いまさらきけないCUDA高速化（2024/12/19）

10.4K Views

December 24, 24

#cuda #gpu #llm #nvidia #高速化 #CUDA #GPU #並列計算 #ソフトウェア開発

スライド概要

生成AIをはじめ、近年では様々な分野でGPUの活用が進んでいます。

開発環境やライブラリも充実してきており、GPUによる並列計算の恩恵を受けやすくなってきています。しかしながら、GPU性能をさらに引き出すためには、その動作原理を深く理解する必要があります。

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

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

@fixstars

スライド一覧

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

またはPlayer版

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

（ダウンロード不可）

関連スライド

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

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

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

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

arm neon

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

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

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

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

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

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

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

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

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

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

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

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

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

各ページのテキスト

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

フィックスターズの強みコンピュータの性能を最大限に引き出す、ソフトウェア高速化のエキスパート集団ハードウェアの知見アルゴリズム実装力各産業・研究分野の知見目的の製品に最適なハードウェアを見抜き、その性能をフル活用するソフトウェアを開発します。ハードウェアの特徴と製品要求仕様に合わせて、アルゴリズムを改良して高速化を実現します。開発したい製品に使える技術を見抜き、実際に動作する実装までトータルにサポートします。 Copyright © Fixstars Group 4

サービス概要お客様専任のエンジニアが直接ヒアリングを行い、高速化を実現するために乗り越えるべき課題や問題を明確にしていきます。高速化のワークフローお客様オリジナルソースコードのご提供高速化したコードコンサルティング高速化サポート先行技術調査アルゴリズムの改良・開発レポートやコードへのQ&A 性能評価・ボトルネックの特定ハードウェアへの最適化実製品への組込み支援レポート作成 Copyright © Fixstars Group 5

サービス提供分野半導体産業機器金融自動車 ● NAND型フラッシュメモリ向けファームウェア開発 ● 次世代AIチップの開発環境基盤生命科学 ● Smart Factory実現への支援 ● マシンビジョンシステムの高速化 ● 自動運転の高性能化、実用化 ● ゲノム解析の高速化 ● 次世代パーソナルモビリティの研究開発 ● 医用画像処理の高速化 Copyright © Fixstars Group ● デリバティブシステムの高速化 ● HFT(アルゴリズムトレード)の高速化 ● AI画像診断システムの研究開発 6

画像処理アルゴリズム開発高速な画像処理需要に対して、経験豊富なエンジニアが責任を持って製品開発をご支援します。お客様の課題ご支援内容高度な画像処理や深層学習等のアルゴリズムを開発できる人材が社内に限られているアルゴリズム調査・改変課題に合ったアルゴリズム・実装手法を調査製品実装に向けて適切な改変を実施機能要件は満たせそうだが、ターゲット機器上で性能要件までクリアできるか不安深層学習ネットワーク精度の改善様々な手法を駆使して深層学習ネットワークの精度を改善製品化に結びつくような研究ができていない論文調査・改善活動論文調査から最先端の手法の探索性能向上に向けた改善活動を継続 Copyright © Fixstars Group 8

10.

GPU向け高速化高性能なGPUの本来の性能を十分に引き出し、ソフトウェアの高速化を実現します。お客様の課題ご支援内容 GPUで計算してみたが期待した性能が出ない GPU高速化に関するコンサルティング GPU/CPUを組み合わせた全体として最適な設 CPU・GPU混在環境でのシステム設計計がしたいアルゴリズムのGPU向け移植原価を維持したまま機能を追加するため、もう少し処理を速くしたい GPUプログラム高速化品質確保のため、精度を上げたく演算量は増えるが性能は維持したい Copyright © Fixstars Group 継続的な精度向上 9

11.

12.

なぜGPUを使うのか • CPUと比べて • • ピーク性能の高さ電力効率の良さ浮動小数点数演算性能メモリバンド幅 CPU: AMD Ryzen 9 9950X 5.84 [TFLOPS] GPU: NVIDIA GeForce RTX 4070 SUPER 35.48 [TFLOPS] • TDP 価格 87.5 [GB/s] 170 [W] ¥109,000~ 504.2 [GB/s] 220 [W] ¥95,000~ その他のアクセラレータと比べて • • 入手性・価格性能比の良さプログラミングの容易さ Copyright © Fixstars Group 11

13.

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

14.

15.

[beta]

例題: 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 Group

14

16.

17.

18.

ホストメモリからデバイスメモリへデータを転送 • 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 Group 17

19.

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 Group // 切り上げ 18

20.

[beta]

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 Group

19

21.

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

22.

23.

スレッドの階層構造 • CUDAではスレッド間に階層構造がある • Warp: 同時に命令が発行されるスレッドをまとめたもの • • Thread Block: いくつかのスレッドをまとめたもの • • • 現行アーキテクチャでは1ブロックあたり最大1024スレッド同一ワープに属するスレッドは必ず同一スレッドブロックに属する Grid: いくつかのスレッドブロックをまとめたもの • • 現行アーキテクチャでは32スレッドカーネル呼び出しは1つのグリッドで処理される階層構造上で近いスレッド同士はより密に同期や通信を行うことができる Copyright © Fixstars Group 22

24.

同期・通信: 同一グリッド内スレッド間の同期 • cooperative_group::sync(grid_group) による同期 • 制約に注意 • • • グリッド中のすべてのスレッドが並行実行されている必要があるカーネル起動時に Cooperative Launch API を使用する必要があるコストも大きいためグリッド単位の同期は避けたほうが良いことが多いスレッド間のデータ交換 • グローバルメモリを用いたデータ共有 Copyright © Fixstars Group 24

25.

同期・通信: 同一スレッドブロック内スレッド間の同期 • cooperative_group::sync(grid_group) による同期 • __syncthreads() による同期スレッド間のデータ交換 • グローバルメモリを用いたデータ共有 • シェアードメモリを使ったデータ共有 • グローバルメモリよりはかなり高速にやり取りができる (詳細は後述) Copyright © Fixstars Group 25

26.

同期・通信: 同一ワープ内スレッド間の同期 • cooperative_group::sync(grid_group) による同期 • __syncthreads() による同期 • __syncwarp() による同期スレッド間のデータ交換 • グローバルメモリを用いたデータ共有 • シェアードメモリを使ったデータ共有 • Warp Shuffle を用いたデータ交換 • • レジスタからレジスタに直接値を渡す上の2つよりさらに低コスト Copyright © Fixstars Group 26

27.

28.

29.

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

30.

31.

32.

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

33.

34.

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

35.

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

36.

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

37.

ストリーム • ストリームとホスト間で同期をとることもできる • ストリームごとに別のタイミングで同期をとることができる • デバイス全体での同期よりこちらを使うほうが便利なことが多い kernel1<<<1, 1, 0, stream1>>>(); kernel2<<<1, 1, 0, stream2>>>(); cudaStreamSynchronize(stream1); foo(); CPU Synchronize GPU kernel1() foo() kernel2() Copyright © Fixstars Group 37

38.

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

39.

40.

41.

42.

Compute Capability (CC) • デバイスの仕様を表す値 • • • • おおむね大きいほど新しい新しいものがそれ以前の機能をすべて含むとは限らない CCが同じであればチップの規模が違うのみ（コア数・メモリ帯域など）デバイスと対応するCCの例: CC アーキテクチャデバイスの例 7.0 Volta Tesla V100 など 7.5 Turing GeForce RTX 20xx など 8.0 Ampere NVIDIA A100 など 8.6 Ampere GeForce RTX 30xx など 8.9 Ada Lovelace GeForce RTX 40xx など 9.0 Hopper NVIDIA H100 など Copyright © Fixstars Group 42

43.

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 Group 43

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

44.

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

45.

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

46.

47.

48.

49.

50.

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

51.

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

52.

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

53.

54.

55.

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

56.

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

57.

レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • FADDのレイテンシは4とする • 並行実行しているワープ数が1の場合: 4 ops / 16 cycles • 並行実行しているワープ数が4の場合: 16 ops / 19 cycles 0x00: FADD R1, R2 0x01: FADD R1, R3 0x02: FADD R1, R4 0x03: FADD R1, 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 Group R1 += R5 R1 += R3 R1 += R5 R1 += R4 R1 += R4 R1 += R5 R1 += R5 58

58.

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

59.

60.

61.

62.

63.

64.

65.

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

66.

67.

メモリバンク • シェアードメモリはバンクを用いてに管理されている • • バンクは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 Group 68

68.

メモリバンク • まとめて処理できるアクセスの例 • 素直なシーケンシャルアクセス 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 Group 69

69.

メモリバンク • まとめて処理できるアクセスの例 • バンクが重複しないランダムアクセス 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 Group 70

70.

メモリバンク • まとめて処理できるアクセスの例 • ブロードキャスト: バンクが重なっても同じアドレスなら問題ない 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 Group 71

71.

メモリバンク • まとめて処理できないアクセスの例 • ストライドアクセス: この場合は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 Group 72

72.

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

73.

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

74.

75.

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

76.

77.

78.

アルゴリズムの概略左から右方向のスキャン 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 Group 79

79.

80.

81.

82.

83.

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

84.

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

85.

86.

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

87.

理論性能で比べる • テーブル引きと計算どちらが速い? • • • • 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 Group 88

88.

89.

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

90.

91.

全体の評価 • 既存実装との性能比較 • • • 比較対象: 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 Group 92

https://github.com/dhernandez0/sgm

92.

93.

94.

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

95.