c++ - 如何从线性内存创建 cudaTextureObject_t？

Question

我无法使引用线性内存的无绑定(bind)纹理正常工作——结果始终是零/黑色读取。我的初始化代码:

缓冲区:

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() 从绑定(bind)到线性纹理内存的一维纹理中读取，并使用 tex1D 从绑定(bind)到 CUDA 数组的一维纹理中读取。这适用于 CUDA 纹理引用和对象传递的 CUDA 纹理。

这两个 API 的区别在于坐标参数。绑定(bind)到线性内存的纹理只能在纹理坐标中寻址(因此 text1Dfetch() 中的整数坐标参数)，而数组同时支持纹理和归一化坐标(因此 tex1D 中的浮点坐标参数) )。