如何使用Cuda tex2D访问像素的每个通道

Question

我正在学习cuda纹理记忆。现在，我得到了一个opencv Iplimage，并且得到了它的图像数据。然后，将纹理绑定到此uchar数组，如下所示：

Iplimage *image = cvCreateImage(cvSize(width, height), IPL_DEPTH_8U, 3);
unsigned char* imageDataArray = (unsigned char*)image->imagedata;

texture<unsigned char,2,cudaReadModeElementType> tex;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 0, 
                                                          cudaChannelFormatKindUnsigned); 
cudaArray *cuArray = NULL;
CudaSafeCall(cudaMallocArray(&cuArray,&channelDesc,width,height));

cudaMemcpy2DToArray(cuArray,0,0,imageDataArray,image->widthstep,
    width * sizeof(unsigned char), height, cudaMemcpyHostToDevice);
cudaBindTextureToArray(texC1_cf,cuArray_currentFrame, channelDesc);

现在，我启动内核，并希望访问该图像的每个像素，每个通道。这就是我感到困惑的地方。

我使用以下代码获取像素坐标（X，Y）：

int X = (blockIdx.x*blockDim.x+threadIdx.x);
int Y = (blockIdx.y*blockDim.y+threadIdx.y);

以及如何访问此（X，Y）的每个通道？return下面的代码是什么？

tex2D(tex, X, Y);

除此之外，您能告诉我如何使用纹理访问数组的纹理内存以及此变换的外观吗？

Answer 1

要将 3 通道 OpenCV 图像绑定到 cudaArray 纹理，您必须创建一个宽度等于 的 cudaArray image->width * image->nChannels，因为通道是由 OpenCV 交错存储的。

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();

cudaArray *cuArray = NULL;
CudaSafeCall(cudaMallocArray(&cuArray,&channelDesc,width * image->nChannels,height));

cudaMemcpy2DToArray(cuArray,0,0,imageDataArray,image->widthstep, width * image->nChannels * sizeof(unsigned char), height, cudaMemcpyHostToDevice);

cudaBindTextureToArray(texC1_cf,cuArray_currentFrame, channelDesc);

现在，要在内核中单独访问每个通道，您只需将 x 索引乘以通道数，然后添加所需通道的偏移量，如下所示：

unsigned char blue = tex2D(tex, (3 * X) , Y);
unsigned char green = tex2D(tex, (3 * X) + 1, Y);
unsigned char red = tex2D(tex, (3 * X) + 2, Y);

第一个是蓝色的，因为 OpenCV 存储具有通道序列 BGR 的图像。

至于当您尝试使用texture<uchar3,..>;tex2D访问时出现的错误 CUDA 仅支持创建 1,2 和 4 元素向量类型的 2D 纹理。不幸的是，不支持 ONLY 3，这非常适合绑定 RGB 图像，并且是一个非常理想的功能。