関西GPGPU勉強会#4 Hackathonの参加報告

関西GPGPU勉強会#4 Hackathonに参加してきた。そこでの成果を簡単にまとめる。このブログのタイトルのon Fermiはすでに実情にそぐわなくなっているのはご愛嬌という事で。

勉強会前日から腹痛にみまわれていたので、Hackathonには午後から参加した。実働時間が少なかったのでプログラムの作成ではなく、ちょっとしたベンチマークとチューニングをした。具体的には

1.  K20とTITANの速度比較
2.  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でチューニングする時には必ず前後で計算時間の確認をされたい。

C2050/2070でECCをON/OFFする。(CUDA4.0RC2の場合)

NVIDIAのGPUでECCのON/OFFを制御するnvidia-smiのオプションがCUDA4.0RC2では変更になった。CUDA3.2以前の方はC2050/2070でECCをON/OFFする。を参照してほしい。

ステータスの確認は

nvidia-smi -L

でGPUの番号と型番の対応が表示される。以前と同様にdeviceQueryコマンドで表示されるものと違う場合があるので注意しよう。ECC Modeなどの詳細な情報は

nvidia-smi -q

を使用する。-dオプションをつけて

nvidia-smi -q -d ECC

とするとECC関連の表示となる。ECCのON/OFFに関しては

GPU 0:1:0
Ecc Mode
Current : Disabled
Pending : Disabled

の部分に注目する。Currentが現在の状態でPendingがリブート後の状態を現している。現在はECCがDisabled、つまりOFFとなっている。

GPU 0のECCのサポートをオンにしたい場合はrootになるかsudoを使って

nvidia-smi -i 0 -e 1

を実行する。-iでGPUの番号を指定し-eで0(OFF)か1(ON)を指定する。コマンドを実行して

Enabled ECC support for GPU 0:1:0.
Reboot required.

と表示されればGPU 0はリブート後にECCがONになる。この状態で

nvidia-smi -q -d ECC

を実行すると

GPU 0:1:0
Ecc Mode
Current : Disabled
Pending : Enabled

となり、PendingがDisabledからEnabledに変更されているのが確認できるだろう。

UbuntuにCUDAドライバーをインストールする。

CUDAドライバーをLinuxにインストールするときにはXを停止しなければならない。FedoraやRedHatなどではrootで

/sbin/init 3

を実行すればいいが、Ubuntuの場合にはこれでは停止しない。

UbuntuでXを停止させる方法は、使用しているDisplay Manegerに依存する。
gdmの場合:

/etc/init.d/gdm stop

kdm(Kubuntu)の場合:

/etc/init.d/kdm stop

となる。

C2050/2070でECCをON/OFFした場合のパフォーマンスの変化を調べる。

C2050/2070でECCをONにした時のパフォーマンスの変化を調べた。場の収束計算で3次元FFTやベクトルの内積を計算する部分が大きな割合を占めるためにメモリアクセスが律速となるプログラムをベンチマークとして使用した。

結果は

ECC 　　 OFF　　　　ON
時間　　 9.6s　　　11.9s

となり明らかな違いが観察された。ECCをONにするとOFFの時の8割程度の速度しか出なくなる。これはECC用のメモリを別に用意しているわけではなくユーザーメモリを利用しているためである。

メモリアクセスが計算の律則になっていて、ECCを必要としない場合、もしくはメモリをギリギリまで必要とする場合にはECCをオフにしていた方がパフォーマンスが向上する。

C2050/2070でECCをON/OFFする。

NVIDIAのFermiのアーキテクチャではECCがサポートされた。金融などの要求に対応するためであるが、ECCをONにすると使用可能なユーザーメモリが12.5%減少する。また、メモリアクセスのパフォーマンスにも影響が出る事が懸念される。そこでECCのON/OFFを切り替える方法を紹介する。NVIDIA Forumsの"How to disable/enable ECC on C2050?"を参考にした。

ECCのON/OFFを切り替えるのはnvidia-smiを使用する。コマンド使用後にリブートが必要となる。まずは

nvidia-smi -r

で現在のステータスを確認しよう。ECCがサポートされていないGPUの場合

ECC is not supported by GPU 0

と表示され、サポートされている場合には

ECC configuration for GPU 1:
Current: 0
After reboot: 0

と表示される。ここでGPUに割り当てられる番号はdeviceQueryコマンドで表示されるものと違う場合があるので注意しよう。この表示の場合、Current（現在)はECCサポートなしで0、リブート後も0となっている。

GPU 1のECCのサポートをオン(1)にしたい場合

nvidia-smi -g 1 -e 1

とする。-gはGPUを指定し、-eはリブート後のECCサポートのオン(1)オフ(0)を指定する。つまりオフにしたいときは

nvidia-smi -g 1 -e 0

とすればいい。コマンドが正しく実行されたかを確認するためには

nvidia-smi -g 1 -e 1

を実行した後に

nvidia-smi -r

で確認してみれば

ECC configuration for GPU 1:
Current: 0
After reboot: 1

となっているはずだ。この状態でリブートすればECCがオンになる。

はじめに

Fermi CoreのTeslaが発売されました。Tesla C1060からFermiに移行するときの注意点などを書き連ねて行く予定です。