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.