如何在 CUDA 内核中正确操作 CV_16SC3 Mat

Question

我正在使用 OpenCV 编写 CUDA 程序。我有一个给定大小的空垫子（例如 1000x800），我使用数据类型 CV_16SC3 将其显式转换为 GPUMat。希望在 CUDA 内核中以这种格式操作图像。然而，尝试操纵垫子似乎无法正常工作。

我调用我的 CUDA 内核如下：

    my_kernel <<< gridDim, blockDim >>>( (unsigned short*)img.data, img.cols, img.rows, img.step);

我的示例内核如下所示

__global__ void my_kernel( unsigned short* img, int width, int height, int img_step)
{
    int x, y, pixel;
    y = blockIdx.y * blockDim.y + threadIdx.y;
    x = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (y >= height)
        return;
    
    if (x >= width)
        return;

    pixel = (y * (img_step)) + (3 * x);
    
    img[pixel] = 255; //I know 255 is basically an uchar, this is just part of my test
    img[pixel+1] = 255
    img[pixel+2] = 255;

}

我期待这个小内核样本将所有像素写入白色。然而，在从 GPU 再次下载 Mat 并使用 imshow 对其进行可视化后，并非所有像素都是白色的，并且存在一些奇怪的黑线，这让我相信我正在以某种方式写信给无效的内存地址。

我的猜测如下。 OpenCV 文档指出 cv::mat::data returns 一个 uchar 指针。然而，我的 Mat 有一个数据类型“16U”（据我所知是短无符号的）。这就是为什么在内核启动时我将指针转换为 (unsigned short*)。但显然这是不正确的。

我应该如何正确地进行才能在我的内核中读取和写入尽可能短的 Mat 数据？

Answer 1

首先，输入图像类型应该是short而不是unsigned short，因为Mat的类型是16SC3（而不是16UC3）。

现在，由于图像步长以字节为单位且数据类型为 short，因此计算像素索引（或地址）时应考虑到它们的字节宽度差异。有 2 种方法可以解决此问题。

方法一：

__global__ void my_kernel( short* img, int width, int height, int img_step)
{
    int x, y, pixel;
    y = blockIdx.y * blockDim.y + threadIdx.y;
    x = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (y >= height)
        return;
    
    if (x >= width)
        return;

    //Reinterpret the input pointer as char* to allow jump in bytes instead of short
    char* imgBytes = reinterpret_cast<char*>(img);
    
    //Calculate row start address using the newly created pointer
    char* rowStartBytes = imgBytes + (y * img_step); // Jump in byte
    
    //Reinterpret the row start address back to required data type.
    short* rowStartShort = reinterpret_cast<short*>(rowStartBytes);
    
    short* pixelAddress = rowStartShort + ( 3 * x ); // Jump in short
    
    //Modify the image values
    pixelAddress[0] = 255; 
    pixelAddress[1] = 255;
    pixelAddress[2] = 255;

}

方法二：

将输入图像步长除以所需数据类型的大小 (short)。将步骤作为内核参数传递时可能会完成。

my_kernel<<<grid,block>>>( img, width, height, img_step/sizeof(short));

我已经使用方法2很长时间了。这是一种快捷方式，但后来当我查看某些图像处理库的源代码时，我意识到实际上方法 1 更可移植，因为类型的大小在不同平台上可能会有所不同。

如何在 CUDA 内核中正确操作 CV_16SC3 Mat

How to correctly manipulate a CV_16SC3 Mat in a CUDA Kernel

opencv

cuda

方法一：

方法二：