NVCCの吐くコードを観察する．
ネット探してもなかったんで．
元のコード
#include <stdio.h> #define N 10 __global__ void add(int *a, int *b, int *c) { int tid = threadIdx.x; if (tid < N) c[tid] = a[tid] + b[tid]; } int main(void) { int a[N], b[N], c[N]; int *dev_a, *dev_b, *dev_c; cudaMalloc((void**)&dev_a, N*sizeof(int)); cudaMalloc((void**)&dev_b, N*sizeof(int)); cudaMalloc((void**)&dev_c, N*sizeof(int)); for(int i = 0; i < N; i++){ a[i] = -i; b[i] = i * i; } cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice); add<<<1,N>>>(dev_a, dev_b, dev_c); cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost); for(int i = 0; i < N; i++){ printf("%d + %d = %d\n", a[i], b[i], c[i]); } cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; }
nvcc(Mac)の吐くコード（一部改変）
#include <stdio.h> #define N 10 int main(void) { int a[N], b[N], c[N]; int *dev_a, *dev_b, *dev_c; cudaMalloc((void**)&dev_a, N*sizeof(int)); cudaMalloc((void**)&dev_b, N*sizeof(int)); cudaMalloc((void**)&dev_c, N*sizeof(int)); for(int i = 0; i < N; i++){ a[i] = -i; b[i] = i * i; } cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice); cudaConfigureCall(1, 10) ? ((void)0) : add(dev_a, dev_b, dev_c); cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost); for(int i = 0; i < N; i++){ printf("%d + %d = %d\n", a[i], b[i], c[i]); } cudaFree(dev_a); cudaFree(dev_b); cudaFree(dev_c); return 0; } static const unsigned long long fatbinData[] __attribute__ ((section ("__NV_CUDA,__nv_fatbin")))= { 0x00100001ba55ed50ull,0x0000000000000d99ull,0x0000004001000002ull,0x00000000000003b0ull, /* カーネルコード省略 */ 0x0a0a207d0a3a5f53ull,0x0000000000000000ull }; static const struct { int m; int v; const unsigned long long* d; char* f; } __fatDeviceText __attribute__ ((aligned (8))) __attribute__ ((section ("__NV_CUDA,__fatbin")))= { 0x466243b1, 1, fatbinData, 0}; void __device_stub__Z3addPiS_S_(int *__par0, int *__par1, int *__par2){ if (cudaSetupArgument((void *)(char *)&__par0, sizeof(__par0), (size_t)0UL) != cudaSuccess) return; if (cudaSetupArgument((void *)(char *)&__par1, sizeof(__par1), (size_t)4UL) != cudaSuccess) return; if (cudaSetupArgument((void *)(char *)&__par2, sizeof(__par2), (size_t)8UL) != cudaSuccess) return; { volatile static char *__f; __f = ((char *)((void ( *)(int *, int *, int *))add)); (void)cudaLaunch(((char *)((void ( *)(int *, int *, int *))add))); }; } void add( int *__cuda_0,int *__cuda_1,int *__cuda_2) { __device_stub__Z3addPiS_S_( __cuda_0,__cuda_1,__cuda_2); } static void __sti____cudaRegisterAll_39_tmpxft_00008aa5_00000000_4_main_cpp1_ii_00197a21(void) __attribute__((__constructor__)); static void __sti____cudaRegisterAll_39_tmpxft_00008aa5_00000000_4_main_cpp1_ii_00197a21(void) { __cudaFatCubinHandle = __cudaRegisterFatBinary((void*)&__fatDeviceText); atexit(__cudaUnregisterBinaryUtil); __cudaRegisterFunction(__cudaFatCubinHandle, (const char*)((void ( *)(int *, int *, int *))add), (char*)"_Z3addPiS_S_", "_Z3addPiS_S_", -1, (uint3*)0, (uint3*)0, (dim3*)0, (dim3*)0, (int*)0); }
NVCCの吐くコードでわかったこと
- fatbinData
- CUDAのカーネルコードが入っている．Linuxの場合インラインアセンブラとディレクティブ命令で埋め込まれ，ラベルが貼られている．Macの場合はunsigned long long int型の配列になっている．
- __fatDeviceText
- fatbinDataへのポインタをメンバとしてもつ構造体．
- add
- 元々*.cuの中ではカーネル関数（__global__修飾子で記述された関数）であったが，nvccによってそのまま__device_stub__Z3addPiS_S_関数を呼び出すだけになっている．この関数はmainから直接呼ばれる．
- __device_stub__Z3addPiS_S_
- ここでcudaSetupArgumentとcudaLaunchを呼び出しCUDAのカーネルコードの引数と起動を行う．この関数内にはfatbinDataに関するデータは一切入っていない．よって，別の場所で前もって登録されているCUDAカーネルコードが実行されると考えられる（後述）．
- __sti____cudaRegisterAll_39_tmpxft_00008aa5_00000000_4_main_cpp1_ii_00197a21
- この関数はプロトタイプ宣言に__attribute__( (__constructor__) )と書かれているため，実行時にmain関数の前に呼び出される（GCC拡張*1）．
- __cudaRegisterFatBinary
- カーネルコードを登録する？関数．
- atexit
- stdlib.hに登録されているCの標準ライブラリ関数．プログラムの終了時に第1引数の関数を実行する．
- __cudaUnregisterBinaryUtil
- __cudaRegisterFatBinaryの逆．
- __cudaRegisterFunction
- ここでCUDAのカーネルコードと関数を紐付していると考えられる．注目すべきは引数addである．addは関数だが，関数の中身は関係ない（たぶん）．このaddはどのカーネルコードを呼び出せばいいのかを識別するためのシンボルであり，cudaLaunchのときにこの識別子addを渡すことにより，__cudaregisterFunctionで登録されたカーネルコードが実行される．なので重複しないconst char*型の値なら何でもよい（mainでもOK）．引数の"_Z3addPiS_S_"はfatbinDataから呼び出すべきカーネルコードの先頭を見つけるための索引のようなものである（"_Z3addPiS_S_"自体は関数addのマングリング*2された名前）．実際にfatbinDataの中からこの文字列のASCIIコード列を探索するといくつか出てくる．
カーネルコード呼び出し手順
- addが実行される
- __device_stub__Z3addPiS_S_が実行される
- cudaSetupArgumentで引数を入れる
- cudaLaunchで識別子addにひもづけられたカーネルコードを実行する
cudaLaunchの中身
- addにひもづけられているのは"_Z3addPiS_S_"である．
- fatbinDataから"_Z3addPiS_S_"を探す．
- 見つかった場所からカーネルコードを実行する．
