GPGPU の歴史と応用例 長岡技術科学大学電気電子情報工学専攻出川智啓
今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展 性能の変化 シングルコアからマルチコア GPU の応用例 6
GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ チップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い ) ノート PC に搭載 PCI Ex カードとして販売される GPU には, ビデオメモリと呼ばれる RAM が搭載 7
GPU(Graphics Processing Unit) とは 代表的な製品 NVIDIA GeForce, Quadro, Tesla AMD Radeon, FireProm 代表的な用途 3D グラフィックス処理 3D ゲーム,3DCAD,3DCG 作成 エンコード デコード支援 GPU 上に専用モジュールを搭載していることが多い デスクトップ PC の GUI 処理 Windows Aero が比較的高性能な GPU を要求 8
GPU の性能の遷移 ( 理論演算性能 ) Maxwell GeForce 780 Ti 9 Theoretical GFLOP/s Tesla Fermi GeForce GTX 580 GeForce GTX 480 Kepler GeForce GTX 680 Kepler GeForce GTX TITAN GeForce GTX 280 Tesla M2090 GeForce 8800 GTX Tesla C2050 GeForce 7800 GTX Tesla GeForce 6800 Ultra GeForce FX 5800 Woodcrest Harpertown C1060 Sandy Bridge Pentium 4 Bloomfield Westmere Tesla K40 Tesla K20X Ivy Bridge Apr 01 Sep 02 Jan 04 May 05 Oct 06 Feb 08 Jul 09 Nov 10 Apr 12 Aug 13 Dec 14 GeForce ゲーム用 Quadro CG 用 Tesla GPGPU 用 NVIDIA 社が公開している資料を基に作成
GPU の性能の遷移 ( 理論バンド幅 ) GeForce 780 Ti Maxwell Tesla K40 GeForce ゲーム用 Theoretical GB/s Tesla GeForce GTX 280 Fermi GeForce GTX 480 Tesla K20X Kepler GeForce GTX 680 Tesla M2090 Tesla C2050 Quadro CG 用 Tesla GPGPU 用 GeForce 8800 GTX Tesla C1060 Ivy Bridge GeForce 7800 GTX Sandy Bridge Bloomfield GeForce 6800 GT Woodcrest GeForce FX 5900 Prescott Westmere Harpertown Northwood 2003 2004 2005 2006 2007 2008 2009 2010 2011 2012 2013 10 NVIDIA 社が公開している資料を基に作成
GPU(Graphics Processing Unit) の役割 グラフィックスを表示するために様々な処理を行い, 処理の結果をディスプレイに出力 3 次元グラフィックスの発展に伴って役割が大きく変化 描画情報 CPU 3 次元座標変換 ポリゴンとピクセルの対応付け ピクセル色計算テクスチャ参照 フレームバッファ ( ビデオメモリ ) への書き込み ディスプレイ出力 過去 CPU が 3D 描画の演算を実行 GPUが出力 描画情報 3 次元座標変換 ポリゴンとピクセルの対応付け ピクセル色計算テクスチャ参照 フレームバッファ ( ビデオメモリ ) への書き込み ディスプレイ出力 現在 GPUが演算から出力までの全てを担当 CPUは描画情報の生成やGPUへの情報の引き渡し,GPU の制御を行う ディスプレイコントローラ 画面出力 GPU 画面出力 11
GPU の描画の流れ 1. CPUからGPUへ描画情報を送信 2. 頂点処理 ( 頂点シェーダ ) 座標変換 画面上での頂点やポリゴンの位置 大きさの決定 頂点単位での照明の計算 3. 頂点やポリゴンからピクセルを生成 ( ラスタライザ ) 4. ピクセル処理 ( ピクセルシェーダ ) 画面上のピクセルの色 テクスチャの模様 5. 画面出力 ピクセルの色情報をフレームバッファに書き込み 2. 3. 4. 12
ビデオカードの利点 CPU で描画のための演算を行うと,CPU にかかる負荷が大きい 3 次元画像処理の専用回路を備えたハードウェアを導入 CPU にかかる負荷を減らすことができる 頂点 ピクセルごとに並列処理が可能なため, ハードウェアによる並列処理が可能 13
ビデオカードの欠点 3 次元画像処理の専用回路を備えたハードウェアを導入 新しい描画方法を開発しても,GPU へ実装 製品化されるまで利用できない ユーザが所有している GPU によって, 利用できる機能にばらつきが生じる ある描画手法用の専用回路を実装しても, その描画方法が常に使われる訳ではないので GPU 全体の利用効率が下がる 14
ビデオカードから GPU へ CG の多様化と共に固定機能の実装が困難に 頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装 グラフィックスカード GPU 頂点処理用回路 頂点シェーダユニット ピクセル処理用回路 ピクセルシェーダユニット 15
ビデオカードから GPU へ 頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装 処理によっては利用効率に差が生じる 頂点処理重視の処理 GPU 頂点シェーダユニット ピクセル処理重視の処理 GPU 頂点シェーダユニット 空きユニット ピクセルシェーダユニット ピクセルシェーダユニット 空きユニット 16
ビデオカードから GPU へ 頂点シェーダとピクセルシェーダを統合したユニファイドシェーダへの進化 頂点処理とピクセル処理を切り替えることで利用率を高める 頂点処理重視の処理 GPU ユニファイドシェーダユニット ピクセル処理重視の処理 GPU ユニファイドシェーダユニット 17
ビデオカードから GPU へ 各ピクセルに対して並列に処理を行えるよう, 並列度を高める 単純な処理を行う演算器を大量に搭載 高い並列度で処理を行う GPU の誕生と GPGPU の普及 高性能な 3DCG 画像処理への要求 GPU の高性能化 GPGPU の長所 消費電力あたりの浮動小数点理論演算性能が高い 安価 CPU を使った大規模計算機と比較して相対的に 18
GPU の進化 72 Volta 1 Watt あたりの単精度行列 - 行列積の回数 60 48 36 24 12 0 情報処理センター GPGPU システムに搭載 Tesla Fermi Kepler Maxwell Pascal 2008 2010 2012 2014 2016 2018 年 GPU は発展途上 (2~3 年で世代交代 ) NVIDIA 社プレゼンテーションを基に作成 今プログラムを作っておくと, 勝手に速くなってくれる!* * 多分に誇張的な表現であることに注意 19
GPU のハードウェア構造 CUDA Core( 旧 Streaming Processor, SP) と呼ばれる演算器を多数搭載 Streaming Multiprocessor(SM, SMX) が複数の CUDA Core と SFU, メモリをまとめて管理 SFU(Special Function Unit) 数学関数を計算するユニット 複数の SM が集まって GPU を構成 20
Fermi アーキテクチャ Tesla M2050 の仕様 SM 数 14 CUDA Core 数 448(=32Core/SM 14SM) L1/L2 キャッシュを搭載 ECC( 誤り訂正機能 ) を搭載 21
Fermi アーキテクチャ Tesla M2050 の仕様 * CUDA コア数 ( 単精度 ) CUDA コアクロック周波数 単精度演算ピーク性能 倍精度演算ユニット数 倍精度演算ピーク性能 メモリクロック周波数 メモリバス幅 最大メモリバンド幅 * 1 単精度 CUDA Core を 2 基使って倍精度演算を実行 *http://www.nvidia.co.jp/object/product_tesla_m2050 _M2070_jp.html http://ja.wikipedia.org/wiki/nvidia_tesla 448 Cores 1,150 MHz 1.03 TFLOPS 0* 1 Unit 515 GFLOPS 1.55 GHz 384 bit 148 GB/s 22
Maxwell アーキテクチャ GeForce GTX TITAN X の仕様 SM 数 24 CUDA Core 数 3,072(=128Core/SM 24SM) 23
Maxwell アーキテクチャ GeForce GTX TITAN X の仕様 * CUDA コア数 ( 単精度 ) CUDA コアクロック周波数 単精度演算ピーク性能 3,072 Cores 1,002 MHz 6.14 TFLOPS 倍精度演算ユニット数 0* 1 Units 倍精度演算ピーク性能 192 GFLOPS* 2 メモリクロック周波数 3.5 GHz* 3 メモリバス幅 最大メモリバンド幅 *http://ja.wikipedia.org/wiki/flops http://http://www.geforce.com/hardware/desk top gpus/geforce gtx titan x/specifications 384 bit 336.5 GB/s * 1 http://www.4gamer.net/games/121/g012181/20141225075/ * 2 倍精度演算は単精度演算の性能の 1/32 (1/16 Flop/Core/clock) * 3 DDR(Double Data Rate) 7GHz 相当と書かれている場合もある 24
Pascal アーキテクチャ 2016 年にリリース予定 倍精度演算器を搭載予定 NVLink GPU 同士や GPU と CPU を接続する独自の方式 通信 (CPU メモリ PCI Express メモリ GPU) のボトルネックを解消 (PCI Express3.0 の 5~12 倍 ) 複数の GPU を使って大規模な計算が可能 3D メモリ (High Bandwidth Memory, HBM)* 3 次元積層技術を利用し, メモリの容量と帯域を大幅に増加 最大 32GB, メモリ帯域 1TB/s *http://pc.watch.impress.co.jp/docs/column/kaigai/20150421_698806.html 25
CUDA Compute Unified Device Architecture NVIDIA 社製 GPU 向け開発環境 (Windows,Linux,Mac OS X) 2007 年頃発表 C/C++ 言語 + 独自のGPU 向け拡張 専用コンパイラ (nvcc) とランタイムライブラリ いくつかの数値計算ライブラリ ( 線形代数計算,FFTなど) CUDA 登場以前 グラフィックスプログラミングを利用 足し算を行うために, 色を混ぜる処理を実行 汎用計算のためには多大な労力が必要 26
プログラマブルシェーダを用いた汎用計算 グラフィックス API(DirectX, OpenGL) による描画処理 + シェーダ言語 (HLSL, GLSL) による演算 void gpumain(){ vec4 ColorA = vec4(0.0, 0.0, 0.0, 0.0); vec4 ColorB = vec4(0.0, 0.0, 0.0, 0.0); vec2 TexA = vec2(0.0, 0.0); vec2 TexB = vec2(0.0, 0.0); TexA.x = gl_fragcoord.x; TexA.y = gl_fragcoord.y; TexB.x = gl_fragcoord.x; TexB.y = gl_fragcoord.y; } ColorA = texrect( texunit0, TexA ); ColorB = texrect( texunit1, TexB ); gl_fragcolor = F_ALPHA*ColorA + F_BETA*ColorB; シェーダ言語を用いた配列加算 (c= *a + *b) の例 void main(){ glutinit( &argc, argv ); glutinitwindowsize(64,64);glutcreatewindow("gpgpuhelloworld"); glgenframebuffersext(1, &g_fb); glbindframebufferext(gl_framebuffer_ext, g_fb); glgentextures(4, g_ntexid); // create (reference to) a new texture glbindtexture(opt1, texid); gltexparameteri(opt1, GL_TEXTURE_MIN_FILTER, GL_NEAREST); gltexparameteri(...); glteximage2d(opt1, 0, opt2, width, height, 0, GL_RGBA, GL_FLOAT, 0); ( 以下省略 ) GPU の処理 (GLSL) 各ピクセルに対して実行 CPU の処理 (OpenGL) 27
CUDA による汎用計算 (c= *a + *b) #define N (1024*1024) #define Nbytes (N*sizeof(float)) #define NT 256 #define NB (N/NT) global void init(float *a, float *b, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; a[i] = 1.0; b[i] = 2.0; c[i] = 0.0; } global void add(float *a, float, float *b, float, float *c){ int i = blockidx.x*blockdim.x + threadidx.x; } c[i] = *a[i] + *b[i]; int main(){ float *a,*b,*c; float, ; } cudamalloc((void **)&a, Nbytes); cudamalloc((void **)&b, Nbytes); cudamalloc((void **)&c, Nbytes); =... ; =...; init<<< NB, NT >>>(a,b,c); add<<< NB, NT >>>(a,, b,, c); return 0; 28
GPU の普及の要因 GPU の進展は 15 年程度 普及の速度は驚異的 CPU は数十年かけて進展 CPU も驚異的な速度で進展 様々な高速化技術が導入 GPU が普及している要因は何か? 29
TOP500 List(2014, Nov.) スーパーコンピュータの性能の世界ランキング GPU を搭載したコンピュータは 2 基だけ 計算機名称 ( 設置国 ) アクセラレータ 実効性能 (PFlop/s) 消費電力 (MW) 1 Tianhe 2 (China) Intel Xeon Phi 33.9 17.8 2 Titan (U.S.A.) NVIDIA K20x 17.6 8.20 3 Sequoia (U.S.A.) 17.2 7.90 4 K computer (Japan) 10.5 12.7 5 Mira (U.S.A.) 8.59 3.95 6 Piz Daint (Switzerland) NVIDIA K20x 6.27 2.33 7 Stampede (U.S.A.) Intel Xeon Phi 5.17 4.51 8 JUQUEEN (Germany) 5.01 2.30 9 Vulcan (U.S.A.) 4.29 1.97 10 (U.S.A.) 3.58 1.50 http://www.top500.org/ より引用 30
理論ピーク性能と実効性能 Floating Point Operations per Second 1 秒あたりに浮動小数の演算が何回できるか 理論ピーク性能 プロセッサの数 ( プロセッサ上に実装された演算器の数 ) や動作周波数から求める理論的な速度 全ての機能が全て同時に使えれば という理論的な値 実効性能 ( 実行性能 ) ある問題に対してプログラムを実行したときに得られた性能 プログラムの中で行っている計算 ( 浮動小数点演算 ) の回数を数え, プログラムの実行にかかった時間を測定して割り算 31
CPU の理論性能 公式 FLOPS = 1 コアの演算性能 [?] コア数 [Core] CPU の動作周波数 [Hz=Clock/sec] 1 コアの演算性能 =1 度に発行出来る浮動小数点演算命令 単位は [Floating Point Operations/Clock /Core] 性能の評価には動作周波数だけでなく,1 コアが 1 クロックで発行できる命令数が重要 32
CPU の性能 FLOPS = 1 コアの演算性能 コア数 CPU の動作周波数 1 コアの演算性能の向上 演算器 ( トランジスタ ) の増加 コア数の増加 トランジスタ数の増加 CPU の動作周波数 回路の効率化や印可電圧の向上 様々な機能を追加 パイプライン処理 スーパースカラ実行 分岐予測等 動作周波数の向上に注力 ( ほぼ全ての処理が速くなる ) 33
CPU の性能の変化 Intel の予告 (Intel Developer Forum 2003) 2007 年頃には 10GHz に達する 34 Intel が公開している資料を基に作成 http://pc.watch.impress.co.jp/docs/2003/0227/kaigai01.htm で見ることができる
CPU の性能の変化 2004 年頃からクロックが停滞 35 Intel が公開している資料を基に作成 ASCII.technologies(Dec 2009) や http://www.gdep.jp/page/view/248 で見ることができる
CPU の性能向上 * * 姫野龍太郎, 絵でわかるスーパーコンピュータ, 講談社 (2012) 電子回路の構成部品 機械式リレー 真空管 トランジスタ IC (Integrated Circuit) LSI (Large Scale Integrated Circuit) CPUとメモリの性能向上 回路の配線を細線化 250nm 180nm 130nm 90nm 65nm 45nm 32nm 線幅が減ると同じ回路を作る際の面積が減少 同じ面積に集積できるトランジスタ数が増加 集積率が上昇 36
ムーアの法則 * インテルの共同設立者ムーアによる経験則 半導体の集積率は 1 年で倍になる 後に 18 ヶ月で 2 倍 に修正姫野龍太郎, 絵でわかるスーパーコンピュータ, Number of Transistors 10 10 10 9 10 8 10 7 10 6 10 5 286 ムーアの法則 (12 ヶ月で倍 ) Intel386 プロセッサ Intel486 プロセッサ *Moore, G.E., Electronics, Vol.38,No.8(1965). http://ja.wikipedia.org/wiki/ ムーアの法則 http://en.wikipedia.org/wiki/moore%27s_law 講談社 (2012) に掲載されている絵を基に作成 ムーアの法則 (18 ヶ月で倍 ) インテル Pentium プロセッサ インテル Pentium 4 プロセッサ インテル Pentium III プロセッサ インテル Pentium II プロセッサ デュアルコアインテル Itanium 2 プロセッサ インテルItanium 2プロセッサインテルItaniumプロセッサ 37 10 4 10 3 8080 8008 4004 8086 1970 1975 1980 1985 1990 1995 2000 2005 2010 year
ポラックの法則 * *http://ja.wikipedia.org/wiki/ ポラックの法則 http://en.wikipedia.org/wiki/pollack%27s_rule 2 倍のトランジスタを使っても, プロセッサの性能はその平方根倍 (1.4 倍 ) 程度にしか伸びない 消費電力は 2 倍, 性能は 1.4 倍 一つの CPU に複数のプロセッサ ( コア ) を搭載 消費電力を上げずに 理論的な 性能を倍に プログラムの作り方に工夫が必要 38
プロセッサの性能向上 性能向上 半導体回路の細線化 動作周波数向上 消費電力が低下 低下分の電力をトランジスタのスイッチングに利用 39
プロセッサの性能向上 性能向上 半導体回路の細線化 2 倍のトランジスタを使っても性能は 1.4 倍程度にしか伸びない 絶縁部が狭くなり漏れ電流が発生, 電力が低下しない 動作周波数向上 消費電力が低下 消費電力の増加によって発熱量が増加, 空冷の限界 低下分の電力をトランジスタのスイッチングに利用 40
プロセッサの性能向上 性能向上 半導体回路の細線化 コア数の増加 2 倍のトランジスタを使っても性能は 1.4 倍程度にしか伸びない 絶縁部が狭くなり漏れ電流が発生, 電力が低下しない 動作周波数向上 消費電力が低下 消費電力の増加によって発熱量が増加, 空冷の限界 低下分の電力をトランジスタのスイッチングに利用 41
CPU の性能向上 FLOPS = 1 コアの演算性能 コア数 CPU の動作周波数 1 コアの演算性能の向上 演算器 ( トランジスタ ) の増加 コア数の増加 トランジスタ数の増加 CPU の動作周波数 回路の効率化や印可電圧の向上 コンパイラの最適化を利用 複数のコアを使うようにプログラムを書かないと速くならない 劇的な性能向上は期待できない 42
GPU を使うという選択 GPU 普及の要因の一つは CPU クロックの頭打ち クロックを下げてマルチコア化した CPU への対応が必要 なぜ GPU という選択か? CPU 用プログラムの並列化でもいいのでは? 消費電力の低減 数値計算や高性能計算 (HPC) の業界が GPU に注目 スーパーコンピュータの性能向上 高機能な CPU を大量に使うと消費電力が問題に 高機能な制御用プロセッサと, 計算を実行する低性能なアクセラレータの組み合わせ 43
Green500(2014, Nov.) AMD 社の GPU が 1 位 日本の次世代機 睡蓮 が 2 位 * NVIDIA 社の GPU が 3 位以降にランクイン 計算機名称アクセラレータ GFLOPS/W 消費電力 (kw) 1 L CSC AMD FirePro S9150 5.27 57.15 2 Suiren PEZY SC 4.95 37.83 3 TSUBAME KFC NVIDIA K20x 4.45 35.39 4 Storm1 NVIDIA K40m 3.96 44.54 5 Wilkes NVIDIA K20 3.63 52.62 6 idataplex NVIDIA K20x 3.54 54.60 7 HA PACS NVIDIA K20x 3.52 78.77 8 Cartesius Accelerator Island NVIDIA K20m 3.46 44.40 9 Piz Daint NVIDIA K20x 3.19 1,753.66 10 romeo NVIDIA K20x 3.13 81.41 44 *4 月 2 日に 1 位相当の値を記録 http://news.mynavi.jp/articles/ 2015/04/02/kek_suiren/ http://www.green500.org/ より引用
アクセラレータ コンピュータの特定の機能や処理能力を向上させるハードウェア CPU で行っていた処理を専用ハードウェアが担当 動画像のエンコード デコード等 コンピュータシミュレーションでは CPU の代わりに計算を実行するハードウェアを指す 画像処理装置 (Graphics Processing Unit) メニーコアプロセッサ 45
アクセラレータ ( メニーコアプロセッサ ) PEZY SC 株式会社 PEZY Computing の 1,024 コアの低消費電力型メニーコアプロセッサ 1024 コア, 動作周波数 733MHz 理論演算性能 単精度 3.0 TFLOPS 倍精度 1.5 TFLOPS 46
アクセラレータ ( メニーコアプロセッサ ) Intel Xeon Phi OS を搭載しており, 接続しているワークステーションとは独立して動かすことが可能 61 コア CPU(1GHz), メモリ 8GB の Linux サーバ 理論演算性能 ( 単精度 ) 約 1 TFLOPS CPU からの制御が必要なアクセラレータとは異なる アーキテクチャが Intel CPU と同じであるため, コンパイルし直すだけで動作する 新モデルを投入予定 * 72コア, メモリ16GB 理論演算性能 3.0 TFLOPS *http://news.mynavi.jp/articles/2014/11/17/sc14/ 47
まとめ GPU の特徴 低性能の演算器を大量に搭載 (~3000 コア ) GPU が使われる理由 理論演算性能が高い メモリとチップ間の帯域も広い 省電力と高性能を両立 今後の計算機の主流になると考えられる 将来に対する投資 GPU だけでなく, 制御用 CPU+ 計算用アクセラレータという思想は今後しばらく主流であり続ける 48
ロボットによる心臓外科手術 ロボットアームと内視鏡を使った心臓外科手術 人工心臓を使わず, 心臓を動かしたまま行う 49
ロボットによる心臓外科手術 https://www.youtube.com/watch?v=g3liyn_oceo 50
ロボットによる心臓外科手術 心臓は複雑な形状で, かつ周期的に脈動 毎秒何十枚と撮られる映像をリアルタイムで処理しながらロボットアームを制御 心臓を 2 次元の画像に変換 その画像を基にロボットが動く道筋を計算 実際に 3 次元の動きに変換 心臓の動きに合わせてロボットアームを制御 51
住宅設備機器開発 住宅設備機器開発のための混相流シミュレーション 池端昭夫, 吉田慎也, 住宅設備機器開発のための混相流シミュレーション, 第 27 回数値流体力学シンポジウム講演予稿集, E05-4, 2013(on USB Flash Drive). の結果を著者が撮影 52
車載システム Google, NVIDIA, Audi などによる Open Automotive Alliance 自動車への Android プラットフォーム搭載促進を目指す 53
NVIDIA 社の自動車関連ソリューション NVIDIA DRIVE PX 高機能運転支援システム NVIDIA DRIVE CX 自動車の高度情報化 54
Jetson Tegra TK1 世界初の組み込みスーパーコンピューター NVIDIA Tegra K1を採用 Keplarアーキテクチャ 192 CUDA コア NVIDIA 4 Plus 1 クアッドコア ARM Cortex A15 CPU Linux for Tegra が動作 55
月面探査 Google Lunar XPRIZE Google による国際宇宙開発レース 日本の民間月面探査チーム HAKUTO が参加 ロケット発射時のシミュレーション 着陸などのシミュレーションや実際の制御に GPU を用いた自動運転技術を利用 56