電子情報通信学会研究会 組込みシステム研究会 (IPSJ-EMB)

2010年1月28日

超並列マルチコアGPUを用いた高速演算処理の実用化
NVIDIA Solution Architect
馬路 徹© 2009 NVIDIA Corporation

目次
なぜ 今GPUコンピューティングか?
CPUの性能向上速度が減速“性能 向上= 並列処理”にGPUコンピューティングが応える
CUDAシステムアーキテクチャによる超並列処理の実現、普及

NVIDIA GPUアーキテクチャの進歩
３Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷
統合化シェーダエンジンへのCUDAの組込み

次世代GPUコンピューティング、Fermiアーキテクチャ
プロセッサ / メモリーアーキテクチャ
GigaThreadスレッド制御機構による更なる高性能化
統合化プログラム開発環境

まとめ© 2009 NVIDIA Corporation

CUDAはPerformance = Parallelismを実現
Performance = Parallelismへの期待
このParallelismを実現するために必要なことは?
相当な数の高効率プロセッサ（4個、8個のオーダではない）
並列処理を抽象化できるプログラミング・システム“CUDA (Compute Unified Device Architecture”
GPUは数百個の多機能、高速プロセッサを内蔵
マルチスレッド・アーキテクチャがこの超並列マルチコアを高効率に活用
CUDA は普及しているC言語の拡張からスタートし、現在はFortran、OpenCL やDirect Computeまで拡張されている
CUDA は並列プログラミングの詳細な記述からプログラマーを解放© 2009 NVIDIA Corporation最初に「C with CUDA Extension」を開発

標準CコードCPUからコールしGPUで実行する関数

CUDAコード
ｸﾞﾛｰﾊﾞﾙ･ｽﾚｯﾄﾞIDをそのままｲﾝﾃﾞｯｸｽに活用

void saxpy_serial(int n, float alpha, float*x, float *y){for (int i=0; i<n; ++i)y[i] = alpha * x[i] + y[i];}// Invoke serial saxpy() kernel
saxpy_serial(n, 2.0, x, y);__global__ void saxpy_parallel(int n, float alpha, float *x, float *y){int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) y[i] = alpha * x[i] + y[i];}// Invoke parallel saxpy() kernel (256 threads per block)int nblocks = (n + 255)/256;
saxpy_parallel <<< nblocks, 256 >>>(n, 2.0, x, y);

階層的スレッド分割数（ブロック数／グリッド）© 2009 NVIDIA Corporation（スレッド数／ブロック）今日のGPU コンピューティングの普及
市場展開モーメンタム
CUDA実装のGPUが市場に1億個以上
出荷60,000以上のGPU コ
ンピューティング開発者
Windows, Linux 及びMacOS プラットフォー
ムのサポート
200以上の大学でGPU コンピューティン
グの講座

GPU コンピューティング・アプリ

C+CUDA
Extension

OpenCL

DirectComputeKhronos
AppleMicrosoftFORTRANJava andPython
plannedNVIDIA GPU
with the CUDA Parallel Computing Architecture© 2009 NVIDIA Corporation
OpenCL is trademark of Apple Inc. used under license to the Khronos Group Inc.

CUDA Not 2x or 3x, Speed-ups are 20x to 150x

Results withTelsa 8 Series
relative to146X

36X

Interactive
visualization of
volumetricwhite matter
connectivity

Ionic placement
for molecular
dynamics
simulation onGPU

149X
Financialsimulation ofLIBOR modelwith swaptions

19X

17X

100X

Transcoding
HD video
stream toH.264

Simulation inMatlab using.mex file CUDA
function

Astrophysics Nbody simulation47X

20X

24X

30X

GLAME@lab:
An M-script API
for linear
Algebraoperations onGPU

Ultrasoundmedicalimaging forcancerdiagnostics

Highlyoptimized
object oriented
moleculardynamics

Cmatch exact
string matching
to find similarproteins andgenesequencescomputeperformanceusing CPU
exclusively© 2009 NVIDIA Corporation

並みの性能向上ではない
2-3倍の性能向上は単なる「高性能化」
顧客の基本的なワークフローに変化は無い

5-10倍の性能向上は「画期的」
装置のアップグレードの価値は十分にある（一部または大部分の）アプリ・ソフトを書き換える意味はある

100倍以上の性能向上は「世界観を変える！」
プラットフォームの取替えの価値はある
アプリケーションのアーキテクチャまでを見直す意味がある
今まで実用的に不可能であったアプリの開発が可能になる
科学技術において「新発見までの時間」を短縮する画期的な変化をもたらす© 2009 NVIDIA Corporation

GPUコンピューティングが導入されたＯＳ
身近になるＧＰＵコンピューティング
Mac OS X Snow Leopard© 2009 NVIDIA Corporation

目次
なぜ 今GPUコンピューティングか?
CPUの性能向上速度が減速“性能 向上= 並列処理”にGPUコンピューティングが応える
CUDAシステムアーキテクチャによる超並列処理の実現、普及

NVIDIA GPUアーキテクチャの進歩
３Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷
統合化シェーダエンジンへのCUDAの組込み

次世代GPUコンピューティング、Fermiアーキテクチャ
プロセッサ / メモリーアーキテクチャ
GigaThreadスレッド制御機構による更なる高性能化
統合化プログラム開発環境

まとめ© 2009 NVIDIA Corporation

NVIDIA GPUによる3Dグラフィックスの進歩
ﾌﾟﾛｸﾞﾗﾏﾋﾞﾘﾃｨｰの大幅向上 (統合化ｼｪｰﾀﾞｰｴﾝｼﾞﾝ) → 汎用演算器

1995

1999

2002

2003

2004

2005

NV1
1 Million
Transistors

GeForce 256
DirectX 7
22 Million
Transistors

GeForce4
Direct X 8
63 Million
Transistors

GeForce FX
DirectX 9
130 Million
Transistors

GeForce 6
DirectX 9c222 Million
Transistors

GeForce 7
DirectX 9c
302 Million
Transistors

Vertex(FP)+Pixel Shaderﾌﾟﾛｸﾞﾗﾏﾋﾞﾘﾃｨｰの向上© 2009 NVIDIA Corporation

Pixel
Shade

Vertex(FP)+Pixel(FP) Shader

2006
GeForce 8
DirectX 10 (Vista)
681 Million
Transistors

統合化Shader
CUDA2008
GeForce
GTX2001.4 Billion
Transistors

20年間続いたグラフィックス・アーキテクチャ
各機能に固定された用途のハードウエア
Vertex

座標変換及び光源処理
DirectX8よりプログラマブル頂点シェーダーとなる

Triangle

三角形、点、線セットアップ

Pixel© 2009 NVIDIA Corporation

フラット・シェーディング、テクスチャーマッピング等
DirectX8よりプログラマブル・ピクセルシェーダーとなる

RasterOPeration

ブレンディング、Ｚバッファー、アンチエイリアシング

Memoryメモリ

DirectX 10 (Vista) 以降のGPU
統合型プログラマブル・シェーダエンジン
ＧＰＵはマルチスレッドを処理する超並列マルチコアプロセッサとなった。
頂点シェーダ、ジオメトリ・シェーダ及びピクセル・シェーダは上記により実行される。
頂点シェーダHost

ジオメトリシェーダ

ピクセルシェーダ

Input AssemblerSP

Thread Execution Manager

Thread Processors

Thread Processors

Thread Processors

Thread Processors

Thread Processors

Thread Processors

Thread Processors

Thread Processors

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

Parallel
Data
Cache

SMLoad/store

SP: Stream Processor
SM: Stream Multiprocessor© 2009 NVIDIA Corporation

Global Memory

Parallel
Data
CacheParallel
Data
Cache

Parallel
DataCache

Parallel
DataCache第1世代のSM (Streaming Multiprocessor)• Streaming Multiprocessor (SM)SM– 8 Streaming Processors (SP)

命令フェッチ

各SPは32b単精度浮動小数点、整数演算サポート

I命令キャッシュ– 2 Super Function Units (SFU)

マルチスレッド命令ディスパッチ

三角関数、平方根等サポート– 全てのSPは同一命令を実行開始（途中分岐あり)

シェアド・メモリ

SIMT (Single instruction Multiple Thread)• マルチスレッド命令ディスパッチ
– 1 から768 スレッドがアクティブ– 32スレッド単位(warp) でSIMT命令実行– 巨大なローカル・レジスタファイルRF
8,192 Registers / SMハードウエア・コンテキストスイッチを容易化• 16 KB シェアード・メモリ© 2009 NVIDIA Corporation

SFU

SP 0

RF 0

RF 4

SP 4

SP 1

RF 1

RF 5SP 5SP 2RF 2

RF 6

SP 6SP 3

RF 3

RF 7

SP 7SFU定数キャッシュテキスチャーフェッチ

NVIDIA GPU コア数の変遷
Unified Shader

Number of programmable shaders / dieFermi Architecture1024

Tesla Architecture

512
256

8800GTX 128

128

Vertex Shader + Pixel Shader

64
32
1684Ti500

FX5600 1+26200

GeForce
5XXX

GeForce
6XXX

7900GT
7800GT
6+165+12 7600GT

GeForce
3XXX

GeForce
4XXX

8+16
6+16
5+12

4+8
3+47500GT

3+41+27300SE
7050SE2+21+4Ti4000 0+2211+46800GT
6800
FX5900 3+86800XT6200
Ti4800 2+4 FX5700 3+28800GS

GeForce
7XXX

96240
GTX285
GTX260 216
9800GTX 128 GTX250
128
GTX240 112

9600GT

648600GTS

329500GT

328500GT

169400GT

168300GT

8FERMI
C1060

240

C870

128

- Vertical axis shows the number of cores, but not the performance- Only major products are shown
GeForce
8XXX

GeForce
9XXX

GeForce
GT2XX

TeslaFermi

Programmable Shader
Unified Shader
CUDA, OpenCL, DirectX Compute etc© 2009 NVIDIA Corporation

512GPU-CPU 性能及びメモリーバンド幅の 差は拡大
メモリーバンド幅ピーク値
GB/sec1000

8x double precision
ECCL1, L2 Caches
1 TF Single Precision
4GB Memory© 2009 NVIDIA Corporation

100

NVIDIA GPU
X86 CPU目次
なぜ 今GPUコンピューティングか?
CPUの性能向上速度が減速“性能 向上= 並列処理”にGPUコンピューティングが応える
CUDAシステムアーキテクチャによる超並列処理の実現、普及

NVIDIA GPUアーキテクチャの進歩
３Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷
統合化シェーダエンジンへのCUDAの組込み

次世代GPUコンピューティング、Fermiアーキテクチャ
プロセッサ / メモリーアーキテクチャ
GigaThreadスレッド制御機構による更なる高性能化
統合化プログラム開発環境

まとめ© 2009 NVIDIA Corporation

‘Fermi’ アーキテクチャのご紹介
DRAM I/F

DRAM I/F

スーパコンピュータの魂を持つGPU

DRAM I/F

DRAM I/F

DRAM I/F

Giga Thread

HOST I/F

DRAM I/F© 2009 NVIDIA Corporation

L2

30億個のトランジスター

コア数を倍増（512 コア）
倍精度浮動小数点演算ピーク性能が8倍GPUとして初めてECCを導入
L1 及び L2 キャッシュを内蔵

約2倍のメモリーバンド幅（GDDR5）
最大1 Terabyte の GPU メモリ
複数Kernel同時実行、C++サポートSM (Streaming Multiprocessor)アーキテクチャ

Instruction Cache
Scheduler Scheduler
Dispatch

SM当たり32 CUDAコア (総数512個)

Dispatch

Register File
Core Core Core Core

倍精度浮動小数点演算ピーク性能は8倍
単精度演算ピーク性能の50%Core Core Core Core
Core Core Core Core
Core Core Core Core

2個の Thread Scheduler
2ワープを2組の16個のCUDA Core, 16個のLoad/Store Unit, 4個のSFUに同時にディスパッチするCore Core Core Core
Core Core Core Core
Core Core Core Core
Core Core Core Core
Load/Store Units x 16
Special Func Units x 4

シェアードメモリ、L1キャッシュとして使用
する 64 KB RAM （構成可変）Interconnect Network
64K Configurable
Cache/Shared Mem
Uniform Cache© 2009 NVIDIA Corporation

CUDA コア・アーキテクチャ

Instruction Cache
Scheduler Scheduler
Dispatch

最新の浮動小数点演算規格 IEEE 754-2008 準拠

Dispatch

Register File

（最新CPUをも凌駕）Core Core Core Core
Core Core Core Core

Fused multiply-add (FMA) 命令を倍精度、単精度でサポート（積和の最終段でラウンディングするため、各段で各々行うより精度が向上）

Core Core Core Core

CUDA Core
Dispatch Port
Operand Collector

Core Core Core Core
Core Core Core Core
Core Core Core Core
Core Core Core Core

新規設計の全命令32bサポート整数ALU
64-bit 及びそれ以上の精度に対しても最適化設計

FP UnitINT Unit

Core Core Core Core
Load/Store Units x 16

Result Queue

Special Func Units x 4
Interconnect Network
64K Configurable
Cache/Shared Mem
Uniform Cache© 2009 NVIDIA Corporation

統合化された64ビットメモリー空間
Ｃ／C++ ポインターのフルサポート
Non-unified Address Space

C言語やC++言語のポインタはターゲットとするアドレス空間が必ずしもLocal

コンパイル時に確定せず、ランタイムに動的に決まるケースがあるため、このポインタを完全に実装することが困難な状況となっていました。しか*p_local

し、統合化されたアドレス空間をもつFermiではこの問題もなくなりました。

Shared*p_shared

Global032-bit*p_global

Unified Address Space
Local

SharedGlobal0© 2009 NVIDIA Corporation

40-bit*pキャッシュメモリ階層構造
オンチップ・シェアードメモリとともに本格的なキャッシュ階層構造を有する初めてのGPU

Fermi Memory Hierarchy
Thread

各SM (32 cores)にL1キャッシュ（48KBまたは16KB）
メモリバンド幅の改善及びアクセス遅延の低減

Shared Memory

L1 Cache

ユニファイド L2 キャッシュ(768 KB)L2 CacheM
AR
D

GPU上の全コアにわたり高速に、コーヒレンント・データをシェア© 2009 NVIDIA Corporation拡張され、高速化されたメモリ・インタフェース
GDDR5 メモリーインタフェース

HOST I/F

DRAM I/F

大きなデータセットを取り扱うことが可能
現在40ｂアドレス。命令は64ｂをサポート

DRAM I/F

最大1 Terabyte のGPUメモリ

DRAM I/FGDDR3の2倍の速度© 2009 NVIDIA Corporation

Giga Thread

DRAM I/F

64b partitioning x 6

DRAM I/F

384ビットのメモリーインタフェース

DRAM I/F

L2

ECC (Error Checking and Correction)
DRAMのECCによるデータエラーの除去
GDDR5 メモリーのECCサポート

内部の主要な記憶素子もECC で保護
レジスターファイル、L1キャッシュ、L2キャッシュ

Single-Error Correct Double-Error Detect (SECDED)サポート
２ビットの誤りはソフト処理（再度実行等）© 2009 NVIDIA Corporation

IEEE 754-2008 規格準拠の高精度演算
IEEE 754-2008 results

Multiply-Add (MAD): D = A*B + C;A

64-bit double precision32-bit single precision

IEEE 754-2008 rounding© 2009 NVIDIA Corporation+C=D

Fused Multiply-Add (FMA): D = A*B + C;AIEEE 754-2008 Fused Multiply-Add (FMA)
D = A*B + C;
No loss of precision
IEEE divide & sqrt use FMA

B= Product (truncate digits)

full-speed denormal operands & results
NaNs, +/- Infinity

nearest even, zero, +inf, -inf

×=×
Product+C=DB(retain all digits)(no loss of precision)

GigaThreadTM Hardware Thread Scheduler (HTS)
階層的に何千ものアクティブなスレッドを管理
コンテキストスイッチが10倍高速
HTS

複数kernelの同時実行© 2009 NVIDIA Corporation

GigaThread Hardware Thread Scheduler
複数Kernel同時実行 + 高速コンテキストスイッチ
Kernel 1

Kernel 1

TimeKernel 2

Kernel 2
nel

Kernel 2

Kernel 2
Kernel 3
Kernel 5

Kernel 3
Kernel 4
Kernel 5

Serial Kernel Execution© 2009 NVIDIA Corporation

Parallel Kernel Execution

Ker4

GigaThread Streaming Data Transfer Engine
デュアルDMA エンジン
CPUGPU 及び GPUCPU データ転送の同時実行
CPU と GPU演算と完全にオーバラップ可能
SDT

処理の流れ:
Kernel 0© 2009 NVIDIA CorporationCPUSDT0GPU

SDT1

Kernel 1

CPU

SDT0

GPU

SDT1

Kernel 2CPU

SDT0GPU

SDT1

Kernel 3CPU

SDT0GPUSDT1

各世代の比較
GPU

G80

GT200

Fermi

集積トランジスタ数
CUDAコア数
倍精度浮動小数点演算能力

6億8100万個
128×

単精度浮動小数点演算能力

128 MAD演算/クロック
1

14億個
240
30 FMA演算/クロック
240 MAD演算/クロック
130億個
512
256 FMA演算/クロック
512 MAD演算/クロック
2224

共有メモリ/SM

16KB

16KB

L1キャッシュ/SM

×××××
32ビット

×××
32ビット

48KB/16KB(構成可能)
16KB/48KB(構成可能)
768KB○
最大1664ビット

ワープスケジューラ(SMあたり搭載数)
特殊関数ユニット(SFU) /SM

L2キャッシュ/SM
ECCメモリのサポート
平行実行カーネル数
ロード/ストアのアドレス幅© 2009 NVIDIA Corporation

NVIDIA Nexus IDE（統合化開発環境）
業界初の超並列アプリ開発用 IDE (Integrated Development Environment)
C, C++, OpenCL, DirectCompute 及びDirectX と OpenGLの両グラフィックスAPIをサポート

完全にVisual Studioに組込まれた 開発環境
(CPU + GPU)コプロセッシング・アプリの開発を効率化
両プロセッサにまたがるソース・デバッグ、性能解析•

両プロセッサにわたるイベント、データのキャプチャ© 2009 NVIDIA Corporation

NVIDIA Nexus IDE（統合化開発環境）画面© 2009 NVIDIA Corporation

まとめ
CPUのIPL性能向上は減速している
Performance = Parallelism が今後の性能確保の要
統合化シェーダエンジンによりGPUは汎用の超並列プロセッサとなる
CUDA GPU コンピューティングはGPUの並列性能を最大限に引き出し、並列プログラミングの抽象化によりプログラミングも容易にする
次世代Fermiアーキテクチャは更にコンピューティング機能、性能を向上した© 2009 NVIDIA Corporation

Thank you for your attention© 2009 NVIDIA Corporation