texture : unsigned char and FilterModeLinear how to use an unsigned char texture with linear interpo

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

let assume I have a texture :

texture<unsigned char,2,cudaReadModeElementType> tex;

I don’t want to have float element in the texture but unsigned char.

I want to have a linear interpolation between my datas.

The programming guide say : cudaFilterModeLinear is only valid for returned value of floating-point type.

So I define the texture like this :

tex.adressMode[0]= cudaAddressModeClamp;

tex.adressMode[1]= cudaAddressModeClamp;

tex.filterMode=cudaFilterModeLinear;//to have the interpolation

tex.channelDesc= cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);//to return value of floating-point type.

tex.normalized=0;

Now the return value from texture is a floating-point type so normally I can use cudaFilterModeLinear.

I also bind the texture with an array (of unsigned char)

cudaBindTexture2D(0,tex,array,cudaCreateChannelDesc(8,0,0,0,cudaChannelFormatKindUnsigned),width,height,pitch);

And finaly in the kernel i perform some computations and i copy them into an other array (of unsigned char)

__global__ myKernel(...)

...

d_array2[index] = tex2D(tex,x,y);

...

I don’t get the correct values.

Someone know how to use filterModeLinear with unsigned char?

Thanks

2

Any idee?

cudaFilterModeLinear is only valid for returned value of floating-point type.

How I can do if I don’t want to have float element in the texture but unsigned char (for speed reasons) and I need an linear interpolation?
I can’t trust that there is no possible solution to use the linear interpolation with an texture who is binded with an array of unsigned char?

Thanks

3

Pre-interpolate the texture yourself (e.g. 16x oversampling) - and store a bigger texture.

4

Answer to myself :)

I have just find the solution and maybe it could help some people.

I have done a mistake.

texture<unsigned char,2, cudaReadModeElementType cudaReadNormalizedFloat> tex;
(and also tex.normalized=1;)
Now it’s work!!

I can use a texture of unsigned char binded with an array of unsigned char with an linear interpolation (who is is only valid for returned value of floating-point type).

5

Yes, linear filtering only works with texture references that are configured to return normalized floating point results (cudaReadNormalizedFloat).