KeplerのKernel呼び出しオーバーヘッド

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

Tesla K20に代表されるKeplerアーキテクチャのGPU、スペックは旧世代のGPUを突き放しているのだが、実性能の向上が意外と小さかったりするケースがしばしば見受けられる。この原因と考えられるものの一つが、Fermi世代に比べて非常に大きい、Kernel呼び出しのオーバーヘッドである。

目次

オーバーヘッドの測定

測定用に、以下のプログラムを用意した。

#include <cstdio>
#include <cstdlib>

static __global__ void
kernel(int *__restrict__ d_a, int x)
{
  *d_a = x;
}

int
main(int argc, char **argv)
{
  int dev = std::atoi(argv[1]);
  cudaSetDevice(dev);
  int *d_a;
  cudaMalloc(reinterpret_cast<void **>(&d_a),
             sizeof(int));
  for (int i = 0; i < 0x1000000; ++i) {
    kernel<<<1, 1>>>(d_a, i);
  }
  int h_a;
  cudaMemcpy(&h_a, d_a, sizeof(int),
             cudaMemcpyDeviceToHost);
  std::printf("%d\n", h_a);
  return 0;
}

int型1つ分確保されたDevice Memoryにintの値を書き込むだけの小さなKernelを224回実行するだけの、シンプルなプログラムである。ところが、このプログラムを以下の環境で実行してみる。

CPU: クアッドコア Xeon W3520 2.67 GHz
GPU 0: Tesla K20X
GPU 1: Tesla C2050
OS: CentOS 6.4
CUDA: CUDA 5.5
GPUドライバ: バージョン319.21開発用ドライバ
$ time ./libcuda 0
16777215

real	1m7.268s
user	1m2.705s
sys	0m2.560s
$ time ./libcuda 1
16777215

real	0m44.276s
user	0m38.143s
sys	0m2.379s

Fermi世代のTesla C2050が44秒で実行するプログラムに、Kepler世代のTesla K20Xは1分07秒も費やしている。

オーバーヘッドの原因の追究

ループ回数を212回に減らしてこのプログラムをnvprofでプロファイルしてみる。

$ nvprof ./libcuda 0
==14239== NVPROF is profiling process 14239, command: ./libcuda 0
4095
==14239== Profiling application: ./libcuda 0
==14239== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 99.94%  5.8544ms      4096  1.4290us  1.3760us  1.9200us  kernel(int*, int)
  0.06%  3.2640us         1  3.2640us  3.2640us  3.2640us  [CUDA memcpy DtoH]
$ nvprof ./libcuda 1
==14252== NVPROF is profiling process 14252, command: ./libcuda 1
4095
==14252== Profiling application: ./libcuda 1
==14252== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 99.95%  3.6785ms      4096     898ns     888ns  1.8200us  kernel(int*, int)
  0.05%  2.0160us         1  2.0160us  2.0160us  2.0160us  [CUDA memcpy DtoH]

Kernel呼び出しのオーバーヘッドが、約500 nsほど増えているようである。しかし、この500 nsを224倍してもせいぜい8秒の差。GPU側のオーバーヘッド増加だけでなく、デバイスドライバやランタイム側でも呼び出しオーバーヘッドが増加しているようである。

仮説

Impress PC Watch後藤弘茂氏の記事であるが、KeplerアーキテクチャのGPUは、CUDA CoreのスケジューリングをGPU上で行っていたFermi世代と比べて、これをCPU上で行うようなアウトソーシングを行っているようである。Kernel呼び出しのオーバーヘッドの増加は、この記事のスケジューリングについての説明とよくつじつまが合う。

結論

Keplerアーキテクチャ向けのGPGPUプログラムを書く際には、このオーバーヘッド増加を見越した対策が必要な場合も在る。CUDA Core数が飛躍的に増加している点と合わせて、多少効率が悪くなっても、複数の処理をまとめた大きなカーネルを作るような工夫で、性能向上を実現できることも在る点を、頭に入れておきたいところである。


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

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

<comments hideform="false" />


Comments

ノート:KeplerのKernel呼び出しオーバーヘッド

個人用ツール