CUDA и nvcc: использование препроцессора для выбора между float или double
Проблема :
Имея .h, я хочу определить real, чтобы быть double, если компилировать для c / c ++ или для cuda с вычислительной способностью> = 1.3. Если компиляция для cuda с вычислительной способностью <1.3, то определите real для float.
После многих часов я пришел к этому (что не работает)
- Xcode 4 не может найти общедоступные файлы заголовков из зависимости статической библиотеки
- Как-то зарегистрировать мои classы в списке
- #define в Java
- Как преобразовать переменную типа enum в строку?
- Должен ли я использовать #include в заголовках?
# если определено (__ CUDACC__) # предупреждение * создание определений для cuda # если определено (__ CUDA_ARCH__) # предупреждение __CUDA_ARCH__ определено # else # предупреждение __CUDA_ARCH__ НЕ определено # endif # if (__CUDA_ARCH__> = 130) # define real double # предупреждение с использованием double in cuda # elif (__CUDA_ARCH__> = 0) # define real float # предупреждение с использованием float в cuda # предупреждение, как, черт возьми, это напечатано, когда __CUDA_ARCH__ не определен? # else # define real # ошибка, черт возьми, значение __CUDA_ARCH__ и как я могу ее распечатать # endif # else # предупреждение * создание определений для c / c ++ # define real double # предупреждение с использованием double для c / c ++ # endif
когда я компилирую (отметьте флаг -arch)
nvcc -arch compute_13 -Ilibcutil testFloatDouble.cu
я получил
* создание определений для cuda Определяется __CUDA_ARCH__ использование двойного в кудах * создание определений для cuda предупреждение __CUDA_ARCH__ НЕ определено предупреждение с использованием float в cuda как, черт возьми, это напечатанное, если __CUDA_ARCH__ не определено сейчас? Неопределенные символы для архитектуры i386: «myKernel (float *, int)», на который ссылаются: ....
Я знаю, что файлы скомпилируются дважды nvcc. Первый – ОК ( CUDACC определено и CUDA_ARCH > = 130), но что происходит во второй раз? CUDA_DEFINED, но CUDA_ARCH undefined или со значением <130? Зачем ?
Спасибо за ваше время.
Кажется, вы можете объединить две вещи: как различать траектории компиляции хоста и устройства, когда nvcc обрабатывает код CUDA и как различать код CUDA и не CUDA. Между ними есть тонкая разница. __CUDA_ARCH__ отвечает на первый вопрос, а __CUDACC__ отвечает на второй.
Рассмотрим следующий fragment кода:
#ifdef __CUDACC__ #warning using nvcc template __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 *); #else #warning compiling with single precision template void add(float *, float *, float *); #else #warning nvcc host code trajectory #endif #else #warning non-nvcc code trajectory #endif
Здесь у нас есть шаблонное kernel 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 немного краткая, но все это рассматривается в обсуждении траекторий компиляции.
На данный момент единственным практическим решением, которое я вижу, является использование пользовательского определения:
# if (! defined (__ CUDACC__) || определено (USE_DOUBLE_IN_CUDA)) # define real double # предупреждение, определяющее double для cuda или c / c ++ # else # define real float # предупреждение, определяющее float для cuda # endif
а потом
nvcc -DUSE_DOUBLE_IN_CUDA -arch compute_13 -Ilibcutil testFloatDouble.cu
Поскольку он выводит для двух компиляций:
#warning определение double для cuda или c / c ++ #warning определение double для cuda или c / c ++
а также
nvcc -Ilibcutil testFloatDouble.cu
делает
#warning определение float для cuda #warning определение float для cuda