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

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

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

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

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

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

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

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

GPU プログラミング言語

GPU プログラム実行環境

DP GFLOPS per Watt GPU ロードマップ 32 Now 16 8 4 Kepler Dynamic Parallelism Maxwell Unified Memory 2 Fermi FP64 1 0.5 Tesla CUDA 2008 2010 2012 2014 2016

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

Unified Memory 現在 将来

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

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)

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

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

プログラマは並列化に注力 ( アーキテクチャ向け最適化はコンパイラが実施 ) 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

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

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

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

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

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

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

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, );

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

DP GFLOPS per Watt Kepler 32 16 8 4 Kepler Dynamic Parallelism Maxwell Unified Memory 2 Fermi FP64 1 0.5 Tesla CUDA 2008 2010 2012 2014 2016

性能と電力 : Fermi Kepler Fermi (M2090) Kepler (K20X) ピーク演算性能 (DP) 0.665 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

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

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

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

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

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

Shuffle の効果 (scan) スレッドブロック内 scan(prefix sum) 入力 0 1 2 3 出力 0 1 3 7 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

Parallel scan 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 0 1 2 3 4 5 6 7 8 10 12 14 16 18 20 22 0 1 2 3 4 6 8 10 12 15 18 21 24 28 32 36 0 1 2 4 6 9 12 16 20 25 30 36 42 49 56 64 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120

Parallel scan 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 Warp scan Warp scan Warp scan Warp scan 0 1 3 6 4 9 15 22 8 17 27 38 12 25 39 54 6 22 38 54 Warp scan 6 28 66 120 0 1 3 6 10 15 21 28 36 45 55 66 78 91 105 120

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

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

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

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

Time (ms) Atomic Operations 効果事例 6 5 4 3 2 1 0 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

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

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

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

GFLOPS RO Cache の効果事例 (Himeno BMT) 100 90 80 70 60 50 40 30 20 10 0 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

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

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

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

小カーネル同時実行テストコード 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

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

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

データ転送とカーネル実行のオーバーラップ 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

データ転送とカーネル実行のオーバーラップ パイプライン化 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

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

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

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

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

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

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();... }

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