CUDA 和 nvcc:使用预处理器在 float 或 double 之间进行选择 [英] CUDA and nvcc: using the preprocessor to choose between float or double

问题描述

问题:

有一个 .h，如果为 c/c++ 或计算能力 >= 1.3 的 cuda 编译，我想将 real 定义为 double.如果为具有计算能力的 cuda 编译 <1.3 然后定义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 未定义或值 CUDA_ARCH130?为什么?

感谢您的宝贵时间.

解决方案

看来你把两件事混为一谈了——nvcc 处理 CUDA 代码时如何区分主机和设备的编译轨迹，以及如何区分 CUDA 和非CUDA代码.两者之间存在细微差别.__CUDA_ARCH__ 回答第一个问题，__CUDACC__ 回答第二个问题.

考虑以下代码片段:

#ifdef __CUDACC__#警告使用 nvcc模板<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__#警告设备代码轨迹#if __CUDA_ARCH__ >120#warning 双精度编译模板 void add<double>(double *, double *, double *);#别的#warning 单精度编译模板无效添加<float>(float *, float *, float *);#别的#warning nvcc 主机代码轨迹#万一#别的#warning 非 nvcc 代码轨迹#万一

这里我们有一个带有 CUDA 架构依赖实例化的模板化 CUDA 内核，一个由 nvcc 引导的主机代码的单独节，以及一个不由 nvcc 引导的主机代码编译节代码>.其行为如下:

$ ln -s cudaarch.cu cudaarch.cc$ gcc -c cudaarch.cc -o cudaarch.ocudaarch.cc:26:2: 警告:#warning 非 nvcc 代码轨迹$ nvcc -arch=sm_11 -Xptxas="-v" -c cudaarch.cu -o cudaarch.cu.ocudaarch.cu:3:2: 警告:#warning using nvcccudaarch.cu:14:2: warning: #warning 设备代码轨迹cudaarch.cu:19:2: 警告:#warning 单精度编译cudaarch.cu:3:2: 警告:#warning using nvcccudaarch.cu:23:2: 警告:#warning nvcc 主机代码轨迹ptxas 信息:为sm_11"编译入口函数_Z3addIfEvPT_S1_S1_"ptxas 信息:使用 4 个寄存器，12+16 字节 smem$ nvcc -arch=sm_20 -Xptxas="-v" -c cudaarch.cu -o cudaarch.cu.ocudaarch.cu:3:2: 警告:#warning using nvcccudaarch.cu:14:2: warning: #warning 设备代码轨迹cudaarch.cu:16:2: 警告:#warning 双精度编译cudaarch.cu:3:2: 警告:#warning using nvcccudaarch.cu:23:2: 警告:#warning nvcc 主机代码轨迹ptxas 信息:为 'sm_20' 编译入口函数 '_Z3addIdEvPT_S1_S1_'ptxas 信息:使用了 8 个寄存器，44 字节 cmem[0]

这里的要点是:

• __CUDACC__ 定义 nvcc 是否指导编译
• __CUDA_ARCH__在编译主机代码时总是未定义，是否由 nvcc 引导
• __CUDA_ARCH__只为nvcc
引导的编译设备代码轨迹定义

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

The problem:

Having a .h, I want to define real to be double if compiling for c/c++ or for cuda with computing capability >= 1.3. If compiling for cuda with computing capability < 1.3 then define real to be float.

After many hours I came to this (which does not work )

#   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

when I compile (note the -arch flag)

nvcc -arch compute_13  -Ilibcutil testFloatDouble.cu 

I get

* 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: ....

I know that files get compiled twice by nvcc. The first one is OK (CUDACC defined and CUDA_ARCH >= 130) but what happens the second time? CUDA_DEFINED but CUDA_ARCH undefined or with value < 130? Why ?

Thanks for your time.

解决方案

It seems you might be conflating two things - how to differentiate between the host and device compilation trajectories when nvcc is processing CUDA code, and how to differentiate between CUDA and non-CUDA code. There is a subtle difference between the two. __CUDA_ARCH__ answers the first question, and __CUDACC__ answers the second.

Consider the following code snippet:

#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

Here we have a templated CUDA kernel with CUDA architecture dependent instantiation, a separate stanza for host code steeered by nvcc, and a stanza for compilation of host code not steered by nvcc. This behaves as follows:

$ 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]

The take away points from this are:

• __CUDACC__ defines whether nvcc is steering compilation or not
• __CUDA_ARCH__is always undefined when compiling host code, steered by nvcc or not
• __CUDA_ARCH__is only defined for the device code trajectory of compilation steered by nvcc

Those three pieces of information are always enough to have conditional compilation for device code to different CUDA architectures, host side CUDA code, and code not compiled by nvccat all. The nvccdocumentation is a bit terse at times, but all of this is covered in the discussion on compilation trajectories.

这篇关于CUDA 和 nvcc:使用预处理器在 float 或 double 之间进行选择的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！