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

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

View Slide

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

View Slide

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

View Slide

フィックスターズのご紹介

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

CUDAの進化

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

SMアーキテクチャの変化
20

View Slide

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

View Slide

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

View Slide

実験: 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 ;
IADD3 R2, R2, 0x2, RZ ;
LOP3.LUT R8, R8, 0x3, RZ, 0x3c, !PT ;
LOP3.LUT R6, R6, 0x3, RZ, 0x3c, !PT ;
FP32 [GFLOPS] 15794.9
INT32 [GOPS] 7897.5

View Slide

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

View Slide

浮動小数点数演算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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

分岐したスレッド間の通信
• 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;
}

View Slide

分岐したスレッド間の通信
• 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;
}

View Slide

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

View Slide

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

View Slide

挙動が変わるケースの例
• 例: コンパイル時に反復回数がわからないループ
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() なし
合流したりしなかったり

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

• 実際にはこうなっている
• 一時レジスタ（と場合によってはL1キャッシュ）は省略できそう
• Ampereで可能になった
Global Memory
L2 Cache
L1 Cache
Register
Shared Memory
Register
Shared Memory
Register
……

View Slide

• 演算とデータ転送のオーバーラップ
• ロードされたデータを使用する前に同期をとる
• 他にもバリアとの組み合わせなどいろいろな書き方がある
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にデータがロードされている
…

View Slide

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

View Slide

Tensor Core
48

View Slide

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

View Slide

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

View Slide

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

View Slide

WMMA API の利用
• fragment: レジスタ上に配置される一時バッファ
• 型情報として用途などの情報を含む
52
__device__ void kernel(float *D, const half *A, const half *B, const float *C){
wmma::fragment A_frag;
wmma::fragment B_frag;
wmma::fragment 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);
}

View Slide

出力コードを読む: 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;

View Slide

出力コードを読む: 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] ;

View Slide

• 行列のロード部分
• 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 ;

View Slide

• 積和演算部分
• 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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

浮動小数点数の表現
• 深層学習周辺では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

View Slide

入出力の精度と速度
• 対応している精度と速度は世代によって異なる
• 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

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

Cooperative Groups
68

View Slide

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

View Slide

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

View Slide

使用例
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();
}

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

グループの分割
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);

View Slide

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

View Slide

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

View Slide

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

View Slide

• Coalesced Group の場合
• もう少しややこしいコードが出てくる
• 取得元のレーン番号を二分探索で求める
• 最後にブロードキャスト
• スレッドあたり O(log^2 n)
80

View Slide

• 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

View Slide

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

View Slide

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

View Slide

CUDA Graphs
84

View Slide

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

View Slide

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 あり

View Slide

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

View Slide

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

View Slide

グラフの構築
• 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);

View Slide

グラフの構築
• Graph API による構築
• カーネルパラメータの渡し方は cudaLaunchKernel() と似た形
90
__global__ void kernel(int x){ printf("%d¥n", x); }
int kernelParam1 = 42;
void *kernelParams[] = { static_cast(&kernelParam1) };
cudaKernelNodeParams params;
params.func = reinterpret_cast(&kernel);
params.gridDim = dim3(1, 1, 1);
params.blockDim = dim3(1, 1, 1);
params.sharedMemBytes = 0;
params.kernelParams = kernelParams;
params.extra = nullptr;

View Slide

グラフの構築
• ストリームキャプチャによる構築
• ストリームに対する操作から自動的にグラフを構築する
• イベントの 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);

View Slide

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

View Slide

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

View Slide

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

View Slide

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

View Slide

Whole Graph Update
96
A
B
src = in_ptr
dst = out_ptr
コピー先
（インスタンス化済み）
A
B
src = in_ptr
dst = out_ptr
コピー元
（インスタンス化前）

View Slide

Whole Graph Update
97
// インスタンス化済みのグラフ
// graphExec と同じトポロジのグラフを構築
cudaGraph_t graph;
……
// インスタンス化済みグラフの更新
cudaGraphNode_t errorNode;
cudaGraphExecUpdateResult updateResult;
cudaGraphExecUpdate(graphExec, graph, &errorNode, &updateResult);

View Slide

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

View Slide

性能評価: 極端なケース
• 空のカーネルを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

View Slide

性能評価: グラフ更新
• 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]

View Slide

性能評価: 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

View Slide

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

View Slide

おわりに
103

View Slide

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

View Slide

Thank You
お問い合わせ窓口 : [email protected]

View Slide

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

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

fixstars

More Decks by fixstars

Other Decks in Programming

Featured

Transcript