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

看起来你可能会混淆两件事 - 如何在nvcc处理CUDA代码时区分主机和设备编译轨迹，以及如何区分CUDA和非CUDA代码。两者之间存在细微差别。 __CUDA_ARCH回答第一个问题，CUDACC__回答第二个问题。

请考虑以下代码段：

#ifdef __CUDACC__

#warning using nvcc

template

__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 *);

#else

#warning compiling with single precision

template void add(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代码和未由nvccat编译的代码。 nvcc文档有时候有点简洁，但所有这些都在编译轨迹的讨论中有所涉及。