CUDA跨文件调用__device__函数的CMake配置解决方案

问题场景 #

• 需求：实现main.cu中的__global__函数调用Maths/GMaths.cu中的__device__函数，同时将Maths目录下的代码编译为可复用库，main编译为可执行文件
• 手动添加-rdc=true、-dc等CUDA编译旗标后，触发LNK2019未解析外部符号错误

项目结构 #

main.cu
main.cuh
CMakeLists.txt
Maths/
  AMaths.cu
  AMaths.cuh
  GMaths.cu
  GMaths.cuh
  CMakeLists.txt

错误配置与核心代码 #

初始根目录CMakeLists.txt核心配置 #

set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -rdc=true -lcudadevrt -lcublas_device -dc")

核心代码片段 #

• main.cu中的全局函数：

__global__ void caller(float* vec1, float* vec2, float* out) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    GMaths::Add(vec1[i], vec2[i], out[i]);
}

• Maths/GMaths.cu中的设备函数：

namespace GMaths {
    __device__ void Add(float val1, float val2, float &out) {
        out = val1 + val2;
    }
}

触发的链接错误 #

Main.obj : error LNK2019: unresolved external symbol __cudaRegisterLinkedBinary_5bf867cd_7_Main_cu_cd620806 referenced in function "void __cdecl __sti____cudaRegisterAll(void)"

解决方案 #

移除手动添加的所有CUDA编译旗标，改为针对Main可执行目标开启CMake原生的CUDA_SEPARABLE_COMPILATION属性，修改后的根目录CMakeLists核心配置如下：

set_target_properties(Main PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

原理说明 #

CMake的CUDA_SEPARABLE_COMPILATION属性会自动处理CUDA分离编译的全流程：

1. 为目标文件添加-dc编译选项，生成可重定位的设备代码
2. 在链接阶段自动添加-rdc=true并链接cudadevrt等必要依赖库
3. 避免手动全局设置编译旗标导致的编译/链接参数不匹配问题，从而解决符号未解析的错误

内容的提问来源于stack exchange，提问作者SCP-1762