How do I create a cudaTextureObject_t from linear memory?

Question

I cannot get bindless textures referencing linear memory to work -- the result is always a zero/black read. My initialization code:

The buffer:

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);

The texture object:

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));

The kernel:

__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);
    }
}

The texture object is initialized with something sensible (4099) when given to the kernel. The Buffer version works flawlessly.

Why does the texture object return zero/black?

Answer 1

As per the CUDA programming reference guide You need to use tex1Dfetch() to read from one-dimensional textures bound to linear texture memory, and tex1D to read from one-dimensional textures bound to CUDA arrays. This applies to both CUDA texture references and CUDA textures passed by object.

The difference between the two APIs is the coordinate argument. Textures bound to linear memory can only be addressed in texture coordinates (hence the integer coordinate argument in text1Dfetch()), whereas arrays support both texture and normalised coordinates (thus the float coordinate argument in tex1D).