ASCII.technologies 2009年12月号の特集扉より
ASCII.technologiesとアスキー総合研究所、ASCII.jpが合同で開催する技術解説セミナーの第1回「GPUコンピューティング」が3月19日に開催される。その講師を務めるフィックスターズは、PS3のCellやNVIDIA/ATIのGPUなど、演算能力の高いマルチセルプロセッサーを最大限に生かすシステム構築、ソフトウェア開発を手がける技術者集団だ。
本セミナーの前から同社ではGPGPU関連のセミナーやハンズオン（実習形式の勉強会）などを重ねており、業界では知らぬものはない存在――とはいえASCII.jp読者の中にはまだ同社の実績や活動を知らない方もいるだろう。
ここではフィックスターズが執筆担当した、「ASCII.technologies」（アスキー・テクノロジーズ）の2009年12月号に掲載されたGPGPU特集から一部抜粋してご紹介する。なおセミナーの参加者には当該号を基本テキストとして配布する予定だ。
驚異の1TFLOPSオーバーパワーを徹底活用
GPGPUによる並列処理
■ GPGPUとCUDAの基礎知識
GPU（Graphics Processing Unit）の特徴のひとつは、チップ内部のシリコンリソースのほとんどを演算器に割り当てていることにある。よって制御部分が少ないため、同じダイサイズの汎用CPUに比べ、演算処理能力が圧倒的に高くなっている。
またもうひとつの特徴としては、メモリI/O速度が非常に高速であることが挙げられる。VRAMのバンド幅は140GiB/sであり、L1テクスチャキャッシュなら最大480GiB/sの高速なデータ転送が可能だ。これらは汎用CPUのメインメモリへのアクセススピードと比べると、ひと桁以上高速な値となる（図1）。
【図1】 メモリバンド幅のCPUとGPUの比較
このような特徴から、GPUは「データ並列」で動作する「Compute bound」なアプリケーションに対して、極めて高い演算処理速度を発揮する。ここでいうデータ並列とは、多数のデータ要素に対して、複数の計算器がまったく同一のプログラムで処理を行なう並列処理手法のことだ。
また、Compute boundとは、プログラム全体のなかでデータ転送時間に対して演算時間の比率が高い問題のことを指す。たとえば、各データ要素に対して同じプログラムが実行される場合、高度なフロー制御の必要性は低い。そして、メモリI/Oに比べて演算時間の比率が高ければ、多数のスレッドをタスク化して用意しておくことで、多量のデータキャッシュを持たなくてもメモリアクセス遅延を隠蔽できる。GPUのハードウェアアーキテクチャは、まさにそのようなデザインになっているといえる……（続く）
■ プログラミングの基本フロー
CUDAのソースコードがホスト側とデバイス側に分かれるということは、前のパートで述べたとおりだ。ここでは、最初にホスト側プログラミングの流れを見ていくことにしよう。
CUDAにおけるホスト側のプログラミングは、デバイスの制御が中心となる。GPU側で動作するカーネルはCPU側のメモリにアクセスできない。そのため、演算の前にデバイスがアクセスするデータについて、ホスト側のメモリ領域からデバイス側のメモリ領域へあらかじめコピーしておく必要がある。
そこで、まずは使用するデバイスメモリを確保する（リスト1の2行め）。リストの例で、第1引数（deviceMemoryPtr）はデバイスメモリへのポインター、第2引数（size）は確保する領域のサイズとなる。
【図1】 CUDAプログラムの基本フロー
次に、ホスト（CPU）⇔デバイス（GPU）間のデータ通信を行なう（リスト1の5行め）。このとき、第1引数（deviceMemoryPtr）は「コピー先」でデバイスメモリへのポインター、第2引数（hostMemoryPtr）は「コピー元」でホストメモリへのポインターとなる。そして第3引数（size）は、コピーするデータのサイズだ。コピーの方向を示す第4引数の「CudaMemcpyHostToDevice」はマクロ名で、この場合はホストからデバイスへのコピーを意味している。これで、処理すべきデータがデバイスメモリ上に置かれたことになる。
一方、GPU側で処理を行なう演算については、このあと説明するようにカーネルとして演算処理を記述したうえで、リスト1の8行めに示すように呼び出す。この例は、GPUで動作する「kernel_func_name」という関数に引数「arg1, arg2, ………」を与えて、3重ブラケット……（続く）。
リスト1 CUDAにおけるホスト側プログラムの基本プログラム
01 /* デバイスメモリの確保 */
02 cudaMalloc(deviceMemoryPtr, size);
03
04 /* ホストからデバイスへのメモリ領域のコピー */
05 cudaMemcpy(deviceMemoryPtr, hostMemoryPtr, size, cudaMemcpyHostToDevice);
06
07 /* カーネルの呼び出し */
08 kernel_func_name<<<dG, dB>>>(arg1, arg2, …);
09
10 /* デバイスからホストへのメモリ領域のコピー
11 cudaMemcpy(hostMemoryPtr, deviceMemoryPtr, size, cudaMemcpyDeviceToHost);
12
13 /* デバイスメモリのデータ領域の解放 */
14 cudaFree(deviceMemoryPtr);
