CUDA и nvcc: использование препроцессора для выбора между float или double - PullRequest
Новая
       16

CUDA и nvcc: использование препроцессора для выбора между float или double

6 голосов
/ 10 января 2012

Проблема :

Имея .h, я хочу определить real как double при компиляции для c / c ++ или для cuda с вычислительной способностью> = 1.3. Если компилируется для cuda с вычислительной способностью <1.3, тогда определите real как float. </p>

Через много часов я пришел к этому (что не работает) 

#   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? Почему? </p>

Спасибо за ваше время.

Ответы [ 2 ]

27 голосов
/ 11 января 2012

Возможно, вы путаете две вещи: как различать траектории компиляции хоста и устройства, когда 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 иногда немного лаконична, но все это рассматривается в обсуждении траекторий компиляции.

3 голосов
/ 10 января 2012

На данный момент единственным практическим решением, которое я вижу, является использование пользовательского определения: 


#   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
...