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

28.9K Views

October 29, 21

#cuda #gpu #cuda高速化 #高速化シリーズ #CUDA #GPU高速化 #並列計算 #ソフトウェア最適化 #フィックスターズ

スライド概要

GPUは幅広い産業分野において適用範囲が拡大し、近年では特に機械学習での活用が注目されています。

本ウェビナーでは、CUDAプログラミングモデルとハードウェアアーキテクチャについて基礎から解説を行い、高速化の実践例を紹介します。

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

@fixstars

スライド一覧

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

またはPlayer版

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

ダウンロード

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

関連スライド

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

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

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

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

arm neon

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

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

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

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

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

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

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

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

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

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

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

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

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

各ページのテキスト

Fixstars Corporation www.fixstars.com 本日のアジェンダ 15:00-15:10 • フィックスターズの紹介 15:10-16:20 • ＣUDA高速化入門 16:20-16:30 • • なぜGPUなのか？ • CUDAプログラミングモデル • ハードウェアアーキテクチャ • 高速化実践例 Q&A Copyright © Fixstars Corporation 3

Fixstars Corporation www.fixstars.com 会社概要グループ体制株式会社フィックスターズ会社名株式会社フィックスターズ本社所在地東京都港区芝浦3-1-1 msb Tamachi 田町ステーションタワーN 28階設立 2002年8月株式会社Fixstars Autonomous Technologies 株式会社ネクスティエレクトロニクスとのJV 自動運転向けソフトウェア開発に特化 Fixstars Solutions, Inc. 上場区分東証一部（証券コード：3687）代表取締役社長三木聡株式会社Sider 資本金 5億5,401万円（2021年3月現在）当社完全子会社ソースコードレビュー、監視・修正漏れ検知ツール開発社員数（連結） 253名（2020年9月現在）当社完全子会社米国での営業及び開発を担当株式会社Smart Opinion プロディジーメディカル株式会社とのJV 乳がんAI画像診断支援事業を担当キオクシア株式会社主なお客様株式会社日立製作所オスカーテクノロジー株式会社株式会社ネクスティエレクトロニクス連結子会社ソフトウェア自動並列化サービスを提供キヤノン株式会社 Copyright © Fixstars Corporation 5

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

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

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

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

10.

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

11.

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

12.

Fixstars Corporation www.fixstars.com GPU開発の課題と当社のサービス製品企画ハードウェア選定 • 原価はあまり上げたくないものの必要な演算リソースは確保したい • GPUはどんな計算を速くできる？ • アルゴリズムをカスタムしても性能が出るか？アルゴリズム設計・実装 • 前処理・後処理も速くできる？ • 機能を追加するため、もう少し処理を速くしたい性能・精度チューニング品質確保 • GPUで動かしてみたものの期待した性能が出ない • 品質確保のため、精度を上げたく演算量は増えるが性能は維持したい Copyright © Fixstars Corporation • GPU高速化に関するコンサルティング • ボトルネック解析 • アルゴリズムのGPU向け設計・実装 • GPUプログラム高速化 • CPU / GPU混在環境での高速化 • 精度向上 • CUDAプログラミングモデルの理解 • ハードウェアアーキテクチャの理解 12

13.

14.

Fixstars Corporation www.fixstars.com なぜGPUを使うのか • CPUと比べて • • ピーク性能の高さ電力効率の良さ浮動小数点数演算性能 CPU: AMD Ryzen 9 5950X GPU: NVIDIA GeForce RTX 3070 • メモリバンド幅 TDP 価格 2.25 [TFLOPS] 51.2 [GB/s] 105 [W] ¥90,000~ 20.31 [TFLOPS] 448.0 [GB/s] 220 [W] ¥93,000~ その他のアクセラレータと比べて • • 入手性・価格性能比の良さプログラミングの容易さ Copyright © Fixstars Corporation 14

15.

Fixstars Corporation www.fixstars.com なぜGPUが速いのか • 並列計算に特化した構成 • 大量のコア・演算器 • • • • CPU: AMD EPYC 7763: 64 Cores, 32 FLOPs/Core/cycle GPU: NVIDIA A100: 108 SMs, 128 FLOPs/SM/cycle バス幅の広い広帯域メモリもちろん弱点もある • • 並列に処理できない問題には弱い最大メモリ容量が小さい Copyright © Fixstars Corporation 15

16.

17.

[beta]

Fixstars Corporation

www.fixstars.com

例題: saxpy
•

Single-precision ax plus y

•

y←a×x+y

•

CPU向けの実装例:
void saxpy(float *y, const float *x, float a, int n){
for(int i = 0; i < n; ++i){
y[i] = a * x[i] + y[i];
}
}

Copyright © Fixstars Corporation

17

18.

19.

Fixstars Corporation www.fixstars.com ホストメモリとデバイスメモリ • CPUとGPUはそれぞれがメモリを持っている • • 目的に応じて適切なほうを利用する必要に応じて片方から他方へデータをコピーする～200 GB/s ホストメモリ (DDR) CPU ～20 GB/s ～2000 GB/s GPU デバイスメモリ (GDDR/HBM) Copyright © Fixstars Corporation 19

20.

Fixstars Corporation www.fixstars.com ホストメモリからデバイスメモリへデータを転送 • cudaMalloc • • • デバイスメモリ上の領域を確保標準Cにおけるmallocに対応 cudaMemcpy • • デバイスメモリに関係するメモリコピー第4引数で転送の方向を指定 (HostToDevice, DeviceToHost など) float *d_y, d_x; // デバイスメモリの確保 cudaMalloc(&d_x, sizeof(float) * n); cudaMalloc(&d_y, sizeof(float) * n); // ホストメモリ (h_x, h_y) から sizeof(float) * n バイト転送 cudaMemcpy(d_x, h_x, sizeof(float) * n, cudaMemcpyHostToDevice); cudaMemcpy(d_y, h_y, sizeof(float) * n, cudaMemcpyHostToDevice); Copyright © Fixstars Corporation 20

21.

Fixstars Corporation www.fixstars.com GPU上でカーネル（プログラム）を実行 • カーネルの呼び出し • • • スレッド数を指定するスレッドブロック数×ブロックあたりのスレッド数で表現ここではループ1回を1スレッドで処理する const int bdim = 128; const int gdim = (n + bdim – 1) / bdim; kernel<<<gdim, bdim>>>(d_y, d_x, a, n); Copyright © Fixstars Corporation // 切り上げ 21

22.

[beta]

Fixstars Corporation

www.fixstars.com

GPU上で動くカーネルの実装
•

__global__ 修飾された関数として定義

•

定義済み変数から自身のインデックスを取得
•
•
•

blockDim: 現在のカーネル実行におけるブロックサイズ
blockIdx: 自身の属するスレッドブロックのインデックス
threadIdx: 自身のスレッドブロック内におけるインデックス
__global__ void kernel(float *y, const float *x, float a, int n){
const int i = threadIdx.x + blockIdx.x * blockDim.x;
if(i < n)
y[i] = a * x[i] + y[i];
}

Copyright © Fixstars Corporation

22

23.

Fixstars Corporation www.fixstars.com デバイスメモリからホストメモリへデータを転送 • cudaMemcpyで逆方向にコピー // デバイスメモリ (d_y) から sizeof(float) * n バイト転送 cudaMemcpy(h_y, d_y, sizeof(float) * n, cudaMemcpyDeviceToHost); // デバイスメモリの解放 cudaFree(d_x); cudaFree(d_y); Copyright © Fixstars Corporation 23

24.

25.

Fixstars Corporation www.fixstars.com スレッドの階層構造 • CUDAではスレッド間に階層構造がある • Warp: 同時に命令が発行されるスレッドをまとめたもの • • Thread Block: いくつかのスレッドをまとめたもの • • • 現行アーキテクチャでは32スレッド現行アーキテクチャでは1ブロックあたり最大1024スレッド同一ワープに属するスレッドは必ず同一スレッドブロックに属する Grid: いくつかのスレッドブロックをまとめたもの • カーネル呼び出しは1つのグリッドで処理される Copyright © Fixstars Corporation 25

26.

Fixstars Corporation www.fixstars.com スレッドの階層構造 • 階層構造上で近いスレッド同士はより密に同期や通信を行うことができる • 同一グリッド • • 同一スレッドブロック • • • カーネル起動・終了時の同期のみ同じブロックに属するスレッド同士での同期シェアードメモリを用いたデータ共有同一ワープ • • スレッド同士でのより軽量な同期ワープシャッフルによるデータ共有 Copyright © Fixstars Corporation 26

27.

Fixstars Corporation www.fixstars.com メモリの階層構造 • メモリにも階層構造がある • おおむねスレッドの階層構造と対応 Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 27

28.

Fixstars Corporation www.fixstars.com メモリの階層構造: レジスタ • プログラム中の自動変数に対応 • 各種演算命令に直接渡すことができる • 他のスレッドとは共有されない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 28

29.

Fixstars Corporation www.fixstars.com メモリの階層構造: ローカルメモリ • プログラム中の自動変数に対応 • 何らかの理由でレジスタに乗せられないときに使用される • 演算命令に渡す際はいったんレジスタにロードする必要がある • 他のスレッドとは共有されない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 29

30.

Fixstars Corporation www.fixstars.com メモリの階層構造: シェアードメモリ • __shared__ 修飾された変数に対応 • 同一スレッドブロック内の全スレッドで共有される Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 30

31.

Fixstars Corporation www.fixstars.com メモリの階層構造: グローバルメモリ • cudaMalloc などで確保された領域に対応 • デバイス全体で共有される • カーネル停止後も値が保持される Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 31

32.

Fixstars Corporation www.fixstars.com メモリの階層構造: コンスタントメモリ • __constant__ 修飾された変数に対応 • デバイス全体で共有される • カーネルから値を書き換えることができない Grid Thread Block Thread Registers Local Memory Shared Memory Global Memory Copyright © Fixstars Corporation Constant Memory 32

33.

Fixstars Corporation www.fixstars.com ホストとデバイス間の同期 • カーネル呼び出しやデータ転送は基本的に非同期実行 • 明示的もしくは暗黙的に同期を挿入する必要がある • cudaMemcpy など一部のAPIは自動的に同期を挿入する kernel<<<1, 1>>>(); // この時点では kernel() はまだ実行されていないかもしれない foo(); cudaDeviceSynchronize(); // この時点では kernel() の処理は確実に完了している CPU GPU foo() cudaDeviceSynchronize() kernel() Copyright © Fixstars Corporation 33

34.

Fixstars Corporation www.fixstars.com ストリーム • デバイスで実行される処理のキュー • • • 投入した順に処理される同じストリームに投入された処理同士はオーバーラップしない指定されなかった場合はデフォルトストリームが使用される kernel1<<<1, 1>>>(); kernel2<<<1, 1>>>(); cudaDeviceSynchronize(); CPU GPU cudaDeviceSynchronize() kernel1() Copyright © Fixstars Corporation kernel2() 34

35.

Fixstars Corporation www.fixstars.com ストリーム • ストリームは複数作成することができる • 別ストリームに投入された処理同士は並行するかもしれない kernel1<<<1, 1, 0, stream1>>>(); kernel2<<<1, 1, 0, stream2>>>(); cudaDeviceSynchronize(); CPU Synchronize GPU kernel1() kernel2() Copyright © Fixstars Corporation 35

36.

Fixstars Corporation www.fixstars.com プログラミングモデルまとめ • 大量のスレッドの間には階層関係がある • • • メモリにも階層関係がある • • • ワープ・スレッドブロック・グリッド距離に応じて同期や通信の制約が変化するレジスタ・ローカルメモリ・シェアードメモリ・グローバルメモリ速度や共有する必要があるスレッド数など要求に応じて適切な領域を使い分けるデバイス上で動く処理は基本的に非同期実行となる • 細かい同期周りの制御にはストリームを活用する Copyright © Fixstars Corporation 36

37.

38.

Fixstars Corporation www.fixstars.com カーネルが遅い原因と対応策要求されている演算量が多すぎる • アルゴリズムを改善して演算量を減らすメモリトラフィックが多すぎる • アルゴリズムを改善してメモリアクセスを減らす • キャッシュなどのハードウェア機能を活用するリソースを有効活用できていない • ハードウェアの制約を理解して演算器やバスなどの稼働率を上げる Copyright © Fixstars Corporation 38

39.

40.

Fixstars Corporation www.fixstars.com NVIDIA A100 Block Diagram • CC 8.0 • 108 SMs/Chip • 6912 FP32 CUDA Cores • コアを活用できるだけの並行実行可能なタスク (=スレッド) を投入する必要がある • スレッド数が足りないならタスクを分割することも視野に入れる https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf Copyright © Fixstars Corporation 41

https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf

41.

Fixstars Corporation www.fixstars.com Streaming Multiprocessor (SM) • スレッドブロックに対応する • • いくつかのスレッドブロックを並行して処理以下の要素を束ねたもの • • • • • • • CUDA Core Tensor Core LD/ST Unit SFU Register File Cache/Shared Memory Scheduler, Dispatcher Copyright © Fixstars Corporation 42

42.

Fixstars Corporation www.fixstars.com Streaming Multiprocessor (SM) • スレッドブロックに対応する • 以下の要素を束ねたもの • • • • • • • CUDA Core Tensor Core LD/ST Unit SFU Register File Cache/Shared Memory Scheduler, Dispatcher 演算器メモリ Copyright © Fixstars Corporation 43

43.

44.

Fixstars Corporation www.fixstars.com CUDA Core • スレッドに対応する • 何らかの演算を行う • • • レジスタファイルから値を読んで演算を行ってレジスタファイルに書き出す FP32/INT32 • Volta以降でINTコアが分離された • 整数演算と浮動小数点数演算を同時に実行できる Copyright © Fixstars Corporation 45

45.

46.

47.

Fixstars Corporation www.fixstars.com SIMTとWarp • ディスパッチャはワープに対して一つの命令を一度に発行する • • • SIMT: Single Instruction, Multiple Threads スレッドごとに異なる命令を発行することはできない条件分岐の取り扱い • • 分岐によって実行の必要がなくなった命令も発行されうるそのような場合はその命令が無視される Copyright © Fixstars Corporation 49

48.

Fixstars Corporation www.fixstars.com Warp Divergence • 条件分岐によって有効な演算を行わないスレッド (=コア) が発生する • ワープ内での異なる方向への分岐は性能劣化につながる • • Warp Divergence と呼ぶ下の例では B(), C() の処理中にコアが半分遊んでいる A() A(); if(threadIdx.x % 2 == 0){ B(); }else{ C(); } B() C() Warp 0 Warp 1 Copyright © Fixstars Corporation 50

49.

Fixstars Corporation www.fixstars.com Warp Divergence • できるだけ同じワープのスレッドが同じように動くことで効率を改善できる • • 連続するスレッドが同じ方向に分岐するようにする下の例では B(), C() におけるコア稼働率が改善している A() A(); if(threadIdx.x < 4){ B(); }else{ C(); } B() C() Warp 0 Warp 1 Copyright © Fixstars Corporation 51

50.

51.

Fixstars Corporation www.fixstars.com ワープスケジューリング • Processing Block はいくつかの実行中ワープの状態を保持している • • サイクルごとに実行可能なワープをその中から選択して命令を発行する • • 可能であれば物理コア数より多くのスレッドの状態を保持する実行可能: 次に発行される命令が依存している処理がすべて完了している実行可能なワープを絶やさないことが効率改善につながる • • 命令のレイテンシを考慮したプログラムを記述する実行可能なワープの候補 (=状態を保持しているスレッド数) を増やす Copyright © Fixstars Corporation 53

52.

Fixstars Corporation www.fixstars.com レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • • FADDのレイテンシは4とする並行実行しているワープ数が1の場合: 4 ops / 16 cycles 0x00: 0x01: 0x02: 0x03: FADD FADD FADD FADD R1, R1, R1, R1, R2 R3 R4 R5 Warp 0 R1 += R2 R1 += R3 Copyright © Fixstars Corporation R1 += R4 R1 += R5 54

53.

Fixstars Corporation www.fixstars.com レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • FADDのレイテンシは4とする • 並行実行しているワープ数が1の場合: 4 ops / 16 cycles • 並行実行しているワープ数が4の場合: 16 ops / 19 cycles 0x00: 0x01: 0x02: 0x03: FADD FADD FADD FADD R1, R1, R1, R1, R2 R3 R4 R5 Warp 0 Warp 1 Warp 2 Warp 3 R1 += R2 R1 += R3 R1 += R2 R1 += R4 R1 += R3 R1 += R2 R1 += R4 R1 += R3 R1 += R2 Copyright © Fixstars Corporation R1 += R5 R1 += R3 R1 += R5 R1 += R4 R1 += R4 R1 += R5 R1 += R5 55

54.

Fixstars Corporation www.fixstars.com Occupancy • SMがいくつのワープを並行実行できるかを表す指標 • • ブロックサイズ・消費レジスタ数・シェアードメモリサイズから求める • • • • 高ければ高いほどレイテンシを隠蔽しやすいブロックサイズ: SMあたりの並行実行可能なブロック数消費レジスタ数: SMあたりのレジスタファイル数シェアードメモリサイズ: SMあたりのシェアードメモリサイズプロファイラ・CUDA Toolkit 付属のExcelシートなどで求められる Copyright © Fixstars Corporation 56

55.

56.

57.

58.

59.

Fixstars Corporation www.fixstars.com Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 32トランザクション / 1アクセススレッドメモリ 0 1 … 2 3 … 4 … … 5 … 31 … … Copyright © Fixstars Corporation … 61

60.

61.

Fixstars Corporation www.fixstars.com L1キャッシュ • SMごとに用意されている • 明示的に指定したものか読み取り専用のデータへのアクセスに対して使用される読み取り専用かどうかの判定 • コンパイラが判定する • ポインタを const __restrict__ 修飾すると読み取り専用であることを明示できる明示的なL1キャッシュの利用 • 組み込み関数 __ldg() を使用する • *ptr → __ldg(ptr) Copyright © Fixstars Corporation 63

62.

63.

Fixstars Corporation www.fixstars.com メモリバンク • シェアードメモリはバンクを用いてに管理されている • • バンクは4バイトごとに切り替わる同じバンクの異なる領域へのアクセスはまとめて処理できない: バンクコンフリクト Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 … 0x00000100 Copyright © Fixstars Corporation 65

64.

Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • 素直なシーケンシャルアクセス Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 66

65.

Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • バンクが重複しないランダムアクセス Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 67

66.

Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できるアクセスの例 • ブロードキャスト: バンクが重なっても同じアドレスなら問題ない Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 68

67.

Fixstars Corporation www.fixstars.com メモリバンク • まとめて処理できないアクセスの例 • ストライドアクセス: この場合は2回に分割される Bank 27 Bank 28 Bank 29 Bank 30 Bank 31 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 27 28 29 30 31 0x00000000 0x00000080 … 0x00000100 スレッド 0 1 2 3 4 Copyright © Fixstars Corporation 69

68.

Fixstars Corporation www.fixstars.com レジスタファイル • プロセッシングブロックごとに用意された領域 • レジスタ幅は1要素あたり32bit • long, double, ポインタなどの64bit型には2つ使われる自動変数に対する領域割り当て • 自動変数は可能ならレジスタに割り当てられる • 特定のケースで低速なローカルメモリに割り当てられる • • 自動変数がレジスタに収まりきらない場合（レジスタスピル）インデックスアクセスが必要な場合 Copyright © Fixstars Corporation 70

69.

Fixstars Corporation www.fixstars.com ハードウェアまとめ演算器 • 演算器を使い切るためには注意が必要なことがある • • 分岐によって何もしないコアが発生することがあるレイテンシを埋めるだけの命令供給が必要メモリ • アクセスパターン次第で効率が落ちることがある • • • グローバルメモリ: Coalescing, キャッシュ利用シェアードメモリ: バンクコンフリクトローカルメモリの利用にも注意する • コンパイラの出力を確認すると確実 Copyright © Fixstars Corporation 71

70.

71.

Fixstars Corporation www.fixstars.com 問題の概要 • 画像のステレオマッチング: Semi-Global Matching (SGM) • ステレオ画像の視差を計算するアルゴリズム • • • 視差: 片方の画像のある画素が他方の画像で何ピクセルずれたところにあるか近くの物体ほど視差が大きくなること利用して距離を計算できるターゲット環境: Pascal 世代のGPU (GeForce GTX 10xx など) Copyright © Fixstars Corporation 73

72.

Fixstars Corporation www.fixstars.com チューニングする部分 • 動的計画法 (DP) である画素における視差が d [px] としたときのスコアを求める • • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい隣接画素に対して急激な視差の変化があるとコストが大きい 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 𝑠𝑐𝑜𝑠𝑡 𝑥, 𝑑 = 𝑙𝑐𝑜𝑠𝑡 𝑥, 𝑑 + min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 𝑑 − min 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑖) 𝑖 𝑥 Copyright © Fixstars Corporation 74

73.

Fixstars Corporation www.fixstars.com チューニングする部分 • 動的計画法である画素における視差が d [px] としたときのスコアを求める • • • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい隣接画素に対して急激な視差の変化があるとコストが大きいこれを縦横斜めの8方向それぞれについてラインごとに計算する Copyright © Fixstars Corporation 75

74.

[beta]

Fixstars Corporation

www.fixstars.com

アルゴリズムの概略
左から右方向のスキャン

for(int y = 0; y < H; ++y){
int prev_min = 0;
for(int x = 0; x < W; ++x){
int cur_min = INT_MAX;
for(int d = 0; d < D; ++d){
X方向のループは依存性がある
int cost = min({
P2,
scost[y][x-1][d-1] - prev_min + P1,
主要な計算は O(HWD) 回行われる
scost[y][x-1][d+1] - prev_min + P1,
計算処理はかなり軽い
scost[y][x-1][d]
- prev_min
});
scost[y][x][d] =
キャッシュヒットが期待できない
cost + dist(left[y][x], right[y][x-d]);
メモリアクセスも O(HWD) 回
cur_min = min(prev_min, cost);
}
prev_min = cur_min;
}
Y方向のループは完全に独立している
}

•

Copyright © Fixstars Corporation

76

75.

Fixstars Corporation www.fixstars.com 並列化方針の検討 y方向ループ1回を1スレッドで担当する • 並列度が足りない: 数百スレッド程度しか利用できない y方向ループ1回を複数スレッドで担当する • x方向のループは分割できない: 前のループに対する依存性があるため • d方向のループは分割可能 • ただしx方向のループ1回ごとに同期が必要になる Copyright © Fixstars Corporation 77

76.

77.

Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • • d方向の最小値を求める端の値を隣のスレッドに渡す最小値の計算・共有 𝑑 Thread 1 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 Thread 0 𝑥 Copyright © Fixstars Corporation 79

78.

Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • • d方向の最小値を求める端の値を隣のスレッドに渡す端の値の共有 𝑑 Thread 1 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃 1 min 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 𝑖 Thread 0 𝑥 Copyright © Fixstars Corporation 80

79.

Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • • 1ワープ以内になるとより軽量な通信が利用できるグローバルメモリへのアクセス効率の向上 • • スレッドあたりのメモリアクセス量が多くなる 1回のアクセスで4要素までアクセスできる Copyright © Fixstars Corporation 81

80.

Fixstars Corporation www.fixstars.com 並列化方針の検討: d方向ループの分割細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • • 1ワープ以内になるとより軽量な通信が利用できるグローバルメモリへのアクセス効率の向上 1ワープを境に実装が大きく変化する • スレッドあたりのメモリアクセス量が多くなる ⇒ 1ワープ以下の範囲で値を変えつつ試せるように実装する • 1回のアクセスで4要素までアクセスできる Copyright © Fixstars Corporation 82

81.

82.

Fixstars Corporation www.fixstars.com アルゴリズムの検討:局所特徴同士の距離の計算 • 局所特徴についての情報 • • • 局所特徴の表現: 64 bit のビット列局所特徴の距離: 互いに異なるビットの数特徴ベクトルの距離は同じ組の距離が何度も使われる • • 8方向すべての処理で同じ計算を行う既存実装では事前計算してテーブル化されていた • table[y][x][d] = distance(left[y][x], right[y][x - d]) Copyright © Fixstars Corporation 84

83.

Fixstars Corporation www.fixstars.com 理論性能で比べる • テーブル引きと計算どちらが速い? • • • • popcounts/s: 763 [Gops/s] • • • 32 [ops/s/SM] × 28 [SM/s] × 1.481 [GHz] = 1326 [Gops/s] 1要素あたり2回必要なのでその半分 Bytes/s: 484.4 [GB/s] • • テーブル化した場合1要素当たり 1 [byte] 特徴同士の距離は popcount 命令2回で求められる GeForce GTX 1080 Ti (sm_61) を例に試算してみる実測値だとおよそ 340 [GB/s] くらい毎回計算するほうが速そう!! Copyright © Fixstars Corporation 85

84.

85.

Fixstars Corporation www.fixstars.com プロファイル結果 (2) • 本当に距離をテーブル化しないほうが速かったのか? • • 実効メモリ帯域で評価するテーブル引きする場合はメモリトラフィックが Reads = Writes になる • • テーブルサイズが結果バッファのサイズと等しいため 52.049×2 = 104.098 [GB/s] 出せなければテーブル化のほうが遅い • • bandwidthTest での帯域が 92.7 [GB/s] 程度テーブル化する方針では勝てないだろうと考えられる Copyright © Fixstars Corporation 87

86.

87.

Fixstars Corporation www.fixstars.com 全体の評価 • 既存実装との性能比較 • • • 比較対象: Embedded real-time stereo estimation via Semi-Global Matching on the GPU, D. Hernandez-Juarez et al, ICCS 2016. https://github.com/dhernandez0/sgm 実際にはもう一つ大きいカーネルがあるのですがそちらの詳細は省略しています Copyright © Fixstars Corporation 89

https://github.com/dhernandez0/sgm

88.

Fixstars Corporation www.fixstars.com 評価結果 • 2.3-12.5% 程度の高速化 • 演算性能に対してメモリ帯域の細いチップで特に強いフレームレート 300 250 261 232 200 150 今回の実装 100 Hernandez+ 69.7 68.1 50 50.2 45.8 0 GTX 1080 Ti GTX 1050 Ti Copyright © Fixstars Corporation DRIVE PX2 90

89.

90.

Fixstars Corporation www.fixstars.com 全体のまとめ • パフォーマンスチューニングにおいてはハードウェアの知識も重要 • • • カーネルのチューニングにおいては特に演算器とメモリに気を配る • • • 使い方を誤ると数倍の性能劣化なども起こりうるもちろんアルゴリズムも重要で両方からのアプローチが必要演算器を余らせない不得意なアクセスパターンによる性能劣化を防ぐ理論をもとに仮説を立てて実装したものを評価する • • プロファイラによる評価理論ピークと実性能の差を読み取る Copyright © Fixstars Corporation 92

91.

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

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

関連スライド

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

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

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

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

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

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

各ページのテキスト

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