GPUを支える技術読み始めた 第4章 [前半]

今年もあと少しになってきたが、なんとか目標だったもう一本を投稿することができて良かった。

私情だが、先日の社内年末パーティでは年間MVPに選出していただいた。＼(^o^)／
非常に嬉しく思うと同時に、いろんな面でサポートをしてくれたHWチームメンバやバックオフィスに感謝したい。
また、来年は社内だけでなくて社外にも影響を与えられるよう頑張っていきたい。

少し気が早いが、来年度の本ブログの方針として「基礎と応用」というコンセプトで書いていきたと考えている。古典的名著と最新の論文の要約などができたら上出来だろうか。（＾ω＾）
直近だと、DeepLearning × HWに関するの新しめの論文や並列処理技法系の本のまとめを計画している。もしかすると、年内にまだいけるかもしれない。

良い報告もできたところで、早速続きを初めていこうと思う。
4章は個人的には一番楽しみな章でじっくり読んでいる。内容が多いので前半と後半に分割して投稿していく。

4章 GPUの超並列処理 [前半]

GPUの並列処理方式

先の章で並列処理方式について以下のように説明した。

• SIMD: 1つの計算を幾つかのデータに対して並列に実行する
• SIMT: 1つの計算を別々の演算機で並列に実行する

4章では上記2つについてもう少し詳しく解説している。

SIMD方式

以下2つのベクトルXと行列Aがあるとする。
$$\begin{align} \bf{X} &= (a, b, c) \\ \bf{A} &= \left( \begin{array}{ccc} a00 & a01 \\ a10 & a12 \\ a21 & a21 \end{array} \right) \\ \bf{Y} &= \bf{X} \cdot \bf{A} \end{align} \ $$

Yを計算する時、SIMDでは先にXの列要素(a)を各演算機にブロードキャストし、Aの行要素(a00, a10, a20)と計算する。この動作をXの列要素分繰り返すことで計算を完了する。
仮に、Xの要素がシェアードメモリ(後述)など、レイテンシのあるメモリに格納されている場合、各ブロードキャストでサイクルを消費することになる。この本ではブロードキャストには1 ~ 4サイクル、積演算自体には5 ~ 8サイクル必要だと書かれている。

出展: 「GPUを支える技術」，p117，Hisa Ando，2017

---

余談だが、上記のような行列演算ライブラリを実装しようとすると、マトリックスをN次のテンソルに拡張して計算を一般化する必要がある。以前自分で実装したときはビャーネ・ストラウストラップ プログラミング言語 C++ 第4版を参考にしたが、2次以上のテンソル実装はのっていなかった。どこかのタイミングでテンソル積の定義を書いた本を読みたいと思っている。他のライブラリはnumpyの仕様に沿ったものが多い気がするけれど、何か参考にした書籍などがあるのだろうか。
読みたい本が多すぎるなー(´ε｀；)…

---

SIMT方式

上記のSIMTの欠点として、ブロードキャストの間は演算を行えないことと、前の命令に依存関係があるような計算は並列に実行できない事が挙げられる。また、以前説明した可変長データに対して効率が減少するという問題もある。

これらに対してSIMTは1つの演算器が1つのデータを処理するので、データをブロードキャストする必要がなく、かつ命令間の依存関係も発生しない。可変長のデータに対してもその分だけ実行すれば良いので、問題は無いはず。
(もちろん演算機数に対してスレッド数が少ないと効率が落ちる?)

演算機の数とかかるサイクル数を表にするとSIMTはSIMTの演算の転置であり、SIMDはベクトルXの要素数が可変の時に効率が落ちるのが分かる。

出展: 「GPUを支える技術」，p118，Hisa Ando，2017

この本はとても丁寧に書かれている分、各章でかなり記述に重複がある気がする。(;￣ー￣)

１命令でいくつスレッドを実行するかはベンダによって違いがあり、
NVIDIは32スレッド(ワープ)、AMは64スレッド(ウェーブフロント)、ARMは4スレッド(クワッド)を基本単位としている。

NVIDIAを例に上げると、基本的に32スレッド同時に実行するので、80スレッドの処理があると、合計3回計算する必要があるが、最後の1回は16個がダミーデータとなる。全体の演算器の有効データに対する利用率は83%ほど。

$$\begin{align} &\lfloor 80 \div 32 \rfloor = 2 \\ \\ &80\ mod\ 16 = 16 \\ \\ &\frac{80}{96} = 0.833 \\ \end{align}$$

利用率の点で言うと、ARMのようにスレッド数が少ないほうが効率がいいが、何度も呼び出すため命令・データフェッチなどのオーバヘッドが増える。用途に応じて最適な割合を選択する必要があるのは、マルチコアCPUでの並列プログラミングと似ている。
3章でもあったように、SIMTはプレディケートコードによって条件分岐ができるが、条件の成立不成立にかかわらず、then, else双方の計算時間がかかる。(後述する)

GPUの構造

NVIDIAとAMD、ARMのGPUの構造をまとめている非常にありがたいパートだ。

NVIDIA Pascal GPUの基礎知識

NVIDIA GP100, GP102などはGPUチップの名前であり、100番は伝統的に科学技術計算用らしい。(ただ、GP102はグラフィック機能もサポートしている)
チップの中にはTPC(Texture Processing Cluster)があり、内部に2つのSM(Streaming Multiprocessor)、さらにその内部に複数の演算器やロード/ストアユニットがある。(2章で既出)

GP100チップ上にはSMが56個搭載されている。正確には60個あるらしいが4個まで不良品があっても出荷できるようにらしい。これは歩留まり率の向上のためで、GPUでは常套手段らしい。P100のWhitepaper があるので、詳しく知りたい方はぜひそちらを参照されたい。

NVIDIA GPUの命令実行のメカニズム

GPUで使うプログラムには以下の２種類がある。

• ホストプログラム
  -> CPU上で実行される
• カーネルプログラム
  -> GPU上で実行される

mainはCPU上にあり一部の関数がGPUを利用する形になる。カーネルプログラムを実行する際は、CPUからその先頭アドレスをGPUに通知する。

出展: 「GPUを支える技術」，p126，Hisa Ando，2017

GPUによるプログラムの実行はスレッドブロックとスレッドという単位で分けられる。
1つのスレッドブロックは複数のスレッドを含み、それらが一つのSMに割当られ計算される。この際、各スレッドは同じプログラムを実行するが、自身の位置などを内部変数として参照できる。(例: blockIdx.x, threadIdx.y)
さらに複数のスレッドブロックをまとめたグリッドという単位がある。グリッド及びスレッドブロックはギガスレッドエンジンでコントロールされる。
具体的な計算フローを以下に示す。

1. CPUがカーネルプログラムを起動
2. N個のスレッドブロックをもつグリッドが実行される
3. N個のスレッドブロックを順番にSMに割り当てる。この際、SMの数よりNが大きい場合は2巡目に再度割り当てを行う
4. SM内でスレッドブロックをワープ(32個のスレッドの塊)に分割し、最大64個までワーププールに格納する
5. ワーププールから実行可能なワープを選択し、連続した2命令をSMで実行する
6. 再度、次に実行するワープを選択して4にもどる

NVIDIA Kepler GPUから、GPUで実行しているカーネルプログラムから直接別のカーネルプログラムを実行するダイナミックパラレリズムという機能も追加された。

SMの構成

細かいSMの構成は、以下のようになっている。

1. FP32演算ユニット x 64 (CUDAコア)
2. FP64演算ユニット x 32 (DPユニット)
3. LD/STユニット x 16
4. Special Function Unit x 16 (逆数、三角関数などを計算)

FP32とFP64演算ユニットは積和演算を実行するので、2演算 / クロックとなる。また、FP32は整数演算にも対応している。整数64bitもDPユニット(CUDAコアではない)でサポートされているのかな？？
GP100チップは56個のSMがあり、64個のCUDAコアを持つため計3,584個のFloat MAC演算ユニットがあることになる。クロック周波数は1,480MHzなので、10.6TFlopsだ。

$$3584 \times 2 \times 1480 = 10,608,640$$

GPUは32bitを基本ビット長として利用しているため、LD/STは基本的に32bitとなっており、FP64は2つのポートを利用して実現している。Pascal GPUからFP32演算器を半分にして計算することでFP16も計算できる。この場合の性能は単純に2倍(21.2TFlops)になる。
Floatは計算方法的に上位と下位に単純にわけられないので、どうやってFP32と2つのFP16を効率的に計算する回路をデザインしているのか興味がある。(^q^)
また、1つのLDポートを利用して2つのFP16をロードできる。ただ、おそらくそれらは連続したアドレスでなければならないはずである。(前の章で128byteを４byte区切りでアクセスできることを説明した)

GPUのメモリシステム

GP100はCPUのようにレジスタファイル、L1、L2、System Memory(DRAM)という記憶階層を持つ。命令体系はRISCと同じでレジスタを演算の入力と出力の両方に利用する。1ワープ分に割り当てられたレジスタのことを1レーンと呼び、1エントリを32bitとして2048エントリ（8KB）ある。1SM(32スレッド)あたりでは65536（262,144KB）となる。(本のp132には1スレッドあたり2048エントリのレジスタとあるが、1ワープあたりとの間違いだろうか??)
基本的には隣のスレッドのレジスタにはアクセスすることができないが、シェアードメモリという独立メモリを利用ると、SM内の異なるスレッド間でデータを共有することができる。また、異なるスレッドのレジスタを読む命令としてshuffle命令というものもあり、こちらはシェアドメモリよりも低レイテンシで利用できる。
SM内で各スレッドに割り当てるレジスタ数はSM内で自由に割合を変更できる。P100の場合、1SMあたり65,536個の32bitレジスタがあり並列実行数は64(ワーププール)なので、1スレッドあたり平均で32エントリである。ただ、スレッド毎に最大で255エントリまで使用するレジスタの割合を変更することができる。これは通常のRISCアーキテクチャに比べて大きな利点である。

$$65536 \div (64 \times 32) = 32 \, \mathrm {entory / average}\\ $$

シェアードメモリ

同一SMで実行されるどのスレッドからもアドレス指定によりアクセスできる独立メモリ。異なるSM間のやり取りはシステムメモリを経由する必要がある。

ロード / ストアユニット

前の章で述べたようにロードストアユニットのバッファは128バイトの幅を持ち、1エントリ(4バイト)単位でアクセスすることができる。仮に、要求されたデータのアドレスが128バイトに含まれていない場合、そのアドレスを含む128バイトを再度バッファにロードする。NVIDIAはこのバッファロードことをリプレイ(Replay)と読んでいる。ロードは128バイトでアラインされている必要があり、100 〜 228などの非アラインアクセスはリプレイの対象となる。リプレイは全ての要求アドレスを満たすまで行われ、最悪のケースは全てのアドレスが128バイトに含まれていないん場合の31回となる。なお、同一アドレスへの書き込みは上書きされるが、その順番は不定。

L1データキャッシュ

NVIDIAが各世代ごとに変更している部分。PascalアーキテクチャではReadonlyとなっており、スヌープなどの処理に対処する必要はなく実装がシンプルになる。SMは命令キャッシュも持っている。

L2データキャッシュ

一般的なCPUとは違い、GPUのL2キャッシュはGDDRメモリと1対1で作られている。これはL1のL2間にクロスバーを挟んでいるということで、メモリサイドキャッシュとも呼ばれる。この構造の問題は1つのGDDRとL2キャッシュペアにアクセスが集中してしまう場合に性能が落ちることだが、ソフトウェア側で要求している連続したアドレスをハッシュして、全てのGDDRにばらばらに配置することで対応している。

デバイスメモリ　(GDDRとHBM)

NVIDIA GPUで利用されているGDDR5は一般的なPCで利用されているDDR4と比べて、約10倍のバンド幅を持っている。

• GDDR5
  • 32bit/cycle
  • 7Gbps
• DDR4
  • 8bit/cyclne
  • 2.4Gbps

ハイエンドGPUはこのGDDR5を12〜16個ほど搭載している。例えば、上記のGDDR5を12個搭載しているGPUのピークバンド幅は336GB/sとなる。

$$4 \times 12 \times 7 = 336 \, \mathrm {GB/s}\\ $$

単純な4Kの入出力だけ考えると6GB/sほどで十分だが、近年の3Dゲーム等では大量のテクスチャデータが利用されており、バンド幅の半分以上がそこに使われることも珍しくないらしい。
近年では、GDDR5Xという改良版(14Gbps)も登場。

GDDR5メモリは非常に高速な信号伝送を行うため、GPUとメモリ間は短い配線で1対1で接続する必要がある。この欠点はデバイスメモリは接続したGDDRメモリの数で固定となってしまうことだ。

上記に対して、TSV(Through Silicon Via)を利用した3D積層DRAM HBMが登場した。HBMは省電力で高バンド幅が得られるが、高額であるためTeslaシリーズのような高価なGPUにしか搭載されていない。

ワープスケジューラ

CPUではパイプライン中にリザルトバイパス等を利用して直接前の命令の演算結果を利用することができるが、GPUは演算結果を一度レジスタファイルに書き込む必要がある(10サイクルほど)。これはワーププールスケジューラがワーププール内の64個のワープから依存関係の無いワープを優先的に実行することで解決される。上記と同様に、システムメモリからのデータロードによる巨大な待ち時間(数100サイクル)もなるべく隠蔽するようになっている。

前の章でもあったように広帯域なP100であったも演算強度は50OP/byteほどであり、50回以上データを使いまわさない限りはメモリネックとなってしまうと考えられる?? (´ε｀;)

プログラマができることとしては、データに依存関係のある計算を引き離したり事前にデータをキャッシュやレジスタに入れることで待ち時間を減らせるらしい。また、NVIDIA GPUはワープ(32スレッド)を一度に実行するので、スレッドブロック中のスレッド数は32の倍数にすると効率的である。

プレディケート実行

NVIDIAのGPUはSIMT内でプレディケート実行をサポートしている。これはプレディケートレジスタに記憶されている値をもとに、次の命令の実行するかどうかを変化させることができる。はじめにsetq命令を任意の条件式で実行し、結果をプレディケートレジスタをセットする。次に命令内にあるプレディケート選択指定に従ってプレディーケートの値を選択してどう実行するかを判断する。命令を実行しない場合、オペランドの読み込みやレジスタファイルへの書き込み、演算の実行などを全て停止する。
このようにワープ内のスレッドが異なる実行をすることをスレッドダイバージェンスと呼ぶ。上述したがif、elseどちらの実行でも同様にサイクルを消費するため、これらの命令数に応じて演算効率は落ちるため、なるべく短くすることが推奨される。

出展: 「GPUを支える技術」，p144，Hisa Ando，2017

本ブログに対するご意見や間違いの指摘などがありましたら、ぜひコメントください。TwitterでもOKです。皆で議論を深めて行けるような場にしていきましょー。