いまさら聞けない！CUDA高速化入門

2021年10月29日開催「いまさら聞けない！CUDA高速化入門」セミナー資料です。

Transcript

1. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation いまさら聞けない！ CUDA高速化入門
2. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 本日のアジェンダ • フィックスターズの紹介 • ＣUDA高速化入門 • なぜGPUなのか？ • CUDAプログラミングモデル • ハードウェアアーキテクチャ • 高速化実践例 • Q&A 2 15:00-15:10 15:10-16:20 16:20-16:30
3. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation フィックスターズのご紹介
4. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 4 会社概要 会社名 株式会社フィックスターズ 本社所在地 東京都港区芝浦3-1-1 msb Tamachi 田町ステーションタワーN 28階 設立 2002年8月 上場区分 東証一部（証券コード：3687） 代表取締役社 長 三木 聡 資本金 5億5,401万円（2021年3月現在） 社員数（連 結） 253名（2020年9月現在） 主なお客様 キオクシア株式会社 株式会社日立製作所 株式会社ネクスティ エレクトロニクス キヤノン株式会社 グループ体制 Fixstars Solutions, Inc. 株式会社Fixstars Autonomous Technologies 株式会社ネクスティ エレクトロニクスとのJV 自動運転向けソフトウェア開発に特化 当社完全子会社 米国での営業及び開発を担当 株式会社Sider 株式会社Smart Opinion 当社完全子会社 ソースコードレビュー、監視・修正漏れ検知ツール開発 プロディジーメディカル株式会社とのJV 乳がんAI画像診断支援事業を担当 オスカーテクノロジー株式会社 連結子会社 ソフトウェア自動並列化サービスを提供 株式会社フィックスターズ
5. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 5 フィックスターズは、コンピュータの性能を最大限に引き出し大量データの高速処理を実現する、 高速化のエキスパート集団です。 低レイヤ ソフトウェア技術 アルゴリズム 実装力 各産業・研究 分野の知見 フィックスターズの強み
6. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ソフトウェア高速化サービス概要 6 お客様のソースコードをご提供いただき、 ソフトウェアの最適化やアルゴリズムの改良を行い高速化したコードをお返しします。 当社 お客様 オリジナルソースコードのご提供 高速化したソースコード コンサルティング 高速化 サポート 性能評価 ボトルネックの特定 アルゴリズムの改良・開発 ハードウェアへの最適化 レポート作成 レポートやコードへのQ&A 実製品への組込み支援
7. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 7 ソフトウェア高速化サービス領域 大量データの高速処理がお客様の製品競争力の源泉となる、 様々な領域でソフトウェア開発・高速化サービスを提供しています。 ・NAND型フラッシュメモリ向けファー ムウェア開発 ・次世代AIチップ向け開発環境基盤開発 Semiconductor ・デリバティブシステムの高速化 ・HFT(アルゴリズムトレード)の高速化 Finance ・自動運転の高性能化、実用化 ・次世代パーソナルモビリティの研究開発 Mobility ・ゲノム解析の高速化 ・医用画像処理の高速化 ・AI画像診断システムの研究開発 Life Science ・Smart Factory化支援 ・マシンビジョンシステムの高速化 Industrial
8. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 画像処理・アルゴリズム開発サービス 8 • お客様の課題 • 高度な画像処理、深層学習等のアルゴリズム開発を行える人材が社内に限られている • 考案中のアルゴリズムで機能要件は満たせそうだが、ターゲット機器上で性能要件まで クリアできるか不安 • 研究開発の成果が製品化にうまく結びつかない • 弊社の支援内容 • 課題に応じたアルゴリズム調査 • 深層学習ネットワーク精度改善、推論高速化手法調査 • 論文調査、実装 https://www.cs.toronto.edu/~frossard/post/vgg16/
9. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation AI・深層学習関連サービス 9 • ディープラーニングの包括的開発技術 • ネットワーク設計からターゲットデバイスでの高速化のノウハウ • 大規模システムからエッジコンピューティングまでの開発実績 ネットワーク設計 データの前処理、データ拡張 精度改善 分散処理による学習高速化 各種DLフレームワーク クラウド・サーバ エッジ モデル圧縮 - 量子化 - 枝刈り - 蒸留 ターゲットデバイスへの ポーティング及び推論高速化 ▪ Visconti, ARM, GPU, DSP ▪ SIMD,NEON,CUDA,TensorRT
10. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation GPU向け高速化サービス 10 • お客様の課題 • GPU 高速化の知見がない • 自力で GPU に乗せてみたものの望む性能が出ない • 弊社の支援内容 • GPU 高速化に関するコンサルティング • ボトルネック調査、GPU プログラムの高速化 • CPU/GPU が混在するヘテロジニアス環境での最適化 10～150 倍の 高速化事例あり
11. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation GPU開発の課題と当社のサービス 11 • 原価はあまり上げたくないものの 必要な演算リソースは確保したい • GPUはどんな計算を速くできる？ • アルゴリズムをカスタムしても性能が 出るか？ • 前処理・後処理も速くできる？ • 機能を追加するため、もう少し処理を 速くしたい • CUDAプログラミングモデルの理解 • ハードウェアアーキテクチャの理解 製品企画 ハードウェア選定 アルゴリズム 設計・実装 性能・精度 チューニング 品質確保 • GPUで動かしてみたものの 期待した性能が出ない • 品質確保のため、精度を上げたく 演算量は増えるが性能は維持したい • GPU高速化に関するコンサルティング • ボトルネック解析 • アルゴリズムのGPU向け設計・実装 • GPUプログラム高速化 • CPU / GPU混在環境での高速化 • 精度向上
12. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation CUDA高速化入門
13. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation なぜGPUを使うのか • CPUと比べて • ピーク性能の高さ • 電力効率の良さ 13 浮動小数点数演算性能 メモリバンド幅 TDP 価格 CPU: AMD Ryzen 9 5950X 2.25 [TFLOPS] 51.2 [GB/s] 105 [W] ¥90,000~ GPU: NVIDIA GeForce RTX 3070 20.31 [TFLOPS] 448.0 [GB/s] 220 [W] ¥93,000~ • その他のアクセラレータと比べて • 入手性・価格性能比の良さ • プログラミングの容易さ
14. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation なぜGPUが速いのか • 並列計算に特化した構成 • 大量のコア・演算器 • CPU: AMD EPYC 7763: 64 Cores, 32 FLOPs/Core/cycle • GPU: NVIDIA A100: 108 SMs, 128 FLOPs/SM/cycle • バス幅の広い広帯域メモリ • もちろん弱点もある • 並列に処理できない問題には弱い • 最大メモリ容量が小さい 14
15. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation CUDAプログラミングモデル
16. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 例題: saxpy • Single-precision ax plus y • y ← a × x + y • CPU向けの実装例: 16 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]; } }
17. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation CUDAを用いたプログラムの流れ • ホストメモリからデバイスメモリへデータを転送 • GPU上でカーネル（プログラム）を実行 • デバイスメモリからホストメモリへデータを転送 17
18. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ホストメモリとデバイスメモリ • CPUとGPUはそれぞれがメモリを持っている • 目的に応じて適切なほうを利用する • 必要に応じて片方から他方へデータをコピーする 18 CPU GPU ホストメモリ (DDR) ～200 GB/s デバイスメモリ (GDDR/HBM) ～2000 GB/s ～20 GB/s
19. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ホストメモリからデバイスメモリへデータを転送 • cudaMalloc • デバイスメモリ上の領域を確保 • 標準Cにおけるmallocに対応 • cudaMemcpy • デバイスメモリに関係するメモリコピー • 第4引数で転送の方向を指定 (HostToDevice, DeviceToHost など) 19 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);
20. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation GPU上でカーネル（プログラム）を実行 • カーネルの呼び出し • スレッド数を指定する • スレッドブロック数×ブロックあたりのスレッド数で表現 • ここではループ1回を1スレッドで処理する 20 const int bdim = 128; const int gdim = (n + bdim – 1) / bdim; // 切り上げ kernel<<<gdim, bdim>>>(d_y, d_x, a, n);
21. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation GPU上で動くカーネルの実装 • __global__ 修飾された関数として定義 • 定義済み変数から自身のインデックスを取得 • blockDim: 現在のカーネル実行におけるブロックサイズ • blockIdx: 自身の属するスレッドブロックのインデックス • threadIdx: 自身のスレッドブロック内におけるインデックス 21 __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]; }
22. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation デバイスメモリからホストメモリへデータを転送 • cudaMemcpyで逆方向にコピー 22 // デバイスメモリ (d_y) から sizeof(float) * n バイト転送 cudaMemcpy(h_y, d_y, sizeof(float) * n, cudaMemcpyDeviceToHost); // デバイスメモリの解放 cudaFree(d_x); cudaFree(d_y);
23. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation スレッドの階層構造 • CUDAではスレッド間に階層構造がある • 近いスレッド同士はより密に通信・同期を行うことができる 23 Warp (32T) … Thread Block (～1024T) … Grid …
24. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation スレッドの階層構造 • CUDAではスレッド間に階層構造がある • Warp: 同時に命令が発行されるスレッドをまとめたもの • 現行アーキテクチャでは32スレッド • Thread Block: いくつかのスレッドをまとめたもの • 現行アーキテクチャでは1ブロックあたり最大1024スレッド • 同一ワープに属するスレッドは必ず同一スレッドブロックに属する • Grid: いくつかのスレッドブロックをまとめたもの • カーネル呼び出しは1つのグリッドで処理される 24
25. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation スレッドの階層構造 • 階層構造上で近いスレッド同士はより密に同期や通信を行うことができる • 同一グリッド • カーネル起動・終了時の同期のみ • 同一スレッドブロック • 同じブロックに属するスレッド同士での同期 • シェアードメモリを用いたデータ共有 • 同一ワープ • スレッド同士でのより軽量な同期 • ワープシャッフルによるデータ共有 25
26. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation • メモリにも階層構造がある • おおむねスレッドの階層構造と対応 メモリの階層構造 26 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
27. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリの階層構造: レジスタ • プログラム中の自動変数に対応 • 各種演算命令に直接渡すことができる • 他のスレッドとは共有されない 27 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
28. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリの階層構造: ローカルメモリ • プログラム中の自動変数に対応 • 何らかの理由でレジスタに乗せられないときに使用される • 演算命令に渡す際はいったんレジスタにロードする必要がある • 他のスレッドとは共有されない 28 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
29. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリの階層構造: シェアードメモリ • __shared__ 修飾された変数に対応 • 同一スレッドブロック内の全スレッドで共有される 29 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
30. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリの階層構造: グローバルメモリ • cudaMalloc などで確保された領域に対応 • デバイス全体で共有される • カーネル停止後も値が保持される 30 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
31. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリの階層構造: コンスタントメモリ • __constant__ 修飾された変数に対応 • デバイス全体で共有される • カーネルから値を書き換えることができない 31 Global Memory Grid Thread Block Thread Registers Local Memory Shared Memory Constant Memory
32. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ホストとデバイス間の同期 • カーネル呼び出しやデータ転送は基本的に非同期実行 • 明示的もしくは暗黙的に同期を挿入する必要がある • cudaMemcpy など一部のAPIは自動的に同期を挿入する 32 kernel<<<1, 1>>>(); // この時点では kernel() はまだ実行されていないかもしれない foo(); cudaDeviceSynchronize(); // この時点では kernel() の処理は確実に完了している kernel() foo() cudaDeviceSynchronize() CPU GPU
33. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ストリーム • デバイスで実行される処理のキュー • 投入した順に処理される • 同じストリームに投入された処理同士はオーバーラップしない • 指定されなかった場合はデフォルトストリームが使用される 33 kernel1<<<1, 1>>>(); kernel2<<<1, 1>>>(); cudaDeviceSynchronize(); kernel1() cudaDeviceSynchronize() CPU GPU kernel2()
34. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ストリーム • ストリームは複数作成することができる • 別ストリームに投入された処理同士は並行するかもしれない 34 kernel1<<<1, 1, 0, stream1>>>(); kernel2<<<1, 1, 0, stream2>>>(); cudaDeviceSynchronize(); kernel2() kernel1() Synchronize CPU GPU
35. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation プログラミングモデルまとめ • 大量のスレッドの間には階層関係がある • ワープ・スレッドブロック・グリッド • 距離に応じて同期や通信の制約が変化する • メモリにも階層関係がある • レジスタ・ローカルメモリ・シェアードメモリ・グローバルメモリ • 速度や共有する必要があるスレッド数など要求に応じて適切な領域を使い分ける • デバイス上で動く処理は基本的に非同期実行となる • 細かい同期周りの制御にはストリームを活用する 35
36. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation ハードウェア
37. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation カーネルが遅い原因と対応策 要求されている演算量が多すぎる • アルゴリズムを改善して演算量を減らす メモリトラフィックが多すぎる • アルゴリズムを改善してメモリアクセスを減らす • キャッシュなどのハードウェア機能を活用する リソースを有効活用できていない • ハードウェアの制約を理解して演算器やバスなどの稼働率を上げる 37
38. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation カーネルが遅い原因と対応策 要求されている演算量が多すぎる • アルゴリズムを改善して演算量を減らす メモリトラフィックが多すぎる • アルゴリズムを改善してメモリアクセスを減らす • キャッシュなどのハードウェア機能を活用する リソースを有効活用できていない • ハードウェアの制約を理解して演算器やバスなどの稼働率を上げる 38
39. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation NVIDIA A100 Block Diagram • CC 8.0 • 108 SMs/Chip • 6912 FP32 CUDA Cores 40 • コアを活用できるだけの並行実行可能なタスク (=スレッド) を投入する必要がある • スレッド数が足りないならタスクを分割することも視野に入れる https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf
40. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Streaming Multiprocessor (SM) • スレッドブロックに対応する • いくつかのスレッドブロックを並行して処理 • 以下の要素を束ねたもの • CUDA Core • Tensor Core • LD/ST Unit • SFU • Register File • Cache/Shared Memory • Scheduler, Dispatcher 41
41. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Streaming Multiprocessor (SM) • スレッドブロックに対応する • 以下の要素を束ねたもの • CUDA Core • Tensor Core • LD/ST Unit • SFU • Register File • Cache/Shared Memory • Scheduler, Dispatcher 42 演算器 メモリ
42. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Processing Block • ワープに対応する • いくつかのワープを並行して処理 • SMからワープをまたがない要素を分割したもの • 各種演算器 • レジスタファイル • スケジューラ・ディスパッチャ 43
43. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation CUDA Core • スレッドに対応する • 何らかの演算を行う • レジスタファイルから値を読んで • 演算を行って • レジスタファイルに書き出す FP32/INT32 • Volta以降でINTコアが分離された • 整数演算と浮動小数点数演算を同時に実行できる 44
44. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Tensor Core • 深層学習向けのアクセラレータ • ワープ単位で協調して小さい行列積を効率よく行う • 世代によって対応する精度・サイズが異なる 45
45. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation その他のユニット • LD/ST (Load/Store) • メモリアクセスを行う • SFU (Special Function Unit) • 特殊関数 (指数関数・三角関数など) の処理を行う • 演算器が少ない分スループットも落ちる 46
46. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation SIMTとWarp • ディスパッチャはワープに対して一つの命令を一度に発行する • SIMT: Single Instruction, Multiple Threads • スレッドごとに異なる命令を発行することはできない • 条件分岐の取り扱い • 分岐によって実行の必要がなくなった命令も発行されうる • そのような場合はその命令が無視される 48
47. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Warp Divergence • 条件分岐によって有効な演算を行わないスレッド (=コア) が発生する • ワープ内での異なる方向への分岐は性能劣化につながる • Warp Divergence と呼ぶ • 下の例では B(), C() の処理中にコアが半分遊んでいる 49 A(); if(threadIdx.x % 2 == 0){ B(); }else{ C(); } A() B() C() Warp 0 Warp 1
48. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Warp Divergence • できるだけ同じワープのスレッドが同じように動くことで効率を改善できる • 連続するスレッドが同じ方向に分岐するようにする • 下の例では B(), C() におけるコア稼働率が改善している 50 A(); if(threadIdx.x < 4){ B(); }else{ C(); } A() B() C() Warp 0 Warp 1
49. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation レイテンシの隠蔽 • 命令のパイプライニングはGPUでも有効 • 依存性のない命令同士を並列実行する • 依存性のない命令の組をどう見つけるか • 近くにある命令との依存性を解析する • 別のスレッドの命令と組み合わせる 51
50. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ワープスケジューリング • Processing Block はいくつかの実行中ワープの状態を保持している • 可能であれば物理コア数より多くのスレッドの状態を保持する • サイクルごとに実行可能なワープをその中から選択して命令を発行する • 実行可能: 次に発行される命令が依存している処理がすべて完了している • 実行可能なワープを絶やさないことが効率改善につながる • 命令のレイテンシを考慮したプログラムを記述する • 実行可能なワープの候補 (=状態を保持しているスレッド数) を増やす 52
51. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • FADDのレイテンシは4とする • 並行実行しているワープ数が1の場合: 4 ops / 16 cycles 53 0x00: FADD R1, R2 0x01: FADD R1, R3 0x02: FADD R1, R4 0x03: FADD R1, R5 R1 += R2 R1 += R3 R1 += R4 R1 += R5 Warp 0
52. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation レイテンシ隠蔽の例 • 依存性のある加算を4回行うプログラム • FADDのレイテンシは4とする • 並行実行しているワープ数が1の場合: 4 ops / 16 cycles • 並行実行しているワープ数が4の場合: 16 ops / 19 cycles 54 0x00: FADD R1, R2 0x01: FADD R1, R3 0x02: FADD R1, R4 0x03: FADD R1, R5 R1 += R2 R1 += R3 R1 += R4 R1 += R5 Warp 0 R1 += R2 R1 += R3 R1 += R4 R1 += R5 Warp 1 R1 += R2 R1 += R3 R1 += R4 R1 += R5 Warp 2 R1 += R2 R1 += R3 R1 += R4 R1 += R5 Warp 3
53. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Occupancy • SMがいくつのワープを並行実行できるかを表す指標 • 高ければ高いほどレイテンシを隠蔽しやすい • ブロックサイズ・消費レジスタ数・シェアードメモリサイズから求める • ブロックサイズ: SMあたりの並行実行可能なブロック数 • 消費レジスタ数: SMあたりのレジスタファイル数 • シェアードメモリサイズ: SMあたりのシェアードメモリサイズ • プロファイラ・CUDA Toolkit 付属のExcelシートなどで求められる 55
54. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリ階層 • 上に行くほど高速な代わりに共有される範囲が狭まる 56 Device Memory L2 Cache SM L1 Cache Shared Memory Processing Block Register File
55. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation デバイスメモリ • グローバルメモリ・ローカルメモリに対応 • GPUのスペックに書かれている容量はこの領域のもの • アクセスパターンによって大きく性能が変わる 57
56. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 58 0 1 2 3 4 5 31 スレッド … メモリ … 1トランザクション / 1アクセス
57. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 59 0 1 2 3 4 5 31 スレッド … メモリ … 2トランザクション / 1アクセス
58. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Coalesce Access • 同一ワープに属するスレッド群のメモリアクセスはまとめて処理される • 範囲が狭ければ1トランザクションにまとめられる 60 0 1 2 3 4 5 31 スレッド … メモリ … 32トランザクション / 1アクセス … … … … … …
59. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation L2 キャッシュ • デバイス中の全SMで共有されている • デバイスメモリへのアクセス時には常に使用される 61
60. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation L1キャッシュ • SMごとに用意されている • 明示的に指定したものか読み取り専用のデータへのアクセスに対して使用される 読み取り専用かどうかの判定 • コンパイラが判定する • ポインタを const __restrict__ 修飾すると読み取り専用であることを明示できる 明示的なL1キャッシュの利用 • 組み込み関数 __ldg() を使用する • *ptr → __ldg(ptr) 62
61. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation シェアードメモリ • SMごとに用意された領域 • L1キャッシュとシェアードメモリの割合は設定で変更可能 • シェアードメモリとして使えるのは 16-96 [KB/SM] 程度 • 残りはL1キャッシュとして使用される 63
62. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリバンク • シェアードメモリはバンクを用いてに管理されている • バンクは4バイトごとに切り替わる • 同じバンクの異なる領域へのアクセスはまとめて処理できない: バンクコンフリクト 64 Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 0x00000100 …
63. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリバンク • まとめて処理できるアクセスの例 • 素直なシーケンシャルアクセス 65 Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 0x00000100 … 0 1 2 3 4 27 28 29 30 31 スレッド
64. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリバンク • まとめて処理できるアクセスの例 • バンクが重複しないランダムアクセス 66 Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 0x00000100 … 0 1 2 3 4 27 28 29 30 31 スレッド
65. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリバンク • まとめて処理できるアクセスの例 • ブロードキャスト: バンクが重なっても同じアドレスなら問題ない 67 Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 0x00000100 … 0 1 2 3 4 27 28 29 30 31 スレッド
66. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation メモリバンク • まとめて処理できないアクセスの例 • ストライドアクセス: この場合は2回に分割される 68 Bank 31 Bank 30 Bank 29 Bank 28 Bank 27 Bank 4 Bank 3 Bank 2 Bank 1 Bank 0 0x00000000 0x00000080 0x00000100 … 0 1 2 3 4 27 28 29 30 31 スレッド
67. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation レジスタファイル • プロセッシングブロックごとに用意された領域 • レジスタ幅は1要素あたり32bit • long, double, ポインタなどの64bit型には2つ使われる 自動変数に対する領域割り当て • 自動変数は可能ならレジスタに割り当てられる • 特定のケースで低速なローカルメモリに割り当てられる • 自動変数がレジスタに収まりきらない場合（レジスタスピル） • インデックスアクセスが必要な場合 69
68. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ハードウェアまとめ 演算器 • 演算器を使い切るためには注意が必要なことがある • 分岐によって何もしないコアが発生することがある • レイテンシを埋めるだけの命令供給が必要 メモリ • アクセスパターン次第で効率が落ちることがある • グローバルメモリ: Coalescing, キャッシュ利用 • シェアードメモリ: バンクコンフリクト • ローカルメモリの利用にも注意する • コンパイラの出力を確認すると確実 70
69. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Copyright © Fixstars Corporation 実践例
70. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 問題の概要 • 画像のステレオマッチング: Semi-Global Matching (SGM) • ステレオ画像の視差を計算するアルゴリズム • 視差: 片方の画像のある画素が他方の画像で何ピクセルずれたところにあるか • 近くの物体ほど視差が大きくなること利用して距離を計算できる • ターゲット環境: Pascal 世代のGPU (GeForce GTX 10xx など) 72
71. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation チューニングする部分 • 動的計画法 (DP) である画素における視差が d [px] としたときのスコアを求める • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい • 隣接画素に対して急激な視差の変化があるとコストが大きい 73 𝑠𝑐𝑜𝑠𝑡 𝑥, 𝑑 = 𝑙𝑐𝑜𝑠𝑡 𝑥, 𝑑 + min 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃1 min 𝑖 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2 − min 𝑖 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑖) 𝑑 𝑥
72. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation チューニングする部分 • 動的計画法である画素における視差が d [px] としたときのスコアを求める • 対応する画素同士の特徴ベクトルの距離が大きいとコストが大きい • 隣接画素に対して急激な視差の変化があるとコストが大きい • これを縦横斜めの8方向それぞれについてラインごとに計算する 74
73. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation アルゴリズムの概略 • 左から右方向のスキャン 75 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){ int cost = min({ P2, scost[y][x-1][d-1] - prev_min + P1, 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]); cur_min = min(prev_min, cost); } prev_min = cur_min; } } X方向のループは依存性がある 主要な計算は O(HWD) 回行われる 計算処理はかなり軽い キャッシュヒットが期待できない メモリアクセスも O(HWD) 回 Y方向のループは完全に独立している
74. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討 y方向ループ1回を1スレッドで担当する • 並列度が足りない: 数百スレッド程度しか利用できない y方向ループ1回を複数スレッドで担当する • x方向のループは分割できない: 前のループに対する依存性があるため • d方向のループは分割可能 • ただしx方向のループ1回ごとに同期が必要になる 76
75. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • d方向の最小値を求める • 端の値を隣のスレッドに渡す 77 𝑑 𝑥 Thread 0 Thread 1 scostの計算
76. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • d方向の最小値を求める • 端の値を隣のスレッドに渡す 78 𝑑 𝑥 Thread 0 Thread 1 最小値の計算・共有 min 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃1 min 𝑖 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2
77. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討: d方向ループの分割 • X方向のループを進めるたびに同期と通信が必要になる • d方向の最小値を求める • 端の値を隣のスレッドに渡す 79 𝑑 𝑥 Thread 0 Thread 1 端の値の共有 min 𝑠𝑐𝑜𝑠𝑡(𝑥 − 1, 𝑑) 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 − 1 + 𝑃1 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑑 + 1 + 𝑃1 min 𝑖 𝑠𝑐𝑜𝑠𝑡 𝑥 − 1, 𝑖 + 𝑃2
78. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討: d方向ループの分割 細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減 粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • 1ワープ以内になるとより軽量な通信が利用できる • グローバルメモリへのアクセス効率の向上 • スレッドあたりのメモリアクセス量が多くなる • 1回のアクセスで4要素までアクセスできる 80
79. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 並列化方針の検討: d方向ループの分割 細かく分割する場合のメリット • スレッド数を増やすことによる Occupancy の向上 • スレッドあたりのレジスタ量の削減 粗く分割する場合のメリット • スレッド間通信などの非本質的な処理の占める割合の減少 • 1ワープ以内になるとより軽量な通信が利用できる • グローバルメモリへのアクセス効率の向上 • スレッドあたりのメモリアクセス量が多くなる • 1回のアクセスで4要素までアクセスできる 81 1ワープを境に実装が大きく変化する ⇒ 1ワープ以下の範囲で値を変えつつ試せるように実装する
80. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation ハイエンドGPU対策 • ハイエンドGPUだとd方向の分割を入れてもスレッド数が足りない • 8方向それぞれが独立なことを利用して複数カーネルを並行させる 82
81. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation アルゴリズムの検討:局所特徴同士の距離の計算 • 局所特徴についての情報 • 局所特徴の表現: 64 bit のビット列 • 局所特徴の距離: 互いに異なるビットの数 • 特徴ベクトルの距離は同じ組の距離が何度も使われる • 8方向すべての処理で同じ計算を行う • 既存実装では事前計算してテーブル化されていた • table[y][x][d] = distance(left[y][x], right[y][x - d]) 83
82. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 理論性能で比べる • テーブル引きと計算どちらが速い? • テーブル化した場合1要素当たり 1 [byte] • 特徴同士の距離は popcount 命令2回で求められる • GeForce GTX 1080 Ti (sm_61) を例に試算してみる • 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] • 実測値だとおよそ 340 [GB/s] くらい • 毎回計算するほうが速そう!! 84
83. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation プロファイル結果 (1) • 横方向の処理のプロファイル結果 • 演算器の稼働率80%弱: うまくリソースを活用できてそう 85
84. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation プロファイル結果 (2) • 本当に距離をテーブル化しないほうが速かったのか? • 実効メモリ帯域で評価する • テーブル引きする場合はメモリトラフィックが Reads = Writes になる • テーブルサイズが結果バッファのサイズと等しいため • 52.049×2 = 104.098 [GB/s] 出せなければテーブル化のほうが遅い • bandwidthTest での帯域が 92.7 [GB/s] 程度 • テーブル化する方針では勝てないだろうと考えられる 86
85. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation プロファイル結果 (3) • 縦方向・斜め方向でも同様の傾向 87
86. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 全体の評価 • 既存実装との性能比較 • 比較対象: Embedded real-time stereo estimation via Semi-Global Matching on the GPU, D. Hernandez-Juarez et al, ICCS 2016. • https://github.com/dhernandez0/sgm • 実際にはもう一つ大きいカーネルがあるのですがそちらの詳細は省略しています 88
87. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 評価結果 • 2.3-12.5% 程度の高速化 • 演算性能に対してメモリ帯域の細いチップで特に強い 89 261 69.7 50.2 232 68.1 45.8 0 50 100 150 200 250 300 GTX 1080 Ti GTX 1050 Ti DRIVE PX2 フレームレート 今回の実装 Hernandez+
88. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation まとめ • 演算とメモリアクセスどちらが重要か見極める • 理論性能から見積もり • 実測で裏付け • コアあたりの効率とチップあたりの効率 • 並列度を下げると演算量は減らしやすい • 一方でリソースが余りやすくなるのでうまくバランスをとる 90
89. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation 全体のまとめ • パフォーマンスチューニングにおいてはハードウェアの知識も重要 • 使い方を誤ると数倍の性能劣化なども起こりうる • もちろんアルゴリズムも重要で両方からのアプローチが必要 • カーネルのチューニングにおいては特に演算器とメモリに気を配る • 演算器を余らせない • 不得意なアクセスパターンによる性能劣化を防ぐ • 理論をもとに仮説を立てて実装したものを評価する • プロファイラによる評価 • 理論ピークと実性能の差を読み取る 91
90. Fixstars Corporation www.fixstars.com Copyright © Fixstars Corporation Thank You お問い合わせ窓口 : contact@fixstars.com