Link kernels together

Question

I have a CUDA kernel in a .cu file and another CUDA kernel in another .cu file. I know that with dynamic parallelism I can call another CUDA kernel from a parent kernel but I'd like to know if there's any way to do this with a child kernel residing in another .cu file.

Answer 1

Yes, you can.

The key is to use separate compilation with device code linking, which is available with nvcc. Since this is already required for usage of dynamic parallelism, there's really nothing new here.

Here's a simple example:

ch_kernel.cu:

#include <stdio.h>

__global__ void ch_kernel(){

  printf("hello from child kernel\n");
}

main.cu:

#include <stdio.h>

extern __global__ void ch_kernel();

__global__ void kernel(){

  ch_kernel<<<1,1>>>();
}

int main(){

  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}

compile with:

nvcc -arch=sm_35 -rdc=true -o test ch_kernel.cu main.cu -lcudadevrt