6

问题：

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

几个小时后，我来到了这个（这不起作用）

# 如果定义（__CUDACC__）

# 警告 * 为 cuda 定义

# 如果定义（__CUDA_ARCH__）
# 警告 __CUDA_ARCH__ 已定义
＃ 别的
# 警告 __CUDA_ARCH__ 未定义
＃ 万一

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

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

当我编译时（注意 -arch 标志）

nvcc -arch compute_13 -Ilibcutil testFloatDouble.cu

我明白了

* 为 cuda 定义
__CUDA_ARCH__ 已定义
在 cuda 中使用双精度

* 为 cuda 定义
警告 __CUDA_ARCH__ 未定义
在 cuda 中使用浮点数发出警告
如果 __CUDA_ARCH__ 现在没有定义，这到底是怎么打印出来的？

架构 i386 的未定义符号：
  “myKernel(float*, int)”，引用自：....

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

谢谢你的时间。

4

2 回答 2

31

似乎您可能混淆了两件事 - 当 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文档有时有点简洁，但所有这些都包含在关于编译轨迹的讨论中。

于 2012-01-10T20:17:16.237 回答
3

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

# if (!defined(__CUDACC__) || 定义(USE_DOUBLE_IN_CUDA))
# 定义真正的双精度
# 警告为 cuda 或 c/c++ 定义双精度
＃ 别的
# 定义真正的浮点数
# 警告为 cuda 定义浮点数
＃ 万一

接着

nvcc -DUSE_DOUBLE_IN_CUDA -arch compute_13 -Ilibcutil testFloatDouble.cu

因为它为两个编译输出：

#warning 为 cuda 或 c/c++ 定义双精度
#warning 为 cuda 或 c/c++ 定义双精度

和

nvcc -Ilibcutil testFloatDouble.cu

做

#warning 为 cuda 定义浮点数
#warning 为 cuda 定义浮点数
于 2012-01-09T23:40:05.687 回答