Texture object fetching in CUDA

Question

I can find many examples online that use CUDA texture references, but not so many that rely on texture objects. I am trying to understand why my code below always fetches 0 rather than my input texture. Am I missing something, or using a wrong setting? I simplified it as much as I could:

#include <stdio.h>

__global__ void fetch(cudaTextureObject_t tex, std::size_t width, std::size_t height) 
{
  for (int j = 0; j < height; j++) {
    for (int i = 0; i < width; i++) {
      float u = (i + 0.5f) / width;
      float v = (j + 0.5f) / height;

      auto p = tex2D<uchar4>(tex, u, v);
      printf("i=%d, j=%d -> u=%3.2f, v=%3.2f, r=%d, g=%d, b=%d, a=%d\n", i, j, u, v, p.x, p.y, p.z, p.w);
      // -> always returns p = {0, 0, 0, 0}
    }
  }
}

int main() {

  constexpr std::size_t width = 2;
  constexpr std::size_t height = 2;

  // creating a dummy texture
  uchar4 image[width*height];
  for(std::size_t j = 0; j < height; ++j) {
    for(std::size_t i = 0; i < width; ++i)
      image[j*width+i] = make_uchar4(255*j/height, 255*i/width, 55, 255);
  }

  cudaArray_t cuArray;
  auto channelDesc = cudaCreateChannelDesc<uchar4>();
  cudaMallocArray(&cuArray, &channelDesc, width, height);
  cudaMemcpy2DToArray(cuArray, 0, 0, image, width*sizeof(uchar4), width*sizeof(uchar4), height, cudaMemcpyHostToDevice);

  struct cudaResourceDesc resDesc;
  memset(&resDesc, 0, sizeof(resDesc));
  resDesc.resType = cudaResourceTypeArray;
  resDesc.res.array.array = cuArray;

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

  cudaTextureObject_t texObj = 0;
  cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

  fetch<<<1, 1>>>(texObj, width, height);

  cudaDeviceSynchronize();

  cudaDestroyTextureObject(texObj);
  cudaFreeArray(cuArray);

  return 0;
}

Answer 1

In your code you specify the texture description as

texDesc.addressMode[0]   = cudaAddressModeBorder;
texDesc.addressMode[1]   = cudaAddressModeBorder;
texDesc.filterMode       = cudaFilterModeLinear;
texDesc.readMode         = cudaReadModeElementType;
texDesc.normalizedCoords = 1;

and the array holding the texture data is defined as

auto channelDesc = cudaCreateChannelDesc<uchar4>();

Quoting the documentation

Linear Filtering

In this filtering mode, which is only available for floating-point textures ......

You have a uchar4 texture. You can't use linear filtering on an integer texture. Either change to a floating point texture type or use another read mode (probably cudaFilterModePoint).