CUDAで2次元連続ウェーブレット変換

CUDAで連続ウェーブレット変換では1次元のウェーブレット変換を行いました. これを2次元に拡張して,画像データのウェーブレット変換ができるようにします.

2次元連続ウェーブレット変換

2次元連続ウェーブレット変換式は以下.

eq_2dcwt.jpg

1次元の時に用いたMexican Hatウェーブレット関数の2次元版は,

eq_2dmexicanhat.jpg

CUDA実装

カーネル呼び出し側のコードは以下.

  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
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
 47
 48
 49
 50
 51
 52
 53
 54
 55
float* CuCWT2DTex(float *img0, int nx, int ny, float s)
{
    // 信号データ
    cudaArray *caSrc;
    float *dW, *hW;
    int size;
 
    // 結果格納用デバイスメモリの確保
    size = nx*ny*sizeof(float);
    cutilSafeCall(cudaMalloc((void**)&dW, size));
    cutilSafeCall(cudaMemset((void*)dW, 0, size));
 
    // 結果格納用ホストメモリの確保
    hW = new float[nx*ny];
    
    int res = 4;
    if(res != 1){
        res = (uint)ceil(res/s);
    }
 
    // 画像データを格納するテクスチャの準備
    // CUDA Arrayの確保とホストからのデータ転送
    cudaChannelFormatDesc cdesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    size = nx*ny*sizeof(float);
    cutilSafeCall(cudaMallocArray(&caSrc, &cdesc, nx, ny));
    cutilSafeCall(cudaMemcpyToArray(caSrc, 0, 0, img0, size, cudaMemcpyHostToDevice));
 
    // テクスチャパラメータ
    g_TexImg.addressMode[0] = cudaAddressModeWrap;
    g_TexImg.addressMode[1] = cudaAddressModeWrap;
    g_TexImg.filterMode = cudaFilterModePoint;
    g_TexImg.normalized = true;    // 正規化されたテクスチャ座標でアクセス
 
    // CUDA Arrayをテクスチャにバインド
    cutilSafeCall(cudaBindTextureToArray(g_TexImg, caSrc, cdesc));
 
    // カーネル実行
    dim3 block(BLOCK_SIZE, BLOCK_SIZE);
    dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y);
 
    cwt2dtex<<< grid, block >>>(dW, nx, ny, 0, 0, 1, 1, s, res);
 
    cutilCheckMsg("Kernel execution failed");    // カーネル実行エラーのチェック
    cutilSafeCall(cudaThreadSynchronize());        // 全てのスレッドが終わるのを待つ
 
    // デバイスからホストへ結果を転送
    size = nx*ny*sizeof(float);
    cutilSafeCall(cudaMemcpy(hW, dW, size, cudaMemcpyDeviceToHost));
 
    // メモリ解放
    cutilSafeCall(cudaFreeArray(caSrc));
    cutilSafeCall(cudaFree(dW));
 
    return hW;
}

画像データはテクスチャメモリに載せています. スレッド数は画像のピクセル数と同じになるようにして, 各スレッドがその座標でのスケールsのウェーブレット変換値を求めます.

カーネルの実装は以下です.

  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
 32
 33
 34
 35
 36
 37
 38
 39
 40
 41
 42
 43
 44
 45
 46
__device__ 
float MexicanHat2D(float x, float y)
{
    x = x*x;
    y = y*y;
    return MEXICAN_HAT_C*(x+y-2)*exp(-(x+y)/2.0);
}
 
__global__ 
void cwt2dtex(float *wt, int nx, int ny, float x0, float y0, float dx, float dy, float s, int res)
{
    // MARK:cwt2dtex
    int i = blockIdx.x*blockDim.x+threadIdx.x;
    int j = blockIdx.y*blockDim.y+threadIdx.y;
 
    if(i < nx && j < ny){
        float x = (x0+i*dx);
        float y = (y0+j*dy);
 
        float w = 0.0;
 
        // ウェーブレット関数の有効範囲
        float kx1 = -s*MEXICAN_HAT_R+x;
        float kx2 =  s*MEXICAN_HAT_R+x;
        float ky1 = -s*MEXICAN_HAT_R+y;
        float ky2 =  s*MEXICAN_HAT_R+y;
 
        kx1 = fmaxf(kx1, 0.0);
        kx2 = fminf(kx2, nx-1.0);
        ky1 = fmaxf(ky1, 0.0);
        ky2 = fminf(ky2, ny-1.0);
 
        // convolution
        float kstep = 1.0/res;
        float kx, ky;
        for(kx = kx1; kx <= kx2; kx += kstep){
            float Tx = (kx-x)/s;
            for(ky = ky1; ky <= ky2; ky += kstep){
                float Ty = (ky-y)/s;
 
                w += tex2D(g_TexImg, kx/nx, ky/ny)*MexicanHat2D(Tx,Ty);
            }
        }
        wt[i+j*nx] = w/(sqrt(s)*res*res);
    }
}

結果

CPU用にも実装し(OpenMPで2スレッド並列で実行),計算時間を比較します. CPU側の実装はウェーブレット関数値の前計算も行っています. 入力データとして,512×512のLena画像をグレースケール化して用いました.

Scalecudacpu
1385242
4395250
813919880
1650276298
321767257270

単位はミリ秒です.各スケールでの変換結果を以下に示します(256×256に縮小しています).

lena.jpglena_wt_s1.jpglena_wt_s4.jpg
lena_wt_s8.jpglena_wt_s16.jpglena_wt_s32.jpg

高速化はできたものの,スケール値を上げたときにタイムアウトエラーが出てプログラムが落ちます. これについてはCUDAカーネル実行のタイムアウトを参照してください. 本来は多数のスレッドで単純な問題を解くのがGPUに適しているので,タイムアウトしないレベルまで問題をさらに分解する方がよいかも知れません.

ソースコード

Visual Studio 2010用のソースコードを以下に置く.

  • 要libpng,libjpeg,CUDA4.0
  • キーボード'1','2','3'でそれぞれCPU,GPU,GPU(テクスチャメモリ使用)を切り替える.その後,'t'キーで再変換.
  • 'j','k'でスケーリングの変更.

添付ファイル: filelena_wt_s16.jpg 1128件 [詳細] filelena_wt_s8.jpg 1127件 [詳細] filelena_wt_s4.jpg 1126件 [詳細] filelena_wt_s32.jpg 1105件 [詳細] filelena_wt_s1.jpg 1158件 [詳細] filelena.jpg 1143件 [詳細] fileeq_2dcwt.jpg 601件 [詳細] filecu_cwt.zip 1277件 [詳細] fileeq_2dmexicanhat.jpg 573件 [詳細]

トップ   編集 凍結 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2011-10-27 (木) 15:09:13 (3254d)