勉強会前日から腹痛にみまわれていたので、Hackathonには午後から参加した。実働時間が少なかったのでプログラムの作成ではなく、ちょっとしたベンチマークとチューニングをした。具体的には
- K20とTITANの速度比較
- Read Only Cacheの効果
の2つの項目について調べた。対象としたプログラムは自作の3D-RISM理論のプログラムである。3D-RISM理論については、GPUと3D-RISMで検索すれば情報は拾えるだろう。
1. K20とTITANの速度比較
まず実行速度の比較を行った。3D-RISMでDNA(12塩基対)の周りの水の分布を計算した。K20c(CUDA5.0)で63.0秒かかるのが、TITAN(CUDA5.5)では47.6秒だった。おおよそ1.3倍の速度向上である。3D-RISMはメモリバンド幅が律速になるプログラムなので、K20cとTITANのメモリバンド幅を調べるとそれぞれ208と288.4 GB/sなので1.39倍となる。このメモリバンド幅の違いが実行時間に影響していると考えるのが妥当だろう。
K20cよりは安くて速いと良い事づく目に見えるが、カーネル実行時のプロセスの保護がされていない可能性がある。次に説明するチューニングを行っている時に他の方との実行がかち合った時に収束状況を示す数値が乱れて収束回数が変化するという現象が発生した。TeslaシリーズであるFermiやK20cでは複数のプロセスを投入しても問題は発生しなかったので複数人で計算に使用する場合には注意されたい。
2. Read Only Cacheの効果
次にGTC Japan 2013の成瀬氏の「CUDAプログラムのKepler向け最適化手法」というセッションで触れられていたRead Only Cacheがどの程度効くかを試してみた。詳細については、セッションの資料やCUDAのマニュアルを参照されたい。チューニングする前のコードを次に示す。
先の3D-RISMの初期化部分の一部でチューニング前は4.4秒かかっていた（nvprofで計測）。実行時のブロック数は256スレッドでグリッド数はx、y供に256であった。またnatuという変数には892が代入されていた。コードを見ればわかるようにdgvとdru、dquは読み出すだけなのでRead Only Cacheの指定をする事で高速化する可能性がある。次にチューニング後のコードを示す。
上記3つの配列の部分を
const double4 * __restrict__ dgv
のようにconstと__restrict__をつけるだけである。他には配列の読み出し部分で__ldg()をつけるという方法もある（CUDA5.5以上）。さてその効果だが、4.4秒が2.5秒に短縮された。約1.8倍の速度向上である。もちろん、これは最も適合した例で、他の部分では1.3倍程度に留まったり、効果が観察されない場合もあった。
もう一つ、このチューニングはTITANとCUDA5.5の組み合わせで行ったが、K20cとCUDA5.0の組み合わせでパフォーマンスが悪化した例もあった（計算時間がおおよそ1.5倍になった）。GTC Japanの成瀬氏のセッション後に質問した時には、悪さはしない（パフォーマンスが悪化する事はない）との事だったが、それはあくまで5.5の場合であるようだ。5.0でチューニングする時には必ず前後で計算時間の確認をされたい。
