スキップしてメイン コンテンツに移動

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サイクル必要だと書かれている。
enter image description here
出展: 「GPUを支える技術」，p117，Hisa Ando，2017


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


SIMT方式

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

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

演算機の数とかかるサイクル数を表にするとSIMTはSIMTの演算の転置であり、SIMDはベクトルXの要素数が可変の時に効率が落ちるのが分かる。
enter image description here
出展: 「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に通知する。

enter image description here
出展: 「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どちらの実行でも同様にサイクルを消費するため、これらの命令数に応じて演算効率は落ちるため、なるべく短くすることが推奨される。
enter image description here
出展: 「GPUを支える技術」，p144，Hisa Ando，2017

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

コメント

コメントを投稿

このブログの人気の投稿

「A Survey of FPGA Based Neural Network Accelerator」の邦訳

A survey of FPGA Based Neural Network Accelerator 今回は2/17日にドワンゴさんで行われたFPGAXでの発表で紹介したサーベイ論文「A Survey of FPGA Based Neural Network Accelerator」の邦訳を掲載することにする。 稚拙ながら発表スライドも以下にあるので興味のある方はぜひ御覧ください。( ´∀｀) https://www.slideshare.net/leapmind/an-introduction-of-dnn-compression-technology-and-hardware-acceleration-on-fpga-88557866 A Survey of FPGA Based Neural Network Accelerator Kaiyuan Guo, Shulin Zeng, Jincheng Yu, Yu Wang, Huazhong Yang https://arxiv.org/abs/1712.08934 Abstruct ニューラルネットは画像やスピーチ、ビデオ認識などの領域に適用され、良い結果を残している。しかし、その計算のストレージや複雑度が、アプリケーションでの活用を難しくしている。CPUでの計算は難しいため、GPUが最初の選択肢となる。 一方で、FPGAを基にしたアクセラレータも研究の対象となってきている。なぜなら、特化したハードウェアの設計はGPUを速度とエネルギー効率で超えるための、有力な選択肢であるためである。様々なFPGAベースのアクセラレータがソフトウェアとハードウェアの最適化手法を用いて提案されてきている。本稿では、それらの全体図とそれらの主要な技術の要約を行う。この調査はFPGAベースのNNアクセラレータについて、そのソフトウェア側からハードウェア側、回路レベルからシステムレベルに渡
コメントを投稿
続きを読む

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

GPUを支える技術読み始めた 第5章[前半] 新年の休みもとうに終わり、皆さんどうお過ごしだろうか。 ちなみに主は年末読む予定であった本や論文を全く消化できずに今に至ってしまった。 年末年始は時間があったので、ひたすら読んでいたはずなのになぜだ。。＿|￣|○ 読むのが遅いのがいけないのか？翻訳するのが遅いのか？？それともそれをまとめるのに時間がかかっているのか。 おそらく全て当てはまるが、時間がかかるものは仕方がないので、少しずつ慣れていくしかないなー。。（-_-;） 余談だが、年末にブログ執筆環境を再構築した。 今まではwebエディタの Classeur を利用していたが、やはり純粋なWebアプリケーションなのでレンダリングやアップデートに問題があった。そこで、今までも利用していたAtomを少しカスタマイズして試してみたが、どうにも動作が重くかつ、vim-modeがいい感じにならずに残念に思っていたところ、vscodeのことを思い出した。どちらもエレクトロンベースだが、vscodeはAtomより全然軽く、markdownプラグインも豊富にあるので、すぐにmarkdown+mathjax環境を構築することができた。いやはや、世の中は便利になったものだ（＾ω＾） せっかくなので、下に執筆環境のスクリーンショットを自慢げに貼ってみようと思う。(markdownは頻繁に見る必要はないので、普段はプレビューは別のタブで開いている) もしかして、こういうことしてるから時間がかかるのかな？？／(^o^)＼ それでは、本題に戻ろう。 今回はGPU支える技術の第5章だ。4章と同じく楽しみにしていた章なのでじっくり読んでいきたい。 なお今回も長い章なので、前半と後半に分けてまとめと感想を書いていく。 第5章 GPUプログラミングの基本[前半] GPUの超並列プロセッサでプログラムを実行するには、超並列で実行でき
コメントを投稿
続きを読む

「VTA: An Open Hardware-Software Stack for Deep Learning」の翻訳と感想

tech_report VTA: An Open Hardware-Software Stack for Deep Learning の翻訳と簡単な感想を書きました(^o^) TVM開発チームには許可を頂いていますが、もし誤訳があったら教えてください。原文は以下です。 https://arxiv.org/abs/1807.04188 1. VTAとは Versatile Tensor Acceleratorの略。 汎用的かつ高速、効率的なHW Deep learningアクセラレータを提供するスタックの総称。 提供されるアクセラレータはプログラム可能であり、柔軟にカスタマイズができる。 VTAは生産性重視で高レベル記述が可能なDLフレームワークと性能重視で低レベルな基盤ハードウェア(FPGAなど)におけるブリッジになることを目的としている。 そのため、アクセラレータ以外にもドライバやJitランタイム、TVMに基づいた最適化済みコンパイラ環境も提供しており、End-to-Endで動作する。 また、FPGAにデプロイするためのインターフェイスや振る舞いHWシュミレーション機能も含まれている。 どのような人や用途に有効なのか HWデザイナとコンピュータアーキテクト 最適化コンパイラのリサーチャ Deep learningのリサーチャ 全体構成 NNVM Intermediate Representation TVM Intermediate Representation VTA JIT Runtime VTA Instruction Set Architecture VTA Hardware Micro-Architecture 2. VTA Hardware Architecture VTAは他の主要なDLアクセラレータと同様に、密線形代数演算を高速に実行できるように設計されている。(Google TP
コメントを投稿
続きを読む