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

Posted 2023-03-23

【中文标题】如何在 CUDA 内核中正确操作 CV_16SC3 Mat

【英文标题】：How to correctly manipulate a CV_16SC3 Mat in a CUDA Kernel

【发布时间】：2021-03-19 14:08:57

【问题描述】：

我在使用 OpenCV 时正在编写一个 CUDA 程序。我有一个给定大小（例如 1000x800）的空垫，我使用 dataytpe CV_16SC3 显式转换为 GPUMat。需要在 CUDA 内核中以这种格式操作图像。但是，尝试操纵 Mat 似乎无法正常工作。

我按如下方式调用我的 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 返回一个 uchar 指针。但是，我的 Mat 有一个数据类型“16U”（据我所知，短无符号）。这就是为什么在内核启动中我将指针转换为 (unsigned short*)。但显然这是不正确的。

我应该如何正确地在内核中读写 Mat 数据？

【参考方案1】：

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

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

方法一：

__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 更具可移植性，因为不同平台的类型大小可能会有所不同。

【讨论】：

非常感谢您的帮助！方法2对我有用。我会记住你提到的可移植性问题。