CUDA を用いた画像処理画像処理を CUDA で並列化基本的な並列化の考え方目標 : 妥当な Naïve コードが書ける最適化の初歩がわかるブロックサイズメモリアクセスパターン

Size: px

Start display at page:

Download "CUDA を用いた画像処理画像処理を CUDA で並列化基本的な並列化の考え方目標 : 妥当な Naïve コードが書ける最適化の初歩がわかるブロックサイズメモリアクセスパターン"

かずまさささおか
7 years ago
Views:

1 CUDA 画像処理入門エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014

2 CUDA を用いた画像処理画像処理を CUDA で並列化基本的な並列化の考え方目標 : 妥当な Naïve コードが書ける最適化の初歩がわかるブロックサイズメモリアクセスパターン

3 RGB Y( 輝度 ) 変換カラー画像からグレイスケールへの変換 Y = R G B

4 CUDA における画像処理の基礎 2 次元メモリ確保 API Pitch を考慮 cudamallocpitch() cudamemcpy2d() 並列化 CUDA の並列度 : 数万以上欲しい Kepler での目安 : CUDA Core 数 x 10 程度 ( 最低限 )

5 PITCH を考慮したメモリレイアウト RGBA(8 bit, uchar4) の配列 index = x + y * pitchinpixel pitchinpixel = pitchinbyte / sizeof(uchar4) width (x, y)

6 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 ( 行 ) コピーする

7 サンプルコード 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);

8 画像処理における並列化の基本基本 : 1 ピクセルに対して 1 スレッドを対応させるピクセル数分スレッドが走る例 : 262,144 (= 512 x 512) スレッドスレッドは処理対象のピクセルを持つ自分の位置 (x, y) を知ることが必要

9 2D での BLOCK THREAD の割り当て 1 Thread : 2 次元でピクセルに対応 Grid 1 Pixel = 1 Thread (x, y) = (Global ID X, Global ID Y) : 2 次元で定義一定のサイズのタイル : 必要数のを 2 次元に敷き詰める

10 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

11 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 = R G B uchar4 value = src[pos]; float Y = 0.299f * value.x f * value.y f * value.z; unsigned char y = (unsigned char)min(255, (int)y); ddst[pos ] = pixel; }

12 カーネル呼び出し (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, );

13 悪い並列化の例 Thread 0 Thread 1 Thread 2 Thread 3 GPU の並列化としては NG 非常に低速並列度が低いメモリアクセスパターンが悪いただし CPU 的発想としてはふつう

14 ここはポイント! コアレス (COALESCED) アクセス Thread : Memory : threadidx.x 連続するスレッドが連続するメモリにアクセスする threadidx.x に対して連続

15 再掲 : 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

16 動かしてみる

17 FAQ : BLOCKDIM の決め方 1. Occupancy ( 占有率 ) を 100 % にする 2. あたりのスレッド数はなるべく小さく 3. 横方向はコアレスアクセスなるべく長くする

18 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 /

19 BLOCKDIM の決め方 (BLOCK の粒度から ) Grid = 4096 Thread の実行例を考えてみる : 256 Thread 1024 Threadで比較 3 SMX / GPU 1 SMXあたり 1 が実行可能とする SMX 0 SMX 1 SMX Thread / t SMX 0 SMX 1 SMX Thread / サイズは小さいほうが得 128 Threads / t

20 BLOCKDIM の決め方 (SMX の構造から ) Warp Scheduler x 4 : 1 clock あたり 4 Warp に対する命令発行のサイズは 128 Thread の倍数が望ましい (128 Thread = 32 Thread/Warp x 4 Warp)

21 タイルは横長がよいタイルの横幅は 32(Warp の幅 ) の倍数がよい 32 より小さい場合 16 もしくは 8 を使う Thread : Memory : threadidx.x

22 blockdim.y RGB Y 変換時のバンド幅 : TESLA K20C blockdim.x Occupancy < 100 % 値 - : バンド幅 - (GB/sec) blockdim.x - - < Tesla - K20c - (ECC - off) -

23 まとめ画像処理における CUDA Pitch を考慮したメモリレイアウト 2 次元の Grid の呼び出し正しい Naïve コード ( カーネル ) の書き方コアレスアクセスピクセルごとにスレッドを割り当てる並列度は数万以上サイズは 128 が適当 ( 単純なカーネルの場合 )

1. マシンビジョンにおける GPU の活用

CUDA 画像処理入門 GTC 213 チュートリアルエヌビディアジャパン CUDA エンジニア森野慎也 1. マシンビジョンにおける GPU の活用 1. 医用画像処理における GPU の活用 CT や MRI から画像を受信して三次元画像の構築をするシステム 2 次元スキャンデータから 3 次元 4 次元イメージの高速生成 CUDA 化により画像処理速度を約 2 倍に高速化 1. CUDA で画像処理

CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン

CUDA を用いた画像処理画像処理を CUDA で並列化基本的な並列化の考え方目標 : 妥当な Naïve コードが書ける最適化の初歩がわかるブロックサイズメモリアクセスパターン