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

Question

问题：

有一个.h，如果为c/c++或计算能力>= 1.3的cuda编译，我想将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？为什么 ？

感谢您抽出时间。

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__总是< /em> 编译主机代码时未定义，由 nvcc 引导或不
• __CUDA_ARCH__ 仅针对由 nvcc 引导的编译设备代码轨迹定义nvcc

这三条信息始终足以对不同 CUDA 架构的设备代码、主机端 CUDA 代码以及根本未由 nvcc 编译的代码进行条件编译。 nvcc 文档有时有点简洁，但所有这些都包含在有关编译轨迹的讨论中。

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