CUDA 5の新機能(1): Device関数の分割コンパイル

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

去る8月16日に配信された、CUDA 5のリリース候補版。今回は、CUDA 5の新機能の中から、Device関数の分割コンパイルについて、記事を書いてみた。

従来のCUDAの制限

従来のCUDAには、「Kernel内で使うDevice関数は、同じソースコード(includeなどは同じソースコードとみなす)内で実装されている必要が有る。」という制限が有った。たとえば、以下のような2つのコードを用意する。

#include <iostream>

extern __device__ int f(int);

__global__ void
kernel(int *dst, const int *src)
{
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  dst[i] = f(src[i]);
}

int
main()
{
  const dim3 grid(30000);
  const dim3 block(1024);
  size_t size = grid.x * block.x;
  int *h_src, *h_dst, *d_src, *d_dst;
  h_src = new int[size];
  for (ptrdiff_t i = 0; i < size; ++i) {
    h_src[i] = i;
  }
  cudaMalloc(reinterpret_cast<void **>(&d_src), size * sizeof(int));
  cudaMemcpy(d_src, h_src, size * sizeof(int), cudaMemcpyHostToDevice);
  delete[] h_src;
  cudaMalloc(reinterpret_cast<void **>(&d_dst), size * sizeof(int));
  cudaEvent_t start, end;
  cudaEventCreate(&start);
  cudaEventCreate(&end);
  cudaEventRecord(start);
  kernel<<<grid, block>>>(d_dst, d_src);
  cudaEventRecord(end);
  cudaEventSynchronize(end);
  float ms;
  cudaEventElapsedTime(&ms, start, end);
  std::cout << ms << " milliseconds" << std::endl;
  double gflops = 2 * size * sizeof(int) / ms * 1.0e-6;
  std::cout << gflops << " GB/sec" << std::endl;
  h_dst = new int[size];
  cudaMemcpy(h_dst, d_dst, size * sizeof(int), cudaMemcpyDeviceToHost);
  cudaFree(d_src);
  cudaFree(d_dst);
  long long int sum = 0;
  for (ptrdiff_t i = 0; i < size; ++i) {
    sum += h_dst[i];
  }
  std::cout << "CHECKSUM: " << sum << std::endl;
  delete[] h_dst;
  return 0;
}
__device__ int
f(int a)
{
  return a + 1;
}

従来のCUDAでは、後者のソースコードはコンパイルできるが、前者のソースコードはコンパイルできない。Kernelに必要なDevice関数「f」が、前者ではなく後者のソースコードで実装されているからである。

CUDA 5でコンパイルする

CUDA 5では、このような分割コンパイルを行うために、NVCCに-rdc=trueというオプションが追加された。オプションの使用にはCompute Capability 2.0以上でのコンパイルが必要であるため、適宜-arch=sm_20などのオプションを追加で指定したい。

具体的には、

$ nvcc -arch=sm_20 -rdc=true a.cu b.cu

のようにすれば、実行バイナリが得られ、実行結果の速度を確認すると、

1.92019 milliseconds
127.987 GB/sec
CHECKSUM: 471859215360000

と、オーバーヘッドに関して、実用上は問題無く実行できているように思われる。

分割して静的ライブラリを作る

-rdc=trueオプションは、ファイルを一つ一つ独立に中間ファイルにコンパイルして、最後にリンクするやり方でも有効に機能する。特にオブジェクトファイルを生成する-cオプションは有用性が高いためか、-rdc=true -cを一つに省略した-dcというオプションが有る。

$ nvcc -arch=sm_20 a.cu -dc
$ nvcc -arch=sm_20 b.cu -dc
$ nvcc -arch=sm_20 a.o b.o

他のOSでは未検証であるが、Linuxにおいては、上記の方法で生成した.oファイルから、静的ライブラリ.aファイルを作ることができる。最終的にKernelに必要な処理がすべて揃うようなリンクをすれば、.aファイルからでも正しいCUDAの実行バイナリが得られる。

$ ar cr a.a a.o
$ ar cr b.a b.o
$ nvcc -arch=sm_20 a.a b.a

一方、動的ライブラリ.soファイルを作ってリンクしようとすると、最終的に実行バイナリを得るのには失敗するようであるが、筆者が正しい方法を見つけていないだけなのかもしれない。今後この方法を確認でき次第、更新していきたいところである。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:CUDA 5の新機能(1): Device関数の分割コンパイル

個人用ツール