CODEWITHSUNDEEP
c++cuda
Maria
Maria

Reputation: 25

Is it possible for invoke a kernel function within an another kernel function in CUDA?

Is it possible we invoke a __global__ function within another __global__ function which is also a kernel(__global__) in CUDA?

for example:

    __global__ void func()
    {
            .
            .
    }

    __global__ void foo()
    {
            .
            .
            func //this is a "func" function that has defination on the kernel
    }

    int main(void)
    {
            .
            .
            func <<<1, 1 >>>()
            foo <<<1, 1 >>>()
    }

And could it be use any function from thrust library in a __global__ function ?

Upvotes: 2

Views: 1444

Answers (1)

The Vivandiere
The Vivandiere

Reputation: 3201

Compute capability 3.5 and newer hardware support what is called Dynamic Parallelism, which gives them the ability to have kernels launched by running kernels on the GPU without requiring any host API calls.

Older hardware supports functions which can be called from kernels (these are denoted as __device__ instead of __global__) and are executed at thread scope only, so no new kernel is launched.

Since Thrust 1.8 was release, a serial execution policy has been introduced, which allows thrust algorithms to be call by threads within an existing running kernel, much like __device__functions. Thrust should also support dynamic parallelism via the thrust::device execution policy on supported hardware.

Upvotes: 3

Related Questions