RGB texture lookup

[yahastu]

EDIT - poking around in the examples, I discovered that the imageDenoising example uses a single array and just packs/unpacks the float3 into a single unsigned int. That seems workable so I’ll just assume that’s standard and go with it. I’ve also got it working using 3 different 32-bit textures if I want more precision…although I’m storing them in 3 different arrays

---

Forgive me if this is covered elsewhere. I’ve read the programming guide and didn’t find the answer there, and when I searched here all I found was other people asking the same question with no answers.

I start with an image buffer. Each channel is stored in row-major format. In other words, if v = channel index, then

img[ vwidthheight + width*y + x ] = img(x,y,v)

I can read this as a 1-channel image by setting it up liek this:

texture<float, 2, cudaReadModeElementType> texRef;

…

// Allocate CUDA array in device memory

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);

cudaArray *cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);

// Copy to device memory some data located at address imgData in host memory

size_t size = sizeof(float) * width * height;

cudaMemcpyToArray(cuArray, 0, 0, imgData, size, cudaMemcpyHostToDevice);

// Set texture parameters

texRef.addressMode[0] = cudaAddressModeClamp;

texRef.addressMode[1] = cudaAddressModeClamp;

texRef.filterMode = cudaFilterModePoint;

texRef.normalized = false;

// Bind the array to the texture

cudaBindTextureToArray(texRef, cuArray, channelDesc);

…

float samp = tex2D(texRef, ix, iy);

Now, I wanted to modify that to work for 3 channels. I have no idea how to describe to CUDA the format of the values…how to tell it I am using RRRR GGG BBB vs RGB RGB RGB, etc. Perhaps that’s my problem. Anyway I used float4 since you’re not allowed to use float3. This is what I did:

texture<float4, 2, cudaReadModeElementType> texRef;

…

// Allocate CUDA array in device memory

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 32, 32, 0, cudaChannelFormatKindFloat);

cudaArray *cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);

// Copy to device memory some data located at address imgData in host memory

size_t size = sizeof(float) * width * height * 3;

cudaMemcpyToArray(cuArray, 0, 0, imgData, size, cudaMemcpyHostToDevice);

// Set texture parameters

texRef.addressMode[0] = cudaAddressModeClamp;

texRef.addressMode[1] = cudaAddressModeClamp;

texRef.filterMode = cudaFilterModePoint;

texRef.normalized = false;

// Bind the array to the texture

cudaBindTextureToArray(texRef, cuArray, channelDesc);

…

float4 samp = tex2D(texRef, ix, iy);

That didn’t work. I’m not exactly sure what part didn’t work because I’m just getting started and haven’t mastered debugging, but I just didn’t get any output. I suppose I could work around this by creating 3 one-dimensional textures for each of the Red, green, and blue channels…but is that the preferred way?