nvccがCUDAコードを処理しているときにホストとデバイスのコンパイル軌跡を区別する方法と、CUDAコードと非CUDAコードを区別する方法の2つを混同しているようです。2つの間に微妙な違いがあります。__CUDA_ARCH__最初の質問に__CUDACC__答え、2番目の質問に答えます。
次のコードスニペットについて考えてみます。
#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
これらの3つの情報は、さまざまなCUDAアーキテクチャへのデバイスコード、ホスト側のCUDAコード、およびまったくコンパイルされていないコードの条件付きコンパイルを行うのに常に十分ですnvcc。nvccドキュメントは時々少し簡潔ですが、これらすべてはコンパイルの軌跡に関する議論でカバーされています。