従来のGPUは定数やテクスチャ用のキャッシュを持っていたが、これらはCPUで言えば命令キャッシュのような読み出しオンリーのキャッシュで、書き込みが行えるデータキャッシュを持っていなかった。このため、高速に読み書きが必要なデータはローカルメモリに配置し、必要に応じてGDDR DRAMで構成されるグローバルメモリとのデータの入れ替えを行うというプログラムを書く必要がある。この点は、PS3に使われているCellプロセサの同様で、ローカルメモリをうまく使えば高い性能が得られるが、プログラミングが難しい、面倒という批判がある。

これに対して、NVIDIAのGF100(開発コード名:Fermi)では各Streaming Multiprocessor(SM)が持つ64KBのローカルメモリをシェアードメモリ部分と1次データキャッシュ部分に分割できるようになった。この分割はシェアードメモリを16KBでキャッシュを48KB、あるいはその逆にシェアードメモリが48KBで16KBをキャッシュとすることができる。

シェアードメモリはグローバルメモリとは独立のメモリ空間のローカルなメモリであり、プログラムで明示的に区別して使用しグローバルメモリとの間でデータの転送を行う必要があるが、1次データキャッシュはグローバルメモリとのコヒーレンスがハードウェアで維持される本物のキャッシュである。一部、推測が入っているが、メモリ系を中心にGF100(Fermi)のブロック図を書くと、次の図のようになる。なお、2次キャッシュは6個あるが、それぞれ異なるメモリアドレスを分担しているので、コヒーレンシ制御は必要ない。

GF100(Fermi)のブロック図(一部筆者の推測を含む)

キャッシュ利用により、非利用時比で50%の性能向上も

質問に答えるNVIDIAフェローのDavid Kirk氏

2010年7月16日に開催されたNVIDIAの「GPU Computing 2010」において、同社フェローのDavid Kirk氏との質疑の機会があったので、キャッシュの効果を示すデータは無いのかと質問してみた。すると、公表されたデータは無いが、設計の意図を説明することはできると述べて、次のように説明してくれた。

CPUのキャッシュは1つのスレッドが多数のキャッシュラインを使うことができ、頻繁に使うデータ群を格納しておくのであるが、GPUのキャッシュは1つのスレッドに1つのキャッシュラインがあるかどうかという程度の量しかないので、データを溜めておく効果は小さい、むしろ、1つのスレッドが計算した値を多くの(同じSMで動く)スレッドに渡すというクロスバスイッチ的な効果を期待しているという。

確かに48KB分をキャッシュとして確保しても、64バイトキャッシュラインとすると768ラインしかない。これに対して、各SMに100ワープを割り当てたとすると、各ワープには32スレッドが含まれるので全体では3200スレッドとなり、約4スレッドに1つのキャッシュラインしかないことになる。したがって、スレッドごとに個別となるデータを1次キャッシュに格納することはできない。しかし、この事情はシェアードメモリとして使っても同じで、やはり全スレッドで共通的に使うデータを格納することになる。

ということで、設計の意図は分かったが、どの程度効果があるのかというデータは分からず仕舞いかと思っていたら、スケジュールのほぼ最後に富士通研究所の成瀬氏が登壇してキャッシュの効果を説明するプレゼンテーションを行った。このセッションはNVIDIAの杉本事業部長の「GPUコンピューティングソリューションのご紹介」と題する発表で、その中で出展各社が10分程度で製品やサービスの説明を行ったのであるが、他社が宣伝的なプレゼンテーションであった中で富士通だけが純技術的なプレゼンテーションであった。

成瀬氏は、Fermiを使うGTX480と1世代前のGTX285グラフィックスカードを使って、理化学研究所(理研)の姫野龍太郎氏が作った流体解析の主要計算部分をベースとした姫野ベンチマークの実行性能を比較した。このベンチマークはメモリバンド幅リミットになるプログラムで、GTX480とGTX285のメモリバンド幅はほぼ同じであり、ほぼ同性能になることが予想される。

シェアードメモリを使って頑張ってチューニングしたCUDAプログラムの場合、GTX285では76.9GFlops、GTX480では74.0GFlopsという性能が得られたという。しかし、姫野ベンチの主要実行部分はオリジナルのCコードでは44行であるが、頑張ってシェアードメモリを使うチューニングを行ったCUDAプログラムは232行と、オリジナルの5.27倍の行数を必要としている。

一方、シェアードメモリを使わずオリジナルコードを単純にCUDA化したプログラムは69行になる。このプログラムでの姫野ベンチの性能は、GTX285が30.8GFlops、GTX480は43.0GFlopsとなったという。キャッシュを持たないGTX 285ではすべてのメモリアクセスがデバイスメモリとなるので、シェアードメモリチューニング版と比較すると約40%に性能が低下している。しかし、GTX 480では自動的にキャッシュが使われるので、シェアードメモリチューニング版の約58%の性能が得られている。つまり、キャッシュが利用できることにより50%近く性能が向上している。

しかし、シェアードメモリチューニングを行ったレベルとは大きな差がある。ということで詳しく調べてみると下図のように、同一のSM上で実行されるスレッドブロックの中の各ワープの実行時間がバラついており、これがキャッシュミスを引き起こしていることが分かったという。1つのスレッドがメモリからデータを読んでくると、他のスレッドはキャッシュに乗ったそのデータを利用できるのであるが、実行時間がバラつくと、遅いスレッドがそのデータを使う前に、速いスレッドが別のデータを読んでそのデータを上書きして消してしまうのが原因という。

スレッドブロック内のワープ間の実行時間のバラつき

ということで、__syncthreads( )関数を呼んでスレッド間の同期をとるようにプログラムを変更すると、同一スレッドブロック内のワープの実行時間が揃い、キャッシュミスが大幅に減ったという。さらにキャッシュミスを減らす若干の修正を加えたプログラムは74行となったが、GTX 480での実行性能は72.3GFlopsとシェアードメモリを使ってチューニングしたプログラムの74.0GFlopsと並ぶ性能になったという。

キャッシュチューニングを行った結果のワープ実行時間

姫野ベンチは1つの例であり、どのプログラムでも同じような効果が得られるというわけではないが、単純にCUDA化してコンパイルしただけで、キャッシュが使われることから50%程度の性能向上が得られ、さらに5行を追加してキャッシュミスを減らすチューニングを行うとシェアードメモリチューニング版のプログラムとほぼ並ぶ性能が得られている。このとき、キャッシュチューニングを行ったプログラムは74行に対して、シェアードメモリチューニングを行ったプログラムは232行であり、オリジナルからの増加分でみると30行と188行という違いとなる。プログラムの開発にかかる手間としては、この追加行数の比の6倍、あるいはそれ以上の違いとなると思われるので、Fermiのキャッシュは開発効率の点でも非常に効果が大きいと言える。

キャッシュチューニングを行った結果の性能とプログラム行数比較

プレゼンテーションの中では発表されなかったが、成瀬氏によるとGTX480のグローバルメモリのレイテンシは400ns程度でGTX285の350ns程度に比べて若干遅くなっている。メモリバンド幅では、多少、勝っているGTX 480の方が姫野ベンチの性能が低いのは、これが影響しているのではないかということであった。また、成瀬氏の実測ではGTX 480の1次キャッシュのアクセスレーテンシは約70ns、2次キャッシュのアクセスレーテンシは約250ns程度とのことで、CPUのキャッシュと比べるとこれでもキャッシュ? という程度の速度であるが、多数ワープを切り替えて実行する超マルチスレッド実行であるのでレイテンシを隠ぺいでき効果が出ているのであろう。