今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展性能の変化シングルコアからマルチコア GPU の応用例 6

Size: px

Start display at page:

Download "今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展性能の変化シングルコアからマルチコア GPU の応用例 6"

るるみわかはら
6 years ago
Views:

1 GPGPU の歴史と応用例長岡技術科学大学電気電子情報工学専攻出川智啓

2 今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展性能の変化シングルコアからマルチコア GPU の応用例 6

3 GPU(Graphics Processing Unit) とは画像処理専用のハードウェア具体的には画像処理用のチップチップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い ) ノート PC に搭載 PCI Ex カードとして販売される GPU には, ビデオメモリと呼ばれる RAM が搭載 7

4 GPU(Graphics Processing Unit) とは代表的な製品 NVIDIA GeForce, Quadro, Tesla AMD Radeon, FireProm 代表的な用途 3D グラフィックス処理 3D ゲーム,3DCAD,3DCG 作成エンコードデコード支援 GPU 上に専用モジュールを搭載していることが多いデスクトップ PC の GUI 処理 Windows Aero が比較的高性能な GPU を要求 8

5 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 社が公開している資料を基に作成

6 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 NVIDIA 社が公開している資料を基に作成

7 GPU(Graphics Processing Unit) の役割グラフィックスを表示するために様々な処理を行い, 処理の結果をディスプレイに出力 3 次元グラフィックスの発展に伴って役割が大きく変化描画情報 CPU 3 次元座標変換ポリゴンとピクセルの対応付けピクセル色計算テクスチャ参照フレームバッファ ( ビデオメモリ ) への書き込みディスプレイ出力過去 CPU が 3D 描画の演算を実行 GPUが出力描画情報 3 次元座標変換ポリゴンとピクセルの対応付けピクセル色計算テクスチャ参照フレームバッファ ( ビデオメモリ ) への書き込みディスプレイ出力現在 GPUが演算から出力までの全てを担当 CPUは描画情報の生成やGPUへの情報の引き渡し,GPU の制御を行うディスプレイコントローラ画面出力 GPU 画面出力 11

8 GPU の描画の流れ 1. CPUからGPUへ描画情報を送信 2. 頂点処理 ( 頂点シェーダ ) 座標変換画面上での頂点やポリゴンの位置大きさの決定頂点単位での照明の計算 3. 頂点やポリゴンからピクセルを生成 ( ラスタライザ ) 4. ピクセル処理 ( ピクセルシェーダ ) 画面上のピクセルの色テクスチャの模様 5. 画面出力ピクセルの色情報をフレームバッファに書き込み

9 ビデオカードの利点 CPU で描画のための演算を行うと,CPU にかかる負荷が大きい 3 次元画像処理の専用回路を備えたハードウェアを導入 CPU にかかる負荷を減らすことができる頂点ピクセルごとに並列処理が可能なため, ハードウェアによる並列処理が可能 13

10 ビデオカードの欠点 3 次元画像処理の専用回路を備えたハードウェアを導入新しい描画方法を開発しても,GPU へ実装製品化されるまで利用できないユーザが所有している GPU によって, 利用できる機能にばらつきが生じるある描画手法用の専用回路を実装しても, その描画方法が常に使われる訳ではないので GPU 全体の利用効率が下がる 14

11 ビデオカードから GPU へ CG の多様化と共に固定機能の実装が困難に頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装グラフィックスカード GPU 頂点処理用回路頂点シェーダユニットピクセル処理用回路ピクセルシェーダユニット 15

12 ビデオカードから GPU へ頂点処理とピクセル処理をユーザが書き換えられるプログラマブルシェーダの実装処理によっては利用効率に差が生じる頂点処理重視の処理 GPU 頂点シェーダユニットピクセル処理重視の処理 GPU 頂点シェーダユニット空きユニットピクセルシェーダユニットピクセルシェーダユニット空きユニット 16

13 ビデオカードから GPU へ頂点シェーダとピクセルシェーダを統合したユニファイドシェーダへの進化頂点処理とピクセル処理を切り替えることで利用率を高める頂点処理重視の処理 GPU ユニファイドシェーダユニットピクセル処理重視の処理 GPU ユニファイドシェーダユニット 17

14 ビデオカードから 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

15 GPU の進化 72 Volta 1 Watt あたりの単精度行列 - 行列積の回数情報処理センター GPGPU システムに搭載 Tesla Fermi Kepler Maxwell Pascal 年 GPU は発展途上 (2~3 年で世代交代 ) NVIDIA 社プレゼンテーションを基に作成今プログラムを作っておくと, 勝手に速くなってくれる!* * 多分に誇張的な表現であることに注意 19

16 GPU のハードウェア構造 CUDA Core( 旧 Streaming Processor, SP) と呼ばれる演算器を多数搭載 Streaming Multiprocessor(SM, SMX) が複数の CUDA Core と SFU, メモリをまとめて管理 SFU(Special Function Unit) 数学関数を計算するユニット複数の SM が集まって GPU を構成 20

17 Fermi アーキテクチャ Tesla M2050 の仕様 SM 数 14 CUDA Core 数 448(=32Core/SM 14SM) L1/L2 キャッシュを搭載 ECC( 誤り訂正機能 ) を搭載 21

18 Fermi アーキテクチャ Tesla M2050 の仕様 * CUDA コア数 ( 単精度 ) CUDA コアクロック周波数単精度演算ピーク性能倍精度演算ユニット数倍精度演算ピーク性能メモリクロック周波数メモリバス幅最大メモリバンド幅 * 1 単精度 CUDA Core を 2 基使って倍精度演算を実行 * _M2070_jp.html Cores 1,150 MHz 1.03 TFLOPS 0* 1 Unit 515 GFLOPS 1.55 GHz 384 bit 148 GB/s 22

19 Maxwell アーキテクチャ GeForce GTX TITAN X の仕様 SM 数 24 CUDA Core 数 3,072(=128Core/SM 24SM) 23

20 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 メモリバス幅最大メモリバンド幅 * top gpus/geforce gtx titan x/specifications 384 bit GB/s * 1 * 2 倍精度演算は単精度演算の性能の 1/32 (1/16 Flop/Core/clock) * 3 DDR(Double Data Rate) 7GHz 相当と書かれている場合もある 24

21 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 * 25

22 CUDA Compute Unified Device Architecture NVIDIA 社製 GPU 向け開発環境 (Windows,Linux,Mac OS X) 2007 年頃発表 C/C++ 言語 + 独自のGPU 向け拡張専用コンパイラ (nvcc) とランタイムライブラリいくつかの数値計算ライブラリ ( 線形代数計算,FFTなど) CUDA 登場以前グラフィックスプログラミングを利用足し算を行うために, 色を混ぜる処理を実行汎用計算のためには多大な労力が必要 26

23 プログラマブルシェーダを用いた汎用計算グラフィックス 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

24 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

25 GPU の普及の要因 GPU の進展は 15 年程度普及の速度は驚異的 CPU は数十年かけて進展 CPU も驚異的な速度で進展様々な高速化技術が導入 GPU が普及している要因は何か? 29

26 TOP500 List(2014, Nov.) スーパーコンピュータの性能の世界ランキング GPU を搭載したコンピュータは 2 基だけ計算機名称 ( 設置国 ) アクセラレータ実効性能 (PFlop/s) 消費電力 (MW) 1 Tianhe 2 (China) Intel Xeon Phi Titan (U.S.A.) NVIDIA K20x Sequoia (U.S.A.) K computer (Japan) Mira (U.S.A.) Piz Daint (Switzerland) NVIDIA K20x Stampede (U.S.A.) Intel Xeon Phi JUQUEEN (Germany) Vulcan (U.S.A.) (U.S.A.) より引用 30

27 理論ピーク性能と実効性能 Floating Point Operations per Second 1 秒あたりに浮動小数の演算が何回できるか理論ピーク性能プロセッサの数 ( プロセッサ上に実装された演算器の数 ) や動作周波数から求める理論的な速度全ての機能が全て同時に使えればという理論的な値実効性能 ( 実行性能 ) ある問題に対してプログラムを実行したときに得られた性能プログラムの中で行っている計算 ( 浮動小数点演算 ) の回数を数え, プログラムの実行にかかった時間を測定して割り算 31

28 CPU の理論性能公式 FLOPS = 1 コアの演算性能 [?] コア数 [Core] CPU の動作周波数 [Hz=Clock/sec] 1 コアの演算性能 =1 度に発行出来る浮動小数点演算命令単位は [Floating Point Operations/Clock /Core] 性能の評価には動作周波数だけでなく,1 コアが 1 クロックで発行できる命令数が重要 32

29 CPU の性能 FLOPS = 1 コアの演算性能コア数 CPU の動作周波数 1 コアの演算性能の向上演算器 ( トランジスタ ) の増加コア数の増加トランジスタ数の増加 CPU の動作周波数回路の効率化や印可電圧の向上様々な機能を追加パイプライン処理スーパースカラ実行分岐予測等動作周波数の向上に注力 ( ほぼ全ての処理が速くなる ) 33

30 CPU の性能の変化 Intel の予告 (Intel Developer Forum 2003) 2007 年頃には 10GHz に達する 34 Intel が公開している資料を基に作成で見ることができる

31 CPU の性能の変化 2004 年頃からクロックが停滞 35 Intel が公開している資料を基に作成 ASCII.technologies(Dec 2009) やで見ることができる

32 CPU の性能向上 * * 姫野龍太郎, 絵でわかるスーパーコンピュータ, 講談社 (2012) 電子回路の構成部品機械式リレー真空管トランジスタ IC (Integrated Circuit) LSI (Large Scale Integrated Circuit) CPUとメモリの性能向上回路の配線を細線化 250nm 180nm 130nm 90nm 65nm 45nm 32nm 線幅が減ると同じ回路を作る際の面積が減少同じ面積に集積できるトランジスタ数が増加集積率が上昇 36

33 ムーアの法則 * インテルの共同設立者ムーアによる経験則半導体の集積率は 1 年で倍になる後に 18 ヶ月で 2 倍に修正姫野龍太郎, 絵でわかるスーパーコンピュータ, Number of Transistors ムーアの法則 (12 ヶ月で倍 ) Intel386 プロセッサ Intel486 プロセッサ *Moore, G.E., Electronics, Vol.38,No.8(1965). ムーアの法則講談社 (2012) に掲載されている絵を基に作成ムーアの法則 (18 ヶ月で倍 ) インテル Pentium プロセッサインテル Pentium 4 プロセッサインテル Pentium III プロセッサインテル Pentium II プロセッサデュアルコアインテル Itanium 2 プロセッサインテルItanium 2プロセッサインテルItaniumプロセッサ year

34 ポラックの法則 * * ポラックの法則 2 倍のトランジスタを使っても, プロセッサの性能はその平方根倍 (1.4 倍 ) 程度にしか伸びない消費電力は 2 倍, 性能は 1.4 倍一つの CPU に複数のプロセッサ ( コア ) を搭載消費電力を上げずに理論的な性能を倍にプログラムの作り方に工夫が必要 38

35 プロセッサの性能向上性能向上半導体回路の細線化動作周波数向上消費電力が低下低下分の電力をトランジスタのスイッチングに利用 39

36 プロセッサの性能向上性能向上半導体回路の細線化 2 倍のトランジスタを使っても性能は 1.4 倍程度にしか伸びない絶縁部が狭くなり漏れ電流が発生, 電力が低下しない動作周波数向上消費電力が低下消費電力の増加によって発熱量が増加, 空冷の限界低下分の電力をトランジスタのスイッチングに利用 40

37 プロセッサの性能向上性能向上半導体回路の細線化コア数の増加 2 倍のトランジスタを使っても性能は 1.4 倍程度にしか伸びない絶縁部が狭くなり漏れ電流が発生, 電力が低下しない動作周波数向上消費電力が低下消費電力の増加によって発熱量が増加, 空冷の限界低下分の電力をトランジスタのスイッチングに利用 41

38 CPU の性能向上 FLOPS = 1 コアの演算性能コア数 CPU の動作周波数 1 コアの演算性能の向上演算器 ( トランジスタ ) の増加コア数の増加トランジスタ数の増加 CPU の動作周波数回路の効率化や印可電圧の向上コンパイラの最適化を利用複数のコアを使うようにプログラムを書かないと速くならない劇的な性能向上は期待できない 42

39 GPU を使うという選択 GPU 普及の要因の一つは CPU クロックの頭打ちクロックを下げてマルチコア化した CPU への対応が必要なぜ GPU という選択か? CPU 用プログラムの並列化でもいいのでは? 消費電力の低減数値計算や高性能計算 (HPC) の業界が GPU に注目スーパーコンピュータの性能向上高機能な CPU を大量に使うと消費電力が問題に高機能な制御用プロセッサと, 計算を実行する低性能なアクセラレータの組み合わせ 43

40 Green500(2014, Nov.) AMD 社の GPU が 1 位日本の次世代機睡蓮が 2 位 * NVIDIA 社の GPU が 3 位以降にランクイン計算機名称アクセラレータ GFLOPS/W 消費電力 (kw) 1 L CSC AMD FirePro S Suiren PEZY SC TSUBAME KFC NVIDIA K20x Storm1 NVIDIA K40m Wilkes NVIDIA K idataplex NVIDIA K20x HA PACS NVIDIA K20x Cartesius Accelerator Island NVIDIA K20m Piz Daint NVIDIA K20x , romeo NVIDIA K20x *4 月 2 日に 1 位相当の値を記録 /04/02/kek_suiren/ より引用

41 アクセラレータコンピュータの特定の機能や処理能力を向上させるハードウェア CPU で行っていた処理を専用ハードウェアが担当動画像のエンコードデコード等コンピュータシミュレーションでは CPU の代わりに計算を実行するハードウェアを指す画像処理装置 (Graphics Processing Unit) メニーコアプロセッサ 45

42 アクセラレータ ( メニーコアプロセッサ ) PEZY SC 株式会社 PEZY Computing の 1,024 コアの低消費電力型メニーコアプロセッサ 1024 コア, 動作周波数 733MHz 理論演算性能単精度 3.0 TFLOPS 倍精度 1.5 TFLOPS 46

43 アクセラレータ ( メニーコアプロセッサ ) Intel Xeon Phi OS を搭載しており, 接続しているワークステーションとは独立して動かすことが可能 61 コア CPU(1GHz), メモリ 8GB の Linux サーバ理論演算性能 ( 単精度 ) 約 1 TFLOPS CPU からの制御が必要なアクセラレータとは異なるアーキテクチャが Intel CPU と同じであるため, コンパイルし直すだけで動作する新モデルを投入予定 * 72コア, メモリ16GB 理論演算性能 3.0 TFLOPS * 47

44 まとめ GPU の特徴低性能の演算器を大量に搭載 (~3000 コア ) GPU が使われる理由理論演算性能が高いメモリとチップ間の帯域も広い省電力と高性能を両立今後の計算機の主流になると考えられる将来に対する投資 GPU だけでなく, 制御用 CPU+ 計算用アクセラレータという思想は今後しばらく主流であり続ける 48

45 ロボットによる心臓外科手術ロボットアームと内視鏡を使った心臓外科手術人工心臓を使わず, 心臓を動かしたまま行う 49

46 ロボットによる心臓外科手術 50

47 ロボットによる心臓外科手術心臓は複雑な形状で, かつ周期的に脈動毎秒何十枚と撮られる映像をリアルタイムで処理しながらロボットアームを制御心臓を 2 次元の画像に変換その画像を基にロボットが動く道筋を計算実際に 3 次元の動きに変換心臓の動きに合わせてロボットアームを制御 51

48 住宅設備機器開発住宅設備機器開発のための混相流シミュレーション池端昭夫, 吉田慎也, 住宅設備機器開発のための混相流シミュレーション, 第 27 回数値流体力学シンポジウム講演予稿集, E05-4, 2013(on USB Flash Drive). の結果を著者が撮影 52

49 車載システム Google, NVIDIA, Audi などによる Open Automotive Alliance 自動車への Android プラットフォーム搭載促進を目指す 53

50 NVIDIA 社の自動車関連ソリューション NVIDIA DRIVE PX 高機能運転支援システム NVIDIA DRIVE CX 自動車の高度情報化 54

51 Jetson Tegra TK1 世界初の組み込みスーパーコンピューター NVIDIA Tegra K1を採用 Keplarアーキテクチャ 192 CUDA コア NVIDIA 4 Plus 1 クアッドコア ARM Cortex A15 CPU Linux for Tegra が動作 55

52 月面探査 Google Lunar XPRIZE Google による国際宇宙開発レース日本の民間月面探査チーム HAKUTO が参加ロケット発射時のシミュレーション着陸などのシミュレーションや実際の制御に GPU を用いた自動運転技術を利用 56

Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx

Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx GPGPU の歴史と応用例長岡技術科学大学電気電子情報工学専攻出川智啓今回の内容 GPU の進化の歴史 GPU のアーキテクチャ GPU の産業応用例 38 GPGPU 実践基礎工学 GPU(Graphics Processing Unit) とは画像処理専用のハードウェア具体的には画像処理用のチップチップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU

今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展 性能の変化 シングルコアからマルチコア GPU の応用例 6

今回の内容 GPU の発展 GPU のアーキテクチャ CPU の発展性能の変化シングルコアからマルチコア GPU の応用例 6