时间戳

首页 > 解决方案 > 在 CUDA 中获取纹理对象

c++ - 在 CUDA 中获取纹理对象

问题描述

我可以在网上找到很多使用 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;
}

标签： c++cudatexturescuda-arrays

解决方案

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

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

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

auto channelDesc = cudaCreateChannelDesc<uchar4>();

引用文档

线性滤波

在这种过滤模式下，它只适用于浮点纹理......

你有uchar4质感。您不能对整数纹理使用线性过滤。更改为浮点纹理类型或使用其他读取模式。

推荐阅读