---
title: CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）
tags:  #gpu #cuda #cuda高速化 #高速化シリーズ #cuda高速化シリーズ  
author: [株式会社フィックスターズ](https://www.docswell.com/user/fixstars)
site: [Docswell](https://www.docswell.com/)
thumbnail: https://bcdn.docswell.com/page/4JMVKG15JW.jpg?width=480
description: 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 　画像処理アルゴリズムの高速化２：https://www.docswell.com/s/fixstars/ZQ81QX-20220929
published: June 23, 22
canonical: https://www.docswell.com/s/fixstars/5RXQJ2-20220623
---
# Page. 1

![Page Image](https://bcdn.docswell.com/page/4JMVKG15JW.jpg)

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


# Page. 2

![Page Image](https://bcdn.docswell.com/page/PJR2XZLZE9.jpg)

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


# Page. 3

![Page Image](https://bcdn.docswell.com/page/PEXD69W17X.jpg)

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


# Page. 4

![Page Image](https://bcdn.docswell.com/page/3EKX1PVM7D.jpg)

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


# Page. 5

![Page Image](https://bcdn.docswell.com/page/L735QXM2E5.jpg)

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


# Page. 6

![Page Image](https://bcdn.docswell.com/page/87DD92667G.jpg)

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


# Page. 7

![Page Image](https://bcdn.docswell.com/page/VJPXZLZZ78.jpg)

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


# Page. 8

![Page Image](https://bcdn.docswell.com/page/2EVLWQWMJQ.jpg)

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


# Page. 9

![Page Image](https://bcdn.docswell.com/page/57GXDWDXEL.jpg)

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


# Page. 10

![Page Image](https://bcdn.docswell.com/page/4EQ81315JP.jpg)

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


# Page. 11

![Page Image](https://bcdn.docswell.com/page/KJ42R1RVE1.jpg)

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


# Page. 12

![Page Image](https://bcdn.docswell.com/page/LE1MQGQ47G.jpg)

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


# Page. 13

![Page Image](https://bcdn.docswell.com/page/GEW3MKMZE2.jpg)

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


# Page. 14

![Page Image](https://bcdn.docswell.com/page/47ZVQZQLE3.jpg)

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


# Page. 15

![Page Image](https://bcdn.docswell.com/page/YJ656Z6M7V.jpg)

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


# Page. 16

![Page Image](https://bcdn.docswell.com/page/GJ5Y6W6QJ4.jpg)

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


# Page. 17

![Page Image](https://bcdn.docswell.com/page/9E2GMQMWER.jpg)

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


# Page. 18

![Page Image](https://bcdn.docswell.com/page/D7YVXWX97M.jpg)

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


# Page. 19

![Page Image](https://bcdn.docswell.com/page/VENM191DJ8.jpg)

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


# Page. 20

![Page Image](https://bcdn.docswell.com/page/Y79N525873.jpg)

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


# Page. 21

![Page Image](https://bcdn.docswell.com/page/G78KY5YZJD.jpg)

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


# Page. 22

![Page Image](https://bcdn.docswell.com/page/L7LZ4Y41ER.jpg)

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


# Page. 23

![Page Image](https://bcdn.docswell.com/page/4EMVKNK5EW.jpg)

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


# Page. 24

![Page Image](https://bcdn.docswell.com/page/PER2XDXZ79.jpg)

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


# Page. 25

![Page Image](https://bcdn.docswell.com/page/P7XD6161JX.jpg)

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


# Page. 26

![Page Image](https://bcdn.docswell.com/page/37KX121MJD.jpg)

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


# Page. 27

![Page Image](https://bcdn.docswell.com/page/LJ35Q4Q275.jpg)

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


# Page. 28

![Page Image](https://bcdn.docswell.com/page/8JDD9Q96JG.jpg)

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


# Page. 29

![Page Image](https://bcdn.docswell.com/page/VEPXZL9ZJ8.jpg)

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


# Page. 30

![Page Image](https://bcdn.docswell.com/page/27VLWQMMEQ.jpg)

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


# Page. 31

![Page Image](https://bcdn.docswell.com/page/5JGXDWGX7L.jpg)

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


# Page. 32

![Page Image](https://bcdn.docswell.com/page/47Q813X5EP.jpg)

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


# Page. 33

![Page Image](https://bcdn.docswell.com/page/KE42R16V71.jpg)

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


# Page. 34

![Page Image](https://bcdn.docswell.com/page/L71MQGV4JG.jpg)

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 &lt; input[threadIdx.x]; ++i){
smem[threadIdx.x] += 1;
}
__syncwarp();
printf(&quot;%08x¥n&quot;, __activemask());
}
__activemask(): アクティブなスレッドの取得
ffffffff
ffffffff
ffffffff
ffffffff
…
ffffffff
ffffffff
ffffffff
ffffffff
000003ff
000003ff
000003ff
000003ff
…
44000000
44000000
88000000
88000000
__syncwarp() あり
全スレッドが同時に到達
__syncwarp() なし
合流したりしなかったり
34
Copyright © Fixstars Group


# Page. 35

![Page Image](https://bcdn.docswell.com/page/G7W3MKLZ72.jpg)

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


# Page. 36

![Page Image](https://bcdn.docswell.com/page/4JZVQZML73.jpg)

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 =&gt; 0, 1, 3, 6, 7
0x0034 =&gt; 2, 4, 5
Copyright © Fixstars Group
36


# Page. 37

![Page Image](https://bcdn.docswell.com/page/YE656ZNMJV.jpg)

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


# Page. 38

![Page Image](https://bcdn.docswell.com/page/GE5Y6W5QE4.jpg)

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


# Page. 39

![Page Image](https://bcdn.docswell.com/page/972GMQNW7R.jpg)

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


# Page. 40

![Page Image](https://bcdn.docswell.com/page/DJYVXWP9JM.jpg)

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


# Page. 41

![Page Image](https://bcdn.docswell.com/page/V7NM195DE8.jpg)

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


# Page. 42

![Page Image](https://bcdn.docswell.com/page/YJ9N52Y8J3.jpg)

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


# Page. 43

![Page Image](https://bcdn.docswell.com/page/GJ8KY56ZED.jpg)

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


# Page. 44

![Page Image](https://bcdn.docswell.com/page/LJLZ4Y917R.jpg)

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


# Page. 45

![Page Image](https://bcdn.docswell.com/page/47MVKNZ57W.jpg)

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


# Page. 46

![Page Image](https://bcdn.docswell.com/page/P7R2XD6ZJ9.jpg)

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


# Page. 47

![Page Image](https://bcdn.docswell.com/page/PJXD61M1EX.jpg)

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


# Page. 48

![Page Image](https://bcdn.docswell.com/page/3JKX12GMED.jpg)

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


# Page. 49

![Page Image](https://bcdn.docswell.com/page/LE35Q4N2J5.jpg)

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


# Page. 50

![Page Image](https://bcdn.docswell.com/page/8EDD9QW6EG.jpg)

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


# Page. 51

![Page Image](https://bcdn.docswell.com/page/V7PXZLXZE8.jpg)

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


# Page. 52

![Page Image](https://bcdn.docswell.com/page/2JVLWQLM7Q.jpg)

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


# Page. 53

![Page Image](https://bcdn.docswell.com/page/5EGXDWXXJL.jpg)

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


# Page. 54

![Page Image](https://bcdn.docswell.com/page/4JQ813857P.jpg)

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


# Page. 55

![Page Image](https://bcdn.docswell.com/page/K742R12VJ1.jpg)

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


# Page. 56

![Page Image](https://bcdn.docswell.com/page/LJ1MQGM4EG.jpg)

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


# Page. 57

![Page Image](https://bcdn.docswell.com/page/GJW3MK3ZJ2.jpg)

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


# Page. 58

![Page Image](https://bcdn.docswell.com/page/4EZVQZVLJ3.jpg)

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


# Page. 59

![Page Image](https://bcdn.docswell.com/page/Y7656Z5MEV.jpg)

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


# Page. 60

![Page Image](https://bcdn.docswell.com/page/G75Y6WYQ74.jpg)

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


# Page. 61

![Page Image](https://bcdn.docswell.com/page/9J2GMQGWJR.jpg)

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


# Page. 62

![Page Image](https://bcdn.docswell.com/page/DEYVXWV9EM.jpg)

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


# Page. 63

![Page Image](https://bcdn.docswell.com/page/VJNM19MD78.jpg)

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


# Page. 64

![Page Image](https://bcdn.docswell.com/page/YE9N52N8E3.jpg)

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


# Page. 65

![Page Image](https://bcdn.docswell.com/page/GE8KY5KZ7D.jpg)

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


# Page. 66

![Page Image](https://bcdn.docswell.com/page/LELZ4YZ1JR.jpg)

Fixstars Group
www.fixstars.com
Structured Sparsity
•
WMMA API ではまだ提供されていない?
•
•
•
ドキュメント中にそれらしい言及がない
ヘッダファイルにもそれらしい定義は見当たらない
CUTLASSではPTXを直書きしている
asm volatile(
&quot;mma.sp.sync.aligned.m16n8k32.row.col.f16.f16.f16.f16 {%0,%1}, &quot;
&quot;{%2,%3,%4,%5}, {%6,%7,%8,%9}, {%10,%11}, %12, 0x0;¥n&quot;
: &quot;=r&quot;(D[0]), &quot;=r&quot;(D[1])
: &quot;r&quot;(A[0]), &quot;r&quot;(A[1]), &quot;r&quot;(A[2]), &quot;r&quot;(A[3]), &quot;r&quot;(B[0]), &quot;r&quot;(B[1]),
&quot;r&quot;(B[2]), &quot;r&quot;(B[3]), &quot;r&quot;(C[0]), &quot;r&quot;(C[1]), &quot;r&quot;(E));
https://github.com/NVIDIA/cutlass/blob/v2.9.0/include/cutlass/arch/mma_sparse_sm80.h
66
Copyright © Fixstars Group


# Page. 67

![Page Image](https://bcdn.docswell.com/page/4JMVKNV5JW.jpg)

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


# Page. 68

![Page Image](https://bcdn.docswell.com/page/PJR2XD2ZE9.jpg)

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


# Page. 69

![Page Image](https://bcdn.docswell.com/page/PEXD61D17X.jpg)

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


# Page. 70

![Page Image](https://bcdn.docswell.com/page/3EKX12XM7D.jpg)

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


# Page. 71

![Page Image](https://bcdn.docswell.com/page/L735Q452E5.jpg)

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


# Page. 72

![Page Image](https://bcdn.docswell.com/page/87DD9QD67G.jpg)

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


# Page. 73

![Page Image](https://bcdn.docswell.com/page/VJPXZL6Z78.jpg)

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


# Page. 74

![Page Image](https://bcdn.docswell.com/page/2EVLWQ3MJQ.jpg)

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


# Page. 75

![Page Image](https://bcdn.docswell.com/page/57GXDW3XEL.jpg)

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


# Page. 76

![Page Image](https://bcdn.docswell.com/page/4EQ81355JP.jpg)

Fixstars Group
www.fixstars.com
グループの分割
tiled_partition
•
スレッドブロックまたはタイルをより細かいタイルに分割
cg::thread_block block = cg::this_thread_block();
auto tile32 = cg::tiled_partition&lt;32&gt;(block);
auto tile16 = cg::tiled_partition&lt;16&gt;(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


# Page. 77

![Page Image](https://bcdn.docswell.com/page/KJ42R19VE1.jpg)

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


# Page. 78

![Page Image](https://bcdn.docswell.com/page/LE1MQGK47G.jpg)

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


# Page. 79

![Page Image](https://bcdn.docswell.com/page/GEW3MKDZE2.jpg)

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


# Page. 80

![Page Image](https://bcdn.docswell.com/page/47ZVQZGLE3.jpg)

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


# Page. 81

![Page Image](https://bcdn.docswell.com/page/YJ656ZYM7V.jpg)

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


# Page. 82

![Page Image](https://bcdn.docswell.com/page/GJ5Y6WGQJ4.jpg)

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


# Page. 83

![Page Image](https://bcdn.docswell.com/page/LE35Q4RZJ5.jpg)

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


# Page. 84

![Page Image](https://bcdn.docswell.com/page/8EDD9QY4EG.jpg)

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


# Page. 85

![Page Image](https://bcdn.docswell.com/page/V7PXZL6VE8.jpg)

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


# Page. 86

![Page Image](https://bcdn.docswell.com/page/2JVLWQ3R7Q.jpg)

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


# Page. 87

![Page Image](https://bcdn.docswell.com/page/5EGXDW36JL.jpg)

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


# Page. 88

![Page Image](https://bcdn.docswell.com/page/4JQ813527P.jpg)

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


# Page. 89

![Page Image](https://bcdn.docswell.com/page/K742R19PJ1.jpg)

Fixstars Group
www.fixstars.com
グラフの構築
•
Graph API による構築
•
カーネルパラメータの渡し方は cudaLaunchKernel() と似た形
cudaGraph_t graph;
cudaGraphCreate(&amp;graph, 0);
// カーネルに渡すパラメータ (後述)
cudaKernelNodeParams params1, params2;
// ノードの生成
cudaGraphNode_t node1, node2;
cudaGraphAddKernelNode(&amp;node1, graph, nullptr, 0, &amp;params1);
cudaGraphAddKernelNode(&amp;node2, graph, nullptr, 0, &amp;params2);
// ノード間の依存関係の定義
cudaGraphAddDependencies(graph, &amp;node1, &amp;node2, 1);
89
Copyright © Fixstars Group


# Page. 90

![Page Image](https://bcdn.docswell.com/page/LJ1MQGKXEG.jpg)

Fixstars Group
www.fixstars.com
グラフの構築
•
Graph API による構築
•
カーネルパラメータの渡し方は cudaLaunchKernel() と似た形
__global__ void kernel(int x){ printf(&quot;%d¥n&quot;, x); }
int kernelParam1 = 42;
void *kernelParams[] = { static_cast&lt;void*&gt;(&amp;kernelParam1) };
cudaKernelNodeParams params;
params.func
= reinterpret_cast&lt;void*&gt;(&amp;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


# Page. 91

![Page Image](https://bcdn.docswell.com/page/GJW3MKDKJ2.jpg)

Fixstars Group
www.fixstars.com
グラフの構築
•
ストリームキャプチャによる構築
•
•
ストリームに対する操作から自動的にグラフを構築する
イベントの Record/Wait で他ストリームもキャプチャ対象になる
cudaGraph_t graph;
cudaGraphCreate(&amp;graph, 0);
cudaStream_t stream;
cudaStreamCreate(&amp;stream);
cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);
kernel&lt;&lt;&lt;1, 1, 0, stream&gt;&gt;&gt;(42);
kernel&lt;&lt;&lt;1, 1, 0, stream&gt;&gt;&gt;(0);
cudaStreamEndCapture(stream, &amp;graph);
91
Copyright © Fixstars Group


# Page. 92

![Page Image](https://bcdn.docswell.com/page/4EZVQZGNJ3.jpg)

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


# Page. 93

![Page Image](https://bcdn.docswell.com/page/Y7656ZY9EV.jpg)

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


# Page. 94

![Page Image](https://bcdn.docswell.com/page/G75Y6WGD74.jpg)

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


# Page. 95

![Page Image](https://bcdn.docswell.com/page/9J2GMQDMJR.jpg)

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


# Page. 96

![Page Image](https://bcdn.docswell.com/page/DEYVXW2PEM.jpg)

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


# Page. 97

![Page Image](https://bcdn.docswell.com/page/VJNM198M78.jpg)

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


# Page. 98

![Page Image](https://bcdn.docswell.com/page/YE9N52KWE3.jpg)

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


# Page. 99

![Page Image](https://bcdn.docswell.com/page/GE8KY5RR7D.jpg)

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


# Page. 100

![Page Image](https://bcdn.docswell.com/page/LELZ4YK2JR.jpg)

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


# Page. 101

![Page Image](https://bcdn.docswell.com/page/4JMVKN29JW.jpg)

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


# Page. 102

![Page Image](https://bcdn.docswell.com/page/PJR2XDM9E9.jpg)

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


# Page. 103

![Page Image](https://bcdn.docswell.com/page/PEXD61V37X.jpg)

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


# Page. 104

![Page Image](https://bcdn.docswell.com/page/3EKX12QN7D.jpg)

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


# Page. 105

![Page Image](https://bcdn.docswell.com/page/L735Q4PZE5.jpg)

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


