CUDA 9 AND MORE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12

Transcription

1 CUDA 9 AND MORE 成瀬彰, シニアデベロッパーテクノロジーエンジニア, 2017/12/12

2 CUDA 9 の概要 VOLTA に対応 ライブラリの高速化 Tesla V100 Volta アーキテクチャ Tensor コア NVLink Independent スレッドスケジューリング cublas ( 主に DL 向け ) NPP ( 画像処理 ) cufft ( 信号処理 ) cusolver COOPERATIVE GROUPS 開発ツールの改善 柔軟なスレッドグループ並列アルゴリズムの抽象化スレッドブロック間の同期 (over SM or GPU) partition sync sync コンパイル時間の短縮 Unified Memory プロファイル NVLink 可視化コンパイラサポート 2

3 VOLTA 対応 3

4 TESLA V100 の概要 Volta アーキテクチャ NVLink と HBM2 の改善 Volta MPS SIMT モデルの改善 Tensor コア Most Productive GPU Efficient Bandwidth Inference Utilization New Algorithms 125 Programmable TFLOPS Deep Learning DL と HPC の両方に最適な GPU 4

5 VOLTA: 米国最大規模スパコンのエンジン Next Milestone In AI Supercomputing AI Exascale Today Performance Leadership 200 PF Accelerated Science ACME DIRAC FLASH GTC HACC LSDALTON NAMD 20 PF NUCCOR NWCHEM QMCPACK RAPTOR SPECFEM XGC 3+EFLOPS Tensor Ops 10X Perf Over Titan 5-10X Application Perf Over Titan 5

6 Relative to Tesla P100 エクサスケール (FP64) に向けて Volta: 米国最大規模スパコンのエンジン HPC ベンチマーク アプリ性能 (P100 V100) System Config Info: 2X Xeon E v4, 2.6GHz, w/ 1X Tesla P100 or V100. V100 measured on pre-production hardware. Summit Supercomputer 200+ PetaFlops ~3,400 Nodes 10 Megawatts 6

7 TENSOR コア混合精度行列計算ユニット 4x4 の行列の積和演算を 1 サイクルで計算する性能 (128 演算 / サクル ) 行列の FMA (Fused Multiply-Add) A 0,0 A 0,1 A 0,2 A 0,3 B 0,0 B 0,1 B 0,2 B 0,3 C 0,0 C 0,1 C 0,2 C 0,3 D = A 1,0 A 1,1 A 1,2 A 1,3 B 1,0 B 1,1 B 1,2 B 1,3 C 1,0 C 1,1 C 1,2 C 1,3 A 2,0 A 2,1 A 2,2 A 2,3 B 2,0 B 2,1 B 2,2 B 2,3 C 2,0 C 2,1 C 2,2 C 2,3 A 3,0 A 3,1 A 3,2 A 3,3 B 3,0 B 3,1 B 3,2 B 3,3 C 3,0 C 3,1 C 3,2 C 3,3 FP16 or FP32 FP16 FP16 FP16 or FP32 D = AB + C 7

8 TENSOR コアの使われ方 16x16 の行列の積和演算を ワープレベル (32 スレッド ) で協調して実行 Warp (32 スレッド ) 32 スレッドで同期 Tensor コアを使い 16x16 行列の行列積和演算を実行 32 スレッドで同期 8

9 TENSOR コアの使い方 device void tensor_op_16_16_16( float *d, half *a, half *b, float *c) { wmma::fragment<matrix_a, > Amat; wmma::fragment<matrix_b, > Bmat; wmma::fragment<matrix_c, > Cmat; wmma::load_matrix_sync(amat, a, 16); wmma::load_matrix_sync(bmat, b, 16); wmma::fill_fragment(cmat, 0.0f); wmma::mma_sync(cmat, Amat, Bmat, Cmat); NVIDIA cublas, cudnn, TensorRT Volta 向けに最適化されたフレームワーク ライブラリ } wmma::store_matrix_sync(d, Cmat, 16, wmma::row_major); CUDA C++ Warp レベル行列演算テンプレート 9

10 CUDA TENSOR コアプログラミング 16x16x16 Warp Matrix Multiply and Accumulate (WMMA) D = A B C FP16 or FP32 FP16 FP16 FP16 or FP32 D = AB + C 10

11 CUDA TENSOR コアプログラミング WMMA: 行列データ型 wmma::fragment<matrix_a, > Amat; A fragment Tensor コア用の行列データ型 各スレッドは 行列の要素の一部を 自分のレジスタに保持 ( 割当は未公開 ) ワープレベル (32 スレッド ) で 行列の全要素を保持すればよいという考え 従来のスレッド単位の行列演算と比べ レジスタ使用量を削減 11

12 CUDA TENSOR コアプログラミング WMMA: ロード命令 wmma::load_matrix_sync(amat, a, stride); A load_matrix_sync Tensor コア行列用のロード命令 ワープ単位で メモリ上の行列要素値を fragment データ型にロード warp 12

13 CUDA TENSOR コアプログラミング WMMA: 行列乗算 wmma::mma_sync(dmat, Amat, Bmat, Cmat); mma_sync Tensor コアを使用して 行列乗算を実行 D = A B C 13

14 CUDA TENSOR コアプログラミング WMMA: ストア命令 wmma::store_matrix_sync(d, Dmat, stride); D load_store_sync Tensor コア行列用のストア命令 ワープ単位で fragment データ型上の行列要素を メモリにストア warp 14

15 TENSOR コアの使い方 device void tensor_op_16_16_16( float *d, half *a, half *b, float *c) { wmma::fragment<matrix_a, > Amat; wmma::fragment<matrix_b, > Bmat; wmma::fragment<matrix_c, > Cmat; wmma::load_matrix_sync(amat, a, 16); wmma::load_matrix_sync(bmat, b, 16); wmma::fill_fragment(cmat, 0.0f); wmma::mma_sync(cmat, Amat, Bmat, Cmat); NVIDIA cublas, cudnn, TensorRT Volta 向けに最適化されたフレームワーク ライブラリ } wmma::store_matrix_sync(d, Cmat, 16, wmma::row_major); CUDA C++ Warp レベル行列演算テンプレート 15

16 PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S PC,S VOLTA INDEPENDENT THREAD SHCEDULING Pascal まで Program Counter (PC) and Stack (S) Warp(32 スレッド ) 毎に PC は 1 つ Volta Convergence Optimizer スレッド毎に PC が存在 個別にスケジューリングが可能 16

17 WARP 同期用ビルドイン関数 CUDA 9 で導入 スレッド同期 syncwarp アクティブなスレッド (PC の同じスレッド ) の取得 activemask スレッド間のデータ交換 all_sync, any_sync, uni_sync, ballot_sync shfl_sync, shfl_up_sync, shfl_down_sync, shfl_xor_sync match_any_sync, match_all_sync (*) 従来の shfl, ballot, any, all は CUDA 9 で deprecated 17

18 WARP 同期プログラミング 想定通りに Warp(32 スレッド ) が同期する保証はない if (threadidx.x < 4) { A; } else { B; } /* 32 スレッド同期実行 */ ここで Warp 内の 32 スレッドが同期している保証はない 分岐のあるプログラムでは Warp は分離 集合して命令実行 if (threadidx.x < 8) { X; } else { Y; } 暗黙の Warp 同期を前提としたプログラミングは危険 特に Volta から (Volta 以前も 安全ではなかった ) 時間 18

19 暗黙的な WARP 同期プログラミング Warp に関して 以下のことを仮定している 1. スレッドは再集合する 2. スレッドはロックステップ実行する if (threadidx.x < 16) A; else B; assert( activemask() = 0xffffffff); if ( activemask() == 0xffffffff) { assert( activemask() = 0xffffffff); } 19

20 暗黙的な WARP 同期プログラミング Warp に関して 以下のことを仮定している 1. スレッドは再集合する 2. スレッドはロックステップ実行する if (threadidx.x < 16) A; else B; assert( activemask() = 0xffffffff); if ( activemask() == 0xffffffff) { assert( activemask() = 0xffffffff); } どちらも True になる保証はない 明示的な Warp 同期が必要そのため CUDA 9 で Warp 同期の Build-in 関数を追加 改変 20

21 例 : 共有メモリを使用したワープ内スレッド間 REDUCTION shmem[tid] = v; v += shmem[tid+16]; shmem[tid] = v; v += shmem[tid+8 shmem[tid] = v; v += shmem[tid+4 shmem[tid] = v v += shmem[tid+2 shmem[tid] = v; v += shmem[tid+1 shmem[tid] = v; これも 安全ではない shmem[tid] = v; v += shmem[tid+16]; shmem[tid] = v; v += shmem[tid+8]; shmem[tid] = v; v += shmem[tid+4]; shmem[tid] = v; v += shmem[tid+2]; shmem[tid] = v; v += shmem[tid+1]; shmem[tid] = v; syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); syncwarp(); 21

22 ライブラリの改善 22

23 CUDA 9: ライブラリの改善 VOLTA 対応 Tensorコアの活用 cublas: Voltaに最適化したGEMMs 全ライブラリ : すぐにVoltaを性能を発揮 スピード cublas: RNNs 向けGEMM 最適化 NPP: 画像処理の高速化 cufft: 様々なサイズのFFT 最適化 DEEP LEARNING Scientific Computing 新アルゴリズム cusolver: マルチ GPU 向け密行列 疎行列ソルバー 密行列固有値解析 nvgraph: 幅優先探索 (BFS) クラスタリング Triangle-Counting グラフ挿入 抽出 インストール CUDA ライブラリだけのパッケージ (without CUDA driver, runtime, etc.) NPP: モジュラー化 23

24 Relative Performance Relative Performance FP32 cublas: GEMMS 性能改善 Volta Tensor コア + CUDA 9 P100 (CUDA 8) V100 (CUDA 9) Mixed Precision (FP16 Input, FP32 compute) P100 (CUDA 8) V100 Tensor Cores (CUDA 9) x x Matrix Size (M=N=K) Matrix Size (M=N=K) 24

25 cublas: cublasgemmex() アルゴリズム選択が可能 (CUDA 8 から ) A B A B A B 18 種類のアルゴリズムから選択可能 CUBLAS_GEMM_ALGO[0:17] CUBLAS_GEMM_DFALT: 自動選択 Tensor コア 3 種類のアルゴリズムから選択可能 CUBLAS_GEMM_ALGO[0:2]_TENSOR_OP CUBLAS_GEMM_DFALT_TENSOR_OP: 自動選択 25

26 Speed up Vs. CUDA 8* CUFFT, NPP cufft NPP CUDA 8 と比べて最大 2 倍の高速化 1D, 2D, 3D IPP と比べて最大 100 倍の性能イメージプロセッシング コンピュータビジョン 1D 2D 3D 2.5x Morphological Ops. 2.0x JPEG 1.5x 1.0x 0.5x 0.0x Geometry Transforms Filters Color Proc. 0x 50x 100x Data Size Speedup Vs. IPP** * V100 and CUDA 9 (r384); Intel Xeon Broadwell, dual socket, E v4@ 2.6GHz, 3.5GHz Turbo with Ubuntu x86_64 with 128GB System Memory * P100 and CUDA 8 (r361); For cublas CUDA 8 (r361): Intel Xeon Haswell, single-socket, 16-core E v3@ 2.3GHz, 3.6GHz Turbo with CentOS 7.2 x86-64 with 128GB System Memory ** CPU system running IPP: Intel Xeon Haswell single-socket 16-core E v3@ 2.3GHz, 3.6GHz Turbo Ubuntu x86_64 with 128GB System Memory 26

27 cusolver: ヤコビ法ベースの固有値ソルバー QR 法と比べて計算量は増えるが並列性が高い 行列サイズ 128~256 までは MKL より高速 QR Jacobi MKL CPU:E5-2690v2, GPU:K40 27

28 cusolver: ヤコビ法ベースの固有値ソルバー バッチ実行 ( 各行列のサイズ :32x32) QR 法と比べて計算量は増えるが並列性が高い CPU:E5-2690v2, GPU:K40 28

29 CUTLASS: FAST LINER ALGEBRA IN CUDA C++ (version 0.1) ユーザの CUDA カーネルから使用できる 高性能な行列積 C++ テンプレート DL アプリの多くは 行列積の組み合わせ CUDA の様々な階層で利用可能 デバイスレベル ブロックレベル ワープレベル スレッドレベル 29

30 CUTLASS の性能 ( 対 CUBLAS) cublas と遜色ない性能を CUDA C++ レベルで実現 データ型 : FP16, FP32, FP64, INT Tensor コア対応 行列データレイアウト : NN, NT, TN, TT 30

31 COOPERATIVE GROUPS 31

32 COOPERATIVE GROUPS スケーラブルで柔軟性の高い スレッド間同期 通信機構 協調動作するスレッドグループの 定義 分割 同期を容易にする スケーラブルなグループサイズ : 数スレッド ~ 全スレッド Thread Block Group 動的なグループの生成 分割が可能 CUDA としてサポート グループサイズにより適切なハードウェアを選択 分割後の Thread Groups Kepler 世代以後の GPU で利用可能 * Note: Multi-Block and Mult-Device Cooperative Groups are only supported on Pascal and above GPUs 32

33 多様なスレッド間同期を簡単に 3 つのスケール スレッドブロック内 シングル GPU 内 (SM 間の同期 ) マルチ GPU 間 (GPU 間の同期 ) 協調動作するスレッドグループを動的に生成し 各グループで同期 スレッドブロック間の同期 partition sync sync sync sync 33

34 カーネル内でのスレッド同期 CUDA 8 まで 34

35 小さいグループ カーネル内でのスレッド同期 CUDA 9 から スレッドブロック 大きいグループ 35

36 COOPERATIVE グループ 5 種類のグループ Coalesced Group Thread Group Thread Block Tile Thread Block グループのメソッド size() スレッド数 thread_rank() スレッドのID sync() スレッド間同期 Grid Group Multi-Grid Group 36

37 COOPERATIVE グループ Thread Block から Thread Block Tile( サブグループ ) を生成 thread_group block = this_thread_block(); this_thread_block() は 自 Thread Block に対応 Thread Block block.sync(); synchthreads() と等価 thread_group tile32 = tiled_partition(block, 32); Thread Block から 32 スレッドのグループを作成 分割後の Thread Groups tile32.sync(); サブグループ内の 32 スレッド間で 同期 thread_group tile4 = tiled_partition(tile32, 4); 再帰的なサブブループ作成も可能 (*) Tile サイズは 32 以下 かつ 2^N に制限 (CUDA 9.0) 37

38 COOPERATIVE グループ同じデバイス関数を サイズの異なるグループで共用できる Thread Block (~1024 スレッド ) Warp (32 スレッド ) g = this_thread_block(); val = reduce(g, shmem, myval); g = tiled_partition(this_thread_block(), 32); val = reduce(g, shmem, myval); device int reduce(thread_group g, int *shmem, int val) { int myrank = g.thread_rank(); for (int i = g.size()/2; i > 0; i /= 2) { shmem[myrank] = val; g.sync(); } val += shmem[myrank ^ i]; g.sync(); } return val; 並列 reduction ( 共有メモリ使用 ) 38

39 THREAD BLOCK TILE ワープ内スレッド間通信 Build-in 関数を使える.shfl().shfl_down().shfl_up().shfl_xor().any().all().ballot().match_any().match_all() template <unsigned size> device int reduce(thread_block_tile<size> g, int val) { for (int i = g.size()/2; i > 0; i /= 2) { val += g.shfl_xor(val, i); } } return val; 並列 reduction (shfl_xor 使用 ) コンパイル時にサイズが分かると高速 thread_group_tile<32> tile32 = tiled_partition<32>(this_thread_block()); thread_group_tile<4> tile4 = tiled_partition<4>(this_thread_block()); 39

40 COALESCED GROUP 同時に同じパスを実行しているスレッドのグループ 40

41 COALESCED GROUP 並列 Array Push ( サイズ不定 ) Y head a b c x y z t0 t7 t0 t7 Thread Block 0 Thread Block 1 41

42 COALESCED GROUP 並列 Array Push ( サイズ不定 ) Y a b c x y z head a b c x y z t0 t7 t0 t7 Thread Block 0 Thread Block 1 42

43 COALESCED GROUP Atomic Aggregation device int atomicagginc(int *head_ptr) { coalesced_group g = coalesced_threads(); int old_head; if (g.thread_rank() == 0) { old_head = atomicadd(head_ptr, g.size()) } int my_head = g.shfl(old_head, 0) + g.thread_rank(); return my_head; } 43

44 COALESCED GROUP 並列 Array Push ( サイズ不定 ) Y head a b c t0 Thread Block 0 t7 44

45 COALESCED GROUP 並列 Array Push ( サイズ不定 ) Y a b c my_head head a b c t0 Thread Block 0 t7 my_head = atomicagginc(head); 45

46 ATOMIC AGGREGATION Build-In 関数でも実装は可能 Cooperative Groups coalesced_group g = coalesced_threads(); int ret; if (g.thread_rank() == 0) { ret = atomicadd(ptr, g.size()) } ret = g.shfl(ret, 0); return ret + g. thread_rank(); Build-In Functions int mask = activemask(); int rank = popc(mask & lanemask_lt()); int leader_lane = ffs(mask) 1; int ret; If (rank == 0) { ret = atomicadd(p, popc(mask)); } ret = shfl_sync(mask, ret, leader_lane); return ret + rank; 記述しやすいのは どちらか? 46

47 専用 API でカーネル起動 GRID GROUP グリッド ( シングル GPU) 内の 全スレッドのグループ global kernel() { grid_group grid = this_grid(); while (...) {... grid.sync(); } } cudalaunchcooperativekernel(...); 起動したカーネルの全スレッドが 同時に active になる必要あり (Persistent Kernel) cudaoccupancymaxactiveblockspermultiprocessor(&numblockspersm, kernel, numthreads, 0); 47

48 MULTI GRID GROUP マルチグリッド ( マルチ GPU) 内の 全スレッドのグループ global kernel() { multi_grid_group multi_grid = this_multi_grid(); while (...) {... grid.sync(); } } 48

49 MULTI GRID GROUP マルチグリッド ( マルチ GPU) 内の 全スレッドのグループ 専用 API でカーネル起動 struct cudalaunchparams params[numdevices]; for (int i = 0; i < numdevices; i++) { params[i].func = (void*) kernel; params[i].griddim = dim3(...); params[i].blockdim = dim3(...); params[i].sharedmem =...; params[i].stream =...; params[i].args =...; } cudalaunchcooperativekernelmultidevice(params, numdevices); 49

50 例 : 粒子シミュレーション Cooperative Groups 無し // threads update particles in parallel integrate<<<blocks, threads, 0, stream>>>(particles);

51 例 : 粒子シミュレーション Cooperative Groups 無し // threads update particles in parallel integrate<<<blocks, threads, 0, s>>>(particles); // Collide each particle with others in neighborhood collide<<<blocks, threads, 0, s>>>(particles); (*) 粒子の位置が移動したら CUDA スレッドへの粒子のマッピングを変えたほうが 高速に処理できる 51

52 例 : 粒子シミュレーション Cooperative Groups 無し // threads update particles in parallel integrate<<<blocks, threads, 0, s>>>(particles); // ここで暗黙的に同期しているので マッピング変更が可能 // Collide each particle with others in neighborhood collide<<<blocks, threads, 0, s>>>(particles); (*) 粒子の位置が移動したら CUDA スレッドへの粒子のマッピングを変えたほうが 高速に処理できる 52

53 GRID GROUP で粒子シミュレーション 2 種類の処理を シングルカーネルで実行 global void particlesim(particle *p, int N) { grid_group g = this_grid(); for (i = g.thread_rank(); i < N; i += g.size()) integrate(p[i]); g.sync() // GPU 全体の同期 } for (i = g.thread_rank(); i < N; i += g.size()) collide(p[i], p, N); cudalaunchcooperativekernel( ) でカーネル起動

54 MULTI-GRID GROUP で粒子シミュレーションシングルカーネルで 大規模な問題をマルチ GPU 実行 global void particlesim(particle *p, int N) { multi_grid_group g = this_multi_grid(); for (i = g.thread_rank(); i < N; i += g.size()) integrate(p[i]); g.sync() // マルチ GPU の全てで同期 } for (i = g.thread_rank(); i < N; i += g.size()) collide(p[i], p, N); cudalaunchcooperativekernelmultidevice( ) で起動 54

55 ロードマップ : COOPERATIVE GROUPS より柔軟なグループ作成 任意ラベルによる グループの分割 (Volta 限定 ) // 計算結果が同じスレッドのグループ int label = foo() % 4; thread_group block = partition(this_thread_block(), label); (*) ランダムなグループは SIMT 実行効率が低下するので 注意が必要 32 より大きなタイル thread_group g = tiled_partition(this_thread_block(), 64); 55

56 ロードマップ : COOPERATIVE GROUPS Collective アルゴリズムのライブラリ Reductions, sorting, prefix sum (scan), 等など. // collective key-value sort using all threads in the block cooperative_groups::sort(this_thread_block(), myvalues, mykeys); // collective scan-based allocate across block int sz = myallocationsize(); // amount each thread wants int offset = cooperative_groups::exclusive_scan(this_thread_block(), sz); 56

57 開発ツール 57

58 多様な開発ツール 58

59 CUDA-MEMCHECK Cooperative Groups 対応 安全ではない Warp 同期プログラミングの検出 (racecheck) 59

60 NVVP: UNIFIED MEMORY プロファイリング CPUページフォールトの発生箇所とソースコードとの対応付け Page Fault Correlation 60

61 NVVP: UNIFIED MEMORY イベントの追加 仮想メモリ関連の挙動の可視化 Memory Thrashing Page Throttling Remote Map 61

62 NVVP: NVLINK トポロジー NVLINK の各リンクの利用率 DGX1 62

63 CUDA 9 の概要 VOLTA に対応 ライブラリの高速化 Tesla V100 Volta アーキテクチャ Tensor コア NVLink Independent スレッドスケジューリング cublas ( 主に DL 向け ) NPP ( 画像処理 ) cufft ( 信号処理 ) cusolver COOPERATIVE GROUPS 開発ツールの改善 柔軟なスレッドグループ並列アルゴリズムの抽象化スレッドブロック間の同期 (over SM or GPU) partition sync sync コンパイル時間の短縮 Unified Memory プロファイル NVLink 可視化コンパイラサポート 63

64 CUDA

65