Texture clamp address mode does not work

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I’m finding that the clamp addressing mode (out of bound access) does not work for my 2D texture. Basically I get garbage, instead of the value that is at the edge of the texture. Has anybody encountered this?

Here is how I created the texture:

texture<unsigned short,2,cudaReadModeNormalizedFloat> texRef_s;

	size_t offset;

	cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned short>();

// set texture parameters

	texRef_s.addressMode[0] = cudaAddressModeClamp;

	texRef_s.addressMode[1] = cudaAddressModeClamp;

	texRef_s.filterMode = cudaFilterModeLinear;

	texRef_s.normalized = false; 

	

	bool status = true;

	cudaError_t CUDAstatus;

	if ((CUDAstatus =  cudaBindTexture2D (&offset, texRef_s, d_proj, desc, columns, rows, pitch)) != cudaSuccess) {

		fprintf(stderr,"Cuda Error : %s %s %d\n", cudaGetErrorString(CUDAstatus), __FILE__, __LINE__);

		status = false;

	}

// d_proj is an array allocated earlier with cudaMallocPitch(&d_proj, &pitch_s, sizeof(unsigned short)*columns, rows);

The texture gets bound correctly (otherwise, the code above would return an error).

Also, the values fetched from the interior of the texture (valid range access) are correct.

Any clues as to why I get garbage values for out-of-bounds texture fetches??

2

Actually, it seems the garbage texture fetches come only from the “bottom” of the 2D array.
(i.e. when the row index is greater than the number of rows in the texture)

I’m still clueless as to what’s going on…

3

Actually, it seems the garbage texture fetches come only from the “bottom” of the 2D array.
(i.e. when the row index is greater than the number of rows in the texture)

I’m still clueless as to what’s going on…

4

Sorry, dumb mistake
I specified the wrong number of rows in
cudaBindTexture2D (&offset, texRef_s, d_proj, desc, columns, rows, pitch)) != cudaSuccess)

5

Sorry, dumb mistake
I specified the wrong number of rows in
cudaBindTexture2D (&offset, texRef_s, d_proj, desc, columns, rows, pitch)) != cudaSuccess)