动态并行 - 编译时对__cudaRegisterLinkedBinary链接错误的未定义引用 - 单独编译

Question

当我尝试编译一个简单的代码时,我遇到了一个问题,那就是C++和Cuda代码以分开的方式编译.

这是我的代码

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 >>>();
}

我编译:

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

这是编译时的错误:

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卡:K20x

任何的想法？

谢谢

Answer 1

这个问题几乎与最近这个问题重复.

除了编译之外,动态并行性还需要可重定位的设备代码链接.

您的nvcc命令行指定了仅编译操作(-rdc=true -c).

g++不做任何设备代码链接.因此,在这样的场景中,当使用g++ 额外的设备代码链接操作进行最终链接操作时,需要步骤.

像这样的东西:

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

使用 CMake 时，设置CUDA_SEPARABLE_COMPILATIONbeforefind_package()启用可重定位设备代码编译和链接：

SET(CUDA_SEPARABLE_COMPILATION ON)
find_package(CUDA QUIET REQUIRED)