2010/7/26
高性能コンピューティング学論
講義資料
新しい並列処理
「GPGPU」について
大島聡史 (東京大学 情報基盤センター)
1
2010/7/26
高性能コンピューティング学論
講義資料
2
目次
• GPGPUとは何か
• GPGPUの概要、はじまりと発展について
• CUDAについて
• GPGPUの普及に大きな影響を与えたCUDA、そのアーキテクチャ
とプログラミング環境について
• 今後のGPUとGPGPU
• GPGPUの将来はどうなるのか、最新のGPU事情
高性能コンピューティング学論
2010/7/26
講義資料
3
自己紹介
• 大島聡史
• 経歴
• 2000年4月～2009年9月 電気通信大学
• 2000年
• 2003年
• 2004年
• 2006年
• 2009年
入学 (情報工学科)
卒研時配属で弓場・本多研究室へ
博士前期課程進学 (並列処理学講座 弓場研究室)
博士後期課程進学 (並列処理学講座 本多研究室)
博士研究員 (高性能コンピューティング学講座 本多研究室)
• 2009年10月～
東京大学 情報基盤センター 助教
2010/7/26
高性能コンピューティング学論
講義資料
GPGPUとは何か
GPGPUの概要、はじまりと発展について
4
2010/7/26
高性能コンピューティング学論
講義資料
5
「GPGPU」？
• General-Purpose computation using GPUs
• 「GPUを用いた汎用演算」
• GPU
• Graphics Processing Unit
• いわゆるビデオカードや統合グラフィックスチップ（上のLSI）
• 代表的な製品：NVIDIA GeForce, ATI (AMD) Radeon, Intel GMA
• 用途＝画面出力全般
• 3Dグラフィックス処理
• 3Dゲーム、3DCAD、3DCG作成
• エンコード・デコード支援
• GPU上に専用チップを搭載していることが多い
• デスクトップ処理
• Windows Aeroが比較的高性能なGPUを
要求したため普及
2010/7/26
高性能コンピューティング学論
講義資料
6
GPUの性能向上と並列処理
• GPUに対する要求＝高速・高解像度・複雑な画像描画
• 2000年以降、特に写実的な3DCGへの要求が高まる
• もちろん、家庭用ゲーム機にもGPUは積まれています
• GPUの高性能化におけるキーワード
• 並列化
• プログラマブル化
高性能コンピューティング学論
2010/7/26
講義資料
7
一般的な3D描画の手順
1. 頂点情報の設定
2. 頂点処理：頂点が空間上のどこに対応するか（カメラ
計算）
3. ラスタライズ：画面上のピクセルとの対応付け
4. ピクセル処理：ピクセル色の決定、テクスチャ適用
① (2, 2)
② (8, 3)
③ (5, 7)
頂点ごと・ピクセルごとに独立計算（並列計算）可能なため、
ハードウェアの並列化による性能向上が行われた
2010/7/26
高性能コンピューティング学論
講義資料
8
GPUのプログラマブル化
• 高度な描画処理手法の開発における問題
• 専用ハードウェアを用いる高度な手法を開発
• ハードウェアに実装されるまで時間が必要
• 対応ハードウェアを所有していないと使えない
• 次々に新たな手法が提案される
ハードウェアと描画手法の開発サイクルがうまく回らない
• プログラマブルシェーダの登場
• 頂点処理やピクセル処理をソフトウェア処理で設定可能にした
• 新しい描画手法をすぐに利用可能になった
• シェーダのバージョン（世代）による実行可能不可能は存在
• ハイエンドGPU＝多数の高速なシェーダユニット、高速大容量メモリ
2010/7/26
高性能コンピューティング学論
講義資料
9
描画処理以外にも使われ始めたGPU
• 初期のプログラマブルシェーダは限定的な処理のみ可能
• GeForce6800(2004年)に搭載されたシェーダ(3.0)で大き
く改善、様々なプログラムが記述可能に
• 複雑な照明処理や陰影処理が実装可能になり、リアルタイム3Dグ
ラフィックス（主に3Dゲーム）の表現力が爆発的に向上
• 描画（照明・陰影）処理に物理計算（数値計算）が導入される
• 単精度浮動小数点演算への対応
高性能コンピューティング学論
2010/7/26
講義資料
10
GPUを用いた汎用演算(GPGPU)
• GPUによる処理を数値計算等に利用する研究が始まる
• GPGPUにおける入出力と演算の位置づけ
CPUからGPUへの転送（テクスチャへのデータセット）
＝入力
CPU
GPU
GPUからCPUへの書き戻し（描画結果の取得）
＝出力
描画
＝演算
高性能コンピューティング学論
2010/7/26
講義資料
11
描画処理を用いた汎用計算の例
• 配列加算
tex1
a
b
c
d
tex2
w x
y
z
a*w
b*x
c*y
• 行列積
• 乗算描画結果の加算
tex2
d*z
方法１：乗算ブレンド描画
方法２：tex1を描画し、
tex2を乗算描画
w1
x1
y1
z1
w2
x2
y2
z2
w3
x3
y2
z3
w4
x4
y4
z4
乗算描画
tex1
a1
b1
c1
d1
a1*w1
a1*x1
a1*y1
a1*z1
a2
b2
c2
d2
a2*w1
a2*x1
a2*y1
a2*z1
a3
b3
c3
d3
…
a4
b4
c4
d4
2010/7/26
高性能コンピューティング学論
講義資料
12
プログラマブルシェーダを用いた汎用計
算の例
• グラフィックス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)
各ピクセルに対し
て実行される
シェーダ言語を用いた配列加算
（vc=a*va + b*vb）の例
CPUの処理
(OpenGL)
2010/7/26
高性能コンピューティング学論
講義資料
13
GPGPUのはじまり：まとめ
• 高性能な描画処理への要求
→GPUの高性能化（高並列、一定のプログラマビリティ）
→高い並列浮動小数点演算性能の汎用計算への利用可能性
→数値計算や可視化などで成果
• GPGPUの特徴（メリット）
• 高い浮動小数点理論演算性能
• CPUが10GFLOPS程度の時代にGeForce6800は120GFLOPS
• 高並列（並列度の高い演算に有用）
• CPUがシングルコア～デュアルコアの時代に数十～百の並列度
• GeForce6800は6(vp)+16(pp)並列
• 安価
• CPUと変わらないレンジ（～10万円）、多くのPCに搭載
• GPGPUの課題
• 直接的でなく扱いにくいプログラミング記述
2010/7/26
高性能コンピューティング学論
講義資料
14
余談
• 私が知る限り、一番はじめの頃は「GPGP」「GP2」
「GPU computing」という呼称が多かった
• 「GPGPU」という名称が普及したが、最近また「GPU
computing」という言葉も使われている（「コンピュー
ティング」を強調したいらしい）
• そのうちまた呼称が変わるのだろうか？
• GPU＝General-Purpose GPU、などの記述を見ると違和
感を感じずにはいられない
2010/7/26
高性能コンピューティング学論
講義資料
15
CUDAについて
GPGPUの普及に大きな影響を与えたCUDA、そのアー
キテクチャとプログラミング環境について
高性能コンピューティング学論
2010/7/26
講義資料
16
黎明期GPGPUの課題
• ハードウェア
• 仕様がよくわからない
• 非公開（グラフィックスAPI＋シェーダ言語
レベルでしか見えない）
• 一応、共有メモリ型並列計算機のようなもの？
• テクスチャとフレームバッファを共有メモリとして、描画ユニットが演算
• プログラミング
• グラフィックスAPI＋シェーダ言語
• 画像処理プログラミングのスキルが必須
ライブラリによる隠蔽
にわかに脚光を浴びたストリーミング言語
高性能コンピューティング学論
2010/7/26
講義資料
mem
PE
PE
PE
PE
PE
PE
PE
PE
PE
PE
PE
PE
17
2010/7/26
高性能コンピューティング学論
講義資料
18
ライブラリによる隠蔽
• GPUプログラミングをライブラリで隠蔽し、既存のプロ
グラミング言語から容易に使えるようにする
• 良く用いられる処理を高速化
• 行列積やFFTなどの数値計算はCPU上でもライブラリがよく利用されて
いる
• GPUを利用してくれるライブラリを作ればGPUの性能を誰でも容易に
活用可能
• ちなみに、私もこのタイプの研究を行った（行っている）
• CPUとGPUを両方使って（片方の方が良い場合は片方だけ使って）計
算する行列計算ライブラリを作った
2010/7/26
高性能コンピューティング学論
講義資料
19
ストリーミング言語
• GPU向けの言語やクラスライブラリを提供
• 多数の入力データ（配列、ストリーム）に次々に演算操
作（カーネル）を施すプログラミングモデル
• 入力配列と出力配列、配列に対する演算を明確に記述する
• 元々GPU用というわけではないが、GPGPUの普及をうけて注目された
• BrookGPU、Brook+、RapidMind、SPRATなど
• 一部の言語はマルチコアCPUやCellにも対応
入力ストリーム
演算操作（カーネル）
出力ストリーム
2010/7/26
高性能コンピューティング学論
講義資料
20
CUDAの登場
• ライブラリもプログラミング言語も一般へ普及せず
• 理由：性能が出せない、使い勝手が良くない、etc.
• 2007年初頭、NVIDIA社がCUDAを公開
• CUDA Unified Device Architecture
• アーキテクチャ＋プログラミング環境
• GPUのアーキテクチャを以前より（とても）詳しく公開
• Cベースのプログラミング環境とドキュメントを無償で公開
• まともなGPU最適化プログラミングが可能に
ユーザが爆発的に増加
高性能コンピューティング学論
2010/7/26
講義資料
21
CUDAアーキテクチャ1
• 物理的な構成の概要
• ※最新版ではキャッシュを搭載するなど変わってきているが、こ
こでは初期型について簡単に解説する
HOST
GPU
Streaming Multiprocessor (MP) ×N
Register
Instruction Unit
×8
SharedMemory
ConstantCache
TextureCache
DeviceMemory
MainMemory
PCI-Express
ScalarProcessor (SP)
2010/7/26
高性能コンピューティング学論
講義資料
22
ハードウェアの特徴
• 階層性のあるハードウェア構成
• 演算器の構成
• 階層性のある演算器配置（SP*8×MP*N）
• 同一MP内のSPは同時に同じ演算しか行えない（SIMD的な構成）
• NVIDIAはSIMTと呼んでいる
• CPUのコアとは趣が異なるので注意
• メモリの構成
• 階層性と局所性のあるメモリ配置
• GPU上に搭載された大容量でグローバルなDRAM：DeviceMemory
• MPローカルの小容量高速共有メモリ：SharedMemory
• 階層的な共有メモリ＋ローカルメモリ
• 全体で共有可能なメモリ＋部分的な共有メモリ＋ローカルメモリ
高性能コンピューティング学論
2010/7/26
講義資料
23
CUDAアーキテクチャ2
• 実行モデル＋メモリ構成の概要
MPに対応
Grid ×α
SPに対応
Register
Thread ×γ
SharedMemory
GlobalMemory
ConstantMemory
Host(CPU, MainMemory)
TextureMemory
Block ×β
LocalMemory
CPUのプロセスやスレッド同様に、Block・Threadは物理的な
数を超えて生成可能
むしろThreadは物理的な数を超えて作成するべき（後述）
高性能コンピューティング学論
2010/7/26
講義資料
24
詳細なメモリ構成
• 特徴の異なる複数種類のメモリ
• 動かすだけならRegisterとGlobalMemoryだけで十分
• GlobalMemoryとSharedMemoryの最適化が重要
名称
Lifetime
共有範囲
速度
容量
GlobalMemory
プログラム GPU全体 高速・高レイテンシ
～*GB
ConstantMemory
プログラム GPU全体 高速・高レイテンシ
＋キャッシュ
64KB
TextureMemory
プログラム GPU全体 高速・高レイテンシ
＋キャッシュ
GlobalMemor
yと共用
SharedMemory
Grid
MP単位
超高速・低レイテンシ
16～KB/MP
Register
Grid
SP単位
超高速・低レイテンシ
8～KB/MP
LocalMemory ※
Grid
SP単位
高速・高レイテンシ
-
※実体はGlobalMemory、レジスタを使いすぎるとLocalMemoryに配置される
高性能コンピューティング学論
2010/7/26
講義資料
25
CUDAアーキテクチャ3
• 処理の流れ
Grid ×α
Register
Thread ×γ
LocalMemory
SharedMemory
GlobalMemory
ConstantMemory
Host(CPU, MainMemory)
TextureMemory
Block ×β
高性能コンピューティング学論
2010/7/26
講義資料
26
CUDAアーキテクチャ4
• もう少し詳しい実行モデル解説
既存のCPU
CUDA
命令H
命令h
命令H
命令G
命令g
命令G
命令F
命令f
命令F
命令E
命令e
命令E
命令D
core 0
命令C
命令d
core 1
命令c
命令D
SP 1 SP 2 SP 3
命令C
命令B
命令b
命令A
命令a
命令A
↓ SP毎（Thread毎）に分岐させたい場合にはマスク処理
↓
変数（Register）を活用することでSP毎に別のメモリを処
理するのが一般的
↓
SP 0
…
命令B
高性能コンピューティング学論
2010/7/26
講義資料
27
CPUとGPUの比較
CPU
(Intel Xeon W3520)
GPU
(NVIDIA GeForceGTX280)
演算器（コア）数
4
240
クロック周波数
2.66GHz
602MHz(SP)
1296MHz(SP)
搭載メモリ容量
-
1GB
メモリ帯域幅
25.6GB/s
141.7GB/s
最大消費電力
130W
236W
理論演算性能
42.56GFlops
933GFlops
CPUとの接続
-
PCIe x16(Gen2)
（主なカタログスペックだけでも多くの差があるが、
中身を知るとそれ以上に違うことがわかるだろう）
2010/7/26
高性能コンピューティング学論
講義資料
28
プログラミング環境としてのCUDA
• CUDA C (RuntimeAPI)
• C/C++ベースの並列化プログラミング言語
• 基本的にC/C++
• 接頭辞による実行対象ハードウェアおよびメモリ配置指定
• 実行対象指定
• __global__
CPUから呼び出しGPU上で実行
• __device__
GPUから呼び出しGPU上で実行
• __host__
CPUから呼び出しCPU上で実行
• 配置指定
• __device__
GlobalMemory
• __shared__
SharedMemory
• __constant__
ConstantMemory
• (TextureMemoryは専用のクラスを用いて扱う)
• バックエンドにGCCやVC++が必要
• （より細かくCUDAを制御可能なDriverAPIもあるが、ここでは
割愛）
2010/7/26
高性能コンピューティング学論
講義資料
29
CUDAプログラミング
• シンプルな配列加算の例
__global__ void gpumain
(float* vc, float* va, float* vb, int nSize, float a, float b){
for(int i=0; i<nSize; i++){
vc[i] = a*va[i] + b*vb[i];
}
}
GPUの処理
grid,threadで指定
された数だけ実行
される
（CPUからGPU上のプロセッサを個別に操作しない）
void main(){
CUT_DEVICE_INIT();
cudaMalloc((void**)&d_va, n);
cudaMalloc((void**)&d_vb, n);
cudaMalloc((void**)&d_vc, n);
cudaMemcpy(d_va, h_va, n, cudaMemcpyHostToDevice);
cudaMemcpy(d_vb, h_vb, n, cudaMemcpyHostToDevice);
gpumain<<< grid, threads >>>(d_vc, d_va, d_vb, nSize, alpha, beta);
cudaMemcpy(h_vc, d_vc, n, cudaMemcpyDeviceToHost);
}
CUDAを用いた配列加算
（vc=a*va + b*vb）の例
CPUの処理
2010/7/26
高性能コンピューティング学論
講義資料
30
CUDAプログラミング
• シンプルな配列加算の例（並列計算版）
__global__ void gpumain
(float* vc, float* va, float* vb, int nSize, float a, float b){
int begin = blockIdx.x*blockDim.x + threadIdx.x;
int step = gridDim.x*blockDim.x;
for(int i=begin; i<nSize; i+=step){
vc[i] = a*va[i] + b*vb[i];
}
（シェーダと比べて）非常にわかりやすくなった
}
CUDAに関する理解は必要
void main(){
CUT_DEVICE_INIT();
cudaMalloc((void**)&d_va, n);
cudaMalloc((void**)&d_vb, n);
cudaMalloc((void**)&d_vc, n);
cudaMemcpy(d_va, h_va, n, cudaMemcpyHostToDevice);
cudaMemcpy(d_vb, h_vb, n, cudaMemcpyHostToDevice);
gpumain<<< grid, threads >>>(d_vc, d_va, d_vb, nSize, alpha, beta);
cudaMemcpy(h_vc, d_vc, n, cudaMemcpyDeviceToHost);
}
GPUの処理
grid,threadで指定
された数だけ実行
される
CUDAを用いた配列加算
（vc=a*va + b*vb）の例
CPUの処理
高性能コンピューティング学論
2010/7/26
講義資料
31
CUDAプログラミング（実行方法）
• 通常のCプログラムと同様にコンパイル・実行が可能
nvcc main.cu
./a.out
• CUDAコンパイラ(NVCC)がGPUカーネルを分離し、CPU
部とGPU部それぞれをコンパイルし、単一の実行ファイ
ルを生成
• CPU部またはGPU部のみをコンパイルすることなども可能
• 中間表現ファイルを出力して解析することも可能
• 中間表現の仕様も公開されているが、実行時にさらに改変されること
などから最適化に使うのは現時点では困難
高性能コンピューティング学論
2010/7/26
講義資料
32
MPIやOpenMPとの実行モデル比較
MPI
OpenMP
CUDA
PE0
main
thread
CPU
PE1
PE2
PE3
thread 0,1,2,3
GPU
Block, Thread
(MP, SP)
×
分散メモリ（プロセス間通信）
共有メモリ
CPU-GPU間分散メモリ
GPU内階層型共有メモリ
高性能コンピューティング学論
2010/7/26
講義資料
33
CUDA最適化プログラミング
• 主な最適化技術の概要
（どのようなプログラムに対して高性能が得られるか）
• 大量のThreadを生成する
• 理想的なBlockあたりThread数は128程度
• GlobalMemoryアクセスのコアレスアクセス
• 複数SPに同時に連続したメモリをアクセスさせる
→まとめられて性能向上
• SharedMemoryのバンクコンフリクト回避
• SharedMemoryにアクセスする際に別々のメモリバンクを叩かせる
→アクセスが衝突せずに性能低下を回避
• ループアンローリング
• 分岐回数が減るため重要
2010/7/26
高性能コンピューティング学論
講義資料
34
大量のThreadを生成する
• SP(Thread)のコンテキスト切り替えコストは非常に小
• メモリアクセスを待つよりコンテキストを切り替えて別のThread
を処理した方が速い
• 大量のThreadを生成することで高性能
• （むしろ、大量のThreadでGlobalMemoryメモリアクセスのレイテ
ンシを隠蔽しないと高い性能が得られない）
• RegisterやSharedMemoryの使用量が多いと多数のThread
を保持できない
• GPUカーネルをシンプルにするべき
高性能コンピューティング学論
2010/7/26
講義資料
35
GlobalMemoryのコアレスアクセス
• 同一MP内の複数SPによるメモリアクセスを揃える（近
づける）とまとめてアクセスできる
• 詳細な条件はGPUの世代によって異なる（緩和が進んでいる）
• アクセスが揃っていない場合
SP0
SP1
SP2
SP3
4回のメモリアクセスが行われる
GlobalMemory
• アクセスが揃っている場合
SP0
SP1
SP2
GlobalMemory
SP3
1回のメモリアクセスに纏められる
高性能コンピューティング学論
2010/7/26
講義資料
36
SharedMemoryのバンクコンフリクト回避
• SharedMemoryは16個・32bitずつのバンクにより構成
• 同一バンクへのアクセスが集中すると性能低下
• 均等なアクセス＝性能低下しない
SharedMemory
• アクセスが集中＝性能低下する
SharedMemory
2-way バンクコンフリクトの例
2010/7/26
高性能コンピューティング学論
講義資料
37
CUDA最適化プログラミングのまとめ
• どのようなプログラムで高い性能が得られるか
• 高い並列度を持っている
• メモリアクセスに規則性がある
• GlobalMemoryのコアレスアクセスやSharedMemoryのバンクコンフリ
クト回避ができる
• テクスチャキャッシュに収まる範囲であれば多少のランダム性も容認
• GlobalMemoryへのアクセスがランダムでも、SharedMemoryに落とせ
ればOK
• CPU-GPU間の通信が頻繁に必要でない
• SP間の同期を繰り返す必要がない
2010/7/26
高性能コンピューティング学論
講義資料
38
GPGPUを試す方法（情報源）
• NVIDIAとAMDの開発者サイトに開発に必要なものと資料
が揃っている
• ダウンロードしてサンプルの実行・改造をすると良い
• NVIDIA
• http://developer.nvidia.com/page/home.html
• CUDA
• AMD
• http://developer.amd.com/pages/default.aspx
• ATI Stream SDK
高性能コンピューティング学論
2010/7/26
講義資料
39
CUDAプログラミングを試す方法
• 準備
• ハードウェア
• 普通のPC
• CUDA対応GPUが無くてもエミュレーションで試すことは可能だが、１万円
未満のGPUでも良いので積んでおきたい
• CUDA関連ソフトウェア
• CUDA Driver
• CUDA Toolkit
• CUDA SDK code samples
• その他
• Windows・Linux・MacOSXに対応
• gccかVisualStudioが必要
• 個人的にはWindowsよりLinuxの方が環境を作りやすいが、VisualStudio統合
デバッガが強力なのは捨てがたい
• 一部のメーカー製PCではドライバ導入が大変（最悪、不可能）
2010/7/26
高性能コンピューティング学論
講義資料
今後のGPUとGPGPU
GPGPUの将来はどうなるのか、最新のGPU事情
40
2010/7/26
高性能コンピューティング学論
講義資料
現在のGPGPU
• ハードウェア
• NVIDIAが大きな影響力
• CUDAが大きな原動力
• OpenCLが普及し始めるとAMDも対等に戦えるか？
• AMDはほぼシェーダの開発環境のみ提供していた
• IntelがGPUを出すという噂は常にあるが……
• Larrabee？
• プログラミング環境
• CUDAが中心
• 徐々に高まるOpenCLへの注目
41
2010/7/26
高性能コンピューティング学論
講義資料
42
OpenCL
• Khronosがまとめている並列計算フレームワーク
• OpenGLの並列計算版のようなもの
• Appleが提案
• Apple版CUDA……？
• C言語(C99)の拡張言語
• CUDAに似た接頭辞(__kernel、__globalなど)
• ベクトル型
• マルチコアCPUやCell B/E、DSPなどの並列計算記述にも使えるよ
うに規格化されている
2010/7/26
高性能コンピューティング学論
講義資料
スパコンとGPU1
• GPUスパコンの登場
• TSUBAME1.2（東工大）
• TOP500(2008.11) #29、Tesla S1070 170台(x4)
• PeakFlops 161.82 TFlops / GPU 58.65 TFlops
• Tianhe-1（中国）
• TOP500(2009.11) #5、RadeonHD4870x2 2560台(x2)
• PeakFlops 1206.20 TFlops / GPU * TFlops
• Nebulae（中国）
• TOP500(2010.6) #2、Tesla C2050 4640台(x1)
• PeakFlops CPU 593.77 TFlops + GPU 2390.53 TFlops
• TSUBAME2.0（東工大）
• TOP500(2010.11) 予定、Tesla M2050 4224台(x1)
• PeakFlops CPU 224.69 TFlops + GPU 2175.36 TFlops
43
2010/7/26
高性能コンピューティング学論
講義資料
44
スパコンとGPU2
• TOP500とGPU
• スコアを大きく左右するのは行列積の性能（演算速度・メモリ速
度）と通信性能
• GPUでスコアを出しやすい問題ではある
• GPUスパコンの課題
• 演算効率（性能が出せるか）
• Tianhe-1やNebulaeは50%も出せていない
• 稼働率（使われるか）
• ユーザがGPUプログラムを作成できるか→教育、ライブラリ
• GPUに適している問題ばかりではない
• 故障率（壊れないか）
• 年単位で動かした実績がほとんどない
2010/7/26
高性能コンピューティング学論
講義資料
45
最新GPUの動向
• GPUが様々なアプリケーションに有用であることは広く
知れ渡りつつある
• より様々な分野で利用され始めた
• プロダクトでの利用が始まった
• 動画編集ソフト、科学技術計算ライブラリ
• 倍精度浮動小数点演算性能やメモリの信頼性を求める声
が大きくなってきた
• 倍精度浮動小数点演算性能が可能な（性能が高い）GPUの登場
• Fermiアーキテクチャ、RadeonHD5xxx
• ECCサポートGPUの登場
• Fermiアーキテクチャ版Tesla
2010/7/26
高性能コンピューティング学論
講義資料
46
最新のCUDAアーキテクチャ
• Fermiアーキテクチャの主な更新ポイント
• 高速な倍精度浮動小数点演算に対応（単精度の半分の性能）
• 浮動小数点atomic演算に対応
• BlockあたりThread数が増加（512->1024）
• MPあたりRegister数が増加（8K,16K->32K）
• MPあたりSharedMemory容量が増加（16KB->48KB）
• SharedMemoryのバンク数が増加（16->32）
• キャッシュを搭載（L1$新規搭載、L2$の汎用化）
• DirectXとCUDAの切り替えが高速化
• チューニングに大きく影響
• （アーキテクチャ更新が速いのは善し悪し？）
高性能コンピューティング学論
2010/7/26
講義資料
47
GPU vs メニーコアCPU
• 汎用化するGPU
• よりプログラムを組みやすく（性能を出しやすく）
• 増加するCPUコア
• 2010年7月現在：最大6コア
• 微細化等によりさらなるコア数増加が予定されている
• 5年後には100コア規模に達するという予測も
• 課題：どうやって使う？何に使う？本当に100コアも必要か？（→CPU
コアとGPUコアの統合）
• 汎用化したGPUとメニーコアCPUは似たようなハード
ウェアになる？
• シンプルCPUコアを大量に並べるLarrabeeはCPUとGPUの両方を
取り込む……予定だった
• 私見：「シンプル超多数コアのGPU」＋「複雑さを維持しコア数
も抑えるCPU」
2010/7/26
高性能コンピューティング学論
講義資料
48
今後（近い将来）のGPU/GPGPU像
• コア数・理論演算性能のさらなる増加
• CPUとの統合が進む
• IntelのCore iに統合されているGPUはGPGPUに使えないレベル
• 本命：AMDのFusion
• 疑問：メモリはどうなる？プログラミングモデルに影響は？
• HPC向けGPUとグラフィックスGPUが分かれる可能性
• そもそも要求が異なる
• 共通のベースアーキテクチャを持つ異なる者
• プログラミング環境はどうなる
• CUDA=NVIDIA GPUアーキテクチャ
• 競争が起きる必要がある->OpenCL？
※勝手な予想を含む
2010/7/26
おまけ
その他
高性能コンピューティング学論
講義資料
49
2010/7/26
高性能コンピューティング学論
講義資料
50
SIMD型アクセラレータ
• GPU以外にもCPUと比べて浮動小数点演算を並列・高速
に扱えるハードウェアは存在する（していた）
• ClearSpeed（英ClearSpeed Technology社）
• TSUBAMEに搭載されているアクセラレータ
• 事実上GPUに駆逐された（高価だったのが最大の問題か）
• Cell Broadband Engine（SCE、ソニー、IBM、東芝）
• PS3に搭載されていることで有名なプロセッサ
• HPC向けの強化版や東芝のノートPC搭載SpursEngineなども
• 一時期GPGPUとともに盛り上がった（次の製品が出ないのが問題か）
• GRAPE
• 重力多体問題専用の計算機
• 現在もバージョンアップが続いており天文学の研究で活用されている
• 最近、TOP500のLittle Green500で一位になった(GRAPE-DR)
2010/7/26
高性能コンピューティング学論
講義資料
51
おわりに
• GPGPUについて、およびCUDAについて（広く浅く）紹
介した
• 興味を持った場合は実際にプログラムを作成してみて欲
しい
• GPUプログラミングコンテストも開催されているので是
非挑戦して欲しい
• （恐らく今年度も開催する、はず）
• 非常に身近でHOTな（CPUと同等以上に発展著しい）
ハードウェア・ソフトウェアなので、挑戦する価値あり
• 情報を追いかけるだけでも楽しいかもしれない？
2010/7/26
高性能コンピューティング学論
講義資料
内容に関する質問・問い合わせ先
大島聡史(東京大学 情報基盤センター 助教)
[email protected]
52
© Copyright 2024 ExpyDoc
