Does cudaAddressModeBorder work with non-normalized coordinates?

Question

I'm using texture objects to access a PGM image pixels. my desire is to have texture fetch the pixel value in the given coordinate, or 0 if i'm out of boundaries.

this is my texture description:

unsigned char *device_input=NULL;
size_t input_pitch;
checkCudaErrors(cudaMallocPitch(&device_input, &input_pitch, sizeof(unsigned char)*IMAGE_WIDTH, IMAGE_HEIGHT));
checkCudaErrors(cudaMemcpy2D(device_input, input_pitch, image, sizeof(unsigned char)*IMAGE_WIDTH, sizeof(unsigned char)*IMAGE_WIDTH, IMAGE_HEIGHT, cudaMemcpyHostToDevice));

cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypePitch2D;
resDesc.res.pitch2D.devPtr = device_input; // 
resDesc.res.pitch2D.pitchInBytes =  input_pitch;
resDesc.res.pitch2D.width = IMAGE_WIDTH;
resDesc.res.pitch2D.height = IMAGE_HEIGHT;
resDesc.res.pitch2D.desc = cudaCreateChannelDesc<unsigned char>();

cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords=false;
texDesc.addressMode[0]=cudaAddressModeBorder;
texDesc.addressMode[1]=cudaAddressModeBorder;

cudaTextureObject_t tex;
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);

however, inside my kernel:

tex2D<unsigned char>(tex_inputImage,-100,-100)

which is obviously outside the boundaries of the image returns the value at image[0,0] instead of the value 0.

same goes for:

tex2D<unsigned char>(tex_inputImage,IMAGE_WIDTH+1,IMAGE_HEIGHT+1)

returns the value at image[IMAGE_WIDTH,IMAGE_HEIGHT] instead of 0.

please note that by using normalized Coordinates the cudaAddressModeBorder works as expected, but I don't want to use normalized Coordinates. according to nvidia's programming guide (Here), cudaAddressModeBorder is supported by non-normalized coordinates.

am I doing something wrong ?

Answer 1

Here is the answer to my own question:

The program ran on a machine with driver version of 319.32, apparently the driver had a bug that treated cudaAddressModeBorder like cudaAddressModeClamp when using normal coordinates (More on the problem here - check the last couple of replies).

the bug is fixed in version 319.49 and cudaAddressModeBorder works as expected with both normalized and non-normalized coordinates.