CUDA和nvcc：使用预处理器在float和double之间进行选择

Question

问题： 我有一个.h文件，如果在c/c++或者CUDA的计算能力 >= 1.3的情况下编译，希望将real定义为double。如果是在CUDA的计算能力 < 1.3的情况下编译，则将real定义为float。

经过多个小时的尝试，我写出了以下代码，但是它不起作用：

# if defined(__CUDACC__)

# warning * making definitions for cuda

# if defined(__CUDA_ARCH__) # warning __CUDA_ARCH__ is defined # else # warning __CUDA_ARCH__ is NOT defined # endif

# if (__CUDA_ARCH__ >= 130) # define real double # warning using double in cuda # elif (__CUDA_ARCH__ >= 0) # define real float # warning using float in cuda # warning how the hell is this printed when __CUDA_ARCH__ is not defined? # else # define real # error what the hell is the value of __CUDA_ARCH__ and how can I print it # endif

# else # warning * making definitions for c/c++ # define real double # warning using double for c/c++ # endif

当我使用以下命令进行编译（注意-arch标志）：

nvcc -arch compute_13 -Ilibcutil testFloatDouble.cu

我得到以下输出：

* making definitions for cuda __CUDA_ARCH__ is defined using double in cuda

* making definitions for cuda warning __CUDA_ARCH__ is NOT defined warning using float in cuda how the hell is this printed if __CUDA_ARCH__ is not defined now?

Undefined symbols for architecture i386: "myKernel(float*, int)", referenced from: ....

我知道文件会被nvcc编译两次。第一次编译是正确的（CUDACC已定义且CUDA_ARCH> = 130），但第二次发生了什么？CUDA_DEFINED已定义，但CUDA_ARCH未定义或值小于130？为什么？

谢谢你的时间。

- cibercitizen1

Answer 1

看起来你可能混淆了两件事——当nvcc处理CUDA代码时如何区分主机和设备编译轨迹，以及如何区分CUDA和非CUDA代码。这两者之间有微妙的区别。__CUDA_ARCH__回答了第一个问题，__CUDACC__回答了第二个问题。

请考虑以下代码片段：

#ifdef __CUDACC__
#warning using nvcc

template <typename T>
__global__ void add(T *x, T *y, T *z)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    z[idx] = x[idx] + y[idx];
}

#ifdef __CUDA_ARCH__
#warning device code trajectory
#if __CUDA_ARCH__ > 120
#warning compiling with double precision
template void add<double>(double *, double *, double *);
#else
#warning compiling with single precision
template void add<float>(float *, float *, float *);
#else
#warning nvcc host code trajectory
#endif
#else
#warning non-nvcc code trajectory
#endif

在这里，我们有一个带有CUDA架构相关实例化的模板化CUDA内核，一个由nvcc控制的主机代码分段，以及一个用于编译未受nvcc控制的主机代码的分段。它的行为如下：

$ ln -s cudaarch.cu cudaarch.cc
$ gcc -c cudaarch.cc -o cudaarch.o
cudaarch.cc:26:2: warning: #warning non-nvcc code trajectory

$ nvcc -arch=sm_11 -Xptxas="-v" -c cudaarch.cu -o cudaarch.cu.o
cudaarch.cu:3:2: warning: #warning using nvcc
cudaarch.cu:14:2: warning: #warning device code trajectory
cudaarch.cu:19:2: warning: #warning compiling with single precision
cudaarch.cu:3:2: warning: #warning using nvcc
cudaarch.cu:23:2: warning: #warning nvcc host code trajectory
ptxas info    : Compiling entry function '_Z3addIfEvPT_S1_S1_' for 'sm_11'
ptxas info    : Used 4 registers, 12+16 bytes smem

$ nvcc -arch=sm_20 -Xptxas="-v" -c cudaarch.cu -o cudaarch.cu.o
cudaarch.cu:3:2: warning: #warning using nvcc
cudaarch.cu:14:2: warning: #warning device code trajectory
cudaarch.cu:16:2: warning: #warning compiling with double precision
cudaarch.cu:3:2: warning: #warning using nvcc
cudaarch.cu:23:2: warning: #warning nvcc host code trajectory
ptxas info    : Compiling entry function '_Z3addIdEvPT_S1_S1_' for 'sm_20'
ptxas info    : Used 8 registers, 44 bytes cmem[0]

这里要点如下：

• __CUDACC__ 定义了是否由 nvcc 控制编译
• 在由 nvcc 或非 nvcc 控制的主机代码编译时，__CUDA_ARCH__ 总是未定义的
• 只有在由 nvcc 控制的设备代码编译轨迹中，__CUDA_ARCH__ 才会被定义

这三个信息足以实现针对不同 CUDA 架构的设备代码的条件编译，主机端 CUDA 代码以及完全不由 nvcc 编译的代码。虽然 nvcc 文档有时有点简略，但所有这些都包含在编译轨迹的讨论中。

- talonmies

---

非常好的贡献，我会保留。不幸的是，对于我的实际问题，我需要检查它是否可行。我们正在移植/扩展一个现有的c/c++库（很多.h和.c文件），用双精度编写，以便可以用于cuda（浮点数和双精度）。我们将double替换为real，并希望有一个一致的real定义，可以是double或float，取决于情况。我们必须考虑如何机械地转换当前的函数头为模板，最重要的是，这是否适用于想要使用纯C的用户。谢谢。 - cibercitizen1

另外，我没有意识到你对add()的定义是在#ifdef __CUDACC__内部。但它也应该对c/c++代码可用以供使用。 - cibercitizen1

最后一点 - 不行。根据定义，使用kernel<<<>>>语法调用cuda内核的设备代码和将调用它的主机代码必须使用nvcc编译。而且内核定义必须对编译的主机和设备轨迹都可用，因为主机端需要一个内部生成的入口存根函数来调用内核。如果你想从常规的C或C++代码中调用内核，你要么需要在.cu文件中使用C/C++包装器，要么使用驱动程序API加载cubin或JIT编译PTX。 - talonmies

我相信当nvcc解析任何代码时，包括主机代码时，__CUDA_ARCH__都会被定义。在主机代码中似乎被定义为0。 - Jared Hoberock

@JaredHoberock：我尊重您在这方面的丰富经验，但我记得在CUDA 3.0 beta期间，当nvcc中的轨迹发生了一些变化时，我与Tim Murray和其他一些人进行了讨论，我们共同得出结论，即在编译主机代码时__CUDA_ARCH__未定义。当然，我们很可能是错的。 - talonmies

Answer 2

目前我看到的唯一实用的解决方案是使用自定义定义：

#   if (!defined(__CUDACC__) ||  defined(USE_DOUBLE_IN_CUDA)) 
#       define real double
#       warning defining double for cuda or c/c++
#   else
#       define real float
#       warning defining float for cuda
#   endif

然后运行：

nvcc -DUSE_DOUBLE_IN_CUDA -arch compute_13  -Ilibcutil testFloatDouble.cu

得到以下两个编译过程的输出：

#warning defining double for cuda or c/c++
#warning defining double for cuda or c/c++

和

nvcc  -Ilibcutil testFloatDouble.cu 

得到以下输出：

#warning defining float for cuda
#warning defining float for cuda

- cibercitizen1