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

Question

在使用OpenCV时，我正在编写一个CUDA程序。我有一个给定大小(例如1000x800(的空Mat，我用dataytpe 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返回一个uchar指针。然而，我的Mat有一个数据类型"；16U"；(据我所知，缩写为unsigned(。这就是为什么在内核启动时，我将指针投射到(unsigned short*(。但显然这是不正确的。

我应该如何正确地读取和写入内核中的Mat数据

Answer 1

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

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

方法1：

__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;
}

方法2：

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

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

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