読者です 読者をやめる 読者になる 読者になる

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コード列を探索するといくつか出てくる.

カーネルコード呼び出し手順

  1. addが実行される
  2. __device_stub__Z3addPiS_S_が実行される
  3. cudaSetupArgumentで引数を入れる
  4. cudaLaunchで識別子addにひもづけられたカーネルコードを実行する

cudaLaunchの中身

  1. addにひもづけられているのは"_Z3addPiS_S_"である.
  2. fatbinDataから"_Z3addPiS_S_"を探す.
  3. 見つかった場所からカーネルコードを実行する.

*1:http://www4.big.or.jp/~kanai/MT/archives/000507.html

*2:http://en.wikipedia.org/wiki/Name_mangling