title: How do I create a cudaTextureObject_t from linear memory?
body: <p>I cannot get bindless textures referencing linear memory to work -- the result is always a zero/black read. My initialization code:</p>
  
  <p>The buffer:</p>
  
  <pre><code>int const num = 4 * 16;
  int const size = num * sizeof(float);
  cudaMalloc(buffer, size);
  
  auto b = new float[num];
  for (int i = 0; i &lt; num; ++i)
  {
      b[i] = i % 4 == 0 ? 1 : 1;
  }
  cudaMemcpy(*buffer, b, size, cudaMemcpyHostToDevice);
  </code></pre>
  
  <p>The texture object:</p>
  
  <pre><code>cudaTextureDesc td;
  memset(&amp;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(&amp;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, &amp;resDesc, &amp;td, nullptr));
  </code></pre>
  
  <p>The kernel:</p>
  
  <pre><code>__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 &lt; imageW) &amp;&amp; (y &lt; imageH))
      {
          // write output color
          uint i = y * imageW + x;
          //auto f = make_float4(buffer[0], buffer[1], buffer[2], buffer[3]);
          auto f = tex1D&lt;float4&gt;(texture, 0);
          d_output[i] = to_uchar4(f * 255);
      }
  }
  </code></pre>
  
  <p>The texture object is initialized with something sensible (4099) when given to the kernel. The Buffer version works flawlessly.</p>
  
  <p>Why does the texture object return zero/black?</p>