Slide 1

Size: px

Start display at page:

Download "Slide 1"

ありさじゅふく
4 years ago
Views:

1 電子情報通信学会研究会組込みシステム研究会 (IPSJ-EMB) 2010 年 1 月 28 日超並列マルチコア GPU を用いた高速演算処理の実用化 NVIDIA Solution Architect 馬路徹

2 目次なぜ今 GPU コンピューティングか? CPUの性能向上速度が減速性能向上 = 並列処理にGPUコンピューティングが応える CUDAシステムアーキテクチャによる超並列処理の実現普及 NVIDIA GPU アーキテクチャの進歩 3Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷統合化シェーダエンジンへのCUDAの組込み次世代 GPU コンピューティング Fermi アーキテクチャまとめプロセッサ / メモリーアーキテクチャ GigaThreadスレッド制御機構による更なる高性能化統合化プログラム開発環境

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

4 最初に 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); 階層的スレッド分割数 ( ブロック数 / グリッド ) ( スレッド数 / ブロック )

今日の GPU コンピューティングの普及 OpenCL is trademark

5 今日の GPU コンピューティングの普及 OpenCL is trademark of Apple Inc. used under license to the Khronos Group Inc. 市場展開モーメンタム CUDA 実装の GPU が市場に 1 億個以上出荷 60,000 以上の GPU コンピューティング開発者 Windows, Linux 及び MacOS プラットフォームのサポート 200 以上の大学で GPU コンピューティングの講座 C + CUDA Extension GPU コンピューティングアプリ OpenCL Khronos Apple Direct Compute Microsoft FORTRAN NVIDIA GPU with the CUDA Parallel Computing Architecture Java and Python planned

6 CUDA Not 2x or 3x, Speed-ups are 20x to 150x Results with Telsa 8 Series relative to compute performance using CPU exclusively 146X 36X 19X 17X 100X Interactive visualization of volumetric white matter connectivity Ionic placement for molecular dynamics simulation on GPU Transcoding HD video stream to H.264 Simulation in Matlab using.mex file CUDA function Astrophysics N- body simulation 149X 47X 20X 24X 30X Financial simulation of LIBOR model with swaptions GLAME@lab: An M-script API for linear Algebra operations on GPU Ultrasound medical imaging for cancer diagnostics Highly optimized object oriented molecular dynamics Cmatch exact string matching to find similar proteins and gene sequences

7 並みの性能向上ではない 2-3 倍の性能向上は単なる高性能化顧客の基本的なワークフローに変化は無い 5-10 倍の性能向上は画期的装置のアップグレードの価値は十分にある ( 一部または大部分の ) アプリソフトを書き換える意味はある 100 倍以上の性能向上は世界観を変える! プラットフォームの取替えの価値はあるアプリケーションのアーキテクチャまでを見直す意味がある今まで実用的に不可能であったアプリの開発が可能になる科学技術において新発見までの時間を短縮する画期的な変化をもたらす

8 GPU コンピューティングが導入された OS 身近になる GPU コンピューティング Mac OS X Snow Leopard

9 目次なぜ今 GPU コンピューティングか? CPUの性能向上速度が減速性能向上 = 並列処理にGPUコンピューティングが応える CUDAシステムアーキテクチャによる超並列処理の実現普及 NVIDIA GPU アーキテクチャの進歩 3Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷統合化シェーダエンジンへのCUDAの組込み次世代 GPU コンピューティング Fermi アーキテクチャまとめプロセッサ / メモリーアーキテクチャ GigaThreadスレッド制御機構による更なる高性能化統合化プログラム開発環境

NVIDIA GPU による 3D グラフィックスの進歩フロクラマヒリティーの大幅向上 (

1999 GeForce 256 DirectX 7 22 Million Transistors 2002

FX DirectX 9 130 Million Transistors 2004 GeForce 6

DirectX 9c 302 Million Transistors 2006 GeForce 8

10 NVIDIA GPU による 3D グラフィックスの進歩フロクラマヒリティーの大幅向上 ( 統合化シェーターエンシン ) 汎用演算器 1995 NV1 1 Million Transistors 1999 GeForce 256 DirectX 7 22 Million Transistors 2002 GeForce4 Direct X 8 63 Million Transistors 2003 GeForce FX DirectX Million Transistors 2004 GeForce 6 DirectX 9c 222 Million Transistors 2005 GeForce 7 DirectX 9c 302 Million Transistors 2006 GeForce 8 DirectX 10 (Vista) 681 Million Transistors 2008 GeForce GTX Billion Transistors フロクラマヒリティーの向上 Vertex(FP) +Pixel Shader Vertex(FP) +Pixel(FP) Shader 統合化 Shader CUDA Pixel Shade

Pixel Raster OPeration フラットシェーディングテクスチャーマッピング等 DirectX8

11 20 年間続いたグラフィックスアーキテクチャ各機能に固定された用途のハードウエア Vertex Triangle 座標変換及び光源処理 DirectX8 よりプログラマブル頂点シェーダーとなる三角形点線セットアップ Pixel Raster OPeration フラットシェーディングテクスチャーマッピング等 DirectX8 よりプログラマブルピクセルシェーダーとなるブレンディング Z バッファーアンチエイリアシング Memory メモリ

12 DirectX 10 (Vista) 以降の GPU 統合型プログラマブルシェーダエンジン GPU はマルチスレッドを処理する超並列マルチコアプロセッサとなった頂点シェーダジオメトリシェーダ及びピクセルシェーダは上記により実行される Host 頂点シェーダジオメトリシェーダピクセルシェーダ SP Input Assembler 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 Parallel Data Cache Parallel Data Cache Parallel Data Cache Parallel Data Cache SM SP: Stream Processor SM: Stream Multiprocessor Load/store Global Memory

第 1 世代の SM (Streaming Multiprocessor) Streaming Multiprocessor (SM) 8 Streaming Processors (SP) 各 SP は 32b 単精度浮動小数点整数演算サポート 2 Super Function Units (SFU) 三角関数平方根等サポート全ての SP は同一命令を実行開始 ( 途中分岐あり )

13 第 1 世代の SM (Streaming Multiprocessor) Streaming Multiprocessor (SM) 8 Streaming Processors (SP) 各 SP は 32b 単精度浮動小数点整数演算サポート 2 Super Function Units (SFU) 三角関数平方根等サポート全ての SP は同一命令を実行開始 ( 途中分岐あり ) SIMT (Single instruction Multiple Thread) マルチスレッド命令ディスパッチ 1 から768 スレッドがアクティブ 32スレッド単位 (warp) でSIMT 命令実行巨大なローカルレジスタファイルRF 8,192 Registers / SM ハードウエアコンテキストスイッチを容易化 16 KB シェアードメモリ S F U SP 0 RF 0 SP 1 RF 1 SP 2 RF 2 SP 3 RF 3 SM 命令フェッチ I 命令キャッシュマルチスレッド命令ディスパッチシェアドメモリ RF 4 RF 5 RF 6 RF 7 定数キャッシュ SP 4 SP 5 SP 6 SP 7 S F U テキスチャーフェッチ

14 1024 NVIDIA GPU コア数の変遷 Number of programmable shaders / die Ti GeForce 3XXX Vertex Shader + Pixel Shader Ti Ti GeForce 4XXX Programmable Shader FX FX FX GeForce 5XXX 6800GT XT GeForce 6XXX 7900GT GT GT GT SE SE 1+4 GeForce 7XXX 8800GTX GS GTS GT GT 8 GeForce 8XXX 9800GTX GT GT GT 16 GeForce 9XXX Unified Shader Tesla Architecture GTX GTX GTX GTX GeForce GT2XX C C Tesla Unified Shader CUDA, OpenCL, DirectX Compute etc Fermi Architecture FERMI Vertical axis shows the number of cores, but not the performance - Only major products are shown Fermi

15 GPU-CPU 性能及びメモリーバンド幅の差は拡大メモリーバンド幅ピーク値 GB/sec TF Single Precision 4GB Memory 8x double precision ECC L1, L2 Caches 100 NVIDIA GPU X86 CPU

16 目次なぜ今 GPU コンピューティングか? CPUの性能向上速度が減速性能向上 = 並列処理にGPUコンピューティングが応える CUDAシステムアーキテクチャによる超並列処理の実現普及 NVIDIA GPU アーキテクチャの進歩 3Dグラフィックスの進歩及び統合化シェーダエンジンへの変遷統合化シェーダエンジンへのCUDAの組込み次世代 GPU コンピューティング Fermi アーキテクチャまとめプロセッサ / メモリーアーキテクチャ GigaThreadスレッド制御機構による更なる高性能化統合化プログラム開発環境

GPUとして初めてECCを導入 Giga Thread L2 DRAM I/F L1 及び L2 キャッシュを内蔵約 2

17 Fermi アーキテクチャのご紹介スーパコンピュータの魂を持つ GPU DRAM I/F DRAM I/F 30 億個のトランジスターコア数を倍増 (512 コア ) HOST I/F DRAM I/F 倍精度浮動小数点演算ピーク性能が8 倍 GPUとして初めてECCを導入 Giga Thread L2 DRAM I/F L1 及び L2 キャッシュを内蔵約 2 倍のメモリーバンド幅 (GDDR5) DRAM I/F DRAM I/F 最大 1 Terabyte の GPU メモリ複数 Kernel 同時実行 C++ サポート

16 個の CUDA, 16 個の Load/Store Unit, 4 個の SFU に同時にディスパッチするシェアードメモリ L1 キャッシュとして使用する 64 KB RAM ( 構成可変 )

18 SM (Streaming Multiprocessor) アーキテクチャ Instruction Cache Scheduler Scheduler SM 当たり 32 CUDA コア ( 総数 512 個 ) Dispatch Dispatch Register File 倍精度浮動小数点演算ピーク性能は 8 倍単精度演算ピーク性能の 50% 2 個の Thread Scheduler 2 ワープを 2 組の 16 個の CUDA, 16 個の Load/Store Unit, 4 個の SFU に同時にディスパッチするシェアードメモリ L1 キャッシュとして使用する 64 KB RAM ( 構成可変 ) Load/Store Units x 16 Special Func Units x 4 Interconnect Network 64K Configurable Cache/Shared Mem Uniform Cache

CUDA コアアーキテクチャ Instruction Cache Scheduler Scheduler 最新の浮動小数点演算規格 IEEE 754-2008 準拠 ( 最新 CPU をも凌駕 ) Dispatch Dispatch Register File Fused multiply-add (FMA) 命令を倍精度単精度でサポート ( 積和の最終段でラウンディングするため CUDA

19 CUDA コアアーキテクチャ Instruction Cache Scheduler Scheduler 最新の浮動小数点演算規格 IEEE 準拠 ( 最新 CPU をも凌駕 ) Dispatch Dispatch Register File Fused multiply-add (FMA) 命令を倍精度単精度でサポート ( 積和の最終段でラウンディングするため CUDA Dispatch Port 各段で各々行うより精度が向上 ) Operand Collector 新規設計の全命令 32b サポート整数 ALU 64-bit 及びそれ以上の精度に対しても最適化設計 FP Unit INT Unit Result Queue Load/Store Units x 16 Special Func Units x 4 Interconnect Network 64K Configurable Cache/Shared Mem Uniform Cache

20 統合化された 64 ビットメモリー空間 C/C++ ポインターのフルサポート Non-unified Address Space Local *p_local Shared C 言語やC++ 言語のポインタはターゲットとするアドレス空間が必ずしもコンパイル時に確定せずランタイムに動的に決まるケースがあるためこのポインタを完全に実装することが困難な状況となっていましたしかし統合化されたアドレス空間をもつFermiではこの問題もなくなりました *p_shared Global 0 32-bit *p_global Unified Address Space Local Shared Global 0 40-bit *p

21 キャッシュメモリ階層構造オンチップシェアードメモリとともに本格的なキャッシュ階層構造を有する初めての GPU Fermi Memory Hierarchy Thread 各 SM (32 cores) に L1 キャッシュ (48KB または 16KB) メモリバンド幅の改善及びアクセス遅延の低減 Shared Memory L1 Cache ユニファイド L2 キャッシュ (768 KB) GPU 上の全コアにわたり高速にコーヒレンントデータをシェア L2 Cache M A R D

22 拡張され高速化されたメモリインタフェース GDDR5 メモリーインタフェース GDDR3 の 2 倍の速度最大 1 Terabyte の GPU メモリ大きなデータセットを取り扱うことが可能現在 40bアドレス命令は64bをサポート 384 ビットのメモリーインタフェース 64b partitioning x 6 DRAM I/F Giga Thread HOST I/F DRAM I/F L2 DRAM I/F DRAM I/F DRAM I/F DRAM I/F

23 ECC (Error Checking and Correction) DRAM の ECC によるデータエラーの除去 GDDR5 メモリーの ECC サポート内部の主要な記憶素子も ECC で保護レジスターファイル L1 キャッシュ L2 キャッシュ Single-Error Correct Double-Error Detect (SECDED) サポート 2 ビットの誤りはソフト処理 ( 再度実行等 )

24 IEEE 規格準拠の高精度演算 IEEE results 64-bit double precision 32-bit single precision full-speed denormal operands & results NaNs, +/- Infinity IEEE rounding nearest even, zero, +inf, -inf IEEE Fused Multiply-Add (FMA) D = A*B + C; No loss of precision IEEE divide & sqrt use FMA Multiply-Add (MAD): D = A*B + C; A B = Product + C = D (truncate digits) Fused Multiply-Add (FMA): D = A*B + C; A B = (retain all digits) Product + C = D (no loss of precision)

25 GigaThread TM Hardware Thread Scheduler (HTS) 階層的に何千ものアクティブなスレッドを管理コンテキストスイッチが 10 倍高速 HTS 複数 kernel の同時実行

26 GigaThread Hardware Thread Scheduler 複数 Kernel 同時実行 + 高速コンテキストスイッチ Kernel 1 Kernel 1 Kernel 2 Kernel 2 Kernel 2 Kernel 3 Ker 4 Time Kernel 2 Kernel 3 nel Kernel 5 Kernel 4 Kernel 5 Serial Kernel Execution Parallel Kernel Execution

27 GigaThread Streaming Data Transfer Engine デュアル DMA エンジン CPU GPU 及び GPU CPU データ転送の同時実行 CPU と GPU 演算と完全にオーバラップ可能処理の流れ : SDT Kernel 0 CPU SDT0 GPU SDT1 Kernel 1 CPU SDT0 GPU SDT1 Kernel 2 CPU SDT0 GPU SDT1 Kernel 3 CPU SDT0 GPU SDT1

28 各世代の比較 GPU G80 GT200 Fermi 集積トランジスタ数 6 億 8100 万個 14 億個 30 億個 CUDAコア数倍精度浮動小数点演算能力 30 FMA 演算 / クロック 256 FMA 演算 / クロック単精度浮動小数点演算能力 128 MAD 演算 / クロック 240 MAD 演算 / クロック 512 MAD 演算 / クロックワープスケジューラ (SMあたり搭載数) 特殊関数ユニット (SFU) /SM 共有メモリ /SM 16KB 16KB 48KB/16KB ( 構成可能 ) L1キャッシュ /SM 16KB/48KB ( 構成可能 ) L2キャッシュ /SM 768KB ECCメモリのサポート平行実行カーネル数最大 16 ロード / ストアのアドレス幅 32ビット 32ビット 64ビット

NVIDIA Nexus IDE( 統合化開発環境 ) 業界初の超並列アプリ開発用

29 NVIDIA Nexus IDE( 統合化開発環境 ) 業界初の超並列アプリ開発用 IDE (Integrated Development Environment) C, C++, OpenCL, DirectCompute 及び DirectX と OpenGL の両グラフィックス API をサポート完全にVisual Studioに組込まれた開発環境 (CPU + GPU) コプロセッシングアプリの開発を効率化両プロセッサにまたがるソースデバッグ性能解析両プロセッサにわたるイベントデータのキャプチャ

30 NVIDIA Nexus IDE( 統合化開発環境 ) 画面

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

32 Thank you for your attention

Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments

Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments 計算機アーキテクチャ第 11 回マルチプロセッサ本資料は授業用です無断で転載することを禁じます名古屋大学大学院情報科学研究科准教授加藤真平デスクトップジョブレベル並列性スーパーコンピュータ並列処理プログラムプログラムの並列化 for (i = 0; i < N; i++) { x[i] = a[i] + b[i]; } プログラムの並列化 x[0] = a[0] + b[0];