使用nvcc编译.cu文件时__CUDA__未定义的问题排查

你的问题出在对__CUDA__和__CUDA_ARCH__这两个宏的作用范围理解有误：

• __CUDA__是nvcc在处理所有CUDA相关代码（包括主机端和设备端编译阶段）时都会定义的宏；
• __CUDA_ARCH__仅在nvcc编译设备端代码（比如核函数、__device__函数的实例化）时才会定义，在.cu文件的主机端编译阶段（由系统编译器如gcc/clang处理）是未定义的。

你原来的条件编译#if defined(__CUDA__) && defined(__CUDA_ARCH__)会导致：当nvcc处理.cu文件的主机端代码时，因为__CUDA_ARCH__未定义，foo()的定义被完全屏蔽。而核函数是在主机端声明的，编译器需要看到foo()的模板声明才能进行语法检查和后续的设备端实例化准备，因此会提示“函数不存在”。

正确的条件编译写法 #

拆分声明与定义，分别用不同的宏保护：

// foo.cuh
// 声明部分：仅让nvcc可见（cpp编译器不会定义__CUDA__，因此看不到这段代码）
#ifdef __CUDA__
template <typename T>
__device__ void foo();
#endif

// 定义部分：仅在设备端编译阶段可见
#if defined(__CUDA__) && defined(__CUDA_ARCH__)
template <typename T>
__device__ void foo(){
    // 你的函数实现代码
}
#endif

这样做的好处：

• .cu文件在主机端编译时，能看到foo()的模板声明，编译器可以正常处理核函数中的调用；
• 设备端编译时，会加载foo()的定义并完成实例化；
• .cpp文件包含头文件时，因为__CUDA__未定义，完全看不到CUDA相关代码，不会出现编译错误。

内容的提问来源于stack exchange，提问作者jjcasmar