CUDAGpuMatで3ch画像の扱い方

GpuMatの3ch画像の扱いに少し手間取ったのでメモ

GpuMatの構造はBGRBGR・・・とチャンネルが続いていく。 CUDAカーネルによってGpuMatを扱うときは PtrStepSzに変換すると扱いやすい

例えば

using namespace cv;


Scalar black(0, 0, 0);
Mat img(Size(1000, 1000), CV_8UC3, black);
cuda::GpuMat gpuImg(img);
cuda::PtrStepSz<uchar> ptrImg = cuda::PtrStepSz<uchar>(gpuImg.rows, gpuImg.cols * gpuImg.channels(), gpuImg.ptr<uchar>(),gpuImg.step);

とかくことによりPtrStepSz型に変換できる。 PtrStepSzはCUDAカーネルの引数に指定するのに便利な型である。

CUDAカーネル内で３chのPtrStepSzを扱う際は以下のようなマクロを定義すると便利

#define GRU(IMG, X, Y) ((IMG).data[(IMG).step * (Y) + 3 * sizeof(unsigned char) * X + sizeof(unsigned char) * 2])
#define GGU(IMG, X, Y) ((IMG).data[(IMG).step * (Y) + 3 * sizeof(unsigned char) * X + sizeof(unsigned char) * 1])
#define GBU(IMG, X, Y) ((IMG).data[(IMG).step * (Y) + 3 * sizeof(unsigned char) * X + sizeof(unsigned char) * 0])

このマクロを定義した後CUDAカーネル内でたとえば (R,G,B) = (255, 0, 0)の画像を作成するときは

__global__ void myKernel(cuda::PtrStepSz<uchar> ptrImg){
    const int x = blockDim.x * blockIdx.x + threadIdx.x;
    const int y = blockDim.y * blockIdx.y + threadIdx.y;

    if((x < ptrImg.cols) && (y < ptrImg.rows)){
        GRU(ptrImg, x, y) = 255;
        GGU(ptrImg, x, y) = 0;
        GBU(ptrImg, x, y) = 0;
    }
}

によって作成することができる。