これからの並列計算のための GPGPU 連載講座 (VI)
様々な GPGPU プログラミング環境
大島聡史
東京大学情報基盤センター

1

はじめに
本連載では GPGPU プログラミング環境として主に CUDA を取り扱ってきた。CUDA は確かに現在主流の GPGPU プログラミング環境であり、今後も主要な GPGPU プログラミング環境と し て 広 く 利 用 さ れ る こ と が 期 待 さ れ る 。一 方 で CUDA 以 外 に も GPGPU プ ロ グ ラ ミ ン グ 環
境 に 関 す る 研 究 や 開 発 は 行 わ れ て い る 。そ こ で 連 載 第 六 回 (最 終 回) の 今 回 は 、CUDA 以 外 の

GPGPU プログラミング環境について解説する。

2

OpenCL
OpenCL は GPGPU を含む様々な並列計算のためのフレームワークであり、Apple 社によって提案された。現在は標準化団体 Khronos Group の作業部会 OpenCL Working Group によって策定されている [1]。CUDA は GPU ベンダーである NVIDIA 社自身の GPU でのみ利用可能である*1 のに 対 し て 、OpenCL は NVIDIA 社 と AMD 社 の GPU を は じ め 、マ ル チ コ ア CPU や DSP(Digital

Signal Processor) など、様々な並列処理ハードウェアをターゲットとしている。特に NVIDIA 社と並ぶ大手 GPU ベンダーである AMD 社は、自社 GPU 向けのプログラミング環境として OpenCLに力を入れている。
本節では OpenCL のプラットフォームモデルや言語仕様について確認した上で、単純なソー
スコードを例示して OpenCL プログラムの解説を行う。なお、この記事は本連載の過去の記事を読んでいる、すなわち CUDA についてある程度の知識を持っている読者を対象とするため、

CUDA と OpenCL の比較にやや重点を置いた解説を行う。

2.1 OpenCL のプラットフォームモデルとメモリモデル
OpenCL の 対 象 と す る プ ラ ッ ト フ ォ ー ム モ デ ル を 図 1 に 示 す 。OpenCL は 階 層 的 な 並 列 性 を
持 つ ハ ー ド ウ ェ ア か ら な る プ ラ ッ ト フ ォ ー ム モ デ ル を 有 し て い る 。す な わ ち 、Host(ホ ス ト 、
汎用 CPU とメインメモリを搭載した計算機) には (複数の)Compute Device(計算デバイス、具体的には GPU カードなど) を搭載することが可能であり、各 Compute Device は (複数の)Compute

Unit(計算ユニット) によって構成されている。さらに、各 Compute Unit は (複数の)Processing
Element(計算エレメント) によって構成されている。
OpenCL の メ モ リ モ デ ル に つ い て は 図 2 の よ う な 構 成 と な っ て い る 。OpenCL で は CUDA における MP と Block および SP(CUDA Core) と Thread のように、Processing Element に対応する

Work-Item(ワークアイテム) および Compute Unit に対応する Workgroup(ワークグループ) という割り当てがなされている。そのうえで、OpenCL のメモリは複数種類のメモリから構成され*1 仕様が公開されているため、別の環境への移植もおそらく不可能ではない

1

図1

OpenCL のプラットフォームモデルており、Work-Item 単位の Private Memory、Workgroup 単位の Local Memory、Compute Device
単位の Global/Constant Memory、そして Host の Host Memory がある。
こ の よ う に OpenCL の モ デ ル を 見 て み る と 、CUDA と 似 て い る こ と に 気 が つ く だ ろ う 。プ
ラ ッ ト フ ォ ー ム モ デ ル を CUDA と 見 比 べ る と 、Compute Unit と MP、Processing Element とSP が同様の位置づけであることがわかる。またメモリモデルについても、Private Memory とRegister、Local Memory と Shared Memory、Global/Constant Memory と Global Memory そし て Host Memory と Host Memory が 同 様 の 位 置 づ け で あ る こ と が わ か る 。そ の た め 、既 にCUDA の知識を得ているユーザにとっては OpenCL のモデルを理解することはそれほど難しくないだろう。
このように OpenCL と CUDA とを比べてみると非常に似ており、OpenCL も GPU に特化したモデル (フレームワーク) であるように見えるかもしれない。OpenCL ではこのモデルで—ただし構成上の一部のメモリ容量が 0 であることもありえるが— 様々な並列計算ハードウェアに対応できるとしている。また、必要条件を引き下げた OpenCL Embedded Proﬁle も策定されており、こちらは組み込み機器における活用が期待されている。

2.2 OpenCL のプログラミングモデルと言語仕様
つづいて、OpenCL のプログラミングモデルと言語仕様について解説する。

OpenCLは、「通常のC/C++で記述されたOpenCL API関数を含むホストコード」から「OpenCL
専用の記述 (指示子や組み込み関数) を用いて記述された計算カーネル関数」を呼び出すというスタイルをとっている。計算カーネル関数呼び出し時に並列度の指定を行うことで、対象計算
カ ー ネ ル が 複 数 の 計 算 デ バ イ ス 上 の 複 数 の Compute Unit お よ び Processing Element 上 で 実 行される。ホストと計算デバイス間のデータ転送については、別途専用の関数 (API) を用いて記2

図 2 OpenCL のメモリモデル

述する必要がある。

OpenCL の プ ロ グ ラ ム 記 述 は C/C++ 言 語 を 拡 張 し た 記 法 で 行 う 。C/C++ 言 語 と 異 な る 点
は、メモリや関数に専用の指示子をつける点である。例えば、GPU(計算デバイス) に実行させる計算カーネルは、 kernel という指示子を付加した関数という形で記述する。また、Global

Memory に配置する変数については、 global という指示子を付加する。これらの記述例は次節のソースコード例を参照していただきたい。
以上のように、具体的な指示子自体は異なるものの、OpenCL プログラムの動作イメージと記述方法は CUDA と非常に似通っている。一方で OpenCL の特徴としては実行時に計算カーネ
ルを生成する (コンパイルする) ことが挙げられる。CUDA ではプログラムコンパイル時に計算カーネルのコンパイルを行い、プログラム実行時にはコンパイル済みの計算カーネルをロードし実行する。これに対して、OpenCL ではプログラム実行時に計算カーネルのコンパイルを行う機構が用いられる。OpenCL ではこの機構により、プログラム作成 (コンパイル) 環境とは異なる実行環境にプログラムを持って行った場合にも実行時にコンパイラが最適なコードを生成することができるため、常に最高の性能が得られるとしている。なお、コンパイル済みの計算カーネルをロードする機構も備えられている。

3

2.3 OpenCL プログラムの例
本節では、実際に OpenCL プログラムソースコードおよび実行手順を見てみることにする。
動 作 確 認 に は CentOS 5.5 x86 64 お よ び CUDA 3.2RC を 導 入 済 み の TeslaC2050 搭 載 PC を 用 いた。残念ながら対象ハードウェアが手元になかったため、AMD 社製 GPU でも同一のプログラ
ム (ソースコード) が利用できるかの確認はできていない。
ソースコード 1 は OpenCL を用いた単純な配列計算プログラム (入力配列の各要素にスカラー値を掛け合わせ、出力配列として返す) におけるホストコード (CPU 側で動作するプログラム)の例である。このプログラムは実際に動くプログラムであり、省略はしておらず、ある程度のエラー処理も記述してある。本プログラムの流れは以下のようになっている。

1. 16 行目〜24 行目 : GPU の初期化処理
2. 25 行目〜28 行目 : コマンドキューの作成
3. 29 行目〜42 行目 : 計算カーネルのコンパイル・ロード
4. 43 行目〜55 行目 : CPU/GPU それぞれのメモリ準備および CPU から GPU へのデータ転送
5. 56 行目〜63 行目 : GPU 演算開始指示および GPU 演算終了待ち
6. 64 行目〜65 行目 : GPU から CPU へのデータ転送
7. 66 行目〜68 行目 : 計算結果の確認
8. 69 行目以降 : 終了処理
ソースコード 1 arraytest.cpp
1
2
3
4
5
6
7
8

#include <oclUtils.h>#deﬁne DATA LENGTH 16
cl context cxGPUContext;
cl kernel kernel;
cl command queue commandQueue;#deﬁne CHK DO(name,o) ciErrNum=o;if(ciErrNum!=CL SUCCESS){printf(name);printf("Ãfailed\n");return(−1);}
9 #deﬁne CHK ERR(name) if(ciErrNum!=CL SUCCESS){printf(name);printf("Ãfailed\n");return(−1);}

10
11 int main(int argc, char∗∗ argv){
12
cl platform id cpPlatform = NULL;
13
cl uint ciDeviceCount = 0;
14
cl device id ∗cdDevices = NULL;
15
cl int ciErrNum = CL SUCCESS;
16// get platform
17
CHK DO("oclGetPlatformID", oclGetPlatformID(&cpPlatform));18// get devices
19
CHK DO("clGetDeviceIDs1", clGetDeviceIDs(cpPlatform, CL DEVICE TYPE GPU, 0, NULL, &
20
21
22
23
24
25
26
27
28

ciDeviceCount));
cdDevices = (cl device id∗)malloc(ciDeviceCount∗sizeof(cl device id));
CHK DO("clGetDeviceIDs2", clGetDeviceIDs(cpPlatform, CL DEVICE TYPE GPU, ciDeviceCount,
cdDevices, NULL));// get context
cxGPUContext = clCreateContext(0, ciDeviceCount, cdDevices, NULL, NULL, &ciErrNum);CHK ERR("clCreateContext");// create command−queue
cl device id device = oclGetDev(cxGPUContext, 0);commandQueue = clCreateCommandQueue(cxGPUContext, device,
CL QUEUE PROFILING ENABLE, &ciErrNum);CHK ERR("clCreateCommandQueue");

4

29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82 }

// program setup
size t program length;
const char∗ source path = "gpu.cl";char ∗source = oclLoadProgSource(source path, "", &program length);if(!source){printf("oclLoadProgSourceÃfailed(%s)\n", source path);return −2000;}// create the program
cl program cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char ∗∗)&source, &program length, &ciErrNum);CHK ERR("clCreateProgramWithSource");free(source);// build the program
CHK DO("clBuildProgram", clBuildProgram(cpProgram, 0, NULL, "-cl-fast-relaxed-math", NULL,
NULL));// Create Kernel
kernel = clCreateKernel(cpProgram, "arraytest", &ciErrNum);CHK ERR("clCreateKernel");// setup data
cl mem d A;
cl mem d R;
ﬂoat∗ h A data = (ﬂoat∗)malloc(sizeof(ﬂoat)∗DATA LENGTH);
for(int i=0; i<DATA LENGTH; i++)h A data[i] = (ﬂoat)(i+1);
ﬂoat∗ h R data = (ﬂoat∗)malloc(sizeof(ﬂoat)∗DATA LENGTH);
d A = clCreateBuﬀer(cxGPUContext, CL MEM READ ONLY|CL MEM COPY HOST PTR, sizeof(ﬂoat)∗DATA LENGTH, h A data, NULL);
d R = clCreateBuﬀer(cxGPUContext, CL MEM WRITE ONLY, sizeof(ﬂoat)∗DATA LENGTH,NULL, NULL);
ﬂoat value = 2.0f;// set args
clSetKernelArg(kernel, 0, sizeof(cl mem), (void∗)&d R);
clSetKernelArg(kernel, 1, sizeof(cl mem), (void∗)&d A);
clSetKernelArg(kernel, 2, sizeof(cl ﬂoat), (void∗)&value);// run kernel
cl event GPUExecution;
size t localWorkSize[] = {4};size t globalWorkSize[] = {DATA LENGTH};
clEnqueueNDRangeKernel(commandQueue, kernel, 1, 0, globalWorkSize, localWorkSize, 0, NULL, &GPUExecution);
clFlush(commandQueue);// sync
clFinish(commandQueue);// bloking readback
clEnqueueReadBuﬀer(commandQueue, d R, CL TRUE, 0, sizeof(ﬂoat)∗DATA LENGTH, h R data,
0, NULL, NULL);// check result
printf("before:Ã"); for(int i=0; i<DATA LENGTH; i++){printf("Ã%.2f", h A data[i]);}printf("\n");printf("afterÃ:Ã"); for(int i=0; i<DATA LENGTH; i++){printf("Ã%.2f", h R data[i]);}printf("\n");// release mem and event
clReleaseMemObject(d A);
clReleaseMemObject(d R);
clReleaseEvent(GPUExecution);// cleanup
ciErrNum |= clReleaseKernel(kernel);ciErrNum |= clReleaseCommandQueue(commandQueue);ciErrNum |= clReleaseProgram(cpProgram);ciErrNum |= clReleaseContext(cxGPUContext);
CHK ERR("release");free(h A data);
free(h R data);
return 0;

GPU の初期化処理については、PlatformID の取得、Device 数の取得および DeviceID の取得、および Context の作成を行っている。もちろん、ソースコード 2 のように数行のコードを追加す
れば搭載されている全 GPU の情報を確認して選択するなどの処理も可能である。やや手順が煩雑ではあるが特別難しい処理ではなく、基本的に一度書いたものを使い回せる記述である。

5

終了処理についても、特筆すべき点はない。
ソースコード 2
1
2
3
4
5
6
7
8
9
10
11
12
13

arraytest.cpp(デバイスの確認)// get context と create command−queue の間に挿入する
size t nDeviceBytes;
CHK DO("clGetContextInfo", clGetContextInfo(cxGPUContext, CL CONTEXT DEVICES, 0,
NULL, &nDeviceBytes));ciDeviceCount = (cl uint)nDeviceBytes/sizeof(cl device id);
if(ciDeviceCount == 0){printf("noÃdevicesÃ(returnÃcodeÃ%i)\n", ciErrNum);return −1;}// check all devices (get device and print the device name)for(unsigned int i = 0; i < ciDeviceCount; ++i){// 変数device にデバイスの情報が得られるため、これを用いて使用する GPU を決めれば良いcl device id device = oclGetDev(cxGPUContext, i);printf("DeviceÃ%d:Ã", i);
oclPrintDevName(LOGBOTH, device); // GPU 名を表示する例
printf("\n");}メモリ確保やデータ転送については、clSetKernelArg による引数設定などにやや手間はかかるが、適切なパラメタとともに API を呼び出すのみで良い。GPU 演算開始指示および終了待ちについても同様に API を用いる。このプログラムではデータ転送に同期転送 (ブロッキング転送) を使用しているが、若干の記述変更で非同期転送も可能であることを付け加えておく。
計算デバイスに処理を行わせる上で注意が必要なのが clEnqueueNDRangeKernel に渡す第 5
引 数 お よ び 第 6 引 数 (localWorkSize と globalWorkSize)、計 算 カ ー ネ ル の 並 列 度 で あ る 。CUDA
では計算カーネルの実行時に Thread と Block の数を指定すると、総数で Thread 数 ×Block 数の並列度で計算が行われた。一方 OpenCL では総数と Work-Item の数、すなわち CUDA におけるThread 数 ×Block 数 と Thread 数 を 指 定 す る 。CUDA プ ロ グ ラ ム を 単 純 に 模 倣 し よ う と す る と勘 違 い す る こ と も あ る (筆 者 は 勘 違 い を し て し ば ら く 悩 ん だ) の で 注 意 が 必 要 で あ る 。な お 各

WorkSize に指定可能な値にはいくつか制限があるが、詳細についてはリファレンスを参照していただきたい。
以上の処理については、本連載の第 2 回で解説した CUDA プログラム (CUDA C) と比べると複雑で手間がかかる記述であり、感覚的には CUDA C よりもむしろ CUDA Driver API に近い。
しかしいずれも処理の流れ自体は CUDA に類似しており、CUDA プログラミングの知識や経験と OpenCL の API 資料さえあれば理解・記述することができるだろう。
一方で OpenCL と CUDA の目立った違いとしては、コマンドキューの作成や計算カーネルのコンパイルが挙げられる。
コマンドキューはホストから計算デバイスに命令を送るための仕組みである。その名の通りキューであり、各計算デバイスは割り当てられたキューに入れられた処理を次々とこなしてい
くことになる。今回は単純なサンプルプログラムであるため、計算カーネルの実行と計算結果のリードバックを行っているのみである。
ち な み に こ の 連 載 中 で は 紹 介 し な か っ た が 、CUDA に も コ マ ン ド キ ュ ー に 近 い も の と し てStream(ストリーム) がある。Stream は GPU による計算と CPU-GPU 間のデータ転送をオーバーラップする際に必要なため、多くのアプリケーションにおいて利用する意義がある。Stream に
ついてはプログラミングガイドやサンプルを参照していただきたい。
カーネルのコンパイルに関する記述は、これまでの CUDA C と類似したプログラム記述と比べて手間のかかる処理であると言わざるを得ない。手順としては以下の通りである。

6

1. ヘッダーファイル(今回のサンプルでは用いていない)やソースファイルをoclLoadProgSource
関数でセット

2. clCreateProgramWithSource 関数でセットしたソースを読み込む
3. clBuildProgram 関数でコンパイルおよびリンク
4. clCreateKernel 関数でカーネル関数を設定
今回のように単一ソースコードからなる単純な計算カーネルを利用する場合はともかく、計算
カーネルの規模が大きくな った場合にはこれらコンパイルに関する一連の処理も伸びてしま
い、さらなる手間となることが懸念される。
つづいて、計算カーネルをソースコード 3 に示す。CUDA と同様に、この計算カーネル (関数)が計算デバイス上の計算エレメントによって一斉に実行されることになる。さらに、具体的な記 述 方 法 こ そ 異 な る も の の 、各 Work-Item ご と に 一 意 に 得 ら れ る ID を 用 い て そ れ ぞ れ 異 な る
データを処理するという基本的な考えも CUDA と同様である。このプログラムは配列の１要素にスカラー値を掛け合わせているだけであるが、デバイスレベルで一意な ID を得ることができる関数 get global id を用いることで各 Work-Item がそれぞれ 1 つずつの配列要素を計算し、全体として配列に対する計算を達成している。
ソースコード 3 gpu.cl
1 #deﬁne DATA LENGTH 16
2
kernel void arraytest( global ﬂoat ∗R,3
int i;
4i = get global id(0);5R[i] = A[i] ∗ value;
6 }global ﬂoat ∗A, ﬂoat value){

計算カーネルに見られる指示子の違いについては、関数や変数に対する指示子が CUDA と若干異なるのはともかく、関数の引数に対しても指示子が必要な点は CUDA との明確な違いである。CUDA は引数に指示子が無いため記述がより簡潔である反面、メモリの扱いがコンパイラ
任 せ の 部 分 が あ り 、特 に 初 期 の CUDA コ ン パ イ ラ で は 変 数 が 実 際 に ど の メ モ リ に 配 置 さ れ る
かわからない旨の警告がしばしば見受けられた。(そして、実行時エラーに見舞われることが
多 々 あ っ た 。) 一 方 で OpenCL は 変 数 の メ モ リ 配 置 を よ り 細 か く 指 定 で き る こ と に な る た め 、プログラマが意図したとおりに記述・最適化が行いやすいケースもあるだろう。
最後に、実際に OpenCL プログラムをコンパイル・実行する際の手順について確認しておく。

CUDA プログラムをコンパイルする際には専用のコンパイラ nvcc を利用した。一方で OpenCLについては専用のコンパイラが存在せず、図 3 に示すように gcc(g++) 等のコンパイラを用いてライブラリ指定を行いつつ通常のコンパイルを行えばよい。もちろん、実行時に計算カーネル
をコンパイルする都合上、実行時に適切な位置に計算カーネルのソースコードが配置されている必要がある。また、計算カーネルに文法エラーなどがあった場合にはプログラム中のカーネ
ルコンパイルでエラーすることになる。
以上、駆け足ではあるが OpenCL の概要と使い方を説明した。OpenCL も CUDA 同様に、正しく動くプログラムを作成するだけならそう難しくはないが、CUDA(CUDA C) と比べると記述
量が多く手間がかかってしまう。CUDA Driver API に対する CUDA C、もしくは OpenGL における GLUT のようなプログラム記述を容易にするライブラリ (ラッパー) が普及すれば、OpenCLを利用するプログラマの負担も減るだろう。また、正しく動くだけならともかく、高速なプロ

7

$ ls
arraytest.cpp gpu.cl$ g++ -O3 -m64 -o arraytest arraytest.cpp -lOpenCL \-I/path_to_cudasdk_3.1/OpenCL/common/inc -I/path_to_cudasdk_3.1/shared/inc \-L/path_to_cudasdk_3.1/OpenCL/common/lib -L/path_to_cudasdk_3.1/shared/lib \-loclUtil_x86_64 -lshrutil_x86_64$ lsarraytest arraytest.cpp gpu.cl$ ./arraytest
before: 1.00 2.00 3.00 4.00 5.00 6.00 7.00 8.00 9.00 10.00 (以下省略)
after : 2.00 4.00 6.00 8.00 10.00 12.00 14.00 16.00 18.00 20.00 (以下省略)
図 3 OpenCL プログラムのコンパイルと実行の例

グラムを作成するには知識と手間が必要である。さらに、同じプログラムが複数のハードウェ
ア で 実 行 で き る と は い え 、同 じ プ ロ グ ラ ム が 複 数 の ハ ー ド ウ ェ ア で 高 速 に 動 作 す る と は 限 らず、ハードウェアに合わせた最適化プログラミングが必要になることに留意せねばならない。

3

その他の GPGPU プログラミング環境

3.1 既存の並列化プログラミング環境を利用した GPGPU プログラミング環境
GPGPU が登場するよりも前から、様々な並列処理プログラムが様々な手段で記述されてきた 。た と え ば C/C++ 言 語 を 用 い た プ ロ グ ラ ミ ン グ に お い て は 、ス レ ッ ド 並 列 化 に pthread やOpenMP が、またプロセス並列化に MPI が長い間利用されてきている。GPGPU プログラミン
グも並列化プログラミングの一種ととらえられることから、既存の並列化プログラミング環境を用いて GPGPU プログラミングが行えれば、アプリケーションプログラマの手間や負担が軽減されることが期待できる。
そこで筆者らは、OpenMP 指示子を用いて記述されたプログラムから CUDA プログラムを作成するプログラム変換ツール OMPCUDA[3] の作成を行っている。OMPCUDA は、OpenMP 指
示子の挿入された対象プログラムから parallel for 指示子 (for ループの並列化を意味する指示子)
で指定された for ループを抽出し、GPU 上で実行されるカーネルに変換したうえで、カーネルが正しく GPU 上で実行できるように CPU-GPU 間のデータ転送等の処理を挿入する、という処理を行う。OpenMP においてはループの各処理を各スレッドが並列に実行するが、OMPCUDA
においてはループの各処理を GPU 上の SP が並列に実行することになる (図 4)。

OMPCUDA の実装は、既存の OpenMP 処理系である OMNI OpenMP compiler(以下 OMNI)[2]を拡張して行っている。OMNI の構造および OMPCUDA との関係を図 5 に示す。OMNI が複数の言語に対するフロントエンドや中間表現を操作するためのツールキットを備えているため、

OMPCUDA は主に中間表現レベルでの実装で機能を達成している。現在のところ OMPCUDA
は入力されたソースコードをシンプルに CUDA 化しており、GPU 上で高い性能が得られるような書き換えは行えていない。またベースとなっている OMPCUDA も最近の OpenMP や C/C++
および Fortran の仕様に対応していない。そのため OMPCUDA は実用的なアプリケーションにおいて利用するには力不足であり、今後さらなる実装が必要である。

8

図4

OpenMP プログラムの実行イメージと OMPCUDA の関係

図 5 OMNI と OMPCUDA の関係

9

OMPCUDA 以 外 に も 、例 え ば Lee ら [5] が 同 様 に OpenMP を 用 い て CUDA を 利 用 す る ソ フ ト
ウェアの開発を行っている。また OpenMP ではないが、指示文を用いた GPGPU プログラミン
グ環境としては PGI 社のアクセラレータコンパイラ [4] や、CAPS 社の HMPP[6] などがある。

3.2 様々なプログラミング環境から GPGPU を利用するためのソフトウェア
CUDA と OpenCL は C/C++ 言語を拡張した言語によるプログラミング環境を提供している。
C/C++ 言語は確かに利用者が多く普及している開発環境であるが、HPC の分野では C/C++
以外に Fortran も多く利用されており、また他の分野では Java, Perl, Python, Ruby, .NET など
様々なプログラミング環境 (言語) が利用されている。そのため、これらの様々な使い慣れたプ
ログラミング環境から CUDA や OpenCL を使いたいという要求も多く、実際に様々なプログラ
ミング環境から GPU を利用するためのソフトウェアが開発されている。本節ではこうしたい
わゆる「言語バインディング」の例をいくつか紹介する。

PyCUDA[7], PyOpenCL[8]PyCUDA および PyOpenCL は Python から CUDA や OpenCL を利用するためのバインディン
グ で あ る 。い ず れ も Python プ ロ グ ラ ム 中 に 文 字 列 と し て CUDA や OpenCL の プ ロ グ ラ ム を 記述しておき、提供されているクラスを介して CPU-GPU 間のデータ転送や GPU の制御を行う。
そのため、Python の言語仕様で記述することができるのはホスト側で動作するプログラムのみとなる。ベースとなるライブラリは C++ で記述されている。

Ruby-OpenCL[9]Ruby-OpenCL も PyOpenCL と同様に、Ruby プログラム中に文字列として OpenCL 計算カー
ネルを記述しておき、ラッパー関数を用いて GPU にアクセスするタイプの言語バインディングである。上記の Python バインディングと同様に、Ruby の言語仕様で記述することができるのはホスト側で動作するプログラムのみとなる。

JCuda[10]JCuda は Java から CUDA を利用するためのバインディングである。上記の Python/Ruby バイ
ンディングとは異なり、別途作成した GPU バイナリ (PTX) を呼び出して利用するためのラッパーと、CUDA が提供する各種ライブラリ (CUBLAS, CUFFT など) を呼び出すための API が主体となっている。

CUDA Fortran[11]CUDA Fortran は Fortran から CUDA を利用するためのバインディングであり、商用コンパイ
ラ で あ る PGI コ ン パ イ ラ に 含 ま れ て い る 。CUDA Fortran は C/C++ と Fortran が 似 て い る (単純に置き換えられる記述が多い) こともあり、実際のプログラム記述もほぼ CUDA C の記述をFortran の仕様で置き換えたものと言える。
以 上 い く つ か の 言 語 バ イ ン デ ィ ン グ を 挙 げ た が 、基 本 的 に は 計 算 カ ー ネ ル の 記 述 に CUDA
や OpenCL を使う必要があり、「様々なプログラミング言語」のみで GPU を活用することは難10

しいのが現状である。逆に言えば、GPU カーネル以外は使い慣れた様々な言語で記述できる、使い慣れた言語で記述したプログラムの中で計算量が多い特定の部分のみを GPU に任せるといった使い方に適していると言えるだろう。

3.3 ストリーミング言語
GPGPU の発展とは独立した平行処理/並列処理のパラダイムとして、ストリーミング処理およびそれを記述するストリーミング言語の研究が行われてきた。GPU による高性能並列計算の 可 能 性 と プ ロ グ ラ ミ ン グ の 難 し さ が 知 ら れ る よ う に な る と 、ス ト リ ー ミ ン グ 言 語 を 用 い た

GPU プログラミングの研究が注目された。残念ながら CUDA が登場したことでストリーミン
グ言語を用いた GPU プログラミングへの注目度は下がってしまったが、本節では GPU に対応したストリーミング言語のいくつかを簡単に紹介する。
ストリーミング処理は、計算対象のデータをストリーム (入力ストリーム/出力ストリーム)、
計算内容をカーネルと定義し、プロセッサがカーネルを次々にストリームへと作用させていくという概念を持っている。これに対応するストリーミング言語も、言語仕様としてストリームとカーネルを明示的に記述する言語仕様を採用している。なお、CUDA もストリーミング言語の一種と見なされることがある。
特に 2005 年頃には GPU に対応した複数のストリーミング言語が登場した。当時注目された
ストリーミング言語の例としては、BrookGPU[12] や RapidMind[14] などが挙げられる。

BrookGPU は米 Stanford University で開発されたオープンソースのストリーミング言語である。図 6 に BrookGPU を用いた配列加算プログラムの例を示す。この例ではカーネル (kfunc) とストリーム (streamRead,streaWrite) が明示的に示されていることがよくわかる。図中のカーネ
ルでは 2 つの変数を加算して別の変数への代入を行っているが、各変数は配列であり、実行時には配列の各要素に対する演算が GPU 上の多数の演算器によって独立に並列処理される。

BrookGPU は カ ー ネ ル を 実 行 す る バ ッ ク エ ン ド と し て DirectX, OpenGL, CPU, CTM(Closeto the Metal、AMD が当時プッシュしていた開発環境) に対応することで、単一ソースコードから様々なハードウェアに対応した。後に AMD 社は自社 GPU 向けの開発環境として BrookGPU
を元にした Brook+[13] を提供している。

RapidMind は University of Waterloo で開発されたオープンソースのメタプログラミング言語
Sh[15] を元に加 RAPIDMIND 社が開発したストリーミング言語である。図 7 に RapidMind を用いたプログラムの例を示す。RapidMind も BrookGPU と同様にカーネルとカーネルへの入出力を明示的に指定する一方で、カーネルの記述方法はより RapidMind に特有の記述になっている。

PetaFLOPS・ExaFLOPS という高性能な計算環境の時代を迎えて、また GPGPU を含めたヘテロジニアス環境の活用が重要な時代を迎えて、プログラムをどのように記述するか、何を用いて記述するかは改めて難しい問題として認知されている。ストリーミング言語が再度注目されることになるかはわからないが、各計算をどこで行うかやデータの配置・移動をどうするかを記述するという考え方は今後の並列化プログラミング言語に引き継がれるだろう。

以上、第六回 (最終回) の今回は本連載でこれまでに扱ってきた CUDA とは違う様々なプログ
ラミング環境についての解説を行った。現在の GPGPU においては CUDA が非常に大きな影響

11// カーネル
kernel void kfunc (float x<>,float y<>, out float z<>) {z = x + y;}
int main() {float a<100>;
float b<100>;
float c <100>;// 入力ストリームの指定
streamRead(a, data1);
streamRead(b, data2);// カーネルの実行
kfunc(a, b);// 出力ストリームの指定
streamWrite(c, result);
return 0;}

int main() {// カーネル
Program kfunc=BEGIN {In<Value3f>x, y;
Out<Value3f>z;z = x + y;
} END;// 入出力ストリームの指定
Array<1, Value3f> a(512);
Array<1, Value3f> b(512);
Array<1, Value3f> c(512);// カーネルの実行
c = kfunc(a, b);
return 0;}

図7

図 6 BrookGPU のソースコード

RapidMind のソースコード

力を持っているのは事実であり、今後もしばらく大きな影響力を維持することは間違いないだ
ろう。一方で、OpenCL も NVIDIA 社製の GPU 以外でも、さらには GPU 以外でも利用可能であることから注目されており。さらに C/C++ 言語以外の様々な方法で GPGPU を利用できることは、GPGPU がさらに広く利用される上で重要である。
この連載は、年々注目と影響力の高まっている GPGPU に関する情報を本センターのユーザに対して提供することを目的として一年にわたって続けてきた。本稿を執筆している 2010 年 11 月
上旬現在、既に中国のスーパーコンピューター Tianhe-1A (2009 年 11 月の TOP500 で RadeonHDを用いて 5 位を達成した「天河一号」の更新版) が 14,336 個の CPU と 7,168 枚の TeslaM2050 を駆使して LINPACK ベンチマークで 2PFLOPS を突破し、次回の TOP500(2010 年 11 月版) の頂点に立つことが確実だと報じられている。また東工大の TSUBAME も 4,224 枚の TeslaM2050 を搭載した TSUBAME 2.0 にリプレイスされて稼働を開始しており、TOP500 および Green500 の上位
ラ ン ク イ ン が 期 待 さ れ て い る 。こ れ ら の 結 果 は 本 記 事 が 公 開 さ れ る 頃 に は 判 明 し て い る だ ろ
う 。GPGPU は 現 在 も 急 速 に 発 展 し て い る 技 術 で あ り 、今 後 も さ ら に GPU の 性 能 を 活 用 す る
スーパーコンピューター「GPU スパコン」は増加するだろう。本センターは現時点で GPU を搭載した計算資源の提供を行っていないが、現在の動向からすれば本センターが近い将来 GPU ス
パコンを提供することも考えられる。来るべき GPU スパコンに向けて、という意味も込めて、今後も本スパコンニュースにて GPGPU に関する情報を不定期に提供していく予定である。

参考文献[1] OpenCL - The open standard for parallel programming of heterogeneous systems, http://www.khronos.org/opencl/[2] M.Sato, S.Satoh, K.Kusano, and Y.Tanaka. Design of OpenMP Compiler for an SMP Cluster.

12

In EWOMP

99, pp. 32―39, 1999.

[3] 大島聡史, 平澤将一, 本多弘樹. OMPCUDA : GPU 向け OpenMP の実装. HPCS2009 2009 年
ハイパフォーマンスコンピューティングと計算科学シンポジウム, pp.131―138, 2009.

[4] PGI. PGI Accelerator Compilers, http://www.pgroup.com/resources/accel.htm[5] Seyong Lee, Seung-Jai Min, Rudolf Eigenmann. OpenMP to GPGPU: a compiler framework for automatic translation and optimization. Proceedings of the 14th ACM SIGPLAN
symposium on Principles and practice of parallel programming, pp.101-110, 2009.
[6] CAPS. HMPP Workbench, http://www.caps-entreprise.com/fr/page/index.php?id=49\&p\ p=36
[7] PyCUDA, http://mathema.tician.de/software/pycuda[8] PyOpenCL, http://mathema.tician.de/software/pyopencl[9] Ruby-OpenCL, http://ruby-opencl.rubyforge.org/[10] jCuda, http://www.jcuda.org/[11] CUDA Fortran, http://www.pgroup.com/resources/cudafortran.htm[12] Ian Buck, Tim Foley, Daniel Horn, Jeremy Sugerman, Kayvon Fatahalian, Mike Houston,and Pat Hanrahan. Brook for GPUs: Stream Computing on Graphics Hardware. SIGGRAPH
2004, 2004.
[13] AMD. Brook+. SC07 BOF Session presentation, November 2007.
[14] Michael D. McCool. Data-Parallel Programming on the Cell BE and the GPU using the
RapidMind Development Platform. In GSPx Multicore Applications Conference, 2006.
[15] Michael McCool and Stefanus Du Toit. Metaprogramming GPUs with Sh. A K Peters Ltd,2004.

13