CUDA画像処理最適化入門

CUDA画像処理最適化入門
エヌビディアジャパン CUDAエンジニア森野慎也
GTC Japan 2014
CUDA画像処理最適化入門
 最適化
1. システムレベル：ホットスポットの特定と、改善
2. GPU : CUDAカーネルの最適化。
 今日のシナリオ : システムレベルでは、ホットスポットが特定済み。
— CUDAを使って、どれだけ速くなるか？
— CUDAカーネルの最適化で、どれだけの改善があるのか？
 Nsight Visual Studio Edition
— CUDAカーネルのプロファイリング
題材
 Disparity Map
— ステレオカメラからの深さ検出
L
L
R
R
ずれ量から深さを推定
動かしてみる
元画像 (RGBA8)
輝度画像 (float)
Disparity Map (int)
DISPARITY MAP
L
𝑑
R
 R側の演算範囲を移動
 L、Rが、よく一致する
ずれ量 d を探す
𝑟
𝑏𝑚 𝑥, 𝑦, 𝑑 =
R
L
𝐿 𝑋, 𝑌 − 𝑅 𝑋 − 𝑑, 𝑌
𝑋,𝑌 ∈ 𝐵𝑙𝑜𝑐𝑘(𝑥,𝑦)
 bm(x, y, d)が最小となるdを求め、D(x, y)とする。
BLOCK MATCHING
R
𝑏𝑚 =
S
R
L
ー
𝑑
L
 ブロック内部の画像が似ていると、bmの値は小さくなる。
CPUで普通に実装すると…
 全部で5重ループ
for (x, y : 出力点座標) {
// 2重ループ
for (dTmp : ずらし量 ) {
// ループ
bm = 0;
for (X, Y : ブロック内部の点の座標)
// 2重ループ
bm += SAD(X, Y, dTmp)
if (bm < bmMin) {
dMin = dTmp; bmMin = bm;
}
}
}
CUDA NAÏVE実装
 出力点(x, y)を並列化。カーネル内部は、3重ループ。
__global__
void ssdNaiveKernel(int *disparity, int disparityPitch, …) {
/* 出力点の(X, Y)の二重ループを並列化 */
int gidx = blockDim.x * blockIdx.x + threadIdx.x, gidy = blockDim.y * blockIdx.y + threadIdx.y;
if ((gidx < roiWidth) && (gidy < roiHeight)) {
float minSad = FLT_MAX; int minDisp = 0;
int imgX = gidx + numDisparities, imgY = gidy;
/* ずらし量で、ループ */
for (int iDisp = 0; iDisp < numDisparities; ++iDisp) {
float sad = 0.f;
/* ブロック内部の点の演算。二重ループ */
for (int iy = 0; iy < blockSize; ++iy) {
int offset = imgX + (imgY + iy) * imgPitch;
for (int ix = 0; ix < blockSize; ++ix)
sad += fabs(dR[offset + ix - iDisp] - dL[offset + ix]);
} // iy
if (sad < minSad) {
minSad = sad; minDisp = iDisp;
}
} // iDisp
disparity[gidx + gidy * disparityPitch] = minDisp;
}
}
性能比較
実装
実行時間
性能向上
CPU (1 core)
970 ms
-
CPU (OpenMP, 8 threads)
320 ms
x 3.0
GPU (Naïve) *
28 ms
x 34.6
 CPU → CUDAに載せ替えるだけでも、速い！
— CUDA : 開発を加速するツールとしても有用。
CPU : Core i7 3870 (4 core), GPU : Tesla K20c (ECC off), Image : 384 x 288 pixels
GPUは、カーネルの実行時間のみ計測
最適化へのアプローチ
 データアクセスの効率化
— Naive版を高速化
— 本セッションでのフォーカス。
 またの機会に取り上げます。
— アルゴリズムの検討
— システムレベルの最適化
— IPCの向上 (IPC = Instruction Per Clock)
演算性能 VS メモリバンド幅
演算性能(Single)
メモリバンド幅
演算性能 /バンド幅
(FLOP / B)
Tesla M2090
(Fermi)
1.33 TFLOPS
177.6 GB/sec
6.97
Tesla K40
(Kepler)
4.29 TFLOPS
288 GB/sec
13.9
比率
x 3.2
x 1.6
-
GPU
 演算性能に対して、メモリバンド幅が十分広くはない。
 Fermi → Keplerで、さらに、顕著になった。
プロファイリング
メモリアクセスのモニタリング実行
NSIGHT : メモリアクセスのモニタリング
 Naïve版をプロファイリング : L2↔SMX のアクセス量が多い
データアクセスの効率化
 低速なパス(L2⇔SM)の利用を減らし、
高速なパス(Shared Memory)を利用する。
Global RO
Tex
高速 (数TB / sec)
PCIe
Host(PC) DRAM
Global Memory
(GPU DRAM)
Global Memory Access
L2 Cache
Shared Memory
Tex L1
演算
Register File
(変数)
CUDA Cores
SM
低速 (~ 10 GB/sec)
Keplerでは、通常、Global Memory Access時のL1 Cacheは、使用されていない。
SHARED MEMORYを使う
R
R−L
L
S
Thread
 1 Line分づつ、読み込み。
 隣り合ったスレッドで、隣り合ったShared Memoryをアクセス。
 値の再利用性が高い。(Blockの横幅分/Thread : 数十回)
実装例 (一部抜粋)
for (int iDisp = 0; iDisp < numDisparities; ++iDisp) {
float sad = 0.f;
for (int iy = 0; iy < blockSize; ++iy) {
__shared__ float shL[128 + 32];
__shared__ float shR[128 + 32];
int posInL = imgX + (imgY + iy) * imgPitch;
if (imgX < width) {
shL[threadIdx.x] = dL[posInL];
shR[threadIdx.x] = dR[posInL - iDisp];
}
else {
shL[threadIdx.x] = shR[threadIdx.x] = 0.f;
}
/* 同期。Shared Memoryの書き込み後には、必須。 */
__syncthreads();
(以下略)
SHARED MEMORY の利用
 L2 : 7.3 GB → 445 MB まで減った！
性能比較
実行時間
L2
アクセス量
CPU (1 core)
970 ms
-
-
-
CPU (OpenMP*)
320 ms
-
-
x 3.0
CUDA Naïve
28 ms
7.27 GB
-
x 34
CUDA1 (Shared Mem)
11 ms
445 MB
15.1 GB
x 88
実装
Shared Mem
性能向上
アクセス量
*OpenMP, 8 threads
CPU : Core i7 3870 (4 core), GPU : Tesla K20c (ECC off), Image : 384 x 288 pixels
もっと、データアクセスを減らす！
 Shared Memoryへのデータ読み込み回数が多い
— R側イメージのピクセルをずらすごとに、読み込み直し。
 Shared Memoryへの読み込み後、複数の「ずれ量」に対して計算。
複数のずれ量を、まとめて計算
R
R−L
L
S
ずれ量
※ 1 Thread分だけ表示
 Rにずれ量を足しつつ、複数のずれ量に対するSADの値を計算
 計算するずれ量の分、再利用回数が増える。
性能比較
実行時間
L2
アクセス量
CPU (1 core)
970 ms
-
-
-
CPU (OpenMP*)
320 ms
-
-
x 3.0
CUDA Naïve
28 ms
7.27 GB
-
x 34
CUDA1 (Shared Mem)
11 ms
445 MB
15.1 GB
x 88
CUDA2 (Shared Mem2)
5.0 ms
60.9 MB
7.7 GB
x 194
実装
Shared Mem
性能向上
アクセス量
*OpenMP, 8 threads
CPU : Core i7 3870 (4 core), GPU : Tesla K20c (ECC off), Image : 384 x 288 pixels
計算方法も見直す
-S line0
-S line1
S(Block)
+S line5
+S line6
 Block Macthingの算出結果の再利用。
— 1 ラインづつ下にBlockを移動。最上列を引き、最下列を足す。
性能比較
実行時間
L2
アクセス量
CPU (1 core)
970 ms
-
-
-
CPU (OpenMP*)
320 ms
-
-
x 3.0
CUDA Naïve
28 ms
7.27 GB
-
x 34
CUDA1 (Shared Mem)
11 ms
445 MB
15.1 GB
x 88
CUDA2 (Shared Mem2)
5.0 ms
60.9 MB
7.7 GB
x 194
CUDA3 (演算見直し)
3.6 ms
10.3 MB
552 MB
x 269
実装
Shared Mem
性能向上
アクセス量
*OpenMP, 8 threads
CPU : Core i7 3870 (4 core), GPU : Tesla K20c (ECC off), Image : 384 x 288 pixels
まとめ
 データアクセスを減らす！
— Nsight Visual Studio Editionによりプロファイリング
— Naïve版では、L2キャッシュからのデータ移動が大きい。
— アクセス量の削減により、5.7倍 ( x 34 → x 269 ) 高速になった。
 もちろん、以下も重要。またの機会に。
— システムレベルの最適化(次のセッションで、一部扱います）
— アルゴリズムレベルの最適化
— IPCの向上

Download Report