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

(1)

補足

MPIプログラムのコンパイル，実行



標準の環境ではmpic++やmpiexecを実行できない

 _{OSがmpic++やmpiexecの場所を把握していないことが原因} 

_{OSが実行ファイルなどを探す場所（PATH）を設定}

 _{PATH（パス）にmpic++やmpiexecがあるディレクトリを追加}  _{PATHは.bashrcに記述}  _{.bashrcはホームディレクトリ（grouseにログインしたときの} ディレクトリ）に置く ‐bash‐3.2$ mpic++ ‐bash: mpic++: command not found ‐bash‐3.2$ mpiexec ‐bash: mpiexec: command not found ‐bash‐3.2$

(2)

補足

MPIプログラムのコンパイル，実行



ホームディレクトリでls ‐a を実行して

 _{.bashrcが見つかる場合}  エディタで.bashrcを開いて編集  _{.bashrcが見つからない場合}  ホームディレクトリに.bashrcを新しく作成  _{.(ドット)から始まるファイルはlsで表示されない}  _{lsに‐aオプションを付けて実行すると全てのファイルを表示}

(3)

補足

MPIプログラムのコンパイル，実行



_{.bashrcに記述（追記）する内容}



_{.bashrcに記述した内容を有効化}

 _{sourceコマンドを利用}  _{source 設定ファイル名として実行} MPIROOT=/opt/mpi/openmpi/gcc PATH=$MPIROOT/bin:$PATH LD_LIBRARY_PATH=$MPIROOT/lib:$LD_LIBRARY_PATH MANPATH=$MPIROOT/share/man:$MANPATH export PATH export LD_LIBRARY_PATH export MANPATH

(4)

補足

MPIプログラムのコンパイル，実行



_{.bashrcをストリーミング配信サイトからダウンロードし}

て利用するには

1.ストリーミング配信サイト（GPGUP実践基礎工学第8回）にアクセス 2.bashrcをダウンロードしてホームディレクトリに置く  アップロードしているファイル名はbashrc 3.mvコマンドを使ってファイル名を変更  _{mv 変更前ファイル名変更後ファイル名}  _{mv bashrc .bashrc} 4.sourceコマンドで設定を有効化  _{source .bashrc}

(5)

長岡技術科学大学電気電子情報工学専攻出川智啓

(6)

今回の内容



_{GPUの進化の歴史}



_{GPUのアーキテクチャ}

(7)

GPU（Graphics Processing Unit）とは



画像処理専用のハードウェア

 具体的には画像処理用のチップ  チップ単体では販売されていない  PCI‐Exカードで販売（チップ単体と区別せずにGPUと呼ぶことも多い）  マザーボードやノートPCに搭載  _{PCI‐Exカードとして販売されるGPUには，ビデオメモリと呼ばれ} るRAMが搭載

(8)

GPU（Graphics Processing Unit）とは



代表的な製品

 _{NVIDIA GeForce}  _AMD Radeon  _{Intel HD Graphics（内蔵）} 

代表的な用途

 _{3Dグラフィックス処理}  _{3Dゲーム，3DCAD，3DCG作成}  エンコード・デコード支援  GPU上に専用チップを搭載していることが多い  デスクトップPCのGUI処理  _{Windows Aeroが比較的高性能なGPUを要求}

(9)

GPU（Graphics Processing Unit）の役割



グラフィックスを表示するために様々な処理を行い，処

理の結果をディスプレイに出力



_{3次元グラフィックスの発展に伴って役割が大きく変化}

3次元座標変換ポリゴンとピクセルの対応付けピクセル色計算テクスチャ参照フレームバッファ（ビデオメモリ）への書き込みディスプレイ出力 CPU ビデオカード _GPU 3次元座標変換ポリゴンとピクセルの対応付けピクセル色計算テクスチャ参照フレームバッファ（ビデオメモリ）への書き込みディスプレイ出力現在過去  _{CPU が 3D 描画} の演算を実行  _GPUが出力描画情報画面出力  _{GPUが演算から} 出力までの全てを担当  _{CPUは描画情報} の生成やGPUへの情報の引き渡し， GPU の制御を行う描画情報画面出力

(10)

GPUの描画の流れ

1. CPUからGPUへ描画情報を送信

2. 頂点処理（頂点シェーダ）

 座標変換  画面上での頂点やポリゴンの位置・大きさの決定  頂点単位での照明の計算

3. 頂点やポリゴンからピクセルを生成

（ラスタライザ）

4. ピクセル処理（ピクセルシェーダ）

 画面上のピクセルの色  テクスチャの模様

5. 画面出力

 ピクセルの色情報をフレームバッファに書き込み 2. 3. 4.

(11)

ビデオカードの利点



_{CPUで描画のための演算を行うと，CPUにかかる負荷が}

大きい



_{3次元画像処理の専用回路を備えたハードウェアを導入}

○CPUにかかる負荷を減らすことができる ○頂点・ピクセルごとに並列処理が可能なため，ハードウェアによる並列処理が可能

(12)

ビデオカードの欠点



_{3次元画像処理の専用回路を備えたハードウェアを導入}

×新しい描画方法を開発しても，GPUへ実装・製品化されるまで利用できない ×ユーザが所有しているGPUによって利用できる機能にばらつきが生じる ×ある描画手法用の専用回路を実装しても，その描画方法が常に使われる訳ではないので_{GPU全体の利用効率が下がる}

(13)

ビデオカードから

GPUへ



_{CGの多様化と共に固定機能の実装が困難に}



頂点処理とピクセル処理をユーザが書き換えられるプロ

グラマブルシェーダの実装

頂点処理用回路ピクセル処理用回路ビデオカード頂点シェーダユニットピクセルシェーダユニット GPU

(14)

レンダリングパイプライン処理

頂点情報光源情報視野変換陰影計算材質情報投影変換クリッピングビューポート変換走査変換合成テクスチャ出力画像投影像を画素へ変換整数演算とメモリアクセス形状データの画面への投影像実数演算

(15)

レンダリングパイプライン処理

頂点情報光源情報視野変換陰影計算材質情報投影変換クリッピングビューポート変換走査変換合成テクスチャ出力画像形状データの画面への投影像ハードウェアで処理（固定機能）実数演算実数演算を行うハードウェアは高価だった

(16)

レンダリングパイプライン処理

頂点情報光源情報視野変換陰影計算材質情報投影変換クリッピングビューポート変換走査変換合成テクスチャ出力画像ハードウェアで処理（固定機能）ハードウェアで処理（固定機能）

(17)

レンダリングパイプライン処理

頂点情報光源情報視野変換陰影計算材質情報投影変換クリッピングビューポート変換走査変換合成テクスチャ出力画像ピクセルシェーダ頂点シェーダ

(18)

ビデオカードから

GPUへ



描画する画像によって頂点処理とピクセル処理の負荷

が変化

 処理によっては利用効率に差が発生し，利用効率が低下 GPU 頂点シェーダユニットピクセルシェーダユニット頂点処理重視の処理 GPU 頂点シェーダユニットピクセルシェーダユニットピクセル処理重視の処理空きユニット空きユニット

(19)

ビデオカードから

GPUへ



頂点シェーダとピクセルシェーダを統合したユニファイド

シェーダへの進化

 頂点処理とピクセル処理を切り替えることで利用率を向上 GPU ユニファイドシェーダユニット頂点処理重視の処理ピクセル処理重視の処理 GPU ユニファイドシェーダユニット

(20)

ビデオカードから

GPUへ



各ピクセルに対して並列に処理実行できるように進化

 単純な処理を行う演算器を大量に搭載  高い並列度で処理を実行 

_{GPUの誕生とGPGPUの普及}

 高性能な3DCG画像処理への要求→GPUの高性能化 

_GPUの長所

 消費電力あたりの浮動小数点理論演算性能が高い  _{GPU単体の消費電力は高い}  （相対的に）安価  CPUだけで同等の計算能力を達成するより安価

(21)

Teslaアーキテクチャ

*

の構造



_{Tesla C1060の仕様}

 _SM数 ₃₀  _{CUDA Core数} _{240(=8 Core/SM×30 SM)}  キャッシュを搭載せず *CUDAのサポートから外れます

(22)

Teslaアーキテクチャの構造



_{Tesla C1060の仕様}

CUDAコア数（単精度） 240 Cores CUDAコアクロック周波数 1,296 MHz 単精度演算ピーク性能 _622*1 _(933*2_) GFLOPS 倍精度演算ユニット数 _30*3 _Units 倍精度演算ピーク性能 _78 GFLOPS メモリクロック周波数 _800 MHz メモリバス幅 _512 bit 最大メモリバンド幅*4 _102 GB/s *1_{単精度演算ピーク性能 = コアクロック周波数×コア数×命令の同時発行数（2）} *2_{CUDA CoreとSFUが同時に命令を発行できれば1296 MHz×240×}₃ *3_{一つのSMに倍精度演算器が一つ搭載} *4_{最大メモリバンド幅=メモリクロック周波数×メモリバス幅/8×2(Double Data Rate)}

(23)

Fermiアーキテクチャの構造



_{Tesla M2050の仕様}

 _SM数 ₁₄  _{CUDA Core数} _{448(=32 Core/SM×14 SM)}  _{L1/L2 キャッシュを搭載}  _{ECC（誤り訂正機能）を搭載}

(24)

Fermiアーキテクチャの構造



_{Tesla M2050の仕様}

CUDAコア数（単精度） 448 Cores CUDAコアクロック周波数 1,150 MHz 単精度演算ピーク性能 _1.03 TFLOPS 倍精度演算ユニット数 _0*1 _Unit 倍精度演算ピーク性能 _515 GFLOPS メモリクロック周波数 _1.55 GHz メモリバス幅 _384 bit 最大メモリバンド幅 _148 GB/s *1_{単精度CUDA Coreを2基使って倍精度演算を実行}

(25)

Keplerアーキテクチャの構造



_{Tesla K20c/mの仕様}

 _SMX数 ₁₃

 _{Streaming Multiprocessor eXtreme (?)}

(26)

Keplerアーキテクチャの構造



_{Tesla K20c/mの仕様}

CUDAコア数（単精度） 2,496 Cores CUDAコアクロック周波数 706 MHz 単精度演算ピーク性能 _3.52 TFLOPS 倍精度演算ユニット数 _832*1 _Units 倍精度演算ピーク性能 _1.17 TFLOPS メモリクロック周波数 _2.6 GHz メモリバス幅 _320 bit 最大メモリバンド幅 _208 GB/s *1_{64基/SMX×13基}

(27)

Maxwellアーキテクチャ



_{GeForce GTX TITAN Xの仕様}

 _SM数 ₂₄

(28)

Maxwellアーキテクチャ



_{GeForce GTX TITAN Xの仕様}

* CUDAコア数（単精度） 3,072 Cores CUDAコアクロック周波数 1,002 MHz 単精度演算ピーク性能 _{6.14 TFLOPS} 倍精度演算ユニット数 _0*1 _Units 倍精度演算ピーク性能 _192 GFLOPS*2 メモリクロック周波数 _3.5 GHz*3 メモリバス幅 _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相当と書かれている場合もある} http://http://www.geforce.com/hardware/desk top‐gpus/geforce‐gtx‐titan‐x/specifications *http://ja.wikipedia.org/wiki/FLOPS

(29)

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

(30)

Voltaアーキテクチャ



_{Pascalの後継}



詳しい情報は不明



アメリカの次世代スーパーコンピュータへ採用予定

 オークリッジ国立研究所 _SUMMIT _{150~300PFLOPS}  ローレンス・リバモア研究所 SIERRA _{100PFLOPS以上}  地球シミュレータと同等の演算性能を1ノードで実現  現在Top500 2位のスーパーコンピュータと同じ電力で5~10 倍高速，サイズは1/5

(31)

GPUの普及の要因



_{GPUの進展は１５年程}

 普及の速度は驚異的 

_{CPUは数十年かけて進展}

 _{CPUも驚異的な速度で進展}  様々な高速化技術を導入 

_{GPUが普及している要因は何か？}

(32)

TOP500 List（2015, Jun.）



スーパーコンピュータの性能の世界ランキング

 _{GPUを搭載したコンピュータは2基だけ} http://www.top500.org/より引用計算機名称（設置国）アクセラレータ 実効性能[PFlop/s] /ピーク性能[PFlop/s] 消費電力[MW] 1 Tianhe‐2 (China) Intel Xeon Phi 33.9/54.9 17.8 2 Titan (U.S.A.) NVIDIA K20x 17.6/27.1 8.20 3 Sequoia (U.S.A.) − 17.2/20.1 7.90 4 K computer (Japan) − 10.5/11.3 12.7 5 Mira (U.S.A.) − 8.59/10.1 3.95 6 Piz Daint (Switzerland) NVIDIA K20x 6.27/7.79 2.33

7 Shaheen II(Saudi Arabia) 5.54/7.24 2.83

8 Stampede (U.S.A.) Intel Xeon Phi 5.17/8.52 4.51 9 JUQUEEN (Germany) − 5.01/5.87 2.30 10 Vulcan (U.S.A.) − 4.29/5.03 1.97

(33)

CPUの性能向上サイクル

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

(34)

CPUの性能向上サイクル

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

(35)

CPUの性能向上サイクル

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

(36)

CPUの性能向上



_FLOPS =

_{1コアの演算性能}

× コア数

× CPUの動作周波数



_{1コアの演算性能の向上}

 演算器（トランジスタ）の増加 

コア数の増加

 トランジスタ数の増加 

_{CPUの動作周波数}

 回路の効率化や印可電圧の向上劇的な性能向上は期待できないコンパイラの最適化を利用複数のコアを使うようにプログラムを書かないと速くならない

(37)

GPUを使うという選択



_{GPU普及の要因の一つはCPUクロックの頭打ち}

 クロックを下げてマルチコア化したCPUへの対応が必要 

なぜGPUという選択か？

 _{CPU用プログラムの並列化でもいいのでは？} 

消費電力の低減

 数値計算や高性能計算（HPC）の業界がGPUに注目  スーパーコンピュータの性能向上  高機能なCPUを大量に使うと消費電力が問題に  高機能な制御用プロセッサと，計算を実行する低性能なアクセラレータの組み合わせ

(38)

Green500（2015, Jun.）



日本の次世代機がTOP3を独占



_{AMD社のGPUが4位}



_{NVIDIA社のGPUが5位以降を占める}

計算機名称アクセラレータ GFLOPS/W 消費電力[kW] 1 Shoubu PEZY‐SC 7.03 50.32

2 Suiren Blue PEZY‐SC 6.84 28.25 3 Suiren PEZY‐SC 6.22 32.59 4 ‐ AMD FirePro S9150 5.27 57.15 5 TSUBAME‐KFC NVIDIA K20x 4.25 39.83 6 XStream NVIDIA K80 4.11 190.0 7 Storm1 NVIDIA K40m 3.96 44.54 8 Wilkes NVIDIA K20 3.63 52.62 9 Taurus NVIDIA K80 3.61 58.01 10 iDataPlex NVIDIA K20x 3.54 54.60

(39)

まとめ



_GPUの特徴

 低性能の演算器を大量に搭載（～3000コア） 

_{GPUが使われる理由}

 理論演算性能が高い  メモリとチップ間の帯域も広い  省電力と高性能を両立  今後の計算機の主流になると考えられる 

将来に対する投資

 _{GPUだけでなく，制御用CPU+計算用アクセラレータという思想} は今後しばらく主流であり続ける

(40)

(41)

情報処理センター

GPGPUシステム



_{16台の計算サーバで構成}

 _{NVIDIA Tesla M2050を搭載}  各サーバに4機ずつ計64機 _grouse tesla01 ‐ tesla04 M2050×16 tesla05 ‐ tesla08 M2050×16 tesla09 – tesla12 M2050×16 tesla13 – tesla16 M2050×16 外部ネットワーク

(42)

(43)

grouseへのログイン

(44)

ターミナルの起動

(45)

(46)

実行イメージ

grouse 処理処理結果出力

(47)

(48)

実行イメージ

grouse ファイル保存ハードディスク共有共有

tesla01 ‐ tesla04 tesla05 ‐ tesla08 tesla09 – tesla12 tesla13 – tesla16 共有

(49)

tesla??へのログイン

もう一つターミナルを起動し，下のターミナルでログイン

$ ssh␣–l␣ユーザID␣tesla?? (??には01~16の番号を入力)

(50)

tesla??でのコンパイルとプログラム実行

grouseでの作業用

tesla??でのコンパイルと実行用

(51)

tesla??でコンパイルとプログラム実行



コンパイル



_{nvcc˽ソースファイル名}

 _{nvccがコンパイル対象とするソースファイルの拡張子は}_.cu  エラーがなければa.outという実行ファイルが作成される 

実行



_./a.out



_{nvccはgccと同じオプションを持っていることが多い}



対応オプションの一覧は‐hオプションで確認

 _{nvcc ‐h}

(52)

実行イメージ

grouse 処理処理結果出力キー入力，マウス入力 $ ssh␣tesla16 tesla16 キー入力処理処理結果

(53)

実行イメージ

grouse キー入力，マウス入力 $ ssh␣tesla16 $ nvcc ??.cu キー入力処理結果ハードディスク ??.cu コンパイル処理結果出力 tesla16

(54)

実行イメージ

grouse 処理結果出力キー入力，マウス入力 $ ssh␣tesla16 $ nvcc ??.cu $ ./a.out キー入力処理結果ハードディスク a.out 実行 tesla16

(55)

実行イメージ

grouse キー入力，マウス入力 $ ssh␣tesla16 $ nvcc ??.cu $ ./a.out キー入力ハードディスク a.out 実行 tesla16 GPU0 GPU1 GPU2 GPU3 処理結果出力処理結果

(56)

情報処理センターでの

CUDAの使い方

1.

grouseやtesla??で開発する場合

 _{grouseやtesla??にログイン}  ソースファイルを作成し，tesla??上でコンパイル・実行  grouseはGPUを搭載していないため実行できない（コンパイルは可能） 2.

研究室のPC等，情報処理センター外で開発する場合

 研究室のPCでソースファイルを作成  _{WinSCPなどでファイルをgrouseにアップロード}  _{grouseを経由してtesla??にログインした後，コンパイル・実行}

(57)

ターミナルのみの利用

（特に演習室外から利用する場合）

(58)

ターミナルのみの利用

（特に演習室外から利用する場合）

2. sshでtesla??にログイン

$ ssh␣–l␣ユーザID␣tesla?? (??には01~16の番号を入力)

(59)

grouseへのファイル転送（WinSCP）



_{Unixコマンドscp（secure copy）のWindows GUIクラ}

イアント



_{Secure Shell (ssh) に含まれるsshの機能を利用し}

て安全性の高いファイル転送を行う

Host名 grouse

統合アカウントの

ユーザ名とパスワード

login

(60)

Windows上のソースファイルをコピー

grouse ハードディスク WinSCPでコピーハードディスク共有共有ソースファイル

tesla01 ‐ tesla04 tesla05 ‐ tesla08 tesla09 – tesla12 tesla13 – tesla16 共有

(61)

ログインしているサーバの確認



_{grouseにログインしているかtesla??にログインしてい}

るか分からなくなったら

(62)

GPUの選択



計算機がGPUを複数搭載している場合

 _{CUDAで利用するGPUを選択} 

_{CUDA APIを利用したGPUの選択}



_{cudaSetDevice()命令}

int main(void){ 変数宣言 _{//ここでは標準でGPU0が使われる} GPUやCPUを使った処理 // cudaSetDevice(3); //ここからGPU3が使われる ... return 0; }

(63)

レポート課題

1（提出期限は2学期末）



_{grouseにログインして，以降のスライドに示す3個のプロ}

グラムを作成，コンパイル，実行せよ

 _{helloworld.cu}  _{hellokernel.cu}  _{hellothread.cu} 

各プログラムで指定された事項について考察せよ

(64)

レポート課題

1（提出期限は2学期末）



通常のC言語の

プログラム

との違いは何か

 違いがあるとすればどの箇所か  違いがないとすればその理由はなぜか #include<stdio.h> int main(void){ printf("hello world¥n"); return 0; } helloworld.cu

(65)

レポート課題

1（提出期限は2学期末）



通常のC言語の

関数

との違いは何か

 違いがあるとすればどの箇所か（どのように違うか） #include<stdio.h> __global__ void kernel(){} int main(void){ kernel<<<1,1>>>(); printf("hello world¥n"); return 0; } _{hellokernel.cu}

(66)

レポート課題

1（提出期限は2学期末）



三重括弧（<<<1,1>>>）内の数字を変えると実行結果は

どのように変化するか



_{cudaDeviceSynchronize();をコメントアウトすると実}

行結果はどのように変化するか，その理由はなぜか

#include<stdio.h> __global__ void hello(){ printf(“Hello Thread¥n”);//printfを実行するには，コンパイルの際に } //‐arch=sm_20オプションが必要 int main(void){ hello<<<1,1>>>(); cudaDeviceSynchronize(); return 0; } _{hellothread.cu}

(67)

レポートの書式



必ず表紙を付けること

 授業名，課題番号，学籍番号，氏名，提出日に加えて課題に要した時間を書く 

課題内容，プログラム，実行結果（ターミナルの出力結果），

考察で構成

 プログラムを実行したtesla??およびGPUの番号も明記すること 

_{pdf形式に変換してメールで提出}

 宛先 _{degawa at vos.nagaokaut.ac.jp}  メール題目 _{GPGPU実践基礎工学課題１（氏名）}