Texture question A 2 dimensional texture question

Accelerated Computing CUDA CUDA Programming and Performance
1

I have running some tests with the texture memory

texture<float,2> textReference;

__global__ void textTest ( float* Out  ) 

{

	*Out = tex2D(textReference,0.0f,0.f);			

}

void meV1(int height, int width )

{

	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

	float test[2][2];

	test[0][0] = 1.f;

	test[1][0] = 1.f;

	test[0][1] = 1.f;

	test[1][1] = 1.f;

	cudaArray* cuArray;

	width = 2;

	height= 2;

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

	cudaMemcpyToArray(cuArray, 0, 0, &test, width*height*sizeof(float), cudaMemcpyHostToDevice);

	textReference.addressMode[0] = cudaAddressModeWrap;

	textReference.addressMode[1] = cudaAddressModeWrap;

	textReference.filterMode = cudaFilterModeLinear;

	cudaBindTextureToArray(textReference, cuArray, channelDesc);

		....

	textTest<<<1,1>>>(output);

		....

}

But the value of output is 0.25 instead of 1.00 , anyone knows why?

Thanks!

2

Try to add 0.5f to the coordinate of the 2D texture fetching in linear mode.
If you read the appendice of the programming guide the formula given to calculate the interpolation begins by subtracting 0.5 in the coordinate.
so when you write (0.0 , 0.0 ) cuda read it at (-0.5 , -0. 5), the 4 nearest neighboors are (-1,-1) (-1,0) (0,-1) which result in zero because of clamp mode and your point (0,0) which result in 0.25*1.f
hope i am right , it works for me

3

That work, but it is amazing since, as you say, when using clamp mode it will make sense but I am using Wrap Mode.

Thanks anyway I will try to use this information carefully but if anyone has an explanation of why the wrap mode doesn’t work as it suppose to… will also be very helpful.

Thanks!

4

Sorry i didn’t see the warp mode.

the programming manual says “cudaAddressModeWrap is only supported for normalized texture coordinates” and you haven’t defined it ?

5

Ouch! I didn’t saw that. Thanks man!