Texture memory-tex2D basics

Question

While using texture memory I have come across the following code:-

uint f = (blockIdx.x * blockDim.x) + threadIdx.x;
uint c = (blockIdx.y * blockDim.y) + threadIdx.y;

uint read = tex2D( refTex, c+0.5f, f+0.5f);

My question is why do we add 0.5f to both c and f? This confuses me.. thankyou

Answer 1

In graphics, a texture is a set of samples that describes the visual appearance of a surface. A sample is a point. That is, it has no size (as opposed to a pixel that has a physical size). When using samples to determine the colors of pixels, each sample is positioned in the exact center of its corresponding pixel. When addressing pixels with whole number coordinates, the exact center for a given pixel becomes its whole number coordinate plus an offset of 0.5 (in each dimension).

In other words, adding 0.5 to texture coordinates ensures that, when reading from those coordinates, the exact value of the sample for that pixel is returned.

However, it is only when filterMode for the texture has been set to cudaFilterModeLinear that the value that is read from a texture varies within a pixel. In that mode, reading from coordinates that are not in the exact center of a pixel returns values that are interpolated between the sample for the given pixel and the samples for neighboring pixels. So, adding 0.5 to whole number coordinates effectively negates the cudaFilterModeLinear mode. But, since adding 0.5 to the texture coordinates takes up cycles in the kernel, it is better to simply turn off the interpolation by setting filterMode to cudaFilterModePoint. Then, reading from any coordinate within a pixel returns the exact texture sample value for that pixel, and so, texture samples can be read directly by using whole numbers.

When using cudaFilterModePoint, if any floating point math is involved in calculating the texture coordinates, care must be taken to ensure that floating point inaccuracies don't cause the texture coordinates to fall outside the range for the intended target pixel.

Also, as the comments mention, there might be a problem in your code. Adding 0.5f to the texture coordinates implies that the cudaFilterModeLinear mode is being used, but that mode returns a float, not an int.

Answer 2

Depending on the texture properties, the value returned by tex2D may be linearly interpolated. In that case, the "indices" f and c should not be integers, but continuous values between the limits of each dimension.

What is a bit odd in this example is that the return value is an integer, which will make any linear interpolant piecewise constant anyway.

For more detailed information, see Section 3.2.10 of the CUDA Programming Guide.