基于cuArray的double类型纹理对象问题：读取结果始终为0

你的代码读取结果始终为0，核心有三个问题：

1. 纹理读取函数与资源类型不匹配 #

tex1Dfetch仅适用于绑定到线性设备内存的纹理，而你使用的是cudaArray类型的纹理资源，必须使用对应数组资源的读取函数（如tex1D）。此外，tex1Dfetch的索引是整数偏移，你传入的0.4会被截断为0，但即使传入整数，函数类型不匹配也会导致读取错误。

2. CUDA数组通道格式与数据类型不匹配 #

你创建cudaArray时用了：

cudaCreateChannelDesc(32,32,0,0,cudaChannelFormatKindUnsigned);

这表示数组每个元素是由两个32位无符号整数组成的uint2，但你实际要存储的是double类型。直接用cudaMemcpyToArray复制double数据到这种格式的数组，会导致数据布局完全错乱，无法正确解析为uint2进而还原为double。

3. 纹理坐标模式未正确设置 #

默认情况下，纹理的normalizedCoords为true（归一化坐标模式），此时坐标范围是[0,1)对应整个数组的元素。如果你想使用绝对索引（比如直接读取第0个元素），需要手动将normalizedCoords设为false，否则传入的小数值会被当作归一化坐标处理，无法定位到正确的元素。

修正后的完整代码 #

#include <vector>
#include <iostream>
#include <cstdio>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess)
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void my_print(cudaTextureObject_t texObject)
{
    // 使用tex1D读取cudaArray纹理，绝对索引0读取第一个元素
    double dval = tex1D<double>(texObject, 0);
    printf("%.6f\n", dval);

    // 如果使用归一化坐标，打开下面的代码，0.0对应第一个元素，0.1对应第二个
    // double dval_norm = tex1D<double>(texObject, 0.0);
    // printf("normalized coord: %.6f\n", dval_norm);
}

int main()
{
    double i = 0.35;
    int numel = 10;

    std::vector<double> h_data(numel, i);

    // 1. 创建对应double类型的cudaArray
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<double>();
    cudaArray_t cArray;
    gpuErrchk(cudaMallocArray(&cArray, &channelDesc, numel, 0));

    // 2. 设置纹理描述：关闭归一化坐标，使用点过滤， clamp寻址
    cudaTextureDesc td{};
    td.filterMode = cudaFilterModePoint;
    td.addressMode[0] = cudaAddressModeClamp;
    td.readMode = cudaReadModeElementType;
    td.normalizedCoords = false; // 使用绝对索引，关闭归一化

    // 3. 配置纹理资源
    cudaResourceDesc resDesc{};
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cArray;

    cudaTextureObject_t texObject;
    gpuErrchk(cudaCreateTextureObject(&texObject, &resDesc, &td, NULL));

    // 4. 直接从主机复制数据到cudaArray（无需额外设备内存）
    gpuErrchk(cudaMemcpyToArray(cArray, 0, 0, h_data.data(), numel*sizeof(double), cudaMemcpyHostToDevice));

    // 5. 启动核函数
    my_print<<<1,1>>>(texObject);
    gpuErrchk(cudaDeviceSynchronize());

    // 清理资源
    gpuErrchk(cudaDestroyTextureObject(texObject));
    gpuErrchk(cudaFreeArray(cArray));

    return 0;
}

额外说明 #

如果你需要使用cudaAddressModeWrap（循环寻址），只需将td.addressMode[0]改为cudaAddressModeWrap即可，但注意：

• 循环寻址仅在归一化坐标模式下生效，此时需要将td.normalizedCoords设为true
• 读取时使用归一化坐标，比如tex1D<double>(texObject, 1.2)会被wrap为0.2，对应数组的第2个元素（索引1）

内容的提问来源于stack exchange，提问作者DocWho