CUDA 中的纹理对象获取

Question

我可以在网上找到很多使用 CUDA 纹理引用的示例，但依赖纹理对象的示例并不多。我试图理解为什么我下面的代码总是获取 0 而不是我的输入纹理。我是不是遗漏了什么，或者使用了错误的设置？我尽可能地简化了它：

#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

在您的代码中，您将纹理描述指定为

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

保存纹理数据的数组定义为

auto channelDesc = cudaCreateChannelDesc<uchar4>();

引用 documentation

Linear Filtering

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

你有一个 uchar4 纹理。您不能对整数纹理使用线性过滤。要么更改为浮点纹理类型，要么使用其他读取模式。

CUDA 中的纹理对象获取

Texture object fetching in CUDA

c++

textures

cuda

cuda-arrays