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

by

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.

More Related Contents:

  1. What is the purpose of using multiple “arch” flags in Nvidia’s NVCC compiler?
  2. CUDA: How to use -arch and -code and SM vs COMPUTE
  3. How to disable a specific nvcc compiler warnings
  4. Unspecified launch failure on Memcpy
  5. Different CUDA versions shown by nvcc and NVIDIA-smi
  6. Thrust inside user written kernels
  7. Modifying registry to increase GPU timeout, windows 7
  8. Understanding CUDA grid dimensions, block dimensions and threads organization (simple explanation) [closed]
  9. How are 2D / 3D CUDA blocks divided into warps?
  10. What can I do against ‘CUDA driver version is insufficient for CUDA runtime version’?
  11. What is a bank conflict? (Doing Cuda/OpenCL programming)
  12. nvidia-smi Volatile GPU-Utilization explanation?
  13. Can anyone provide sample code demonstrating the use of 16 bit floating point in cuda?
  14. Any particular function to initialize GPU other than the first cudaMalloc call?
  15. Polymorphism and derived classes in CUDA / CUDA Thrust
  16. CUDA: How many concurrent threads in total?
  17. CUDA apps time out & fail after several seconds – how to work around this?
  18. Why is NVIDIA Pascal GPUs slow on running CUDA Kernels when using cudaMallocManaged
  19. How do I select which GPU to run a job on?
  20. multi-GPU basic usage
  21. How to use 2D Arrays in CUDA?
  22. How can I make tensorflow run on a GPU with capability 2.x?
  23. How to implement device side CUDA virtual functions?
  24. How to create a CUDA context?
  25. Why does my CUDA kernel crash (unspecified launch failure) with a different dataset size?
  26. CUDA determining threads per block, blocks per grid
  27. How does CUDA assign device IDs to GPUs?
  28. printf() in my CUDA kernel doesn’t result produce any output
  29. Multiply Rectangular Matrices in CUDA
  30. “invalid configuration argument ” error for the call of CUDA kernel?

Leave a Comment