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

1

はじめに
CUDA プログラムの最適化には、既存の CPU 向けの最適化とは異なる様々な要素が存在する。連載第三回である今回は、CUDA プログラムの性能に大きな影響を与えるいくつかの項目に注目し、CUDA における最適化プログラミングの方法を解説する。

CUDA の開発環境は不定期に更新が行われており、またハードウェアの世代によって利用可能な機能には違いがある。本稿執筆時点における開発環境 (CUDA Toolkit および CUDA SDK)の最新バージョンは 3.0 であり、最新のハードウェア世代ナンバー (Compute Capability) は 2.0である*1 。CUDA プログラムの最適化においてはこれらのバージョン番号が大きな影響を持つ。
最新の環境 (開発環境およびハードウェア環境) ほど新たな最適化技術が利用可能になったり、以前の環境では性能的に問題があるプログラムでも高速に実行可能となっていたりすることがある。開発環境のバージョンは新たな開発環境をインストールすることで上げることが可能である一方、ハードウェア環境のバージョンは GPU に強く依存する (GPU を変更しないと上げることができない) ので注意が必要である。
本稿で用いている実験環境は表 1 の通りである。本稿では多くの環境で共通に利用可能な最適化方法を紹介するよう努めるが、各方法により得られる最適化効果が環境ごとに異なる可能性があることをあらかじめ断っておく。
表1

CPU

XeonW3520 (4 コア, 2.67GHz)

GPU

GeForceGTX285 (240 コア, 1.48GHz, Compute Capability 1.3)

OS

CentOS 5.4 x86 64

開発環境

2

実験環境

CUDA Toolkit 3.0, CUDA SDK 3.0

対象問題と測定範囲および性能測定方法

2.1 対象問題と測定範囲
今回は単純な行列積 (行列 A = 行列 B × 行列 C) を最適化対象問題とする。実装の内容は主にCUDA Programming Guide 等を参考にして独自に実装したものである。今回実験に用いた全プログラムは筆者の web サイト (http://www.cspp.cc.u-tokyo.ac.jp/ohshima/) にて公開する予定である。
*1 Compute Capability 2.0 の GPU は GeForceGTX480/470 および TeslaC2010 のみであり本稿執筆時点では発売直後のため用いていない1

今回は単純化のため問題を以下のように定める：• 問題サイズ (行列の一辺の長さ) は 4096 に固定• データ型は単精度浮動小数点形式 (ﬂoat 型)• メモリの配置は一次元配列• 行列の転置については考えない
こ の 問 題 設 定 で 単 純 に 実 装 し た 場 合 、行 列 B に 対 す る メ モ リ ア ク セ ス は 連 続 に 、行 列 Cに 対 す る メ モ リ ア ク セ ス は 不 連 続 に な る 。性 能 測 定 範 囲 (実 行 時 間 測 定 範 囲) に つ い て は 、GPU 上 で の 計 算 の み で は な く CPU-GPU 間 の デ ー タ 転 送 時 間 を 含 め て 実 行 時 間 を 測 定 し 、
2 × SIZE × SIZE × SIZE ÷ time(ただし SIZE は行列の一辺の長さ) の式にて性能 FLOPS 値を算出することにする。
上記の問題設定にしたがい、GPU 上の単一ブロック・単一スレッドを用いて単純な三重ルー
プ処理による行列積を行ったところ、わずか 5MFLOPS という低い性能しか得られなかった。
また単純な並列化の例として、64 ブロック*64 スレッドの 4096 並列を用いて各行の計算を並列実行させてみた (図 1) ところ、1GFLOPS 程度の性能しか得られなかった。一方、CUDA に付属の数値計算ライブラリ CUBLAS を用いて行列積を実行したところ、約 315GFLOPS の性能が得られた*2 。本稿は CUBLAS の性能に追いつき追い越すことが趣旨ではなく、CUDA プログラミ
ングの最適化技術を紹介することが趣旨であるが、CUBLAS の性能を目指して様々な高速化手法を適用していくことにする。

図1*2

単純な並列実行の例

残念ながら CUBLAS はソースコードが公開されていないため詳細な実装を知ることはできない

2

2.2 性能測定方法
CPU プ ロ グ ラ ム (C 言 語 プ ロ グ ラ ム) の 実 行 時 間 測 定 方 法 と し て は 、時 刻 取 得 関 数 、特 に 高
精度な gettimeofday 関数を用いた方法が知られている。今回は性能評価対象に CPU-GPU 間のデ ー タ 転 送 を 含 め て い る た め 、CPU か ら GPU へ デ ー タ の セ ッ ト を 行 う 前 と GPU か ら CPU へ
デ ー タ の 書 き 戻 し を 行 っ た 後 に gettimeofday 関 数 を 呼 び 出 し て 差 分 を と る こ と で 性 能 FLOPS
値を求めている。
ちなみに単純な時刻取得関数の利用では、CPU から GPU へのデータセット単体や、GPU 上
での計算単体での実行時間を測定することができない。
たとえば、CPU から GPU に対して演算開始指示を行うためには

funcname<<<xyz,xyz>>>(args1,args2);という形式の呼び出し記述を用いる。この前後の行で gettimeofday を行い差分を取っても GPUに対する演算開始指示は非同期であるため、すなわちプログラムは CPU から GPU への演算開
始指示が終了した時点で呼び出し指示の次の行へ遷移してしまうため、GPU 上での計算時間を取得することはできない。このように、CUDA プログラムにおいてはいくつかの処理が非同期処理であることに気をつける必要がある*3 。GPU による処理の時間を確実に測定するためには、cudaThreadSynchronize 関数を呼び出して GPU による処理が確実に終了するのを待つ必要がある。

GPU による演算時間を測定する別の方法としては、4 章で紹介する CUDA プロファイラを利用する方法がある。CUDA プロファイラを利用すれば、プログラム実行終了後に CPU-GPU 間の転送時間や GPU によるカーネル関数実行時間をそれぞれ個別に確認することができる。

3

CUDA 最適化プログラミング
本章では、以下の 5 つの最適化手法について解説する。

1. SIMD 化 (GlobalMemory への並列連続メモリアクセス)
2. ConstantMemory・TextureMemory の活用
3. SharedMemory の活用
4. ループの展開 (計算担当範囲の調整)
5. ホスト側のチューニング

3.1 SIMD 化 (GlobalMemory への並列連続メモリアクセス)
本節では GlobalMemory の活用方法について解説する。

GlobalMemory は GPU 上の全てのプロセッサで共有し読み書き可能なため様々な用途に利用
可 能 な メ モ リ で あ り 、連 続 ア ク セ ス 時 の 転 送 速 度 は 高 速 で あ る 一 方 で レ イ テ ン シ が 大 き く ラ
ンダムアクセス性能が低いという特徴を持つ。GlobalMemory の最適化における基本的な戦略*3 逆に非同期処理を利用すれば、CPU と GPU で同時に処理を行い CPU と GPU 両方の演算能力を利用できる

3

は、GPU 上の多数のプロセッサ (SP) による連続領域並列アクセスを行うことである。このメ
モリアクセスはコアレスなメモリアクセスと呼ばれており、CUDA の性能最適化において非常に重要である。
図 1 のプログラムでは、行列 B については各スレッドが独立に行列積の各行を計算しており、また行列 C についてはメモリが不連続な方向への連続アクセスが必要である。この割り当て方の場合、CPU(C 言語) プログラムであれば、行列 B についてはスレッド毎に連続メモリアクセ
スを行うことができる。一方 CUDA では、同一 MP 内の複数の SP にまたがって連続メモリアクセスを行うことが重要である。しかし今回のプログラムでは、各ブロック内のスレッドは不連
続な GlobalMemory アクセスを多く行ってしまっている。
そこで、ブロック内のスレッドが連続メモリアクセスを行うようにプログラムを変更する。
図 2 はブロック単位で行列積の各行の計算を担当し、行内の計算においては各スレッドが各列の 計 算 を 担 当 す る よ う に 変 更 し た 例 で あ る 。各 ブ ロ ッ ク 毎 の ス レ ッ ド 数 を W と し た と き 、ブロック数は横に 問題サイズ/W、縦に 問題サイズ となるようにプログラムを記述している。これにより行列 B に対するメモリアクセスが連続になり性能が向上する。

図 2 SIMD 化 (スレッドによる連続メモリアクセス) の例

各ブロック毎のスレッド数 W を変化させて性能を比較したところ、表 2 に示すように最大で41.1GFLOPS まで性能が向上した。このように、適切な連続領域アクセス (コアレスなメモリ
アクセス) を行うことが GlobalMemory を用いた高速化の基本である。より正確には、同一 MP
内の複数 SP が適切な GlobalMemory アクセスを行う場合には、複数のメモリアクセスがまとめて実行されるため性能が向上する。CUDA の世代 (開発環境およびハードウェアのバージョン)
が 進 む に し た が っ て コ ア レ ス な メ モ リ ア ク セ ス を 行 え る 条 件 が 緩 和 さ れ て き て お り 、最 近 のGPU では不連続でも一定領域内に収まっている GlobalMemory アクセスはまとめて実行される

4

ことが多い*4 。GlobalMemory は容量が大きく GPU 全体で共有されるため使用する機会が多い
メモリである。不用意に不連続なメモリアクセスを繰り返して性能低下を招かぬように気をつけることが重要である。
表 2 SIMD 化時の性能 (GFLOPS)

W(SIMD 長)

8

16

32

64

128

256

512

性能

8.4

17.2

32.5

41.1

40.6

40.7

38.3

3.2 ConstantMemory・TextureMemory の活用
本節では ConstantMemory と TextureMemory の活用方法について解説する。

ConstantMemory と TextureMemory は GPU から見ると読み込み専用であり、またキャッシュ効果が得られるメモリである。そのため、頻繁にアクセスする読み取り専用データを配置するのに適している 。ConstantMemory は容量が 64KB に制限 され ているため 格納するデ ータを 厳選する必要があるが、TextureMemory は大容量のメモリを配置できるために様々な問題に適用
し や す い 。さ ら に TextureMemory は 二 次 元 お よ び 三 次 元 の デ ー タ を 扱 っ た り デ ー タ の 補 間 を行うこともできるため、利用する対象データによっては高い性能が期待できる。
行列積においては、行列 B と行列 C が読み取り専用のデータであり、またアクセス回数も多
いため、ConstantMemory や TextureMemory の活用による性能向上が期待できる。今回は利用可能な容量が多い TextureMemory を利用してみることにする。
図 3 に TextureMemory を利用するプログラムを示す。このプログラムは前節の SIMD 化した
行列積プログラムに対して、行列 B のみを TextureMemory に配置するよう書き換えたものである。TextureMemory を使うにはデータセットのための手続きがホスト側に必要なため、ホスト
側プログラムの記述量が若干増加していることがわかる。またデバイス側プログラムについて
も、テクスチャの参照座標を計算する記述が追加されていることがわかる。

TextureMemory を利用した場合の性能は表 3 の通りである。最大 58.9GFLOPS の性能が得られており GlobalMemory のみの場合と比べて性能が向上していることがわかる。なお、行列 C を

TextureMemory に配置した場合や行列 B と行列 C の両方を TextureMemory に配置したプログラ
ムも作成して性能を比較してみたが、残念ながら良い性能は得られなかった。TextureMemory
による性能向上は、各 MP に搭載された TextureCache にデータがキャッシュされることによるも の で あ り 、ど の 程 度 性 能 向 上 が 得 ら れ る か は メ モ リ ア ク セ ス パ タ ー ン に 依 存 し て い る 。行
列 C のみまたは行列 B と C の両方を TextureMemory に配置したプログラムのメモリアクセスパ
タ ー ン よ り も 行 列 B の み を TextureMemory に 配 置 し た 方 が TextureCache を う ま く 活 用 で き たものと考えられる。

*4 特に Compute Capability 1.3 では 1.1/1.2 と比べて条件が緩和されている。詳細なルールについては CUDA

Programming Guide を参照していただきたい。

5

図 3 TextureMemory を用いたプログラム

表 3 TextureMemory 活用時の性能 (GFLOPS)

W(SIMD 長)

8

16

32

64

128

256

512

行列 B のみ TextureMemory に配置

6.6

13.0

24.9

46.7

58.9

57.4

54.5

行列 C のみ TextureMemory に配置

5.4

10.7

18.7

31.1

23.5

23.6

22.6

行列 B と C を TextureMemory に配置

9.0

17.1

29.1

31.0

24.5

25.5

26.7

参考:TextureMemory 未使用

8.4

16.9

31.8

39.9

39.4

39.6

37.3

3.3 SharedMemory の活用
本節では SharedMemory の活用方法について解説する。

SharedMemory は、高速かつレイテンシが低い一方で 1MP あたりの容量がわずか 16KB しかな
い、各ブロック内の全スレッドでデータを共有可能な共有メモリである。SharedMemory の位置づけとしては、既存の CPU におけるキャッシュメモリや Cell B.E. の SPE における LS(Local

Storage) に近いと見なされることがある。しかし SharedMemory は、CPU のキャッシュと異なりプログラマが明示的に操作を記述する必要があり、SPE の LS と異なり共有メモリであり、そのうえキャッシュや LS と比べて容量が小さい。そのため、SharedMemory ならではの活用方法を理解する必要がある。

6

SharedMemory の基本的な活用方法としては、GlobalMemory に対するキャッシュとしての利用が挙げられる。3.1 節で述べたように、GlobalMemory は転送速度自体は高速であるもののレ
イテンシが大きくランダムアクセスに弱いという特徴を持っている。そのため GlobalMemory
上の特定の不連続データを複数回利用する場合には、一度 SharedMemory へコピーしてから利用することで性能を改善することが可能となる。

3.1 節 で SIMD 化 し た プ ロ グ ラ ム を 再 度 見 直 す と 、ラ ン ダ ム ア ク セ ス こ そ 無 い も の の 、各 ブ
ロ ッ ク 内 の ス レ ッ ド が 全 く 同 じ GlobalMemory(行 列 B) 上 の デ ー タ を 参 照 し て い る こ と が わ
か る 。そ こ で 、行 列 B へ の ア ク セ ス を SharedMemory に よ っ て 高 速 化 す る こ と に す る 。図 4 にSharedMemory を用いたプログラムを示す。本プログラムと 3.1 節のプログラムとの違いは、行
列 A の要素に対するアクセスに SharedMemory を用いるか否かのみである。そのためプログラムの変更点も shared 指示子を用いた SharedMemory の宣言および SharedMemory にアクセス
す るための記述 (ks ル ープのステップ数変更 および SharedMemory 内 のデータ にアクセス す る
ための内側ループの追加) のみである。小さなプログラムの修正ながらその効果は大きく、表 4に示すように最大で 75.4GFLOPS を達成している。

図 4 SharedMemory を用いたプログラム

表 4 SharedMemory 活用時の性能 (GFLOPS)

W(SIMD 長)

8

16

32

64

128

256

512

性能

22.1

49.8

38.6

68.8

75.1

75.4

75.3

参考:SharedMemory 未使用

8.4

16.9

31.8

39.9

39.4

39.6

37.3

今回のプログラムは各スレッドから SharedMemory へのアクセスが単純で偏りのないものであったが、SharedMemory の特定の範囲に対して偏ったアクセスがある場合にはあまり高い性能が得られないことがある。これは SharedMemory におけるバンクコンフリクトの問題として知られている。SharedMemory は全体が 16 のメモリバンクに分割されており、特定のバンクにアクセスが集中するとメモリアクセスが待たされてしまい性能が低下するという問題である。

7

そのため SharedMemory へのアクセスパターンによっては参照・格納順序の変更やパディングなどを行うことで初めて性能が向上することがある。また 16KB という限られた空間では適切にデータを配置しきれず SharedMemory を活かせないこともあるので注意が必要である。

3.4 ループの展開 (計算担当範囲の調整)
CPU プログラムの高速化において有用な最適化手法の 1 つにループの展開 (アンローリング)がある。特に GPU は分岐処理に弱いという特徴があるため、アンローリングにより分岐処理(ループの終了判定) が削減されることによる性能向上が期待できる。
表5に、SIMD化されたプログラム、TextureMemoryを用いたプログラム、そして SharedMemoryを用いたプログラムのそれぞれをアンローリングした際の性能および性能比 (それぞれアンロー
リングを行わない場合を 1 とした相対性能) を示す。プログラムの差分は単純なアンローリン
グの有無のみであるため、プログラムの詳細な説明は省略する。実験の結果、残念ながら単純なアンローリングによる性能向上は得られなかった。
表5

アンローリング後の性能 (GFLOPS、括弧内はアンローリング前に対する性能比)

アンローリング段数 (SIMD 長は全て 256)

1

2

3

4

SIMD 化 + アンローリング

40.3

40.6

40.3

40.0

(0.99)(1.00)(0.99)(0.98)

55.3

55.1

55.2

55.3

(0.92)(0.92)(0.92)(0.92)

75.5

58.9

75.4

75.4

(1.00)(0.78)(1.00)(1.00)

TextureMemory(行列 B のみ)+ アンローリング
SharedMemory+ アンローリング

ところで、これまでに挙げたプログラムはいずれもブロック形状が縦に 1 のもの (gridDim.y
が 1 のもの) のみを扱ってきた。一方でブロック形状を変更して各スレッドが担当する行数を増やし、ブロック数を減らすと、プログラム全体で見たループの終了判定回数 (全スレッドが行う終了判定の累計回数) を減少させることができる。これは、三重ループによる行列積において外側のループをアンローリングすること に相当する。

SIMD 化 さ れ た プ ロ グ ラ ム の 計 算 担 当 範 囲 を 調 整 し た プ ロ グ ラ ム の ソ ー ス コ ー ド と 動 作 イ
メ ー ジ を 図 5 に 、SIMD 化 さ れ た プ ロ グ ラ ム 、TextureMemory を 用 い た プ ロ グ ラ ム 、そ し て

SharedMemory を用いたプログラムのそれぞれについて計算担当範囲を調整したプログラムの性能と性能比を表 6 に示す。今回のアンローリングはいずれも高い性能向上が得られており、特に SharedMemory を用いたプログラムについては 3 倍を超える速度向上比を達成している。最大性能が得られたのは SharedMemory+ 範囲拡大 4 倍にした場合で、245.8GFLOPS を達成した。

3.5 ホスト側のチューニング
ここまではデバイス (GPU) 側のチューニングについて紹介してきた。本節ではホスト (CPU,
メインメモリ) 側のチューニング手法を紹介する。

8

図5

計算担当範囲の調整

表 6 問題割り当てを変更した際の性能 (GLOPS、括弧内はアンローリング前に対する性能比)

割り当て範囲拡大率 (SIMD 長は全て 256)

2

4

8

16

SIMD 化 + 範囲拡大

55.8

60.3

78.7

65.6

(1.37)(1.48)(1.93)(1.61)

104.2

104.4

104.4

104.5

(1.73)(1.74)(1.74)(1.74)

146.2

245.8

137.0

-

TextureMemory+ 範囲拡大
SharedMemory+ 範囲拡大

(1.94) (3.26) (1.82)(SharedMemory 16 倍は SharedMemory の容量が 16KB を超えてしまったため測定不可能)

特に簡単で効果が高い手法としては、pinned memory(Page-Locked memory) の使用が挙げられる。pinned memory を使用すると CPU-GPU 間のデータ転送性能を向上させることができる。
使い方は CPU 側でメモリを確保する際にcudaHostAlloc((void**)&g_data, sizeof(data)*length, cudaHostAllocWriteCombined);の よ う に cudaHostAlloc 関 数 を 使 え ば 良 い だ け で あ り 、非 常 に 簡 単 で あ る 。た だ し 、pinned

memory を使いすぎると CPU 側の性能が低下することがあるので注意せねばならない。
例として、3.4 節の SharedMemory+ 範囲拡大 (4 倍) プログラムに対して pinned memory を適用
して性能を測定した。その結果、245.8GFLOPS から 247.0GFLOPS へとわずかながら性能を向9上させることができた。CPU 側のメモリ確保方法を変更するだけで若干の性能向上が得られていることから、データ送受信のサイズが大きかったり通信回数が多くプログラム実行時間に対する CPU-GPU 間のデータ転送時間の占める割合が大きな場合には、pinned memory を有効に活用すべきである。
他 の ホ ス ト 側 の チ ュ ー ニ ン グ 手 法 と し て は 、計 算 と 通 信 の オ ー バ ー ラ ッ プ が 挙 げ ら れ る 。

CUDA で は CPU-GPU 間 のデータ通信 と GPU に よる計算 をオーバー ラップさせ ることが できるため、複数回のカーネル実行を行う場合には適切にオーバーラップ実行することで性能向上が可能となる。しかし、今回は行列積を一度だけ実行して性能を測定しているため、通信と計算をオーバーラップ実行する機会がない。本手法による性能向上については機会があれば改めて紹介したい。
別のオーバーラップの考え方として、CPU と GPU で問題を分割し同時に計算することが考えられる。行列積は CPU でも高速に計算できる問題であり、また問題分割がとても容易である。
そのため、これまでに解説してきたプログラムはいずれも GPU が計算をしている間に CPU は何もせずに待っていたが、対象問題をデータ分割して CPU でも計算すればより高い性能が得られることが期待できる。なおこれは CUDA 特有の手法ではなく、グラフィックス API を用い
た GPGPU でも可能だった手法である。

3.6 CUDA 最適化プログラミング:まとめ
以上、本章では 5 つの最適化手法について解説した。今回解説した最適化手法のいくつかは複数種類を組み合わせて利用することでより高い性能を得ることができる。(実際、いくつかの解説中では特に説明せずに組み合わせて利用していた。) 本章にて解説した最適化手法の中で高い性能が得られた組み合わせとしては、SIMD 化・SharedMemory・カーネル内最内ループのアンローリング・計算担当範囲の調整・pinned memory の活用を組み合わせて実行したケー
スで最大 277.7GFLOPS を達成している。
今回は主に行列 B のメモリアクセスを中心とした最適化を行ってきたが、さらに高い性能を得るためには、行列 C のメモリアクセスについても最適化する必要がある。CUDA Programming

Guide に掲載されている行列積の例では二次元ブロック化によるメモリアクセスを行っているので参考になるだろう。

4

CUDA プロファイラを用いた性能の解析

4.1 CUDA プロファイラ
GPU は CPU と同等以上にメモリアクセスパターンやプロセッサへの計算割り当てによって性能が左右されるハードウェアである。しかし CUDA 最適化プログラミングにおいては、様々な最適化手法を組み合わせた場合にどの手法がどの程度影響を与えているのか、またプログラム
記述の変更がハードウェアの動作にどのように影響しているのかを知ることは容易ではない。

CUDA 最適化プログラミングの手助けになりうるツールとしては、CUDA Toolkit に付属のNVIDIA 社製 CUDA プロファイラが挙げられる。本章では CUDA プロファイラの簡単な使い方を紹介し、また前章のいくつかの実行結果に対するプロファイル出力結果を確認してみること
にする。

10

CUDA プロファイラは CUDA toolkit インストール時に自動的にインストールされる。CUDA
プロファイラを使うために最低限必要な準備は、環境変数 CUDA PROFILE に 1 を設定することである。CUDA PROFILE を設定した状態で CUDA プログラムを実行すると、実行後にカレ
ントディレクトリに cudab proﬁle.log というプロファイル結果が書き込まれたファイルが作られる。出力ファイル名は環境変数 CUDA PROFILE LOG にて変更することも可能である。
図 6 に行列積 プロ グラ ムの 1 つを実行 した 際のプ ロファ イル結果を 示す。デ フォルトで プ ロ
ファイル出力されるデータは CPU 時間、GPU 時間 ()、occupancy(GPU 占有率) のみであるが、環境変数 CUDA PROFILE CONFIG によって設定ファイルを指定することで出力される項目を変更するこ とも でき る。CUDA プロ ファ イラの 詳細な使用 方法は付属 のドキュ メント*5 を参照されたい。
> cat cuda_profile_0.log# CUDA_PROFILE_LOG_VERSION 1.6# CUDA_DEVICE 0 GeForce GTX 285# TIMESTAMPFACTOR fffff72f1365e070
method,gputime,cputime,occupancy
method=[ memcpyHtoD ] gputime=[ 11760.608 ] cputime=[ 11928.000 ]method=[ memcpyHtoD ] gputime=[ 11759.840 ] cputime=[ 11926.000 ]method=[ memcpyHtoD ] gputime=[ 11767.712 ] cputime=[ 11936.000 ]method=[ mmkernel ] gputime=[ 4120903.500 ] cputime=[ 10.000 ] occupancy=[ 0.250 ]method=[ memcpyDtoH ] gputime=[ 11828.768 ] cputime=[ 12003.000 ]
図6

プロファイル出力結果例

この他、CUDA プロファイラの出力を可視化する Visual Proﬁler も提供されており、OS を問
わず (Windows,Linux,MacOSX のいずれでも) 使うことができる。さらに VisualStudio と統合された開発環境 nsight も開発が進んでいる。今後これらについても機会があれば紹介する。

4.2 CUDA プロファイラによるプログラム解析の例
本節ではこれまでに取り上げたプログラムのいくつかについてプロファイル出力結果を確認する。なお、同時にプロファイル可能なパラメタ数が限られているため、複数回実行して取得
し比較を行っている。
まずはじめに、2.1 節で作成した単純な並列実行のプログラム (1.0GFLOPS) と 3.1 節で作成した SIMD 長 8(8.4GFLOPS) の プ ロ グ ラ ム の プ ロ フ ァ イ ル 出 力 結 果 を 比 較 し て み る こ と に す る 。
以下に示すのはそれぞれのプロファイル出力結果である。
(単純な並列実行)
method=[ mmkernel ] gputime=[ 140879904.000 ] cputime=[ 140881600.000 ] occupancy=[ 0.500 ]local_load=[ 0 ] local_store=[ 0 ] gst_request=[ 100663296 ] divergent_branch=[ 4294967295 ]branch=[ 100687899 ] sm_cta_launched=[ 3 ] gld_incoherent=[ 0 ] gld_coherent=[ 4294967295 ]gld_32b=[ 4294967295 ] gld_64b=[ 0 ] gld_128b=[ 0 ] gst_incoherent=[ 0 ]gst_coherent=[ 4294967295 ] gst_32b=[ 4294967295 ] gst_64b=[ 0 ] gst_128b=[ 0 ]instructions=[ 906313776 ] warp_serialize=[ 0 ] cta_launched=[ 7 ]*5 toolkit をインストールしたディレクトリ (デフォルトでは/usr/local/cuda/) 直下の doc ディレクトリにある

CUDA Proﬁler 3.0.txt

11

(SIMD 長 8)
method=[ mmkernel ] gputime=[ 16227636.000 ] cputime=[ 16227876.000 ] occupancy=[ 0.250 ]local_load=[ 0 ] local_store=[ 0 ] gst_request=[ 69906 ] divergent_branch=[ 4294967295 ]branch=[ 286824318 ] sm_cta_launched=[ 69906 ] gld_incoherent=[ 0 ] gld_coherent=[ 4294967295 ]gld_32b=[ 1717993472 ] gld_64b=[ 0 ] gld_128b=[ 0 ] gst_incoherent=[ 0 ]gst_coherent=[ 419432 ] gst_32b=[ 209716 ] gst_64b=[ 0 ] gst_128b=[ 0 ]instructions=[ 2291798310 ] warp_serialize=[ 0 ] cta_launched=[ 209716 ]こ れ ら の プ ロ グ ラ ム を 比 較 す る と 、単 純 な 並 列 実 行 は SIMD 化 し た も の に 比 べ て

gld 32b(GlobalMemory に 対 す る load の 回 数) や gst 32b(GlobalMemory に 対 す る store の 回 数) が非常に多いことがわかる。これは SIMD 化することによって GlobalMemory へのアクセスがまとめて行われるようになったことを反映しているものであると考えられる。
つづいて、3.1 節で作成した SIMD 長 8(8.4GFLOPS) のプログラムと SIMD 長 256(40.7GFLOPS)のプログラムのプロファイル出力結果を比較してみる。以下に示すのは各 SIMD 長 256 におけるプロファイル出力結果である。
(SIMD 長 256)
method=[ mmkernel ] gputime=[ 3337158.750 ] cputime=[ 3337238.000 ] occupancy=[ 1.000 ]local_load=[ 0 ] local_store=[ 0 ] gst_request=[ 17480 ] divergent_branch=[ 2165478906 ]branch=[ 71705145 ] sm_cta_launched=[ 2185 ] gld_incoherent=[ 0 ] gld_coherent=[ 1708726971 ]gld_32b=[ 429522944 ] gld_64b=[ 429522944 ] gld_128b=[ 0 ] gst_incoherent=[ 0 ]gst_coherent=[ 419456 ] gst_32b=[ 0 ] gst_64b=[ 104864 ] gst_128b=[ 0 ]instructions=[ 573064352 ] warp_serialize=[ 0 ] cta_launched=[ 6554 ]SIMD 長 8 と SIMD 長 256 で異なる値として、occupancy(GPU 上の演算器がどの程度使われたか) が挙げられる。本プログラムでは SIMD 長=ブロックあたりのスレッド数であることから、

SIMD 長 8 では GlobalMemory へのアクセスレイテンシをアクティブスレッドの切り替えによって隠蔽し切れておらず、GPU 上の演算器を使い切れていなかったためにこのような差が生じたと考えられる。また、SIMD 長 8 では 32b という 32bit 単位のメモリ読み書きばかりなのに対して、SIMD 長 256 では 64b という 64bit 単位のメモリ読み書きが多くなっており、SIMD 長を大きくすることでメモリの読み書きがまとめて行われるようになったことが確認できる。
つづいて、SIMD 長 256 のプログラム (40.7GFLOPS) と、これに対して 3.4 節で計算割り当て範囲の拡大 (8 倍) を行ったプログラム (78.7GFLOPS) のプロファイル出力結果を比較してみる。
(割り当て範囲 8 倍)
method=[ mmkernel ] gputime=[ 1810732.000 ] cputime=[ 1810776.000 ] occupancy=[ 0.500 ]local_load=[ 0 ] local_store=[ 0 ] gst_request=[ 17536 ] divergent_branch=[ 1100197911 ]branch=[ 8991858 ] sm_cta_launched=[ 274 ] gld_incoherent=[ 0 ] gld_coherent=[ 1122244519 ]gld_32b=[ 429916160 ] gld_64b=[ 53739520 ] gld_128b=[ 0 ] gst_incoherent=[ 0 ]gst_coherent=[ 419840 ] gst_32b=[ 0 ] gst_64b=[ 104960 ] gst_128b=[ 0 ]instructions=[ 269445042 ] warp_serialize=[ 0 ] cta_launched=[ 820 ]割り当て範囲を拡大した結果、occupancy は低下してしまったが、branch(プログラム全体における分岐回数) や instructions(プログラム全体における実行命令数) などの大幅な減少が確認
できており、複数行で同様に行われていたループ処理 (ループ終了判定に関わる分岐命令) がまとめられたことにより性能向上したことが読み取れる。
さらに SharedMemory を用いた場合のプロファイル出力結果と比較を行いたいところであるが、残念ながら Compute Capability 2.0 以上の GPU でないと SharedMemory の読み書き回数は12取得できないため、比較することができなかった。
以 上 の よ う に 、CUDA プ ロ フ ァ イ ラ を 使 う こ と で 詳 細 な 実 行 情 報 取 得 す る こ と が 可 能 と なり、最適化の効果を確認することができる。CUDA プロファイラを用いると様々な情報を確認することができる一方で、現在のところプロファイラ出力のどの値がどの程度性能に影響するかについてはあまり明らかになっていない。そのため単純に最適化の指標として使うことは容易ではないが、意図したとおりに最適化が行えているかを確認する際などに役に立つだろう。
今後は様々な最適化手法や各手法におけるパラメタの選択にプロファイラが有効活用できるよ
うになることが期待される。
以上、第三回の今回は行列積を対象問題としてとりあげて CUDA 向けの最適化プログラミン
グについて紹介した。
今回は対象問題が単純なこともあり多くの最適化手法が少なからず性能向上を得られたが、
実 ア プ リ ケ ー シ ョ ン で は 常 に 性 能 向 上 が 得 ら れ る と は 限 ら な い 。ま た 、行 列 積 の よ う に 一 度

GPU に計算命令を出して終わりではなく、CPU と GPU それぞれが様々な計算を繰り返し行う
アプリケーションや複数台の GPU を使うようなアプリケーションでは、CPU-GPU 間の通信がボトルネックにならないようにアプリケーション全体の最適化を行わなくては高い性能を得ることはできない。アプリケーションの特徴と GPU の特徴を見極めて適切な最適化を行う必要がある。

(次回に続く)13