How to access each channel of a pixel using cuda tex2D

Question

I'm learning cuda texture memory. Now, I got a opencv Iplimage, and I get its imagedata. Then I bind a texture to this uchar array, like below:

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);

Now I lanch my kernel, and I want to access each pixel, every channel of that image. This is where I get confused.

I use this code to get the pixel coordinate (X,Y):

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

And how can I access each channel of this (X,Y)? what's the code below return?

tex2D(tex, X, Y);

Besides this, Can you tell me how texture memory using texture to access an array, and how this transform looks like?

B

Simply: int4 C = tex2D(tex, X, Y); int R = C.r; int G = C.g;... – Peltry

@Peltry But what text2D returns is a unsigned char, there is no proper way to convert uchar to int4. – B

You make an error in texture defining: texture<unsigned char,... will be a monochrome texture of uchar. If you want store a coloured image in texture, use uchar3 or uchar4. – Peltry

@Peltry Is there anything else I should change? I am sorry, but Im really new to cuda. When I change texture<unsigned char,...> to texture<uchar3,..>, I got a error: error : no instance of overloaded function "tex2D" matches the argument list argument types are: (texture<uchar3, 2, cudaReadModeElementType>, int, int) – B

Sorry, there's no uchar3 in available texture types, only uchar4, so you should make a simple program to convert your RGB data into RGBA array for texture binding. – Peltry

Answer 1

To bind a 3 channel OpenCV image to cudaArray texture, you have to create a cudaArray of width equal to image->width * image->nChannels, because the channels are stored interleaved by 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);

Now, to access each channel separately in the kernel, you just have to multiply the x index with number of channels and add the offset of desired channel like this:

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);

First one is blue because OpenCV stores images with channel sequence BGR.

As for the error you get when you try to access texture<uchar3,..> using tex2D; CUDA only supports creating 2D textures of 1,2 and 4 element vector types. Unfortunately, ONLY 3 is not supported which is very good for binding RGB images and is a really desirable feature.

Huehuebner

Great! Your answers in stackoverflow related to cuda are very helpful to me.Thank you! – B