Dynamic Parallelism - undefined reference to __cudaRegisterLinkedBinary linking error while compiling - separate compilation

Question

I got a problem when I try to compile a simple code there are C++ and Cuda code compile in a separated way.

Here's my code

main.cpp:

#include "file.cuh"

int main( void )
{
     test();
     return 0;
}

file.cuh:

void test( void );

file.cu:

#include <cuda.h>
#include <cuda_runtime.h>
#include <cstdio>

#include "file.cuh"

__global__ void printId( void )
{
    printf("Hello from block %d \n", blockIdx.x);
}

__global__ void DynPara( void )
{
    dim3 grid( 2, 1, 1 );
    dim3 block( 1, 1, 1 );

    printId<<< grid, block >>>();
}

void test( void )
{
    dim3 grid( 1, 1, 1 );
    dim3 block( 1, 1, 1 );

    dynPara<<< grid, block >>>();
}

I compile with:

nvcc -arch=sm_35 -lcudadevrt -rdc=true -c file.cu
g++ file.o main.cpp -L<path> -lcudart

And here's the error while compiling:

file.o: In function `__sti____cudaRegisterAll_39_tmpxft_00005b2f_00000000_6_file_cpp1_ii_99181f96()':
tmpxft_00005b2f_00000000-3_file.cudafe1.cpp:(.text+0x1cd): undefined reference to `__cudaRegisterLinkedBinary_39_tmpxft_00005b2f_00000000_6_file_cpp1_ii_99181f96'

os: Red Hat card: K20x

Any idea?

Thanks

Answer 1

This question is pretty much a duplicate of this recent question.

Dynamic parallelism requires relocatable device code linking, in addition to compiling.

Your nvcc command line specifies a compile-only operation (-rdc=true -c).

g++ does not do any device code linking. So in a scenario like this, when doing the final link operation using g++ an extra device code link step is required.

Something like this:

nvcc -arch=sm_35 -rdc=true -c file.cu
nvcc -arch=sm_35 -dlink -o file_link.o file.o -lcudadevrt -lcudart
g++ file.o file_link.o main.cpp -L<path> -lcudart -lcudadevrt

Answer 2

When using CMake, setting CUDA_SEPARABLE_COMPILATION before find_package() enables both relocatable device code compiling and linking:

SET(CUDA_SEPARABLE_COMPILATION ON)
find_package(CUDA QUIET REQUIRED)