How to change the adress mode of a texture?

Hello,

I am trying to change the adress mode of a texture, but I don’t see any effect.

My code:

texture<uint8_t, cudaTextureType2D, cudaReadModeNormalizedFloat> texI;

udaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uint8_t>();

texI.filterMode = cudaFilterModeLinear;
texI.addressMode[0] = cudaAddressModeBorder;
texI.addressMode[1] = cudaAddressModeBorder;


//in the kernel:
for(float f = -10.5; f <= 10.5; f += 1){
  printf("%.1f ", tex2D(texI,f,1.5)*255);
}

No matter, which adress mode I choose, the fetched values are clamped.

What am I doing wrong?

May be https://devtalk.nvidia.com/default/topic/479435/cudaaddressmodeborder/
helps?
Bill

Hi,

Just off the top of my head:

You seem to be using normalized coordinates. Hence they should be in the range of [0, 1] which you don’t seem to be.

With that type of loop you would almost always be on the border…

EDIT: try cudaReadModeElementType instead?

Hello Bill, unfortunately that did not help.

Hello Jimmy, I’m not using normalized coordinates, it’s just the return values that are normalised floats (my texture consists of unsigned 8-bit integers and the hardware casts them into normalized floats). I tried cudaReadModeElementType, but the outcome was the same (apart from the type that was returned).

Does anyone has another idea?

Are you using CUDA 5.0? I vaguely remember that cudaAddressModeBorder was unimplemented in some earlier releases (don’t recall which ones though).

Are you trying the change the address mode while the texture is bound? I believe you need to rebind the texture in that case.

Yes, I am using CUDA 5.
I’m doing the change before binding the texture. Here the reduced code:

texture<uint8_t, cudaTextureType2D, cudaReadModeNormalizedFloat> texI;

void call_kernel( uint8_t* I, uint32_t IPitch)
{
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uint8_t>();

        texI.filterMode = cudaFilterModeLinear;
	texI.addressMode[0] = cudaAddressModeWrap;
	texI.addressMode[1] = cudaAddressModeWrap;

	
        CudaSafeCall( cudaBindTexture2D(NULL, texI, I, channelDesc, TEMPLATE_WIDTH + 20, TEMPLATE_HEIGHT + 20, IPitch) );

	kernel<<<1,1>>>();

	CudaSafeCall( cudaUnbindTexture(texI) );
}

//kernel:
__global__ void kernel(){	
  	for(float f = -10.5; f <= 10.5; f += 1){
  		printf("%.1f ", tex2D(texI,f,1.5)*255);
  	}
}

It does not matter, which address mode I chose, it allways behaves like clamped address mode.

As far as I know, cudaAddressModeBorder works just fine but has the limitation that one cannot set the border value which currently is always zero. That is to say, an out of bounds texture access returns zero. There was an issue with missing documentation regarding this mode prior to CUDA 4.1 (see https://devtalk.nvidia.com/default/topic/502623/clamping-texture-to-zero-adding-a-zero-border)

The code in #6 above appears to use unnormalized texture addressing with cudaModeAddressWrap. Note that according to the CUDA C Programming Guide “cudaAddressModeWrap and cudaAddressModeMirror are only supported for normalized texture coordinates”

Thanks Norbert for refreshing my memory. That is exactly what I was trying to refer to.

I attempted to write a very simple test app using a 2D texture to demonstrates the difference between cudaAddressModeBorder and cudaAddressModeClamp. The result is that I do not see any difference.

My understanding was that an out-of-bounds access with cudaAddressModeBorder should return the border value (which should be fixed at zero), whereas an out-of-bounds access with cudaAddressModeClamp returns the appropriate edge value in such a case. But what I am seeing is that both modes show clamp-to-edge behavior, which puzzles me.

I will follow up on my side.

That is exactly what I experienced.

The “wrap” address mode in that code was to demonstrate, that neither of the address modes made a change. I did not know, that cudaAddressModeWrap and cudaAddressModeMirror are only supported for normalized texture coordinates. Thank you for that!

There are indications that cudaAddressModeBorder may also apply to normalized texture coordinates only. The issue is being investigated and a bug has been filed.

+1 to this bug. Is there any actual bug report website registered developers can see? I can only find the form to submit a bug, but not to browse the tracker.

Seeing this with a Geforce 650Ti using CUDA/C++ VS2010 SP1, CUDA 5, 320.18 NVIDIA driver.

This used to work fine with non-normalized coordinates and totally caught us by surprise.

Sorry for the inconvenience. I have used BORDER mode with non-normalized addressing myself in the past and it worked just fine. The appropriate team is actively working on this issue. As far as I am aware, the bug reporting system allows customers to track only issues that they filed themselves.

Bump. Can we get an official update here?

Are you still seeing the issue with CUDA 5.5 and the accompanying driver?