CUDA e nvcc: usando o pré-processador para escolher entre float ou double

O problema :

Tendo um .h, eu quero definir real para ser duplo se compilar para c / c ++ ou cuda com capacidade de computação> = 1,3. Se compilar para cuda com capacidade de computação <1.3, defina real como float.

Depois de muitas horas eu vim para isso (o que não funciona)

 # se definido (__ CUDACC__)

 # warning * fazendo definições para cuda

 # se definido (__ CUDA_ARCH__)
 # aviso __CUDA_ARCH__ está definido
 # outro
 # warning __CUDA_ARCH__ NÃO está definido
 # fim se

 # if (__CUDA_ARCH__> = 130)
 # define duplo real
 # aviso usando o dobro na cuda
 # elif (__CUDA_ARCH__> = 0)
 # define flutuador real
 # aviso usando float em cuda
 # warning how hell é impresso quando __CUDA_ARCH__ não está definido?
 # outro
 # define real 
 # erro que diabos é o valor de __CUDA_ARCH__ e como posso imprimi-lo
 # fim se

 # outro
 # warning * fazendo definições para c / c ++
 # define duplo real
 # aviso usando o dobro para c / c + +
 # fim se

quando eu compilar (observe o sinalizador -ar)

 nvcc -arch compute_13 -Ilibcutil testFloatDouble.cu 

eu recebo

 * fazendo definições para cuda
 __CUDA_ARCH__ está definido
 usando o dobro na cuda

 * fazendo definições para cuda
 aviso __CUDA_ARCH__ NÃO está definido
 aviso usando float em cuda
 como diabos isso é impresso se __CUDA_ARCH__ não está definido agora?

 Símbolos indefinidos para arquitetura i386:
   "myKernel (float *, int)", referenciado de: ....

Eu sei que os arquivos são compilados duas vezes pelo nvcc. O primeiro é OK ( CUDACC definido e CUDA_ARCH > = 130) mas o que acontece na segunda vez? CUDA_DEFINED mas CUDA_ARCH indefinido ou com valor <130? Por quê ?

Obrigado pelo seu tempo.

Parece que você pode estar confundindo duas coisas – como diferenciar entre as trajetórias de compilation do host e do dispositivo quando o nvcc está processando o código CUDA, e como diferenciar entre o código CUDA e o não-CUDA. Há uma diferença sutil entre os dois. __CUDA_ARCH__ responde a primeira pergunta e __CUDACC__ responde a segunda.

Considere o seguinte trecho de código:

 #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 

Aqui temos um kernel CUDA modelo com instanciação dependente de arquitetura CUDA, uma estrofe separada para código hospedeiro steeered por nvcc , e uma estrofe para compilation de código host não direcionada por nvcc . Isso se comporta da seguinte maneira:

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

Os pontos de tirar disto são:

  • __CUDACC__ define se o nvcc está orientando a compilation ou não
  • __CUDA_ARCH__ é sempre indefinido ao compilar o código do host, orientado por nvcc ou não
  • __CUDA_ARCH__ é definido apenas para a trajetória de código de compilation do código dirigido por nvcc

Essas três informações são sempre suficientes para compilar condicionalmente o código do dispositivo para diferentes arquiteturas CUDA, código CUDA do lado do host e código não compilado pelo nvcc . A documentação do nvcc é um pouco sucinta às vezes, mas tudo isso é abordado na discussão sobre trajetórias de compilation.

No momento, a única solução prática que vejo é usar uma definição personalizada:


 # if (! defined (__ CUDACC__) || definido (USE_DOUBLE_IN_CUDA)) 
 # define duplo real
 # warning definindo double para cuda ou c / c ++
 # outro
 # define flutuador real
 # aviso definindo float para cuda
 # fim se

e depois

 nvcc -DUSE_DOUBLE_IN_CUDA -arch compute_13 -Ilibcutil testFloatDouble.cu

Como ele produz as duas compilações:

 #warning definindo o dobro para cuda ou c / c ++
 #warning definindo o dobro para cuda ou c / c ++

e

 nvcc -Ilibcutil testFloatDouble.cu 

faz

 #warning definindo float para cuda
 #warning definindo float para cuda