GPUコンピューティングの現状と未来

Transcription

1 GPU コンピューティングの現状と未来 成瀬彰, HPC Developer Technology, NVIDIA

2 Summary 我々のゴールと方向性 ゴール実現に向けて進めている技術開発 Unified Memory, OpenACC Libraries, GPU Direct Kepler の機能紹介 Warp shuffle, Memory system Hyper-Q, Dynamic Parallelism

3 Our Goals 電力効率 プログラミング簡易化ポータビリティ 多数アプリをカバー

4 様々な分野で使われる GPU コンピューティング

5 CUDA パラレル コンピューティング プラットフォーム

6 CUDA パラレル コンピューティング プラットフォーム

7 CUDA パラレル コンピューティング プラットフォーム

8 GPU プログラミング言語

9 GPU プログラム実行環境

10 DP GFLOPS per Watt GPU ロードマップ 32 Now Kepler Dynamic Parallelism Maxwell Unified Memory 2 Fermi FP Tesla CUDA

11 Tesla K40 メモリ容量より多くのアプリ CPU クロック電力状況により適切なクロックを選択 6GB 流体解析 地震波解析 レンダリング GPU Boost 12GB

12 Unified Memory 現在 将来

13 Unified Memory void sortfile(file *fp, int N) { char *data = (char*)malloc(n); char *sorted = (char*)malloc(n); fread(data, 1, N, fp); } CPU code cpu_sort(sorted, data, N); use_data(sorted); free(data); free(sorted) GPU code void sortfile(file *fp, int N) { char *data = (char*)malloc(n); char *sorted = (char*)malloc(n); fread(data, 1, N, fp); char *d_data, *d_sorted; cudamalloc(&d_data, N); cudamalloc(&d_sorted, N); cudamemcpy(d_data, data, N, ); gpu_sort<<< >>>(d_sorted, d_data, N); cudamemcpy(sorted, d_sorted, N, ); cudafree(d_data); cudafree(d_sorted); use_data(sorted); free(data); free(sorted) }

14 Unified Memory CPU code void sortfile(file *fp, int N) { char *data = (char*)malloc(n); char *sorted = (char*)malloc(n); fread(data, 1, N, fp); GPU code (UVM) void sortfile(file *fp, int N) { char *data = (char*)malloc(n); char *sorted = (char*)malloc(n); fread(data, 1, N, fp); cpu_sort(sorted, data, N); gpu_sort<<< >>>(sorted, data, N); } use_data(sorted); free(data); free(sorted) } use_data(sorted); free(data); free(sorted)

15 OpenACC: ディレクティブ CPU GPU Program myscience... serial code...!$acc kernels do k = 1,n1 do i = 1,n2... parallel code... enddo enddo!$acc end kernels... End Program myscience オリジナル Fotrran/C コード OpenACC Compiler Hint シンプル : ディレクティブ挿入 パワフル : 少ない労力 コンパイラが並列化 オープン : 多数のベンダのアクセラレータをサポート NVIDIA, AMD, (soon) Intel

16 OpenACC の特徴 親しみやすいプログラミングモデル 多分野への応用 ヘテロジニアスアーキテクチャ オープンスタンダード X86 and ARM AMD, Intel, NVIDIA

17 プログラマは並列化に注力 ( アーキテクチャ向け最適化はコンパイラが実施 ) OpenACC によるアプリ高速化事例事例 (ORNL and Tokyo Tech) (dual-cpu nods vs. CPU+GPU) S3D Combustion NICAM Weather/Climate Tuned top 3 kernels for GPUs (90% of runtime) End result: 2.2X faster with K20X vs. dual AMD node Kepler is 6X faster than Fermi Improved performance of CPU-only version by 50% Tuned top kernels using CUDA, then OpenACC CUDA result: 3.1x faster on GPU vs. CPU node OpenACC result (preliminary ) = 69-77% of CUDA More portable, more maintainable Full OpenACC port in progress Results from Cray/ORNL and Tokyo Tech

18 OpenACC 対応状況 Geology Weather/Climate/ Ocean Plasma & Combustion Fluid Dynamics / Cosmology Quantum Chemistry AWP-ODC CAM-SE Cloverleaf CHIMERA CASTEP EMGS ELAN COSMO Physics GENE PMH bv DELPASS GAMESS CCSD(T) *Seismic CPML* FIM GTC DNS GAUSSIAN SPECFM3D GEOS-5 LULESH MiniGHOST MiniMD Harmonie S3D RAMSES ONETEP HBM UPACS Quantum Espresso ICON NICAM NEMO GYRE NIM PALM-GPU ROMS WRF X-ECHO

19 Speed-up CUBLAS: 逆行列計算 LAPACK 準拠 API cublas<t>getrfbatched() LU 分解 cublas<t>getribatched() 逆行列計算 多数の小サイズ行列用 (*) 行列サイズ :64, 行列数 :1000 CPU-1core (2.8GHz,MKL) GPU (K20,naïve) GPU (K20,cublas)

20 NVIDIA GPUDirect データ移動を最適化する技術ファミリー GPUDirect Shared GPU and System memory ノード内のメモリコピー負荷を削減 GPUDirect Peer-to-Peer ノード内の別 GPUのメモリを直接アクセスノード内のGPU-to-GPUメモリ転送を加速 GPUDirect RDMA ノード間で GPU-to-GPU RDMA 通信

21 1 2 1 GPUDirect Shared GPU and System Memory Without GPUDirect GPU writes to pinned main memory 1 CPU copies main memory 1 to main memory 2 Network driver reads main memory 2 With GPUDirect GPU writes to pinned main memory Network driver reads from main memory CPU Main Mem CPU Main Mem Chip set GPU Chip set GPU Network GPU Memory Network GPU Memory

22 GPUDirect (Peer-to-Peer) ホストメモリを仲介せずにデータ移動 ( ノード内 ) Fermi 以上

23 GPUDirect RDMA ホストメモリを仲介せずにデータ移動 ( ノード間 ) Kepler 以上 System Memory GDDR5 Memory GDDR5 Memory GDDR5 Memory GDDR5 Memory System Memory CPU GPU1 GPU2 GPU2 GPU1 CPU Server 1 PCI-e Network Card Network Network Card PCI-e Server 2

24 GPU-aware MPI MPI 関数だけで GPU-to-GPU 通信を可能に MPI_Send(), MPI_Recv() にデバイスメモリを指定を可能に通信処理の最適化からプログラマを解放 パイプライン転送 (Device Host, Host Host, Host Device) 送信 : Without GPU-aware MPI cudamemcpy( s_buf, s_device, size, ); MPI_Send( s_buf, size, ); 受信 : MPI_Recv( r_buf, size, ); cudamemcpy( r_device, r_buf, size, ); With GPU-aware MPI 送信 : MPI_Send( s_device, size, ); 受信 : MPI_Recv( r_device, size, );

25 GPU-aware MPI Libraries GPU メモリからの送信 受信 多くの集合通信に対応 利用可能な最も良い転送方式を選択 MVAPICH Open MPI IBM Platform Computing Computing IBM Platform MPI Versions: MVAPICH2 1.9 OpenMPI IBM Platform MPI V9.1.2 (Free Community Edition) Reference NVIDIA GPUDirect Technology Overview

26 DP GFLOPS per Watt Kepler Kepler Dynamic Parallelism Maxwell Unified Memory 2 Fermi FP Tesla CUDA

27 性能と電力 : Fermi Kepler Fermi (M2090) Kepler (K20X) ピーク演算性能 (DP) TFLOPS 1.31 TFLOPS ピーク演算性能 (SP) 1.33 TFLOPS 3.95 TFLOPS 最大メモリバンド幅 177 GB/s 250 GB/s TDP 225 Watt 235 Watt x2 x3 x1.4 x1 27

28 Fermi から Kepler へ Kepler Fermi 最大 3 倍の電力あたり性能 28

29 SM(Fermi) SMX(Kepler) Fermi (M2090) Kepler (K20X) CUDA コア コア周波数 1.3GHz 0.73GHz 最大スレッド数 最大スレッドブロック数 ビットレジスタ数 32 K 64 K L2 容量 0.75 MB 1.5 MB x6 x0.6 x1.3 x2 x2 x2 29

30 Kepler で強化された機能 Warp Shuffle Memory System Atomics Operations Read-only Cache Hyper-Q Concurrency Overlapping Dynamic Parallelism 30

31 Warp Shuffle 他スレッドのレジスタの読み出しを可能に 対象 : 同一ワープ内のスレッド (32 スレッド ) 共有メモリ不要のスレッド間データ交換 syncthreads() も不要に Kepler 世代 (CC 3.0 以上 ) から利用可能 31

32 4 種類の関数 idx, up, down, xor a b c d e f g h shfl() shfl_up() shfl_down() shfl_xor() h d f e a c c b a b a b c d e f c d e f g h g h c d a b g h e f Indexed any-to-any Shift right/up to n th neighbour Shift left/down to n th neighbour Butterfly (XOR) exchange

33 Shuffle の効果 (scan) スレッドブロック内 scan(prefix sum) 入力 出力 SMEM (Shared Memory) for (ofst = 1; ofst < BLOCK_SIZE; ofst *= 2) { if (idx >= ofst) smem[idx] += smem[idx - ofst]; syncthreads(); } 4 11 SHFL (Shuffle) for (ofst = 1; ofst < WARP_SIZE; ofst *= 2) { if (idx >= ofst) val += shfl_up(val,ofst,warp_size); } if (idx % WARP_SIZE == WARP_SIZE - 1) smem[idx/warp_size] = val; syncthreads(); if (idx < NUM_WARP) { sum = smem[idx]; for (ofst = 1; ofst < NUM_WARP; ofst *= 2) { if (idx >= ofst) sum += shfl_up(sum,ofst,num_warp); } smem[idx] = sum; } syncthreads(); if (idx/warp_size > 0) val += smem[idx/warp_size - 1]; 33

34 Parallel scan

35 Parallel scan Warp scan Warp scan Warp scan Warp scan Warp scan

36 Time (ms) Shuffle の効果事例 (scan) SMEM Scan (fp32) x3.0 SHFL 3 倍の性能 UP Tesla K20 グリッド形状 (26, 1, 1) ブロック形状 (1024, 1, 1) 1000 回実行 4,096B smem per block 128B 36

37 Shuffle の効果 (reduction) スレッドブロック内 reduction reduction コード例 ( ワープ内 ) SMEM (Shared Memory) SHFL (Shuffle) idx = threadidx.x; for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) { if (idx < mask) smem[idx] += smem[idx ^ mask]; syncthreads(); } for (mask = WARP_SIZE/2 ; mask > 0 ; mask >>= 1) { var = shfl_xor( var, mask, WARP_SIZE ); } 37

38 Time (ms) Shuffle の効果事例 (reduction) Reduction within TB (fp32) x 倍の性能 UP SMEM SHFL Tesla K20 ブロック形状 (1024, 1, 1) グリッド形状 (26, 1, 1) 1000 回実行 4,096B smem per block 128B 38

39 Atomic Operations サポートタイプ データ型の拡張 グローバルメモリ上の Atomic 操作を高速化 複数カーネルに分離していた処理を単一カーネルで 効果確認 16M 要素 reduction データ型は float smem[idx] = input[g_idx]; for (mask = BLOCK_SIZE/2; mask > 0; mask /= 2) { if (idx < mask) { smem[idx] += smem[idx ^ mask]; } syncthreads(); } if (idx == 0) { atomicadd( output, smem[idx] ); } 39

40 Time (ms) Atomic Operations 効果事例 Reduction (Sum, SP, 16M elements) x2.1 x3.7 Fermi(C2075) Kepler(K20) Kepler(K20) with SHFL Fermi から Kepler で 2.1 倍の性能 UP Shuffle 命令併用で 3.7 倍の性能 UP ブロック形状 (1024, 1, 1) ECC off 40

41 Read-Only(RO) Cache SM SMEM Threads L1 Read TEX only TEX Texture API CUDA Arrays 一般的な Read-Only キャッシュとして使用可能 L2 cache Kepler 以降 コンパイラに指示 DRAM 41

42 2 つの使い方 組み込み関数 : ldg() global kernel( int* output, int* input ) {... output[idx] =... + input[idx ldg( &input[idx + delta] + + delta]...; ) +...;... } 型修飾子 : const restrict global kernel( int* output, int* const input int* ) restrict input ) {... output[idx] =... + input[idx + delta] +...;... } 42

43 RO Cache の効果 Himeno BMT 19 ポイント ステンシル テストコード 共有メモリを使用せずに CUDA 化 jacobi_kernel(..., float* p,... ); jacobi_kernel(..., const float* strict p,... ); 43

44 GFLOPS RO Cache の効果事例 (Himeno BMT) Without RO cache (Fermi: C2075) GFLOPS (Himeno BMT) Without RO cache (Kepler: K20) 25% With RO cache (Kepler: K20) 25% 性能 UP Himeno BMT 問題サイズ : L ブロック形状 (128, 2, 1) ECC off 44

45 Hyper-Q Stream Queue Mgmt より多くのカーネルを同時実行可能に Stream Queue Mgmt C R B Q A P Z Y X C B A R Q P Z Y X CUDA Generated Work Grid Management Unit Pending & Suspended Grids 1000s of pending grids Work Distributor Work Distributor 16 active grids 32 active grids SM SM SM SM SMX SMX SMX SMX Fermi Kepler 45

46 Without Hyper-Q (Fermi) Stream 1 Kernel A, B, C A B C P Q R X Y Z Single Hardware Work Queue Stream 2 Stream 3 Kernel P, Q, R Kernel X, Y, Z 最多 16 同時実行 制限 : 同時実行できるのはストリーム端のカーネル 46

47 With Hyper-Q (Kepler) A B C P Q R X Y Z Stream 1 Stream 2 Stream 3 Kernel A, B, C Kernel P, Q, R Kernel X, Y, Z Multiple Hardware Work Queue 最多 32 同時実行 ( 偽の ) ストリーム依存性から開放 47

48 小カーネル同時実行テストコード cudastream_t stream[nstreams]; for (i = 0 ; i < nstreams ; ++i) { // ストリーム生成 cudacreatestream( &stream[i] ); } dim3 gdim( 1, 1, 1 ); dim3 bdim( 1024, 1, 1 ); for (i = 0 ; i < nstreams ; ++i) { // カーネル 1 を投入 kernel_1<<<gdim, bdim, 0, stream[i]>>>(... ); // カーネル 2 を投入 kernel_2<<<gdim, bdim, 0, stream[i]>>>(... ); } // カーネル 3 を投入 kernel_3<<<gdim, bdim, 0, stream[i]>>>(... ); 48

49 小カーネル同時実行テスト (Fermi) 部分的に同時実行 シングルハードウェアキューの制約 Tesla C

50 小カーネル同時実行テスト (Kepler) Tesla K20 全カーネル ( ストリーム ) を同時実行 これまでより簡単に同時実行が可能に 50

51 データ転送とカーネル実行のオーバーラップ 3つの処理をオーバーラップ可能 データ転送 (Host to Device) カーネル実行データ転送 (Device to Host) cudamemcpy( a_dev, a_host, all, cudamemcpyhosttodevice ); kernel_1<<<gdim, bdim>>>( c_dev, a_dev, all ); cudamemcpy( b_dev, b_host, all, cudamemcpyhosttodevice ); kernel_2<<<gdim, bdim>>>( c_dev, b_dev, all ); cudamemcpy( c_host, c_dev, all, cudamemcpydevicetohost ); 51

52 データ転送とカーネル実行のオーバーラップ パイプライン化 cudastream_t stream[nstreams]; for (s = 0 ; s < nstreams ; ++s) { cudacreatestream( &stream[s] ); } s = 0; for (p = 0 ; p < npipeline; ++p) { cudamemcpyasync( a_dev[p], a_host[p], part, cudamemcpyhosttodevice, stream[s] ); kernel_1<<<gdim, bdim, 0, stream[s]>>>( c_dev[p], a_dev[p], part ); cudamemcpyasync( b_dev[p], b_host[p], part, cudamemcpyhosttodevice, stream[s] ); kernel_2<<<gdim, bdim, 0, stream[s]>>>( c_dev[p], b_dev[p], part ); cudamemcpyasync( c_host[p], c_dev[p], part, cudamemcpydevicetohost, stream[s] ); s = (s+1) % nstreams; } パイプラインコードも Hyper-Q で効率化 52

53 オーバーラップ実行テスト (Fermi) パイプライン化前 データ転送とカーネル実行のオーバーラップ無し Tesla C

54 オーバーラップ実行テスト (Fermi) パイプライン後 Tesla C2075 データ転送とカーネル実行 相応の時間でオーバーラップ カーネル実行の間に隙間 54

55 オーバーラップ実行テスト (Kepler) 完全にオーバーラップ Tesla K20 カーネル実行の間に空き無し Hyper-Q の効果 55

56 Dynamic Parallelism とは? GPU からカーネルを起動する仕組み Dynamically 実行時のデータ値に基づくカーネル起動 Simultaneously 複数スレッドから同時に起動 Independently スレッド毎に独自グリッドで起動 CPU GPU CPU GPU Fermi Kepler

57 Dynamic Parallelism の動作イメージ CPU GPU CPU GPU CPU がきめ細かく制御 GPU が自律的にに動作

58 Dynamic Parallelism コードサンプル void rec_func(... ) { }... rec_func(... ); rec_func(... );... global void rec_func(... ) {... if ( blockidx.x == 0 ) { cudastreamcreate( &st0 ); cudastreamcreate( &st1 ); rec_func<<<..., st0 >>>(... ); rec_func<<<..., st1 >>>(... ); cudadevicesynchronize(); cudastreamdestroy( st0 ); cudastreamdestroy( st1 ); } syncthreads();... }

59 まとめ 我々のゴール 電力効率, プログラミング簡易化, 多数アプリその達成のために開発している技術 Unified Memory, OpenACC Libraries, GPU Direct, GPU-aware MPI Kepler の機能 Warp shuffle, Memory system Hyper-Q, Dynamic Parallelism