CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

27K Views

June 23, 22

#gpu #cuda #cuda高速化 #高速化シリーズ #cuda高速化シリーズ #Fixstars Group #CUDA アーキテクチャ #ソフトウェア高速化サービス #画像処理 #AI・深層学習

スライド概要

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

株式会社フィックスターズ

@fixstars

スライド一覧

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

またはPlayer版

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

ダウンロード

ダウンロード(pdf - 3.02MB)

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

kaggle 画像処理機械学習深層学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 63.6K

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

arm neon

株式会社フィックスターズ 38.6K

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

ros2 gpu ロボット自律走行 ros2シリーズ高速化シリーズ

株式会社フィックスターズ 30.8K

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

kaggle 画像処理機械学習ディープラーニングコンペ初心者 kaggleスコアアップシリーズ

株式会社フィックスターズ 30K

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

cuda gpu cuda高速化高速化シリーズ

株式会社フィックスターズ 29K

ROS2自律走行実現に向けて 1 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/10/17）

ros2 自律走行ロボット gpu ros2シリーズ高速化シリーズ

株式会社フィックスターズ 22.7K

各ページのテキスト

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

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

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

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

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

https://www.cs.toronto.edu/~frossard/post/vgg16/

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.

13.

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.

17.

18.

19.

20.

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

https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf

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

https://images.nvidia.com/aem-dam/en-zz/Solutions/geforce/ampere/pdf/NVIDIA-ampere-GA102-GPU-Architecture-Whitepaper-V1.pdf

26.

27.

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.

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.

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.

48.

49.

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

https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf

50.

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.

58.

59.

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.

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

https://arxiv.org/abs/1506.02626

65.

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

https://github.com/NVIDIA/cutlass/blob/v2.9.0/include/cutlass/arch/mma_sparse_sm80.h

67.

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

68.

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.

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.

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.

80.

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.

84.

85.

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.

88.

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.

94.

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.

103.

104.

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

105.

CUDA高速化セミナー vol.2 ～CUDAアーキテクチャの進化～（2022/06/23）

株式会社フィックスターズ

関連スライド

Kaggleスコアアップセミナー～画像系コンペ入門[前編]（2023/08/02）

いまさら聞けない ARM を使ったNEONの基礎と活用事例（2021/08/05）

ROS2自律走行実現に向けて 2 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/11/30）

Kaggleスコアアップセミナー～画像系コンペ入門[後編]（2023/09/26）

いまさら聞けない！CUDA高速化入門 ～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）

ROS2自律走行実現に向けて 1 ～次世代ロボット開発フレームワークROS2のビルドシステム徹底理解～（2022/10/17）

各ページのテキスト

いまさら聞けない！CUDA高速化入門～プログラミングモデルとアーキテクチャの解説、高速化の実践～（2021/10/29）