CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014
CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン
RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587 G + 0.114 B
CUDA における画像処理の基礎 2 次元メモリ確保 API Pitch を考慮 cudamallocpitch() cudamemcpy2d() 並列化 CUDA の並列度 : 数万以上欲しい Kepler での目安 : CUDA Core 数 x 10 程度 ( 最低限 )
PITCH を考慮したメモリレイアウト RGBA(8 bit, uchar4) の配列 index = x + y * pitchinpixel pitchinpixel = pitchinbyte / sizeof(uchar4) width (x, y)
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 の割り当て 1 Thread : 2 次元 でピクセルに対応 Grid 1 Pixel = 1 Thread (x, y) = (Global ID X, Global ID Y) : 2 次元 で定義 一定のサイズのタイル : 必要数の を 2 次元 に敷き詰める
2D での BLOCK THREAD の割り当て blockdim.x * blockidx.x threadidx.x blockdim.y * blockidx.y threadidx.y 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
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 を 2 次元 (x, y 方向 ) に初期化 */ dim3 blockdim(64, 2); /* divroundup() は 切り上げの割り算 */ dim3 griddim(divroundup(width, blockdim.x), divroundup(height, blockdim.y)); RGBToYKernel<<<gridDim, blockdim>>>(ddst, dsrc, );
悪い並列化の例 Thread 0 Thread 1 Thread 2 Thread 3 GPU の並列化としては NG 非常に低速 並列度が低い メモリアクセスパターンが悪い ただし CPU 的発想としてはふつう
ここはポイント! コアレス (COALESCED) アクセス Thread : 0 1 2 3 4 5 6 7 8 Memory : threadidx.x 連続するスレッドが 連続するメモリにアクセスする threadidx.x に対して 連続
再掲 : 2D での BLOCK THREAD の割り当て blockdim.x * blockidx.x threadidx.x blockdim.y * blockidx.y threadidx.y 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
動かしてみる
FAQ : BLOCKDIM の決め方 1. Occupancy ( 占有率 ) を 100 % にする 2. あたりのスレッド数は なるべく小さく 3. 横方向は コアレスアクセス なるべく 長くする
BLOCKDIM の決め方 (OCCUPANCY から ) 項目 値 最大の 数 / SMX 16 最大のThread 数 / SMX 2048 最大のThread 数 / 1024 SMX あたり 2048 Thread 走らせたい Occupancy ( 占有率 ) = 100 % Occupancy = 100 % を満たす あたりのスレッド数は 2048 Thread / 16 = 128 Thread / 2048 Thread / 8 = 256 Thread / 2048 Thread / 4 = 512 Thread / 2048 Thread / 2 = 1024 Thread /
BLOCKDIM の決め方 (BLOCK の粒度から ) Grid = 4096 Thread の実行例を考えてみる : 256 Thread 1024 Threadで比較 3 SMX / GPU 1 SMXあたり 1 が実行可能とする SMX 0 SMX 1 SMX 2 256 Thread / t SMX 0 SMX 1 SMX 2 1024 Thread / サイズは小さいほうが得 128 Threads / t
BLOCKDIM の決め方 (SMX の構造から ) Warp Scheduler x 4 : 1 clock あたり 4 Warp に対する命令発行 のサイズは 128 Thread の倍数が望ましい (128 Thread = 32 Thread/Warp x 4 Warp)
タイルは横長がよい タイルの横幅は 32(Warp の幅 ) の倍数がよい 32 より小さい場合 16 もしくは 8 を使う Thread : 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 Memory : threadidx.x
blockdim.y RGB Y 変換時のバンド幅 : TESLA K20C blockdim.x 1 2 4 8 16 32 64 128 256 512 1024 1 1.4 2.8 5.6 11.2 22.1 43.9 78.5 119.8 119.3 115.4 87.7 Occupancy < 100 % 2 2.6 5.2 10.4 20.6 40.7 77.9 119.8 119.4 115.3 87.4-4 4.8 9.6 19.2 37.8 74.0 119.4 118.2 114.2 87.3 - - 8 8.4 16.7 33.3 69.6 115.0 117.9 111.9 87.1 - - - 16 13.4 26.3 60.6 106.7 115.0 114.3 87.2 - - - - 32 17.7 40.4 81.1 103.9 110.9 86.9 - - - - - 64 20.7 41.7 79.8 99.0 83.5 - - - - - - 128 20.7 41.6 75.6 75.3 - - - - - - - 256 20.7 41.0 60.3 - - - - - - - - 512 20.5 37.6 - - - - - 値 - : バンド幅 - (GB/sec) - - 1024 19.1 blockdim.x - - < 8 - - - - Tesla - K20c - (ECC - off) -
まとめ 画像処理における CUDA Pitch を考慮したメモリレイアウト 2 次元の Grid の呼び出し 正しい Naïve コード ( カーネル ) の書き方 コアレスアクセス ピクセルごとに スレッドを割り当てる 並列度は 数万以上 サイズは 128 が適当 ( 単純なカーネルの場合 )