CUDA 和 nvcc : using the preprocessor to choose between float or double

Question

问题 :

如果为 c/c++ 或具有计算能力 >= 1.3 的 cuda 编译，有一个 .h，我想将 real 定义为 double。如果为计算能力 < 1.3 的 cuda 编译，则将 real 定义为 float。

几个小时后，我来到了这个(这不起作用)

# 如果已定义(__CUDACC__)

# 警告 * 为 cuda 制定定义

# 如果已定义(__CUDA_ARCH__)
# 警告 __CUDA_ARCH__ 已定义
＃ 别的
# 警告 __CUDA_ARCH__ 未定义
＃ 万一

# 如果 (__CUDA_ARCH__ >= 130)
# 定义真正的双
# 在 cuda 中使用 double 的警告
# elif (__CUDA_ARCH__ >= 0)
# 定义真正的浮点数
# 在 cuda 中使用 float 的警告
# 警告当 __CUDA_ARCH__ 未定义时，这到底是怎么打印的？
＃ 别的
# 定义真实
# 错误 __CUDA_ARCH__ 的值到底是什么，我该如何打印它
＃ 万一

＃ 别的
# 警告 * 为 c/c++ 定义
# 定义真正的双
# 在 c/c++ 中使用 double 的警告
＃ 万一

当我编译时(注意 -arch 标志)

nvcc -arch compute_13 -Ilibcutil testFloatDouble.cu

我得到

* 为 cuda 制定定义
__CUDA_ARCH__ 已定义
在 cuda 中使用 double

* 为 cuda 制定定义
警告 __CUDA_ARCH__ 未定义
在 cuda 中使用 float 警告
如果 __CUDA_ARCH__ 现在没有定义，这到底是怎么打印的？

架构 i386 的 undefined symbol :
“myKernel(float*, int)”，引用自:....

我知道文件被 nvcc 编译了两次。第一个没问题( CUDACC 定义和 CUDA_ARCH >= 130)但是第二次会发生什么？
CUDA_DEFINED 但是 CUDA_ARCH 未定义或值 < 130？为什么 ？

谢谢你的时间。

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 是否为转向编译

__CUDA_ARCH__ 在编译主机代码时总是未定义的，由 nvcc 引导或不

__CUDA_ARCH__ 仅定义为nvcc

引导的编译的设备代码轨迹

这三条信息总是足以将设备代码条件编译到不同的 CUDA 架构、主机端 CUDA 代码和根本没有被 nvcc 编译的代码。 nvcc 文档有时有点简洁，但所有这些都包含在关于编译轨迹的讨论中。