CUDA画像処理入門
エヌビディアジャパン CUDAエンジニア 森野慎也
GTC Japan 2014
CUDAを用いた画像処理
画像処理をCUDAで並列化
— 基本的な並列化の考え方 — 目標 : 妥当なNaïveコードが書ける 最適化の初歩がわかる
— ブロックサイズ — メモリアクセスパターンRGB → Y(輝度) 変換
カラー画像から、グレイスケールへの変換
CUDAにおける画像処理の基礎
2次元メモリ確保API
— Pitchを考慮 — cudaMallocPitch()、cudaMemcpy2D() 並列化
— CUDAの並列度 : 数万以上欲しい…PITCHを考慮したメモリレイアウト
RGBA(8 bit, uchar4)の配列
index = x + y * pitchInPixel
width
pitchInPixel = pitchInByte / sizeof(uchar4)
2次元メモリ 確保・転送
cudaError_t
cudaMallocPitch
( void** devPtr, size_t* pitch,
size_t width, size_t height )
— widthバイトのメモリを、height行分、取得する。 — 行は、pitchバイトで整列する。
cudaError_t
cudaMemcpy2D
( void* dst, size_t dpitch,
const void* src, size_t spitch, size_t width, size_t height,
cudaMemcpyKind kind )
— dstで示されるメモリ (dpitchバイトで整列)に、srcで示されるメモリ (spitchバイトで整列) を、width (バイト) x height (行)、コピーする。
サンプルコード
uchar4 *src, *dImage;
size_t spitch, dPitch, dPitchInPixel; // ピッチつきで、メモリをアロケート
cudaMallocPitch(&dImage, *dPitch, width * sizeof(uchar4), height); dPitchInPixel = dPitch / sizeof(uchar4);
// ピッチを変換しつつ、ホスト→デバイスへと、メモリ転送
cudaMemcpy2D(dImage, dPitch, src, sPitch, width * sizeof(uchar4), height, cudaMemcpyHostToDevice);
画像処理における並列化の基本
基本 : 1 ピクセルに対して、 1スレッドを対応させる
— ピクセル数分、スレッドが走る。 例 : 262,144 (= 512 x 512) スレッド スレッドは、処理対象のピクセルを持つ。
— 自分の位置 (x, y) を知ることが必要2DでのBLOCK・THREADの割り当て
Thread : 「２次元」でピクセルに対応。
Block
: 「２次元」で定義。一定のサイズのタイル。
Grid
: 必要数のBlockを「２次元」に敷き詰める。
1 Block
1 Pixel = 1 Thread(x, y) =
(Global ID X, Global ID Y)
2DでのBLOCK・THREADの割り当て
GlobalID は、(x, y
, z
)方向に計算できる
— GlobalID(x) = blockDim.x * blockIdx.x + threadIdx.x — GlobalID(y) = blockDim.y * blockIdx.y + threadIdx.y — GlobalID(z) = blockDim.z * blockIdx.z + threadIdx.z blockDim.x * blockIdx.x
threadIdx.x
blockDim.y * blockIdx.y threadIdx.y
RGB → Y 変換 カーネル
__global__
void RGBToYKernel(uchar4 *dDst, const uchar4 *dSrc, int width, int height, int pitch){ int gidx = blockDim.x * blockIdx.x + threadIdx.x;
int gidy = blockDim.y * blockIdx.y + threadIdx.y;
if ((gidx < width) && (gidy < height)) { int pos = gidx + gidy * pitch;
// Y = 0.299 R + 0.587 G + 0.114 B uchar4 value = src[pos];
float Y = 0.299f * value.x + 0.587f * value.y + 0.114f * value.z; unsigned char y = (unsigned char)min(255, (int)Y);
dDst[pos ] = pixel; }
カーネル呼び出し (GRIDサイズ指定)
/* value、radixで割って、切り上げる */ int divRoundUp(int value, int radix) {
return (value + radix – 1) / radix; }
/* gridDim, blockDimを、２次元(x, y方向)に初期化 */ dim3 blockDim(64, 2);
/* divRoundUp()は、切り上げの割り算 */
dim3 gridDim(divRoundUp(width, blockDim.x), divRoundUp(height, blockDim.y)); RGBToYKernel<<<gridDim, blockDim>>>(dDst, dSrc, …);
悪い並列化の例
GPUの並列化としては、NG。非常に低速。
— 並列度が低い — メモリアクセスパターンが悪い ただし… CPU的発想としてはふつう。 Thread 0 Thread 1 Thread 2 Thread 3 ……
ここはポイント！コアレス(COALESCED)アクセス
連続するスレッドが、連続するメモリにアクセスする。
— threadIdx.xに対して、連続。 Memory : 0 1 2 3 4 5 6 7 8 … threadIdx.x … Thread :再掲 : 2DでのBLOCK・THREADの割り当て
GlobalID は、(x, y
, z
)方向に計算できる
— GlobalID(x) = blockDim.x * blockIdx.x + threadIdx.x — GlobalID(y) = blockDim.y * blockIdx.y + threadIdx.y — GlobalID(z) = blockDim.z * blockIdx.z + threadIdx.z blockDim.x * blockIdx.x
threadIdx.x
blockDim.y * blockIdx.y threadIdx.y
FAQ : BLOCKDIMの決め方
1. Occupancy (占有率) を 100 %にする
2. Blockあたりのスレッド数は、なるべく小さく。
BLOCKDIMの決め方 (OCCUPANCY から)
SMXあたり、2048 Thread走らせたい。
— Occupancy (占有率) = 100 %
Occupancy = 100 % を満たす、Blockあたりのスレッド数は、
2048 Thread / 16 Block = 128 Thread / Block 2048 Thread / 8 Block = 256 Thread / Block 2048 Thread / 4 Block = 512 Thread / Block 2048 Thread / 2 Block = 1024 Thread / Block
項目 値
最大のBlock数 / SMX 16 最大のThread数 / SMX 2048 最大のThread数 / Block 1024
BLOCKDIMの決め方(BLOCKの粒度から)
Grid = 4096 Thread の実行例を考えてみる
— Block : 256 Thread、1024 Threadで比較
— 3 SMX / GPU、1 SMXあたり 1 Blockが実行可能とする 1024 Thread / Block Block Block Block Block SMX 0 SMX 1 SMX 2 256 Thread / Block SMX 0 Block Block Block Block Block Block Block Block Block Block Block Block Block Block Block SMX 1 SMX 2 Block t t
BLOCKDIMの決め方 (SMXの構造から)
Warp Scheduler x 4 :
— 1 clockあたり、4 Warpに対する命令発行
— Blockのサイズは、 128 Thread の倍数が望ましい。 (128 Thread = 32 Thread/Warp x 4 Warp)
