title: Texture object fetching in CUDA
body: <p>I can find many examples online that use CUDA texture references, but not so many that rely on texture objects. I am trying to understand why my code below always fetches <code>0</code> rather than my input texture. Am I missing something, or using a wrong setting? I simplified it as much as I could:</p>
  <pre><code>#include &lt;stdio.h&gt;
  
  __global__ void fetch(cudaTextureObject_t tex, std::size_t width, std::size_t height) 
  {
    for (int j = 0; j &lt; height; j++) {
      for (int i = 0; i &lt; width; i++) {
        float u = (i + 0.5f) / width;
        float v = (j + 0.5f) / height;
  
        auto p = tex2D&lt;uchar4&gt;(tex, u, v);
        printf(&quot;i=%d, j=%d -&gt; u=%3.2f, v=%3.2f, r=%d, g=%d, b=%d, a=%d\n&quot;, i, j, u, v, p.x, p.y, p.z, p.w);
        // -&gt; always returns p = {0, 0, 0, 0}
      }
    }
  }
  
  int main() {
  
    constexpr std::size_t width = 2;
    constexpr std::size_t height = 2;
  
    // creating a dummy texture
    uchar4 image[width*height];
    for(std::size_t j = 0; j &lt; height; ++j) {
      for(std::size_t i = 0; i &lt; width; ++i)
        image[j*width+i] = make_uchar4(255*j/height, 255*i/width, 55, 255);
    }
  
    cudaArray_t cuArray;
    auto channelDesc = cudaCreateChannelDesc&lt;uchar4&gt;();
    cudaMallocArray(&amp;cuArray, &amp;channelDesc, width, height);
    cudaMemcpy2DToArray(cuArray, 0, 0, image, width*sizeof(uchar4), width*sizeof(uchar4), height, cudaMemcpyHostToDevice);
  
    struct cudaResourceDesc resDesc;
    memset(&amp;resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;
    resDesc.res.array.array = cuArray;
  
    struct cudaTextureDesc texDesc;
    memset(&amp;texDesc, 0, sizeof(texDesc));
    texDesc.addressMode[0]   = cudaAddressModeBorder;
    texDesc.addressMode[1]   = cudaAddressModeBorder;
    texDesc.filterMode       = cudaFilterModeLinear;
    texDesc.readMode         = cudaReadModeElementType;
    texDesc.normalizedCoords = 1;
  
    cudaTextureObject_t texObj = 0;
    cudaCreateTextureObject(&amp;texObj, &amp;resDesc, &amp;texDesc, NULL);
  
    fetch&lt;&lt;&lt;1, 1&gt;&gt;&gt;(texObj, width, height);
  
    cudaDeviceSynchronize();
  
    cudaDestroyTextureObject(texObj);
    cudaFreeArray(cuArray);
  
    return 0;
  }
  </code></pre>