CUDA高速化セミナー vol.2 ~CUDAアーキテクチャの進化~(2022/06/23)

9.6K Views

June 23, 22

スライド概要

CUDA高速化セミナーシリーズの第二回として、CUDA高速化の基礎知識となるGPUマイクロアーキテクチャとCUDAツールキットについて解説します。

GPUに関連する開発業務を行っており、コンピュータアーキテクチャに興味がある方にオススメの内容となっています。

<講演内容>
・Volta / Turing / Ampere マイクロアーキテクチャとCUDAツールキットの解説

<過去資料>
・vol.1 画像処理アルゴリズムの高速化: https://www.docswell.com/s/fixstars/K24MYM-20220527
・vol.2 CUDAアーキテクチャの進化: https://www.docswell.com/s/fixstars/5RXQJ2-20220623
・vol.3 ソフトウェア高速化と深層学習: 
https://www.docswell.com/s/fixstars/5DEJQD-20220728
・vol.4 TensorRT化のワークフロー事例紹介: https://www.docswell.com/s/fixstars/524MGM-20220825
・vol.5  画像処理アルゴリズムの高速化2:https://www.docswell.com/s/fixstars/ZQ81QX-20220929

profile-image

フィックスターズは、コンピュータの性能を最大限に引き出すソフトウェア開発のスペシャリストです。車載、産業機器、金融、医療など、幅広い分野での開発経験があります。また、ディープラーニングや機械学習などの最先端技術にも力を入れています。 並列化や最適化技術を駆使して、マルチコアCPU、GPU、FPGA、量子アニーリングマシンなど、さまざまなハードウェアでソフトウェアを高速化するサービスを提供しています。さらに、長年の経験から培ったハードウェアの知識と最適化ノウハウを活かし、高精度で高性能なアルゴリズムの開発も行っています。       ・開催セミナー一覧:https://www.fixstars.com/ja/seminar   ・技術ブログ :https://proc-cpuinfo.fixstars.com/

シェア

埋め込む »CMSなどでJSが使えない場合

関連スライド

各ページのテキスト
1.

Fixstars Group www.fixstars.com CUDA 高速化セミナー vol.2 CUDAアーキテクチャの進化 Copyright © Fixstars Group Copyright © Fixstars Group

2.

Fixstars Group www.fixstars.com 発表者紹介 • 冨田 明彦(とみた あきひこ) • 平櫛 貴章(ひらぐし たかあき) ソリューションカンパニー 営業企画執行役 Fixstars Autonomous Technologies エグゼクティブエンジニア 2008年に入社。金融、医療業界において、 ソフトウェア高速化業務に携わる。その 後、新規事業企画、半導体業界の事業を 担当し、現職。 2015年に新卒で入社。幅広い産業領域で CPU / GPU を用いたパフォーマンスチュ ーニング業務に携わる。 2 Copyright © Fixstars Group

3.

Fixstars Group www.fixstars.com 本日のAgenda フィックスターズの紹介 (15分) • 会社紹介 • 高速化のためにCUDAアーキテクチャの進化を知る CUDAアーキテクチャの進化(60分) • SM アーキテクチャ • Tensor Core • Cooperative Groups • CUDA Graphs Q&A / 告知 3 Copyright © Fixstars Group

4.

Fixstars Group www.fixstars.com フィックスターズのご紹介 Copyright © Fixstars Group Copyright © Fixstars Group

5.

Fixstars Group www.fixstars.com フィックスターズの強み コンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団 低レイヤ ソフトウェア技術 アルゴリズム 実装力 各産業・研究 分野の知見 5 Copyright © Fixstars Group

6.

Fixstars Group www.fixstars.com ソフトウェア高速化サービス (概要) お客様のソースコードをご提供いただき、 最適化やアルゴリズムの改良を施して高速化してお返しします オリジナルソースコードのご提供 当社 高速化したソースコード コンサルティング 高速化 お客様 サポート 先行技術調査 アルゴリズムの改良・開発 レポートやコードへのQ&A 性能評価 ハードウェアへの最適化 実製品への組込み支援 ボトルネックの特定 レポート作成 Copyright © Fixstars Group 6

7.

Fixstars Group www.fixstars.com ソフトウェア高速化サービス 様々な領域でソフトウェア高速化サービスを提供しています 大量データの高速処理は、お客様の製品競争力の源泉となっています Semiconductor Industrial ・NAND型フラッシュメモリ向けファー ・Smart Factory化支援 ムウェア開発 ・マシンビジョンシステムの高速化 ・次世代AIチップ向け開発環境基盤開発 Mobility Life Science ・自動運転の高性能化、実用化 ・ゲノム解析の高速化 ・次世代パーソナルモビリティの研究開発 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Finance ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 7 Copyright © Fixstars Group

8.

Fixstars Group www.fixstars.com 画像処理・アルゴリズム開発サービスの例 • お客様の課題 • 高度な画像処理や深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件まで クリアできるか不安 • 製品化に結びつくような研究ができていない • 弊社の支援内容 • 課題に応じたアルゴリズム調査 • 深層学習ネットワーク精度改善、推論高速化手法調査 • 論文調査、実装 出展:https://www.cs.toronto.edu/~frossard/post/vgg16/ Copyright © Fixstars Group 8

9.

Fixstars Group www.fixstars.com AI・深層学習関連サービス • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績 ネットワーク設計 データの前処理、データ拡張 精度改善 分散処理による学習高速化 各種DLフレームワーク ターゲットデバイスへの ポーティング及び推論高速化 ■ ARM, GPU, DSP ■ SIMD,NEON,CUDA,TensorRT モデル圧縮 - 量子化 - 枝刈り - 蒸留 クラウド・サーバ エッジ Copyright © Fixstars Group 9

10.

Fixstars Group www.fixstars.com GPU向け高速化サービス • お客様の課題 • • GPU 高速化の知見がない 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • • • GPU 高速化に関するコンサルティング ボトルネック調査、GPU プログラムの高速化 CPU/GPU が混在するヘテロジニアス環境での最適化 Copyright © Fixstars Group 10~150 倍の 高速化事例あり 10

11.

Fixstars Group www.fixstars.com 高速化のためにCUDAアーキテクチャの進化を知る 1. GPU搭載の新製品へ機能を追加し 処理を高速化したい • • 新アーキテクチャの GPU を導入 CUDA Toolkit を最新にアップグレード 2. 高速化をやりきったか? • • もっと性能出せないか? もっと可読性を上げられないか? 3. CUDAアーキテクチャの進化を 知っておこう! Copyright © Fixstars Group 11 11

12.

Fixstars Group www.fixstars.com CUDAの進化 Copyright © Fixstars Group Copyright © Fixstars Group

13.

Fixstars Group www.fixstars.com 今回の概要 • Volta世代以降のCUDAの変更点のおさらい • CUDA Toolkit 9.0 以降の変更点のおさらい • 主に計算カーネルを記述する人向け 13 Copyright © Fixstars Group

14.

Fixstars Group www.fixstars.com Voltaアーキテクチャ • Compute Capability 7.0, 7.2 • • NVIDIA V100, TITAN V, Jetson Xavier など 主な新機能 • • • • 深層学習用アクセラレータ: Tensor Core L1キャッシュとシェアードメモリの拡張 浮動小数点数演算と整数演算の同時発行 Independent Thread Scheduling 14 Copyright © Fixstars Group

15.

Fixstars Group www.fixstars.com Turingアーキテクチャ • Compute Capability 7.5 • • • GeForce RTX 20xx, NVIDIA T4 など ゲーミング・グラフィックス向けのアーキテクチャ 主な新機能 • • レイトレーシング用アクセラレータ: RT Core 第2世代 Tensor Core 15 Copyright © Fixstars Group

16.

Fixstars Group www.fixstars.com Ampereアーキテクチャ (GA100) • データセンター向けAmpere • Compute Capability 8.0 • • NVIDIA A100 など 主な新機能 • • • 第3世代 Tensor Core 非同期コピー・バリア タスクグラフの高速化 16 Copyright © Fixstars Group

17.

Fixstars Group www.fixstars.com Ampereアーキテクチャ (GA102) • デスクトップ向けAmpere • Compute Capability 8.6 • • GeForce RTX 30xx など 主な新機能 • • 第2世代 RT Core 浮動小数点数演算2つの同時発行 17 Copyright © Fixstars Group

18.

Fixstars Group www.fixstars.com 今日の内容 • SMアーキテクチャの変化 • • • 整数演算コアの追加 Independent Thread Scheduling スレッド間通信機能の強化 • Tensor Core • Cooperative Groups • CUDA Graphs 18 Copyright © Fixstars Group

19.

Fixstars Group www.fixstars.com 注意事項 • この資料で言及している内容は CUDA 11.7 時点のものです • 今後のバージョンアップで変化することがあります 19 Copyright © Fixstars Group

20.

Fixstars Group www.fixstars.com SMアーキテクチャの変化 20 Copyright © Fixstars Group Copyright © Fixstars Group

21.

Fixstars Group www.fixstars.com 整数演算コアの追加 • Volta以降 • CUDAコアがFP32コアとINT32コアに分けられた • 非本質的な計算によるリソース圧迫を防ぐ • 非本質的な計算: アドレス計算など https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf Copyright © Fixstars Group 21

22.

Fixstars Group www.fixstars.com アドレス計算 • CUDAコアのアドレッシングモードはあまり多機能ではない • • レジスタ上のアドレス+定数オフセットのみ うっかりするとアドレス計算命令の方が多くなる • INTコアで浮動小数点数演算ユニットの圧迫を回避できる • とはいえまだ注意は必要 • • 浮動小数点数演算の倍以上あるとFP32コアが余る IMADはFPコアで処理される 22 Copyright © Fixstars Group

23.

Fixstars Group www.fixstars.com 実験: INT32 Core を動かしてみる • FFMAとADD, LOP3を交互に動かす • ASUS RTX 2080Ti TURBO で測定 • • Theoretical Peak: 17103.4 [GFLOPS] FP32 [GFLOPS] 15794.9 INT32 [GOPS] 7897.5 足し合わせると理論ピークを越える • FFMA IADD3 FFMA IADD3 FFMA LOP3.LUT FFMA LOP3.LUT R9, R4, R7, R2, R5, R8, R3, R6, R9, R4, R7, R2, R5, R8, R3, R6, R10.reuse, 3 ; 0x2, RZ ; R10.reuse, 3 ; 0x2, RZ ; R10.reuse, 3 ; 0x3, RZ, 0x3c, !PT ; R10.reuse, 3 ; 0x3, RZ, 0x3c, !PT ; INTコアの効果が出ている 23 Copyright © Fixstars Group

24.

Fixstars Group www.fixstars.com 実験: INT32 Core を動かしてみる • スレッドブロック番号の偶奇で分岐 • • • • FFMAのみをひたすら発行するスレッドブロック ADD, LOP3のみをひたすら発行するスレッドブロック 演算の総量は前の例と同じ 同様の傾向となる • • FP32とINT32で異なるスレッド・ワープの命令を供給できる スーパースカラというよりはSMT的な振る舞い 24 Copyright © Fixstars Group

25.

Fixstars Group www.fixstars.com 浮動小数点数演算2つの同時発行 • Ampere (GA102) • INT32コアが浮動小数点数も扱えるように • • FP32のスループットが2倍に またアドレス計算まわりのチューニングが必要に https://images.nvidia.com/aem-dam/en-zz/Solutions/geforce/ampere/pdf/NVIDIAampere-GA102-GPU-Architecture-Whitepaper-V1.pdf Copyright © Fixstars Group 25

26.

Fixstars Group www.fixstars.com L1キャッシュの効率改善 • Volta • L1キャッシュが速くなった • • シェアードメモリより少し遅い程度 シェアードメモリをキャッシュ代わりにしなくて良いケースが増える • const __restrict__ をつける(コンパイラが自動でつけることもある) 26 Copyright © Fixstars Group

27.

Fixstars Group www.fixstars.com Independent Thread Scheduling • Volta • スレッドごとにPCとスタックを持つようになった • 分岐時のスケジューリングの選択肢が増える • • 分岐したスレッド間での通信が可能となる レイテンシの隠蔽にも有利かもしれない 27 Copyright © Fixstars Group

28.

Fixstars Group www.fixstars.com Independent Thread Scheduling • Pascal以前 • • 同一ワープ内の全スレッドがPCを共有 ある命令の処理は必ず同時に行われる A() if(threadIdx.x < 8){ A(); B(); }else{ C(); D(); } E(); B() C() D() E() 28 Copyright © Fixstars Group

29.

Fixstars Group www.fixstars.com Independent Thread Scheduling • Volta以降 • • スレッドごとにPCを持つ 1つ以上のスレッドのPCが指している命令が発行される A() if(threadIdx.x < 8){ A(); B(); }else{ C(); D(); } E(); C() B() D() E() 29 Copyright © Fixstars Group

30.

Fixstars Group www.fixstars.com 分岐したスレッド間の通信 • Pascal以前の場合 • • スレッド0側のループから抜けられない スレッド1側の処理が進まない __shared__ int x; if (threadIdx.x == 0) { x = 1; do { // __nanosleep(0); } while(x); } else { do { // __nanosleep(0); } while(!x); x = 0; } 30 Copyright © Fixstars Group

31.

Fixstars Group www.fixstars.com 分岐したスレッド間の通信 • Volta以降の場合 • • __nanosleepで他スレッドに処理を譲る スレッド1側の処理を進めることができる __shared__ int x; if (threadIdx.x == 0) { x = 1; do { __nanosleep(0); } while(x); } else { do { __nanosleep(0); } while(!x); x = 0; } 31 Copyright © Fixstars Group

32.

Fixstars Group www.fixstars.com ワープ内での同期 • Volta以降 • A() コード上のある地点に同時に到達しない場合がある A() B() B() E() E() C() C() D() 期待される処理順の例 - 全スレッドの E() が同時に発行される D() E() 発生しうる処理順の例 - 全スレッドの E() が同時に発行されるとは限らない 32 Copyright © Fixstars Group

33.

Fixstars Group www.fixstars.com ワープ内での同期 • __syncwarp() で明示的に同期をとる • • マスクで指定したスレッドが到達するのを待つ それ以降は再度分岐するまで同期された状態で進む 33 Copyright © Fixstars Group

34.
[beta]
Fixstars Group

www.fixstars.com

挙動が変わるケースの例
•

例: コンパイル時に反復回数がわからないループ
{0, 1, 2, …, 30, 31} を与える

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

__activemask(): アクティブなスレッドの取得

ffffffff
ffffffff
ffffffff
ffffffff
…
ffffffff
ffffffff
ffffffff
ffffffff

000003ff
000003ff
000003ff
000003ff
…
44000000
44000000
88000000
88000000

__syncwarp() あり
全スレッドが同時に到達

__syncwarp() なし
合流したりしなかったり
34

Copyright © Fixstars Group

35.

Fixstars Group www.fixstars.com ワープ内通信 • 他スレッドと通信を行う場合明示的な同期を挿入する • ワープ内通信を伴う組み込み関数は _sync がつく • • • • __all_sync, __shfl_sync など 同期をとる必要があるスレッド集合をビットマスクで指定 CUDA Toolkit 9.0 以降 _sync なしは Deprecated 使える命令が少し増えました • • Match (Volta) Reduce (Ampere) 35 Copyright © Fixstars Group

36.

Fixstars Group www.fixstars.com Warp Match: match_any • valueに同じ値を渡したスレッドの集合を取得 unsigned int __match_any_sync(unsigned mask, T value); 0 0x00cb 0 0x00cb 1 0x0034 0 0x00cb 1 0x0034 1 0x0034 0 0x00cb 0 0x00cb 0x00cb => 0, 1, 3, 6, 7 0x0034 => 2, 4, 5 Copyright © Fixstars Group 36

37.

Fixstars Group www.fixstars.com Warp Match: match_all • 集合中のスレッド全てがvalueに同じ値を渡したかの判定 unsigned int __match_all_sync(unsigned mask, T value, int *pred); mask value result *pred 0x000f 0 0x0000 0 0x000f 1 0x0000 0 0x000f 1 0x0000 0 0x000f 0 0x0000 0 0x00f0 0 0x00f0 1 0x00f0 0 0x00f0 1 0x00f0 0 0x00f0 1 0x00f0 0 0x00f0 1 37 Copyright © Fixstars Group

38.

Fixstars Group www.fixstars.com Reduce • 整数の総和や最小値・最大値などを求める 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); 38 Copyright © Fixstars Group

39.

Fixstars Group www.fixstars.com Asynchronous Barrier • Five Stages of Synchronization • 同期を arrive と wait の2つに分割する • Arrive • • • • その地点への到達を示す 特に何かを待たずに次の処理へ進む Waitでの待ち受けに使用するトークンを返す Wait • • スレッド群がトークンに対応するArriveへ到達するまで待つ 他スレッドのArrive以前のメモリ操作が観測可能であることが保証される 39 Copyright © Fixstars Group

40.

Fixstars Group www.fixstars.com Asynchronous Barrier barrier::arrive • 処理が到達したことを他スレッドに通知 • この時点ではブロックされない • 処理 (2) へ直ちに進む 処理 (1) barrier::arrive 処理 (2) barrier::wait 処理 (3) 40 Copyright © Fixstars Group

41.

Fixstars Group www.fixstars.com Asynchronous Barrier barrier::wait • 全スレッドがarriveに到達するまで待つ • • 他スレッドは処理 (2) の途中かもしれない 処理 (1) での変更が観測可能になる • 処理 (2) での変更については保証なし 処理 (1) barrier::arrive 処理 (2) barrier::wait 処理 (3) 41 Copyright © Fixstars Group

42.

Fixstars Group www.fixstars.com Asynchronous Barrier: 性能評価 • 220回同期するだけのカーネルの処理時間 • • • Async: barrier.arrive_and_wait() • • • 256スレッド, 1スレッドブロック NVIDIA A10G HWアクセラレーションあり: -arch sm_86 HWアクセラレーションなし: -arch sm_75 Sync: __syncthreads() Async (w/ HW accel.) Time [ms] Async (w/o HW accel.) 127.4 577.8 Sync 23.8 42 Copyright © Fixstars Group

43.

Fixstars Group www.fixstars.com Asynchronous Data Copy • メモリコピーの非同期実行を可能にする • ありがちなパターン • • • グローバルメモリからシェアードメモリにコピー シェアードメモリ上のデータを使って計算 これを繰り返す Register Shared Memory Shared Memory Global Memory Shared Memory Global Memory Copyright © Fixstars Group …… 43

44.

Fixstars Group www.fixstars.com Asynchronous Data Copy • 実際にはこうなっている • • 一時レジスタ(と場合によってはL1キャッシュ)は省略できそう Ampereで可能になった Register Register Register L1 Cache Shared Memory Shared Memory …… L2 Cache Global Memory Copyright © Fixstars Group

45.

Fixstars Group www.fixstars.com Asynchronous Data Copy • 実際にはこうなっている • • 一時レジスタ(と場合によってはL1キャッシュ)は省略できそう Ampereで可能になった Register Register Register L1 Cache Shared Memory Shared Memory …… L2 Cache Global Memory Copyright © Fixstars Group

46.

Fixstars Group www.fixstars.com 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にデータがロードされている … Copyright © Fixstars Group

47.

Fixstars Group www.fixstars.com まとめ • 徐々にできることが増えている • • • • 演算器の増加 より柔軟なスケジューリング スレッド間通信の高機能化 移行に伴い気をつけるべき点もある • • • ワープ内の同期 ボトルネックとなる点の変化 新機能=高速とは限らない 47 Copyright © Fixstars Group

48.

Fixstars Group www.fixstars.com Tensor Core 48 Copyright © Fixstars Group Copyright © Fixstars Group

49.

Fixstars Group www.fixstars.com Tensor Core • 深層学習向けのアクセラレータ • 小さい行列の積を効率よく求める • 入力の精度によって性能が変わる https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf Copyright © Fixstars Group 49

50.

Fixstars Group www.fixstars.com Tensor Core を用いた演算の流れ • Tensor Core はワープ単位で協調して使用する • 入力行列をレジスタにロードする • 行列積を求める • 出力行列をメモリにストアする 50 Copyright © Fixstars Group

51.

Fixstars Group www.fixstars.com WMMA API • WMMA: Warp Matrix Multiply and Accumulate • • • ワープ単位で協調して行列積を行うためのAPI 実質的に Tensor Core を利用するためのAPI 利用可能な行列サイズの制約 • • 実ハードウェアのそれよりやや大きい Tensor Core の世代差 (命令当たりの演算量) を吸収する • • 小さい行列を処理する命令を複数回呼んで少し大きい行列を処理 逆は効率の低下につながってしまう 51 Copyright © Fixstars Group

52.
[beta]
Fixstars Group

www.fixstars.com

WMMA API の利用
•

fragment: レジスタ上に配置される一時バッファ
•

型情報として用途などの情報を含む

__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);
}
52
Copyright © Fixstars Group

53.

Fixstars Group www.fixstars.com 出力コードを読む: PTX • WMMA API の呼び出しと対応する • 一塊で 16x16x16 だけ処理している 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; 53 Copyright © Fixstars Group

54.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 行列のロード部分 • 8x8部分行列ごとに2要素ロード /* 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] ; 54 Copyright © Fixstars Group

55.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 行列のロード部分 • • Bのロードは8x8部分行列内での転置が入る Bが列優先であれば転置は省略される /* 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 MOVM.16.MT88 MOVM.16.MT88 MOVM.16.MT88 R0, R0 ; R23, R23 ; R22, R22 ; R24, R24 ; 55 Copyright © Fixstars Group

56.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 R8, R4, R8, R28, R18, R18, R20, R20, R0, R23, R22, R24, R8 R4 R8 R4 R0 R23 R22 R24 ; ; ; ; R18 R20 R8,9 R4,5 R19 R21 R10,11 R6,7 56 Copyright © Fixstars Group

57.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 R8, R4, R8, R28, R18, R18, R20, R20, R0, R23, R22, R24, R8 R4 R8 R4 R0 R23 R22 R24 ; ; ; ; R18 R20 R8,9 R4,5 R19 R21 R10,11 R6,7 57 Copyright © Fixstars Group

58.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 R8, R4, R8, R28, R18, R18, R20, R20, R0, R23, R22, R24, R8 R4 R8 R4 R0 R23 R22 R24 ; ; ; ; R18 R20 R8,9 R4,5 R19 R21 R10,11 R6,7 58 Copyright © Fixstars Group

59.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_75) • 積和演算部分 • 4回の16x8x8行列積に分割される HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 HMMA.1688.F32 R8, R4, R8, R28, R18, R18, R20, R20, R0, R23, R22, R24, R8 R4 R8 R4 R0 R23 R22 R24 ; ; ; ; R18 R20 R8,9 R4,5 R19 R21 R10,11 R6,7 59 Copyright © Fixstars Group

60.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_86) • 積和演算部分 • 2回の16x8x16行列積に分割される R24 R26 R25 R27 HMMA.16816.F32 R8, R4.reuse, R24, R8 ; HMMA.16816.F32 R12, R4, R26, R12 ; R4 R6 R8,9 R12,13 R5 R7 R10,11 R14,15 60 Copyright © Fixstars Group

61.

Fixstars Group www.fixstars.com 出力コードを読む: SASS (sm_86) • 積和演算部分 • 2回の16x8x16行列積に分割される R24 R26 R25 R27 HMMA.16816.F32 R8, R4.reuse, R24, R8 ; HMMA.16816.F32 R12, R4, R26, R12 ; R4 R6 R8,9 R12,13 R5 R7 R10,11 R14,15 61 Copyright © Fixstars Group

62.

Fixstars Group www.fixstars.com 浮動小数点数の表現 • 深層学習周辺ではIEEE754以外の形式が用いられることがある • • 指数部と仮数部に用いるビット数が異なる 第3世代 Tensor Core で TF32, BF16 を追加サポート 符号部 指数部 仮数部 FP32: IEEE754 Single 1 8 23 TF32: TensorFloat-32 1 8 10 FP16: IEEE754 Half 1 5 10 BF16: BFloat16 1 8 7 62 Copyright © Fixstars Group

63.

Fixstars Group www.fixstars.com 入出力の精度と速度 • 対応している精度と速度は世代によって異なる • FP32コアによるFFMAとの性能比: 入力 出力 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 1st Gen. 2nd Gen. Copyright © Fixstars Group 3rd Gen. 63

64.

Fixstars Group www.fixstars.com Structured Sparsity • Ampere (GA100) で追加 • ニューラルネットの重みは寄与度の低い要素を多く含む • これらの要素を0とみなして演算量を削減する: pruning Han et al. (2015) Learning both Weights and Connections for Efficient Neural Networks Copyright © Fixstars Group 64

65.

Fixstars Group www.fixstars.com Structured Sparsity • 制約付きでハードウェアによるサポートを行う • 2:4 Sparsity: 連続する4要素のうち2要素が0である • • 行と列の内積を求めるときに0との積和をスキップする 必要な積和の回数が半減する⇒実効性能が倍になる 65 Copyright © Fixstars Group

66.
[beta]
Fixstars Group

www.fixstars.com

Structured Sparsity
•

WMMA API ではまだ提供されていない?
•
•

•

ドキュメント中にそれらしい言及がない
ヘッダファイルにもそれらしい定義は見当たらない

CUTLASSではPTXを直書きしている
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

66
Copyright © Fixstars Group

67.

Fixstars Group www.fixstars.com まとめ • Tensor Core: 行列演算用アクセラレータ • • • ワープ単位で協調して小さい行列同士の積を求める 世代と精度によって性能が大きく変化する CUDA C++ からは WMMA API 経由で操作する • • 実ハードより少し粗い単位で処理を行う 一部の処理はまだ提供されていない 67 Copyright © Fixstars Group

68.

Fixstars Group www.fixstars.com Cooperative Groups 68 Copyright © Fixstars Group Copyright © Fixstars Group

69.

Fixstars Group www.fixstars.com Cooperative Groups • 複数スレッドで協調する処理を記述するためのライブラリ • • • グリッド・スレッドブロック・ワープ…… CUDA Toolkit 9.0 で導入された CUDAの標準ライブラリとして提供される • • 実は従来から頑張れば似たようなことは出来た 黒魔術的な実装で将来の変更におびえなくてよくなる 69 Copyright © Fixstars Group

70.

Fixstars Group www.fixstars.com スレッドブロック • cooperative_groups::thread_block • できること • • • • スレッドブロックサイズの取得 (dim3もしくは1次元化されたもの) スレッドブロック内における自身のスレッド番号の取得 グリッド内における自身のスレッドブロック番号の取得 スレッドブロック内の全スレッドでの同期 70 Copyright © Fixstars Group

71.

Fixstars Group www.fixstars.com 使用例 __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(); } 71 Copyright © Fixstars Group

72.

Fixstars Group www.fixstars.com グリッド • cooperative_groups::grid_group • できること • • • • グリッド内における自身のスレッド番号の取得 グリッド内における自身のスレッドブロック番号の取得 グリッド内のスレッド数とスレッドブロック数の取得 グリッド内の全スレッドでの同期 • カーネル呼び出し時に特定の条件を満たす必要がある 72 Copyright © Fixstars Group

73.

Fixstars Group www.fixstars.com グリッド内での同期 • 全てのスレッドが同時に起動している必要がある • • Occupancyでスレッド数の上限が決まる cudaLaunchCooperativeKernel でカーネルを起動する 73 Copyright © Fixstars Group

74.

Fixstars Group www.fixstars.com タイル (Thread Block Tile) • スレッドブロックをさらに分割したもの • • • 実際はワープサイズ以下の2冪個のスレッドをまとめたもの スレッド数はコンパイル時定数である必要がある できること • • • 自身のインデックスの取得 グループ内での同期 各種ワープ内通信 (シャッフル・ボーティング・マッチング) 74 Copyright © Fixstars Group

75.

Fixstars Group www.fixstars.com Coalesced Group • 同一ワープ内の任意のスレッド群からなるグループ • 典型的にはアクティブなスレッド群から構成する cg::coalesced_group active = cg::coalesced_threads(); • タイルでできることとほとんど同じことができる • 処理によっては追加のコストがかかることがある 75 Copyright © Fixstars Group

76.

Fixstars Group www.fixstars.com グループの分割 tiled_partition • スレッドブロックまたはタイルをより細かいタイルに分割 cg::thread_block block = cg::this_thread_block(); auto tile32 = cg::tiled_partition<32>(block); auto tile16 = cg::tiled_partition<16>(tile32); labeled_partition, binary_partition • ラベルとして同じ値を渡したスレッドからなるグループを生成 cg::thread_block block = cg::this_thread_block(); auto group = cg::labeled_partition(block, block.thread_rank() % 4); 76 Copyright © Fixstars Group

77.

Fixstars Group www.fixstars.com Data Manipulation • グループ内のスレッド同士で協調して計算を行う • • • 現状ではワープ内 (タイルと Coalesced Group) のみ アーキテクチャごとに適切な実装が選択される できること • • Reduce Scan 77 Copyright © Fixstars Group

78.

Fixstars Group www.fixstars.com Reduce • 𝑦 = 𝑥0 op 𝑥1 op … op 𝑥𝑛−1 • • • • 𝑖番目のスレッドからの入力を𝑥𝑖 とする opは結合的な2項演算 総和の計算などに用いる 可能なら __reduce_**_sync が使用される • • Ampere (sm_80) 以降のハードウェアである op が特定の演算 (cg::plus<int> など) である 78 Copyright © Fixstars Group

79.

Fixstars Group www.fixstars.com Reduce: HWアクセラレーションなし • タイルの場合 • __shfl_xor を用いたコードが出力される • • Kepler以降で使われていたものと同様 スレッドあたり O(log n) 79 Copyright © Fixstars Group

80.

Fixstars Group www.fixstars.com Reduce: HWアクセラレーションなし • Coalesced Group の場合 • もう少しややこしいコードが出てくる • • • 取得元のレーン番号を二分探索で求める 最後にブロードキャスト スレッドあたり O(log^2 n) 80 Copyright © Fixstars Group

81.

Fixstars Group www.fixstars.com Reduce: HWアクセラレーションなし • Coalesced Group の場合 • 単純なタイルより遅くなる • mask=0xffffffffの時は別コードが使用されるため例外 実装 処理時間 [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 81 Copyright © Fixstars Group

82.

Fixstars Group www.fixstars.com Scan • 𝑦𝑖 = 𝑥0 op 𝑥1 op … op 𝑥𝑖 (inclusive_scan) • 𝑦𝑖 = 𝑥0 op 𝑥1 op … op 𝑥𝑖−1 (exclusive_scan) • 特にハードウェアアクセラレーションはない • アクセラレーションなしのreduceと似たようなコードになる 82 Copyright © Fixstars Group

83.

Fixstars Group www.fixstars.com まとめ • Cooperative Groups • スレッド間の協調をポータブルに記述できる • • グループの粒度の差異を吸収する アーキテクチャ・ツールキットの変化を吸収する 83 Copyright © Fixstars Group

84.

Fixstars Group www.fixstars.com CUDA Graphs 84 Copyright © Fixstars Group Copyright © Fixstars Group

85.

Fixstars Group www.fixstars.com CUDA Graphs • 多くのジョブは複数のタスクの組み合わせによって構成される • • • • カーネル呼び出し データ転送 ホスト側処理 …… • タスク間の依存性をグラフで表現する • その一連の処理をまとめて実行する 85 Copyright © Fixstars Group

86.

Fixstars Group www.fixstars.com CUDA Graphs のメリット • タスク投入に起因するオーバーヘッドを低減できる • • 細かいタスクを多く含む場合に特に有効 Ampereだとハードウェアアクセラレーションもある GPU A B C D Launch B Launch C Launch D CUDA Graphs なし CPU Launch A A GPU B C D CUDA Graphs あり CPU Launch Graph 86 Copyright © Fixstars Group

87.

Fixstars Group www.fixstars.com グラフの構成要素 • ノード • • • • • • • カーネル起動 CPU処理の実行 メモリ確保・解放・コピー イベント処理 (Record, Wait) セマフォ操作 (Signal, Wait) 子グラフの実行 ノード間の依存関係 87 Copyright © Fixstars Group

88.

Fixstars Group www.fixstars.com グラフのライフサイクル • 構築したグラフを使いまわすことで性能を稼ぐ グラフの構築 インスタンス化 実行 更新 88 Copyright © Fixstars Group

89.

Fixstars Group www.fixstars.com グラフの構築 • Graph API による構築 • カーネルパラメータの渡し方は cudaLaunchKernel() と似た形 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); 89 Copyright © Fixstars Group

90.
[beta]
Fixstars Group

www.fixstars.com

グラフの構築
•

Graph API による構築
•

カーネルパラメータの渡し方は cudaLaunchKernel() と似た形

__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;
90
Copyright © Fixstars Group

91.

Fixstars Group www.fixstars.com グラフの構築 • ストリームキャプチャによる構築 • • ストリームに対する操作から自動的にグラフを構築する イベントの Record/Wait で他ストリームもキャプチャ対象になる 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); 91 Copyright © Fixstars Group

92.

Fixstars Group www.fixstars.com インスタンス化 • 実行可能なグラフオブジェクトを生成する • 最適化・エラーチェックなどもこのタイミングで行われる // グラフの構築 cudaGraph_t graph; …… // グラフのインスタンス化 cudaGraphExec_t graphExec; cudaGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0); 92 Copyright © Fixstars Group

93.

Fixstars Group www.fixstars.com グラフの実行 • 指定したストリームに対してグラフを構成するタスク群を投入 cudaStream_t stream; cudaGraphExec_t graphExec; cudaGraphLaunch(graphExec, stream); 93 Copyright © Fixstars Group

94.

Fixstars Group www.fixstars.com グラフの更新 • 軽微な変更は再インスタンス化せずに行うことができる • • • 操作対象となるメモリアドレスの書き換えなど インスタンス化に起因するコストを抑える 2通りの方法がある • • Whole graph update Individual node update 94 Copyright © Fixstars Group

95.

Fixstars Group www.fixstars.com Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる コピー先 (インスタンス化済み) コピー元 (インスタンス化前) src = nullptr src = in_ptr A A dst = nullptr dst = out_ptr B B 95 Copyright © Fixstars Group

96.

Fixstars Group www.fixstars.com Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる コピー先 (インスタンス化済み) コピー元 (インスタンス化前) src = in_ptr src = in_ptr A A dst = out_ptr dst = out_ptr B B 96 Copyright © Fixstars Group

97.

Fixstars Group www.fixstars.com Whole Graph Update • 同じトポロジの構築済みグラフからパラメータをコピーする • トポロジが異なる場合などコピーできない場合はエラーとなる // インスタンス化済みのグラフ cudaGraphExec_t graphExec; // graphExec と同じトポロジのグラフを構築 cudaGraph_t graph; …… // インスタンス化済みグラフの更新 cudaGraphNode_t errorNode; cudaGraphExecUpdateResult updateResult; cudaGraphExecUpdate(graphExec, graph, &errorNode, &updateResult); 97 Copyright © Fixstars Group

98.

Fixstars Group www.fixstars.com Individual Node Update • 特定ノードのパラメータのみを更新 • • 構築時に得られたノードのハンドルで更新対象を指定 cudaGraphExec**NodeSetParams() // 構築時にノードのハンドルを控えておく cudaGraphNode_t node; cudaGraphAddKernelNode(&node, graph, nullptr, 0, &baseParams); // インスタンス化済みのグラフ cudaGraphExec_t graphExec; // インスタンス化済みグラフの更新 cudaGraphExecKernelNodeSetParams(graphExec, node, &modifiedParams); 98 Copyright © Fixstars Group

99.

Fixstars Group www.fixstars.com 性能評価: 極端なケース • 空のカーネルを1024回直列に投入する処理を1024回実行 • カーネル呼び出し1024回ごとに1回同期を入れる 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 高速化率 99 Copyright © Fixstars Group

100.

Fixstars Group www.fixstars.com 性能評価: グラフ更新 • 100個のカーネルのうち最初と最後のものを書き換える • • 入力ポインタ・出力ポインタのみを書き換えるようなシナリオ これを1024回繰り返した時の処理時間 RTX 2080 Ti (sm_75) 更新なし Whole Graph Update Individual Node Update NVIDIA A10G (sm_86) 81.58 [ms] 104.70 [ms] 152.30 [ms] 227.60 [ms] 89.18 [ms] 106.11 [ms] 100 Copyright © Fixstars Group

101.

Fixstars Group www.fixstars.com 性能評価: TensorRT MobileNet-v2 on NVIDIA A10G • カーネルあたりの負荷が小さいほど相対的に効果的 10 1.16 1.14 1.12 1.1 1.08 1.06 1.04 1.02 1 0.98 0.96 0.94 9 8 処理時間 [ms] • バッチサイズを変えて試行 7 6 5 4 3 2 1 0 N=1 N=4 N=16 N=64 Copyright © Fixstars Group 高速化率 • CUDA Graphs なし CUDA Graphs あり speedup 101

102.

Fixstars Group www.fixstars.com まとめ • 複数のタスクをまとめて投入することでオーバーヘッドを低減 • • • 小さいタスクが多い場合特に効果的 Ampereではハードウェアアクセラレーションも効く 既存ライブラリと併用することもできる • ストリームキャプチャによるグラフ構築 102 Copyright © Fixstars Group

103.

Fixstars Group www.fixstars.com おわりに 103 Copyright © Fixstars Group Copyright © Fixstars Group

104.

Fixstars Group www.fixstars.com おわりに • CUDA環境はまだ拡張が続けられている • • • うまく活用することで様々な恩恵を受けられる • • • • ハードウェアの世代更新 ソフトウェア (CUDA Toolkit) のバージョンアップ 処理速度の向上 従来は実装できなかったアルゴリズムの実装 コードの可読性向上 きちんと理解して適材適所で活用しましょう 104 Copyright © Fixstars Group

105.

Fixstars Group www.fixstars.com Thank You お問い合わせ窓口 : contact@fixstars.com Copyright © Fixstars Group