Textures in CUDA

Hi Everyone,

In the algorithm I’m trying to implement, I want to take 2 rgba images, process them on the GPU, then give back another rgba image as output. So far I’ve been trying to convert the images into textures so that my kernel can read the data from textures- my idea being to read in the array of unsigned char that represents my image data into the texture so that for each uchar4 element of the texture, x = r, y = g, z = b, w = a. However looking at the output its become clear to me that my texture data may not be holding the rgba data in the way I expected it to be. My code so far looks something like:

cudaChannelFormatDesc fmt = cudaCreateChannelDesc<uchar4>();	




	CUDA_SAFE_CALL(cudaMemcpy2DToArray(d_l_img,0,0,l_img,WIDTH*sizeof(unsigned char)*4,WIDTH,HEIGHT,cudaMemcpyHostToDevice));

	CUDA_SAFE_CALL(cudaMemcpy2DToArray(d_r_img,0,0,r_img,WIDTH*sizeof(unsigned char)*4,WIDTH,HEIGHT,cudaMemcpyHostToDevice));

CUDA_SAFE_CALL(cudaBindTextureToArray(l_tex, d_l_img));

CUDA_SAFE_CALL(cudaBindTextureToArray(r_tex, d_r_img));

where d_l_img and d_r_img are cudaArrays, WIDTH is the width of the image, HEIGHT is the height of the image, and l_tex and r_tex are the textures which are declared as:

texture<uchar4,2,cudaReadModeNormalizedFloat> l_tex;

texture<uchar4,2,cudaReadModeNormalizedFloat> r_tex;

Although I’ve looked up on past posts, I’m not sure what the consensus is regarding using chars rather than floats for textures- is it feasible? I would really like to be able to access the rgba data in my algorithm as numbers between 0-255 rather than as numbers between 0-1. It will make my life much easier.

Sorry if I’m missing a vital point regarding textures- the documentation on them isn’t as complete as I’d like and I realize I may be confused in many areas. If something’s not clear, please let me know. I really appreciate the help, thanks!