CUDA高速化セミナーvol.2 ~CUDAアーキテクチャの進化~

2022年6月23日に開催された「CUDA高速化セミナーvol.2　～CUDAアーキテクチャの進化～」の当日資料です。

Transcript

1. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDA 高速化セミナー vol.2 CUDAアーキテクチャの進化
2. Fixstars Group www.fixstars.com Copyright © Fixstars Group 発表者紹介 2 • 冨田 明彦（とみた あきひこ） ソリューションカンパニー 営業企画執行役 2008年に入社。金融、医療業界において、 ソフトウェア高速化業務に携わる。その 後、新規事業企画、半導体業界の事業を 担当し、現職。 • 平櫛 貴章（ひらぐし たかあき） Fixstars Autonomous Technologies エグゼクティブエンジニア 2015年に新卒で入社。幅広い産業領域で CPU / GPU を用いたパフォーマンスチュ ーニング業務に携わる。
3. Fixstars Group www.fixstars.com Copyright © Fixstars Group 本日のAgenda フィックスターズの紹介 (15分) • 会社紹介 • 高速化のためにCUDAアーキテクチャの進化を知る CUDAアーキテクチャの進化（60分） • SM アーキテクチャ • Tensor Core • Cooperative Groups • CUDA Graphs Q&A / 告知 3
4. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group フィックスターズのご紹介
5. Fixstars Group www.fixstars.com Copyright © Fixstars Group フィックスターズの強み コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団 低レイヤ ソフトウェア技術 アルゴリズム 実装力 各産業・研究 分野の知見 5
6. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、 最適化やアルゴリズムの改良を施して高速化してお返しします 当社 お客様 オリジナルソースコードのご提供 高速化したソースコード コンサルティング 高速化 サポート 先行技術調査 性能評価 ボトルネックの特定 アルゴリズムの改良・開発 ハードウェアへの最適化 レポート作成 レポートやコードへのQ&A 実製品への組込み支援 6
7. Fixstars Group www.fixstars.com Copyright © Fixstars Group ソフトウェア高速化サービス 様々な領域でソフトウェア高速化サービスを提供しています 大量データの高速処理は、お客様の製品競争力の源泉となっています ・NAND型フラッシュメモリ向けファー ムウェア開発 ・次世代AIチップ向け開発環境基盤開発 Semiconductor ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 Finance ・自動運転の高性能化、実用化 ・次世代パーソナルモビリティの研究開発 Mobility ・ゲノム解析の高速化 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Life Science ・Smart Factory化支援 ・マシンビジョンシステムの高速化 Industrial 7
8. Fixstars Group www.fixstars.com Copyright © Fixstars Group 画像処理・アルゴリズム開発サービスの例 出展：https://www.cs.toronto.edu/~frossard/post/vgg16/ • お客様の課題 • 高度な画像処理や深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件まで クリアできるか不安 • 製品化に結びつくような研究ができていない • 弊社の支援内容 • 課題に応じたアルゴリズム調査 • 深層学習ネットワーク精度改善、推論高速化手法調査 • 論文調査、実装 8
9. Fixstars Group www.fixstars.com Copyright © Fixstars Group AI・深層学習関連サービス 9 • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績 ネットワーク設計 データの前処理、データ拡張 精度改善 分散処理による学習高速化 各種DLフレームワーク クラウド・サーバ エッジ モデル圧縮 - 量子化 - 枝刈り - 蒸留 ターゲットデバイスへの ポーティング及び推論高速化 ▪ ARM, GPU, DSP ▪ SIMD,NEON,CUDA,TensorRT
10. Fixstars Group www.fixstars.com Copyright © Fixstars Group GPU向け高速化サービス 10 • お客様の課題 • GPU 高速化の知見がない • 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • GPU 高速化に関するコンサルティング • ボトルネック調査、GPU プログラムの高速化 • CPU/GPU が混在するヘテロジニアス環境での最適化 10～150 倍の 高速化事例あり
11. Fixstars Group www.fixstars.com Copyright © Fixstars Group 高速化のためにCUDAアーキテクチャの進化を知る 11 1. GPU搭載の新製品へ機能を追加し 処理を高速化したい 2. 高速化をやりきったか？ 3. CUDAアーキテクチャの進化を 知っておこう！ • 新アーキテクチャの GPU を導入 • CUDA Toolkit を最新にアップグレード • もっと性能出せないか？ • もっと可読性を上げられないか？ 11
12. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDAの進化
13. Fixstars Group www.fixstars.com Copyright © Fixstars Group 今回の概要 • Volta世代以降のCUDAの変更点のおさらい • CUDA Toolkit 9.0 以降の変更点のおさらい • 主に計算カーネルを記述する人向け 13
14. Fixstars Group www.fixstars.com Copyright © Fixstars Group Voltaアーキテクチャ • Compute Capability 7.0, 7.2 • NVIDIA V100, TITAN V, Jetson Xavier など • 主な新機能 • 深層学習用アクセラレータ: Tensor Core • L1キャッシュとシェアードメモリの拡張 • 浮動小数点数演算と整数演算の同時発行 • Independent Thread Scheduling 14
15. Fixstars Group www.fixstars.com Copyright © Fixstars Group Turingアーキテクチャ • Compute Capability 7.5 • GeForce RTX 20xx, NVIDIA T4 など • ゲーミング・グラフィックス向けのアーキテクチャ • 主な新機能 • レイトレーシング用アクセラレータ: RT Core • 第2世代 Tensor Core 15
16. Fixstars Group www.fixstars.com Copyright © Fixstars Group Ampereアーキテクチャ (GA100) • データセンター向けAmpere • Compute Capability 8.0 • NVIDIA A100 など • 主な新機能 • 第3世代 Tensor Core • 非同期コピー・バリア • タスクグラフの高速化 16
17. Fixstars Group www.fixstars.com Copyright © Fixstars Group Ampereアーキテクチャ (GA102) • デスクトップ向けAmpere • Compute Capability 8.6 • GeForce RTX 30xx など • 主な新機能 • 第2世代 RT Core • 浮動小数点数演算2つの同時発行 17
18. Fixstars Group www.fixstars.com Copyright © Fixstars Group 今日の内容 • SMアーキテクチャの変化 • 整数演算コアの追加 • Independent Thread Scheduling • スレッド間通信機能の強化 • Tensor Core • Cooperative Groups • CUDA Graphs 18
19. Fixstars Group www.fixstars.com Copyright © Fixstars Group 注意事項 • この資料で言及している内容は CUDA 11.7 時点のものです • 今後のバージョンアップで変化することがあります 19
20. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group SMアーキテクチャの変化 20
21. Fixstars Group www.fixstars.com Copyright © Fixstars Group 整数演算コアの追加 • Volta以降 • CUDAコアがFP32コアとINT32コアに分けられた • 非本質的な計算によるリソース圧迫を防ぐ • 非本質的な計算: アドレス計算など 21 https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf
22. Fixstars Group www.fixstars.com Copyright © Fixstars Group アドレス計算 • CUDAコアのアドレッシングモードはあまり多機能ではない • レジスタ上のアドレス+定数オフセットのみ • うっかりするとアドレス計算命令の方が多くなる • INTコアで浮動小数点数演算ユニットの圧迫を回避できる • とはいえまだ注意は必要 • 浮動小数点数演算の倍以上あるとFP32コアが余る • IMADはFPコアで処理される 22
23. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実験: INT32 Core を動かしてみる • FFMAとADD, LOP3を交互に動かす • ASUS RTX 2080Ti TURBO で測定 • Theoretical Peak: 17103.4 [GFLOPS] • 足し合わせると理論ピークを越える • INTコアの効果が出ている 23 FFMA R9, R9, R10.reuse, 3 ; IADD3 R4, R4, 0x2, RZ ; FFMA R7, R7, R10.reuse, 3 ; IADD3 R2, R2, 0x2, RZ ; FFMA R5, R5, R10.reuse, 3 ; LOP3.LUT R8, R8, 0x3, RZ, 0x3c, !PT ; FFMA R3, R3, R10.reuse, 3 ; LOP3.LUT R6, R6, 0x3, RZ, 0x3c, !PT ; FP32 [GFLOPS] 15794.9 INT32 [GOPS] 7897.5
24. Fixstars Group www.fixstars.com Copyright © Fixstars Group 実験: INT32 Core を動かしてみる • スレッドブロック番号の偶奇で分岐 • FFMAのみをひたすら発行するスレッドブロック • ADD, LOP3のみをひたすら発行するスレッドブロック • 演算の総量は前の例と同じ • 同様の傾向となる • FP32とINT32で異なるスレッド・ワープの命令を供給できる • スーパースカラというよりはSMT的な振る舞い 24
25. Fixstars Group www.fixstars.com Copyright © Fixstars Group 浮動小数点数演算2つの同時発行 • Ampere (GA102) • INT32コアが浮動小数点数も扱えるように • FP32のスループットが2倍に • またアドレス計算まわりのチューニングが必要に 25 https://images.nvidia.com/aem-dam/en-zz/Solutions/geforce/ampere/pdf/NVIDIA- ampere-GA102-GPU-Architecture-Whitepaper-V1.pdf
26. Fixstars Group www.fixstars.com Copyright © Fixstars Group L1キャッシュの効率改善 • Volta • L1キャッシュが速くなった • シェアードメモリより少し遅い程度 • シェアードメモリをキャッシュ代わりにしなくて良いケースが増える • const __restrict__ をつける（コンパイラが自動でつけることもある） 26
27. Fixstars Group www.fixstars.com Copyright © Fixstars Group Independent Thread Scheduling • Volta • スレッドごとにPCとスタックを持つようになった • 分岐時のスケジューリングの選択肢が増える • 分岐したスレッド間での通信が可能となる • レイテンシの隠蔽にも有利かもしれない 27
28. Fixstars Group www.fixstars.com Copyright © Fixstars Group Independent Thread Scheduling • Pascal以前 • 同一ワープ内の全スレッドがPCを共有 • ある命令の処理は必ず同時に行われる 28 if(threadIdx.x < 8){ A(); B(); }else{ C(); D(); } E(); A() B() C() D() E()
29. Fixstars Group www.fixstars.com Copyright © Fixstars Group Independent Thread Scheduling • Volta以降 • スレッドごとにPCを持つ • 1つ以上のスレッドのPCが指している命令が発行される 29 if(threadIdx.x < 8){ A(); B(); }else{ C(); D(); } E(); A() B() C() D() E()
30. Fixstars Group www.fixstars.com Copyright © Fixstars Group 分岐したスレッド間の通信 • Pascal以前の場合 • スレッド0側のループから抜けられない • スレッド1側の処理が進まない 30 __shared__ int x; if (threadIdx.x == 0) { x = 1; do { // __nanosleep(0); } while(x); } else { do { // __nanosleep(0); } while(!x); x = 0; }
31. Fixstars Group www.fixstars.com Copyright © Fixstars Group 分岐したスレッド間の通信 • Volta以降の場合 • __nanosleepで他スレッドに処理を譲る • スレッド1側の処理を進めることができる 31 __shared__ int x; if (threadIdx.x == 0) { x = 1; do { __nanosleep(0); } while(x); } else { do { __nanosleep(0); } while(!x); x = 0; }
32. Fixstars Group www.fixstars.com Copyright © Fixstars Group ワープ内での同期 • Volta以降 • コード上のある地点に同時に到達しない場合がある 32 A() B() C() D() E() A() B() C() D() E() E() 期待される処理順の例 - 全スレッドの E() が同時に発行される 発生しうる処理順の例 - 全スレッドの E() が同時に発行されるとは限らない
33. Fixstars Group www.fixstars.com Copyright © Fixstars Group ワープ内での同期 • __syncwarp() で明示的に同期をとる • マスクで指定したスレッドが到達するのを待つ • それ以降は再度分岐するまで同期された状態で進む 33
34. Fixstars Group www.fixstars.com Copyright © Fixstars Group 挙動が変わるケースの例 • 例: コンパイル時に反復回数がわからないループ 34 __global__ void kernel(const int *input){ __shared__ volatile int smem[32]; for(int i = 0; i < input[threadIdx.x]; ++i){ smem[threadIdx.x] += 1; } __syncwarp(); printf("%08x¥n", __activemask()); } {0, 1, 2, …, 30, 31} を与える __activemask(): アクティブなスレッドの取得 ffffffff ffffffff ffffffff ffffffff … ffffffff ffffffff ffffffff ffffffff 000003ff 000003ff 000003ff 000003ff … 44000000 44000000 88000000 88000000 __syncwarp() あり 全スレッドが同時に到達 __syncwarp() なし 合流したりしなかったり
35. Fixstars Group www.fixstars.com Copyright © Fixstars Group ワープ内通信 • 他スレッドと通信を行う場合明示的な同期を挿入する • ワープ内通信を伴う組み込み関数は _sync がつく • __all_sync, __shfl_sync など • 同期をとる必要があるスレッド集合をビットマスクで指定 • CUDA Toolkit 9.0 以降 _sync なしは Deprecated • 使える命令が少し増えました • Match (Volta) • Reduce (Ampere) 35
36. Fixstars Group www.fixstars.com Copyright © Fixstars Group Warp Match: match_any • valueに同じ値を渡したスレッドの集合を取得 36 0 0 1 0 1 1 0 0 0x00cb 0x00cb 0x0034 0x00cb 0x0034 0x0034 0x00cb 0x00cb 0x00cb => 0, 1, 3, 6, 7 0x0034 => 2, 4, 5 unsigned int __match_any_sync(unsigned mask, T value);
37. Fixstars Group www.fixstars.com Copyright © Fixstars Group Warp Match: match_all • 集合中のスレッド全てがvalueに同じ値を渡したかの判定 37 unsigned int __match_all_sync(unsigned mask, T value, int *pred); 0 1 1 0 0 0 0 0 0x000f 0x000f 0x000f 0x000f 0x00f0 0x00f0 0x00f0 0x00f0 mask value 0x0000 0x0000 0x0000 0x0000 0x00f0 0x00f0 0x00f0 0x00f0 result 0 0 0 0 1 1 1 1 *pred
38. Fixstars Group www.fixstars.com Copyright © Fixstars Group Reduce • 整数の総和や最小値・最大値などを求める 38 unsigned __reduce_add_sync(unsigned mask, unsigned value); unsigned __reduce_min_sync(unsigned mask, unsigned value); unsigned __reduce_max_sync(unsigned mask, unsigned value); int __reduce_add_sync(unsigned mask, int value); int __reduce_min_sync(unsigned mask, int value); int __reduce_max_sync(unsigned mask, int value); unsigned __reduce_and_sync(unsigned mask, unsigned value); unsigned __reduce_or_sync(unsigned mask, unsigned value); unsigned __reduce_xor_sync(unsigned mask, unsigned value);
39. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Barrier • Five Stages of Synchronization • 同期を arrive と wait の2つに分割する • Arrive • その地点への到達を示す • 特に何かを待たずに次の処理へ進む • Waitでの待ち受けに使用するトークンを返す • Wait • スレッド群がトークンに対応するArriveへ到達するまで待つ • 他スレッドのArrive以前のメモリ操作が観測可能であることが保証される 39
40. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Barrier barrier::arrive • 処理が到達したことを他スレッドに通知 • この時点ではブロックされない • 処理 (2) へ直ちに進む 40 barrier::arrive 処理 (1) 処理 (2) barrier::wait 処理 (3)
41. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Barrier barrier::wait • 全スレッドがarriveに到達するまで待つ • 他スレッドは処理 (2) の途中かもしれない • 処理 (1) での変更が観測可能になる • 処理 (2) での変更については保証なし 41 barrier::arrive 処理 (1) 処理 (2) barrier::wait 処理 (3)
42. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Barrier: 性能評価 • 220回同期するだけのカーネルの処理時間 • 256スレッド, 1スレッドブロック • NVIDIA A10G • Async: barrier.arrive_and_wait() • HWアクセラレーションあり: -arch sm_86 • HWアクセラレーションなし: -arch sm_75 • Sync: __syncthreads() 42 Async (w/ HW accel.) Async (w/o HW accel.) Sync Time [ms] 127.4 577.8 23.8
43. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Data Copy • メモリコピーの非同期実行を可能にする • ありがちなパターン • グローバルメモリからシェアードメモリにコピー • シェアードメモリ上のデータを使って計算 • これを繰り返す 43 Global Memory Shared Memory Register Shared Memory Global Memory Shared Memory ……
44. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Data Copy • 実際にはこうなっている • 一時レジスタ（と場合によってはL1キャッシュ）は省略できそう • Ampereで可能になった Global Memory L2 Cache L1 Cache Register Shared Memory Register Shared Memory Register ……
45. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Data Copy • 実際にはこうなっている • 一時レジスタ（と場合によってはL1キャッシュ）は省略できそう • Ampereで可能になった Global Memory L2 Cache L1 Cache Register Shared Memory Register Shared Memory Register ……
46. Fixstars Group www.fixstars.com Copyright © Fixstars Group Asynchronous Data Copy • 演算とデータ転送のオーバーラップ • ロードされたデータを使用する前に同期をとる • 他にもバリアとの組み合わせなどいろいろな書き方がある auto block = cg::this_thread_block(); __shared__ float smem[BLOCK_SIZE]; // 非同期データ転送の開始 cg::memcpy_async(block, smem, src, sizeof(float) * BLOCK_SIZE); // 何らかのsmemに依存しない処理 … // 非同期データ転送の完了待ち cg::wait(block); // データ転送完了: smemにデータがロードされている …
47. Fixstars Group www.fixstars.com Copyright © Fixstars Group まとめ • 徐々にできることが増えている • 演算器の増加 • より柔軟なスケジューリング • スレッド間通信の高機能化 • 移行に伴い気をつけるべき点もある • ワープ内の同期 • ボトルネックとなる点の変化 • 新機能=高速とは限らない 47
48. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group Tensor Core 48
49. Fixstars Group www.fixstars.com Copyright © Fixstars Group Tensor Core • 深層学習向けのアクセラレータ • 小さい行列の積を効率よく求める • 入力の精度によって性能が変わる 49 https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf
50. Fixstars Group www.fixstars.com Copyright © Fixstars Group Tensor Core を用いた演算の流れ • Tensor Core はワープ単位で協調して使用する • 入力行列をレジスタにロードする • 行列積を求める • 出力行列をメモリにストアする 50
51. Fixstars Group www.fixstars.com Copyright © Fixstars Group WMMA API • WMMA: Warp Matrix Multiply and Accumulate • ワープ単位で協調して行列積を行うためのAPI • 実質的に Tensor Core を利用するためのAPI • 利用可能な行列サイズの制約 • 実ハードウェアのそれよりやや大きい • Tensor Core の世代差 (命令当たりの演算量) を吸収する • 小さい行列を処理する命令を複数回呼んで少し大きい行列を処理 • 逆は効率の低下につながってしまう 51
52. Fixstars Group www.fixstars.com Copyright © Fixstars Group WMMA API の利用 • fragment: レジスタ上に配置される一時バッファ • 型情報として用途などの情報を含む 52 __device__ void kernel(float *D, const half *A, const half *B, const float *C){ wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> A_frag; wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> B_frag; wmma::fragment<wmma::accumulator, 16, 16, 16, float> C_frag; wmma::load_matrix_sync(A_frag, A, 16); wmma::load_matrix_sync(B_frag, B, 16); wmma::load_matrix_sync(C_frag, C, 16, wmma::mem_row_major); wmma::mma_sync(C_frag, A_frag, B_frag, C_frag); wmma::store_matrix_sync(D, C_frag, 16, wmma::mem_row_major); }
53. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: PTX • WMMA API の呼び出しと対応する • 一塊で 16x16x16 だけ処理している 53 wmma.load.a.sync.aligned.row.m16n16k16.global.f16 {%r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9}, [%rd7], %r1; wmma.load.b.sync.aligned.row.m16n16k16.global.f16 {%r10, %r11, %r12, %r13, %r14, %r15, %r16, %r17}, [%rd8], %r1; wmma.load.c.sync.aligned.row.m16n16k16.global.f32 {%f1, %f2, %f3, %f4, %f5, %f6, %f7, %f8}, [%rd6], %r1; wmma.mma.sync.aligned.row.row.m16n16k16.f32.f32 {%f9, %f10, %f11, %f12, %f13, %f14, %f15, %f16}, {%r2, %r3, %r4, %r5, %r6, %r7, %r8, %r9}, {%r10, %r11, %r12, %r13, %r14, %r15, %r16, %r17}, {%f1, %f2, %f3, %f4, %f5, %f6, %f7, %f8}; wmma.store.d.sync.aligned.row.m16n16k16.global.f32 [%rd5], {%f9, %f10, %f11, %f12, %f13, %f14, %f15, %f16}, %r1;
54. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 行列のロード部分 • 8x8部分行列ごとに2要素ロード 54 /* load_matrix_sync(A_frag) */ LDG.E.SYS R18, [R12] ; LDG.E.SYS R19, [R12+0x100] ; LDG.E.SYS R20, [R12+0x10] ; LDG.E.SYS R21, [R12+0x110] ; /* load_matrix_sync(C_frag) */ LDG.E.64.SYS R8, [R16] ; LDG.E.64.SYS R10, [R16+0x200] ; LDG.E.64.SYS R4, [R16+0x20] ; LDG.E.64.SYS R6, [R16+0x220] ;
55. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 行列のロード部分 • Bのロードは8x8部分行列内での転置が入る • Bが列優先であれば転置は省略される 55 /* load_matrix_sync(B_frag) */ LDG.E.SYS R0, [R14] ; LDG.E.SYS R23, [R14+0x10] ; LDG.E.SYS R22, [R14+0x100] ; LDG.E.SYS R24, [R14+0x110] ; MOVM.16.MT88 R0, R0 ; MOVM.16.MT88 R23, R23 ; MOVM.16.MT88 R22, R22 ; MOVM.16.MT88 R24, R24 ;
56. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される 56 HMMA.1688.F32 R8, R18, R0, R8 ; HMMA.1688.F32 R4, R18, R23, R4 ; HMMA.1688.F32 R8, R20, R22, R8 ; HMMA.1688.F32 R28, R20, R24, R4 ; R8,9 R10,11 R4,5 R6,7 R18 R19 R20 R21 R0 R22 R23 R24
57. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される 57 HMMA.1688.F32 R8, R18, R0, R8 ; HMMA.1688.F32 R4, R18, R23, R4 ; HMMA.1688.F32 R8, R20, R22, R8 ; HMMA.1688.F32 R28, R20, R24, R4 ; R8,9 R10,11 R4,5 R6,7 R18 R19 R20 R21 R0 R22 R23 R24
58. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される 58 HMMA.1688.F32 R8, R18, R0, R8 ; HMMA.1688.F32 R4, R18, R23, R4 ; HMMA.1688.F32 R8, R20, R22, R8 ; HMMA.1688.F32 R28, R20, R24, R4 ; R8,9 R10,11 R4,5 R6,7 R18 R19 R20 R21 R0 R22 R23 R24
59. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される 59 HMMA.1688.F32 R8, R18, R0, R8 ; HMMA.1688.F32 R4, R18, R23, R4 ; HMMA.1688.F32 R8, R20, R22, R8 ; HMMA.1688.F32 R28, R20, R24, R4 ; R8,9 R10,11 R4,5 R6,7 R18 R19 R20 R21 R0 R22 R23 R24
60. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_86) • 積和演算部分 • 2回の16x8x16行列積に分割される 60 HMMA.16816.F32 R8, R4.reuse, R24, R8 ; HMMA.16816.F32 R12, R4, R26, R12 ; R8,9 R10,11 R12,13 R14,15 R4 R5 R6 R7 R24 R25 R26 R27
61. Fixstars Group www.fixstars.com Copyright © Fixstars Group 出力コードを読む: SASS (sm_86) • 積和演算部分 • 2回の16x8x16行列積に分割される 61 HMMA.16816.F32 R8, R4.reuse, R24, R8 ; HMMA.16816.F32 R12, R4, R26, R12 ; R8,9 R10,11 R12,13 R14,15 R4 R5 R6 R7 R24 R25 R26 R27
62. Fixstars Group www.fixstars.com Copyright © Fixstars Group 浮動小数点数の表現 • 深層学習周辺ではIEEE754以外の形式が用いられることがある • 指数部と仮数部に用いるビット数が異なる • 第3世代 Tensor Core で TF32, BF16 を追加サポート 62 符号部 指数部 仮数部 FP32: IEEE754 Single 1 8 23 TF32: TensorFloat-32 1 8 10 FP16: IEEE754 Half 1 5 10 BF16: BFloat16 1 8 7
63. Fixstars Group www.fixstars.com Copyright © Fixstars Group 入出力の精度と速度 • 対応している精度と速度は世代によって異なる • FP32コアによるFFMAとの性能比: 63 入力 出力 1st Gen. 2nd Gen. 3rd Gen. FP16 FP16 8x 8x 16x FP16 FP32 8x 8x 16x INT8 INT32 N/A 16x 32x INT4 INT32 N/A 32x 64x TF32 FP32 N/A N/A 8x BF16 FP32 N/A N/A 16x Binary INT32 N/A N/A 256x FP64 FP64 N/A N/A 1x
64. Fixstars Group www.fixstars.com Copyright © Fixstars Group Structured Sparsity • Ampere (GA100) で追加 • ニューラルネットの重みは寄与度の低い要素を多く含む • これらの要素を0とみなして演算量を削減する: pruning 64 Han et al. (2015) Learning both Weights and Connections for Efficient Neural Networks
65. Fixstars Group www.fixstars.com Copyright © Fixstars Group Structured Sparsity • 制約付きでハードウェアによるサポートを行う • 2:4 Sparsity: 連続する4要素のうち2要素が0である • 行と列の内積を求めるときに0との積和をスキップする • 必要な積和の回数が半減する⇒実効性能が倍になる 65
66. Fixstars Group www.fixstars.com Copyright © Fixstars Group Structured Sparsity • WMMA API ではまだ提供されていない? • ドキュメント中にそれらしい言及がない • ヘッダファイルにもそれらしい定義は見当たらない • CUTLASSではPTXを直書きしている 66 asm volatile( "mma.sp.sync.aligned.m16n8k32.row.col.f16.f16.f16.f16 {%0,%1}, " "{%2,%3,%4,%5}, {%6,%7,%8,%9}, {%10,%11}, %12, 0x0;¥n" : "=r"(D[0]), "=r"(D[1]) : "r"(A[0]), "r"(A[1]), "r"(A[2]), "r"(A[3]), "r"(B[0]), "r"(B[1]), "r"(B[2]), "r"(B[3]), "r"(C[0]), "r"(C[1]), "r"(E)); https://github.com/NVIDIA/cutlass/blob/v2.9.0/include/cutlass/arch/mma_sparse_sm80.h
67. Fixstars Group www.fixstars.com Copyright © Fixstars Group まとめ • Tensor Core: 行列演算用アクセラレータ • ワープ単位で協調して小さい行列同士の積を求める • 世代と精度によって性能が大きく変化する • CUDA C++ からは WMMA API 経由で操作する • 実ハードより少し粗い単位で処理を行う • 一部の処理はまだ提供されていない 67
68. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group Cooperative Groups 68
69. Fixstars Group www.fixstars.com Copyright © Fixstars Group Cooperative Groups • 複数スレッドで協調する処理を記述するためのライブラリ • グリッド・スレッドブロック・ワープ…… • CUDA Toolkit 9.0 で導入された • CUDAの標準ライブラリとして提供される • 実は従来から頑張れば似たようなことは出来た • 黒魔術的な実装で将来の変更におびえなくてよくなる 69
70. Fixstars Group www.fixstars.com Copyright © Fixstars Group スレッドブロック • cooperative_groups::thread_block • できること • スレッドブロックサイズの取得 (dim3もしくは1次元化されたもの) • スレッドブロック内における自身のスレッド番号の取得 • グリッド内における自身のスレッドブロック番号の取得 • スレッドブロック内の全スレッドでの同期 70
71. Fixstars Group www.fixstars.com Copyright © Fixstars Group 使用例 71 __global__ void block_kernel(){ cg::thread_block g = cg::this_thread_block(); // ブロックを構成するスレッド数の取得 (1次元) // blockDim.x * blockDim.y * blockDim.z unsigned int size = g.num_threads(); // ブロックを構成するスレッド数の取得 (3次元, blockDim) dim3 block_dim = g.dim_threads(); // 自身がブロック内で何番目のスレッドかの取得 (3次元, threadIdx) dim3 thread_idx = g.thread_index(); // 自身がブロック内で何番目のスレッドかの取得 (1次元) // threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y unsigned int rank = g.thread_rank(); // このスレッドブロックがグリッド内で何番目のブロックかの取得 (3次元, blockIdx) dim3 block_idx = g.group_index(); // ブロック内での同期 (__syncthreads()) g.sync(); }
72. Fixstars Group www.fixstars.com Copyright © Fixstars Group グリッド • cooperative_groups::grid_group • できること • グリッド内における自身のスレッド番号の取得 • グリッド内における自身のスレッドブロック番号の取得 • グリッド内のスレッド数とスレッドブロック数の取得 • グリッド内の全スレッドでの同期 • カーネル呼び出し時に特定の条件を満たす必要がある 72
73. Fixstars Group www.fixstars.com Copyright © Fixstars Group グリッド内での同期 • 全てのスレッドが同時に起動している必要がある • Occupancyでスレッド数の上限が決まる • cudaLaunchCooperativeKernel でカーネルを起動する 73
74. Fixstars Group www.fixstars.com Copyright © Fixstars Group タイル (Thread Block Tile) • スレッドブロックをさらに分割したもの • 実際はワープサイズ以下の2冪個のスレッドをまとめたもの • スレッド数はコンパイル時定数である必要がある • できること • 自身のインデックスの取得 • グループ内での同期 • 各種ワープ内通信 (シャッフル・ボーティング・マッチング) 74
75. Fixstars Group www.fixstars.com Copyright © Fixstars Group Coalesced Group • 同一ワープ内の任意のスレッド群からなるグループ • 典型的にはアクティブなスレッド群から構成する • タイルでできることとほとんど同じことができる • 処理によっては追加のコストがかかることがある 75 cg::coalesced_group active = cg::coalesced_threads();
76. Fixstars Group www.fixstars.com Copyright © Fixstars Group グループの分割 tiled_partition • スレッドブロックまたはタイルをより細かいタイルに分割 labeled_partition, binary_partition • ラベルとして同じ値を渡したスレッドからなるグループを生成 76 cg::thread_block block = cg::this_thread_block(); auto tile32 = cg::tiled_partition<32>(block); auto tile16 = cg::tiled_partition<16>(tile32); cg::thread_block block = cg::this_thread_block(); auto group = cg::labeled_partition(block, block.thread_rank() % 4);
77. Fixstars Group www.fixstars.com Copyright © Fixstars Group Data Manipulation • グループ内のスレッド同士で協調して計算を行う • 現状ではワープ内 (タイルと Coalesced Group) のみ • アーキテクチャごとに適切な実装が選択される • できること • Reduce • Scan 77
78. Fixstars Group www.fixstars.com Copyright © Fixstars Group Reduce • 𝑦 = 𝑥0 op 𝑥1 op … op 𝑥𝑛−1 • 𝑖番目のスレッドからの入力を𝑥𝑖 とする • opは結合的な2項演算 • 総和の計算などに用いる • 可能なら __reduce_**_sync が使用される • Ampere (sm_80) 以降のハードウェアである • op が特定の演算 (cg::plus<int> など) である 78
79. Fixstars Group www.fixstars.com Copyright © Fixstars Group Reduce: HWアクセラレーションなし • タイルの場合 • __shfl_xor を用いたコードが出力される • Kepler以降で使われていたものと同様 • スレッドあたり O(log n) 79
80. Fixstars Group www.fixstars.com Copyright © Fixstars Group Reduce: HWアクセラレーションなし • Coalesced Group の場合 • もう少しややこしいコードが出てくる • 取得元のレーン番号を二分探索で求める • 最後にブロードキャスト • スレッドあたり O(log^2 n) 80
81. Fixstars Group www.fixstars.com Copyright © Fixstars Group Reduce: HWアクセラレーションなし • Coalesced Group の場合 • 単純なタイルより遅くなる • mask=0xffffffffの時は別コードが使用されるため例外 81 実装 処理時間 [ms] thread_block_tile<32> 194.2 coalesced_group (mask=0x7fffffff) 991.0 coalesced_group (mask=0xffffffff) 142.4 reduceをスレッドあたり220回実行 GeForce RTX 2080 Ti (sm_75) gdim=68×4, bdim=256
82. Fixstars Group www.fixstars.com Copyright © Fixstars Group Scan • 𝑦𝑖 = 𝑥0 op 𝑥1 op … op 𝑥𝑖 (inclusive_scan) • 𝑦𝑖 = 𝑥0 op 𝑥1 op … op 𝑥𝑖−1 (exclusive_scan) • 特にハードウェアアクセラレーションはない • アクセラレーションなしのreduceと似たようなコードになる 82
83. Fixstars Group www.fixstars.com Copyright © Fixstars Group まとめ • Cooperative Groups • スレッド間の協調をポータブルに記述できる • グループの粒度の差異を吸収する • アーキテクチャ・ツールキットの変化を吸収する 83
84. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group CUDA Graphs 84
85. Fixstars Group www.fixstars.com Copyright © Fixstars Group CUDA Graphs • 多くのジョブは複数のタスクの組み合わせによって構成される • カーネル呼び出し • データ転送 • ホスト側処理 • …… • タスク間の依存性をグラフで表現する • その一連の処理をまとめて実行する 85
86. Fixstars Group www.fixstars.com Copyright © Fixstars Group CUDA Graphs のメリット • タスク投入に起因するオーバーヘッドを低減できる • 細かいタスクを多く含む場合に特に有効 • Ampereだとハードウェアアクセラレーションもある 86 CPU GPU CPU GPU A B C D Launch A Launch B Launch C Launch D A B C D Launch Graph CUDA Graphs なし CUDA Graphs あり
87. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの構成要素 • ノード • カーネル起動 • CPU処理の実行 • メモリ確保・解放・コピー • イベント処理 (Record, Wait) • セマフォ操作 (Signal, Wait) • 子グラフの実行 • ノード間の依存関係 87
88. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフのライフサイクル • 構築したグラフを使いまわすことで性能を稼ぐ 88 グラフの構築 インスタンス化 実行 更新
89. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの構築 • Graph API による構築 • カーネルパラメータの渡し方は cudaLaunchKernel() と似た形 89 cudaGraph_t graph; cudaGraphCreate(&graph, 0); // カーネルに渡すパラメータ (後述) cudaKernelNodeParams params1, params2; // ノードの生成 cudaGraphNode_t node1, node2; cudaGraphAddKernelNode(&node1, graph, nullptr, 0, &params1); cudaGraphAddKernelNode(&node2, graph, nullptr, 0, &params2); // ノード間の依存関係の定義 cudaGraphAddDependencies(graph, &node1, &node2, 1);
90. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの構築 • Graph API による構築 • カーネルパラメータの渡し方は cudaLaunchKernel() と似た形 90 __global__ void kernel(int x){ printf("%d¥n", x); } int kernelParam1 = 42; void *kernelParams[] = { static_cast<void*>(&kernelParam1) }; cudaKernelNodeParams params; params.func = reinterpret_cast<void*>(&kernel); params.gridDim = dim3(1, 1, 1); params.blockDim = dim3(1, 1, 1); params.sharedMemBytes = 0; params.kernelParams = kernelParams; params.extra = nullptr;
91. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの構築 • ストリームキャプチャによる構築 • ストリームに対する操作から自動的にグラフを構築する • イベントの Record/Wait で他ストリームもキャプチャ対象になる 91 cudaGraph_t graph; cudaGraphCreate(&graph, 0); cudaStream_t stream; cudaStreamCreate(&stream); cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal); kernel<<<1, 1, 0, stream>>>(42); kernel<<<1, 1, 0, stream>>>(0); cudaStreamEndCapture(stream, &graph);
92. Fixstars Group www.fixstars.com Copyright © Fixstars Group インスタンス化 • 実行可能なグラフオブジェクトを生成する • 最適化・エラーチェックなどもこのタイミングで行われる 92 // グラフの構築 cudaGraph_t graph; …… // グラフのインスタンス化 cudaGraphExec_t graphExec; cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0);
93. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの実行 • 指定したストリームに対してグラフを構成するタスク群を投入 93 cudaStream_t stream; cudaGraphExec_t graphExec; cudaGraphLaunch(graphExec, stream);
94. Fixstars Group www.fixstars.com Copyright © Fixstars Group グラフの更新 • 軽微な変更は再インスタンス化せずに行うことができる • 操作対象となるメモリアドレスの書き換えなど • インスタンス化に起因するコストを抑える • 2通りの方法がある • Whole graph update • Individual node update 94
95. Fixstars Group www.fixstars.com Copyright © Fixstars Group Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる 95 A B src = nullptr dst = nullptr コピー先 （インスタンス化済み） A B src = in_ptr dst = out_ptr コピー元 （インスタンス化前）
96. Fixstars Group www.fixstars.com Copyright © Fixstars Group Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる 96 A B src = in_ptr dst = out_ptr コピー先 （インスタンス化済み） A B src = in_ptr dst = out_ptr コピー元 （インスタンス化前）
97. Fixstars Group www.fixstars.com Copyright © Fixstars Group Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる 97 // インスタンス化済みのグラフ cudaGraphExec_t graphExec; // graphExec と同じトポロジのグラフを構築 cudaGraph_t graph; …… // インスタンス化済みグラフの更新 cudaGraphNode_t errorNode; cudaGraphExecUpdateResult updateResult; cudaGraphExecUpdate(graphExec, graph, &errorNode, &updateResult);
98. Fixstars Group www.fixstars.com Copyright © Fixstars Group Individual Node Update • 特定ノードのパラメータのみを更新 • 構築時に得られたノードのハンドルで更新対象を指定 • cudaGraphExec**NodeSetParams() 98 // 構築時にノードのハンドルを控えておく cudaGraphNode_t node; cudaGraphAddKernelNode(&node, graph, nullptr, 0, &baseParams); // インスタンス化済みのグラフ cudaGraphExec_t graphExec; // インスタンス化済みグラフの更新 cudaGraphExecKernelNodeSetParams(graphExec, node, &modifiedParams);
99. Fixstars Group www.fixstars.com Copyright © Fixstars Group 性能評価: 極端なケース • 空のカーネルを1024回直列に投入する処理を1024回実行 • カーネル呼び出し1024回ごとに1回同期を入れる 99 RTX 2080 Ti (sm_75) NVIDIA A10G (sm_86) CUDA Graphs なし 1505.6 [ms] 2568.2 [ms] CUDA Graphs あり 879.0 [ms] 1018.0 [ms] 高速化率 1.71x 2.52x
100. Fixstars Group www.fixstars.com Copyright © Fixstars Group 性能評価: グラフ更新 • 100個のカーネルのうち最初と最後のものを書き換える • 入力ポインタ・出力ポインタのみを書き換えるようなシナリオ • これを1024回繰り返した時の処理時間 100 RTX 2080 Ti (sm_75) NVIDIA A10G (sm_86) 更新なし 81.58 [ms] 104.70 [ms] Whole Graph Update 152.30 [ms] 227.60 [ms] Individual Node Update 89.18 [ms] 106.11 [ms]
101. Fixstars Group www.fixstars.com Copyright © Fixstars Group 性能評価: TensorRT • MobileNet-v2 on NVIDIA A10G • バッチサイズを変えて試行 • カーネルあたりの負荷が小さいほど相対的に効果的 101 0.94 0.96 0.98 1 1.02 1.04 1.06 1.08 1.1 1.12 1.14 1.16 0 1 2 3 4 5 6 7 8 9 10 N=1 N=4 N=16 N=64 高速化率 処理時間 [ms] CUDA Graphs なし CUDA Graphs あり speedup
102. Fixstars Group www.fixstars.com Copyright © Fixstars Group まとめ • 複数のタスクをまとめて投入することでオーバーヘッドを低減 • 小さいタスクが多い場合特に効果的 • Ampereではハードウェアアクセラレーションも効く • 既存ライブラリと併用することもできる • ストリームキャプチャによるグラフ構築 102
103. Fixstars Group www.fixstars.com Copyright © Fixstars Group Copyright © Fixstars Group おわりに 103
104. Fixstars Group www.fixstars.com Copyright © Fixstars Group おわりに • CUDA環境はまだ拡張が続けられている • ハードウェアの世代更新 • ソフトウェア (CUDA Toolkit) のバージョンアップ • うまく活用することで様々な恩恵を受けられる • 処理速度の向上 • 従来は実装できなかったアルゴリズムの実装 • コードの可読性向上 • きちんと理解して適材適所で活用しましょう 104
105. Fixstars Group www.fixstars.com Copyright © Fixstars Group Thank You お問い合わせ窓口 : contact@fixstars.com