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

GPGPUの歴史と応用例

今回の内容



_GPUの発展



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



_CPUの発展

 性能の変化  シングルコアからマルチコア 

_{GPUの応用例}

GPU（Graphics Processing Unit）とは



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

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

GPU（Graphics Processing Unit）とは



代表的な製品

 _{NVIDIA GeForce, Quadro, Tesla}  _{AMD Radeon, FireProm} 

代表的な用途

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

GPUの性能の遷移（理論演算性能）

Tesla Fermi Kepler Maxwell Kepler GeForce ゲーム用 Quadro CG用 Tesla GPGPU用 Theoretical   GFLOP/s GeForce FX 5800GeForce 6800 Ultra GeForce 7800 GTX GeForce 8800 GTX GeForce GTX 280 GeForce GTX 480 GeForce GTX 580 GeForce GTX 680 GeForce GTX TITAN GeForce 780 Ti Pentium 4

WoodcrestHarpertown Sandy Bridge

Ivy Bridge Tesla K40 Tesla K20X Tesla M2090 Tesla C2050 Tesla C1060

Apr‐01 Sep‐02 Jan‐04 May‐05 Oct‐06 Feb‐08 Jul‐09 Nov‐10 Apr‐12 Aug‐13 Dec‐14

Westmere Bloomfield

GPUの性能の遷移（理論バンド幅）

Tesla Fermi Maxwell Kepler GeForce ゲーム用 Quadro CG用 Tesla GPGPU用 GeForce FX 5900 GeForce 6800 GT GeForce 7800 GTX GeForce 8800 GTX GeForce GTX 280 GeForce GTX 480 GeForce GTX 680 GeForce 780 Ti Tesla K40 Tesla K20X Tesla M2090 Tesla C2050 Tesla C1060 Northwood Woodcrest Harpertown Sandy Bridge Ivy Bridge Westmere Bloomfield Prescott Theoretical   GB/s

GPU（Graphics Processing Unit）の役割



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

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



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

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

GPUの描画の流れ

1.

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

2.

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

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

3.

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

（ラスタライザ）

4.

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

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

5.

画面出力

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

ビデオカードの利点



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

大きい



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

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

ビデオカードの欠点



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

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

ビデオカードから

GPUへ



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



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

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

頂点処理用回路 ピクセル処理用回路 グラフィックスカード 頂点シェーダユニット ピクセルシェーダユニット GPU

ビデオカードから

GPUへ



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

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

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

ビデオカードから

GPUへ



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

シェーダへの進化

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

ビデオカードから

GPUへ



各ピクセルに対して並列に処理を行えるよう，並列度を

高める

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

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

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

_{GPGPUの長所}

 消費電力あたりの浮動小数点理論演算性能が高い  安価  _{CPUを使った大規模計算機と比較して}相対的に

GPUの進化



_{GPUは発展途上（２～３年で世代交代）}

 今プログラムを作っておくと，勝手に速くなってくれる！* NVIDIA社プレゼンテーションを基に作成 情報処理センター GPGPUシステムに搭載 1   Watt あたりの単精度行列－行 列積の回数 年 0 12 24 36 48 60 72 2008 2010 2012 2014 2016 2018 Tesla Fermi Maxwell Kepler Pascal Volta *多分に誇張的な表現であることに注意

GPUのハードウェア構造



_{CUDA Core（旧Streaming Processor, SP）と呼ばれ}

る演算器を多数搭載



_{Streaming Multiprocessor（SM, SMX）が複数の}

CUDA CoreとSFU，メモリをまとめて管理

 _{SFU(Special Function Unit)}  数学関数を計算するユニット 

複数のSMが集まってGPUを構成

Fermiアーキテクチャ



_{Tesla M2050の仕様}

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

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基使って倍精度演算を実行} *http://www.nvidia.co.jp/object/product_tesla_M2050  _M2070_jp.html http://ja.wikipedia.org/wiki/NVIDIA_Tesla

Maxwellアーキテクチャ



_{GeForce GTX TITAN Xの仕様}

 _{SM数 24}

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

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

CUDA



_{Compute Unified Device Architecture}

 _{NVIDIA社製GPU向け開発環境(Windows,Linux,Mac OS X)}  _{2007年頃発表}  _{C/C++言語＋独自のGPU向け拡張}  専用コンパイラ（nvcc）とランタイムライブラリ  いくつかの数値計算ライブラリ（線形代数計算，FFTなど） 

_{CUDA登場以前}

 グラフィックスプログラミングを利用  足し算を行うために，色を混ぜる処理を実行  汎用計算のためには多大な労力が必要

プログラマブルシェーダを用いた汎用計算



グラフィックス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; } 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) シェーダ言語を用いた配列加算 （c=*a + *b）の例

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; }

GPUの普及の要因



_{GPUの進展は15年程度}

 普及の速度は驚異的 

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

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

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

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

理論ピーク性能と実効性能



_{Floating‐Point Operations per Second}

 _{1秒あたりに浮動小数の演算が何回できるか} 

理論ピーク性能

 プロセッサの数（プロセッサ上に実装された演算器の数）や動 作周波数から求める理論的な速度  「全ての機能が全て同時に使えれば」という理論的な値 

実効性能（実行性能）

 ある問題に対してプログラムを実行したときに得られた性能  プログラムの中で行っている計算（浮動小数点演算）の回数 を数え，プログラムの実行にかかった時間を測定して割り算

CPUの理論性能



公式

 _{FLOPS =  1コアの演算性能 [？]} × コア数 _[Core] × CPUの動作周波数 [Hz=Clock/sec] 

_{1コアの演算性能}

 ＝1度に発行出来る浮動小数点演算命令  単位は[Floating Point Operations/Clock /Core]  性能の評価には動作周波数だけでなく，1コアが1クロックで 発行できる命令数が重要

CPUの性能



_FLOPS = 

_{1コアの演算性能}

× コア数

× CPUの動作周波数



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

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

コア数の増加

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

_{CPUの動作周波数}

 回路の効率化や印可電圧の向上 動作周波数の向上に注力 （ほぼ全ての処理が速くなる） 様々な機能を追加 • パイプライン処理 • スーパースカラ実行 • 分岐予測等

CPUの性能の変化



_{Intelの予告(Intel Developer Forum 2003)}

CPUの性能の変化



_{2004年頃からクロックが停滞}

Intelが公開している資料を基に作成

CPUの性能向上

* 

電子回路の構成部品

 機械式リレー  真空管  トランジスタ  _{IC (Integrated Circuit)}  _{LSI (Large Scale Integrated Circuit)} 

_{CPUとメモリの性能向上}

 回路の配線を細線化  _{250nm→180nm→130nm→90nm→65nm→45nm→32nm→} 

線幅が減ると同じ回路を作る際の面積が減少

同じ面積に集積できるトランジスタ数が増加

集積率が上昇 *_{姫野龍太郎，絵でわかるスーパーコンピュータ，講談社 (2012)}

ムーアの法則

* 

インテルの共同設立者ムーアによる経験則

 半導体の集積率は1年で倍になる  後に「18ヶ月で2倍」に修正 姫野龍太郎，絵でわかるスーパーコンピュータ， 講談社 (2012)に掲載されている絵を基に作成 http://en.wikipedia.org/wiki/Moore%27s_law *Moore, G.E., Electronics, Vol.38,No.8(1965). http://ja.wikipedia.org/wiki/ムーアの法則 year 1970 1975 1980 1985 1990 1995 2000 2005 2010 103 104 105 106 107 108 109 1010 Number of T ransistors 4004 80088080 8086 286 Intel386プロセッサ Intel486プロセッサ インテルPentiumプロセッサ インテルPentium IIプロセッサ インテルPentium IIIプロセッサ インテルPentium 4プロセッサ インテルItaniumプロセッサ インテルItanium 2プロセッサ デュアルコアインテルItanium 2プロセッサ ムーアの法則 （12ヶ月で倍） ムーアの法則 （18ヶ月で倍）

ポラックの法則

* 

_{2倍のトランジスタを使っても，プロセッサの性能はその}

平方根倍（1.4倍）程度にしか伸びない

 消費電力は2倍，性能は1.4倍 

一つのCPUに複数のプロセッサ（コア）を搭載

 消費電力を上げずに“理論的な”性能を倍に  プログラムの作り方に工夫が必要 *http://ja.wikipedia.org/wiki/ポラックの法則 http://en.wikipedia.org/wiki/Pollack%27s_Rule

プロセッサの性能向上

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

プロセッサの性能向上

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

プロセッサの性能向上

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

CPUの性能向上



_FLOPS = 

_{1コアの演算性能}

× コア数

× CPUの動作周波数



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

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

コア数の増加

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

_{CPUの動作周波数}

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

GPUを使うという選択



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

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

なぜGPUという選択か？

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

消費電力の低減

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

計算機名称 アクセラレータ 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

Green500（2014, Nov.）



_{AMD社のGPUが1位}



日本の次世代機「睡蓮」が2位*



_{NVIDIA社のGPUが3位以降にランクイン}

*4月2日に1位相当の値を記録 http://news.mynavi.jp/articles/ 2015/04/02/kek_suiren/

アクセラレータ



コンピュータの特定の機能や処理能力を向上させる

ハードウェア

 _{CPUで行っていた処理を専用ハードウェアが担当}  動画像のエンコード・デコード等 

コンピュータシミュレーションではCPUの代わりに計算を

実行するハードウェアを指す

 画像処理装置(Graphics Processing Unit)  メニーコアプロセッサ

アクセラレータ（メニーコアプロセッサ）



_PEZY‐SC

 株式会社PEZY Computingの1,024コアの低消費電力型メ ニーコアプロセッサ  _{1024コア，動作周波数733MHz}  理論演算性能  単精度 3.0 TFLOPS  倍精度 1.5 TFLOPS

アクセラレータ（メニーコアプロセッサ）



_{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/

まとめ



_GPUの特徴

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

_{GPUが使われる理由}

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

将来に対する投資

 _{GPUだけでなく，制御用CPU+計算用アクセラレータという思想}

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



ロボットアームと内視鏡を使った心臓外科手術

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

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



心臓は複雑な形状で，かつ周期的に脈動



毎秒何十枚と撮られる映像をリアルタイムで処理しなが

らロボットアームを制御

 心臓を2次元の画像に変換  その画像を基にロボットが動く道筋を計算  実際に3次元の動きに変換 

心臓の動きに合わせてロボットアームを制御

住宅設備機器開発

車載システム



_{Google, NVIDIA, AudiなどによるOpen Automotive} 

Alliance

NVIDIA社の自動車関連ソリューション



_{NVIDIA DRIVE PX}

 高機能運転支援システム



_{NVIDIA DRIVE CX}

Jetson Tegra TK1



世界初の組み込みスーパーコンピューター

 _{NVIDIA Tegra K1を採用}  _{Keplarアーキテクチャ}  192 CUDA コア  NVIDIA 4‐Plus‐1™ クアッドコア ARM® Cortex‐A15 CPU  _{Linux for Tegraが動作}

月面探査



_{Google Lunar XPRIZE}

 _{Googleによる国際宇宙開発レース}  日本の民間月面探査チーム「HAKUTO」が参加 

ロケット発射時のシミュレーション



着陸などのシミュレーションや実際の制御にGPUを用い

た自動運転技術を利用