如何从线性内存创建cudaTextureObject_t？

Question

我无法获得引用线性内存的无绑定纹理 - 结果始终为零/黑色读取。我的初始化代码：

缓冲区：

int const num = 4 * 16;
int const size = num * sizeof(float);
cudaMalloc(buffer, size);

auto b = new float[num];
for (int i = 0; i < num; ++i)
{
    b[i] = i % 4 == 0 ? 1 : 1;
}
cudaMemcpy(*buffer, b, size, cudaMemcpyHostToDevice);

纹理对象：

cudaTextureDesc td;
memset(&td, 0, sizeof(td));

td.normalizedCoords = 0;
td.addressMode[0] = cudaAddressModeClamp;
td.addressMode[1] = cudaAddressModeClamp;
td.addressMode[2] = cudaAddressModeClamp;
td.readMode = cudaReadModeElementType;
td.sRGB = 0;

td.filterMode = cudaFilterModePoint;
td.maxAnisotropy = 16;

td.mipmapFilterMode = cudaFilterModePoint;
td.minMipmapLevelClamp = 0;
td.maxMipmapLevelClamp = 0;
td.mipmapLevelBias = 0;

struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = *buffer;
resDesc.res.linear.sizeInBytes = size;
resDesc.res.linear.desc.f = cudaChannelFormatKindFloat;
resDesc.res.linear.desc.x = 32;
resDesc.res.linear.desc.y = 32;
resDesc.res.linear.desc.z = 32;
resDesc.res.linear.desc.w = 32;

checkCudaErrors(cudaCreateTextureObject(texture, &resDesc, &td, nullptr));

内核：

__global__ void
    d_render(uchar4 *d_output, uint imageW, uint imageH, float* buffer, cudaTextureObject_t texture)
{
    uint x = blockIdx.x * blockDim.x + threadIdx.x;
    uint y = blockIdx.y * blockDim.y + threadIdx.y;

    if ((x < imageW) && (y < imageH))
    {
        // write output color
        uint i = y * imageW + x;
        //auto f = make_float4(buffer[0], buffer[1], buffer[2], buffer[3]);
        auto f = tex1D<float4>(texture, 0);
        d_output[i] = to_uchar4(f * 255);
    }
}

当给予内核时，纹理对象被初始化为合理的（4099）。 Buffer版本完美无瑕。

为什么纹理对象返回零/黑？

Answer 1

根据CUDA programming reference guide你需要使用tex1Dfetch()读取绑定到线性纹理内存的一维纹理，并使用tex1D读取绑定到CUDA数组的一维纹理。这适用于CUDA纹理参考和对象传递的CUDA纹理。

两个API之间的区别在于坐标参数。绑定到线性存储器的纹理只能在纹理坐标中寻址（因此text1Dfetch()中的整数坐标参数），而数组支持纹理和标准化坐标（因此tex1D中的浮点坐标参数）。