Bind a uchar3 texture

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

I would like to interpolate my RGB texture but there is trouble with uchar3 in texture.

here is parts of my program:

texture<uchar3, 2,cudaReadModeElementType> ImageInRGB;

cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar3>();
            checkCudaErrors(cudaBindTexture2D(0,&ImageInRGB,u8_Image_Device,&desc,iwidth_,iheight_,iwidth_*3));

            dim3 blocks(ceil((float)iwidth_*iSamplingCoefficient_ / ( BLOCK_SIZE_X)), ceil((float)iheight_*iSamplingCoefficient_ / BLOCK_SIZE_Y));
            dim3 threads(BLOCK_SIZE_X, BLOCK_SIZE_Y);
            UpSamplingKernelRGB<<<blocks,threads>>>(ptImageOut_Device, iwidth_*iSamplingCoefficient_,iheight_*iSamplingCoefficient_,iSamplingCoefficient_);

__global__ void UpSamplingKernelRGB(u_int8_t *ptOut,  int  w, int h,int iUpSamplingValue)
{
.....
         uchar3 data = tex2D(ImageInRGB,ix,iy); ### ERROR in compilation

.....
}

I get that error in my compilation:

error: no instance of overloaded function "tex2D" matches the argument list
            argument types are: (texture<uchar3, 2, cudaReadModeElementType>, int, int)

Is it possible to interpolate uchar3 texture?

2

According to the programming guide:

[url]http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#texture-and-surface-memory[/url]

“a texel … is restricted to the basic integer and single-precision floating-point types and any of the 1-, 2-, and 4-component vector types defined in char, short, int, long, longlong, float, double that are derived from the basic integer and single-precision floating-point types.”

uchar3 does not fit that description

3

So is there a way to interpolate a rgb texture?
Should I have to spread in different color and interpolate them separatly?

4

I don’t know what you mean by interpolate.

If you want to texture against RGB, you could use uchar4, and ignore the 4th component in subsequent calculations (e.g. your own interpolation calculations). This will require the underlying storage to be arranged as uchar4, e.g. RGBA, not uchar3.

If you mean use the built-in interpolation capabilities of the texture hardware, you’ll need to be more specific with an example of what you have in mind.

5

3-channel interleaved images can not be bound to a texture (object).

But see my last answer at
https://devtalk.nvidia.com/default/topic/481747/is-texture-60-float3-3-62-possible-cudaerrorinvalidchanneldescriptor/?offset=5
for a workaround

I treat the images as one-channel images with three times the width and bind a texture to the one-channel image.
But note for accessing pixels outside of the image dimension you won’t get the right value (according to specified border mode) from the texture hardward - so you have to implement yourself. The same applies for interpolation - you have to implement it yourself, the value delivered from texture hardware is wrong.

You can ‘hide’ all that dirty stuff inside short access functions. The additional overhead does not seem to be a problem in practice.

6

Thank you for the answers.

An other question about cudaBinTexture2d, about the height and width parameters.
Are the width always had to be greater than height? Because I had images in landscape mode and the function send me error with these images.