title: Bound CUDA texture reads zero
body: <p>I try to read values from a texture and write them back to global memory.
  I am sure the writing part works, beause I can put constant values in the kernel and I can see them in the output:</p>
  
  <pre><code>__global__ void
  bartureKernel( float* g_odata, int width, int height) 
  {
      unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
      unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
  
      if(x &lt; width &amp;&amp; y &lt; height) {
              unsigned int idx = (y*width + x);
              g_odata[idx] = tex2D(texGrad, (float)x, (float)y).x;
  
      }
  }
  </code></pre>
  
  <p>The texture I want to use is a 2D float texture with two channels, so I defined it as:</p>
  
  <pre><code>texture&lt;float2, 2, cudaReadModeElementType&gt; texGrad;
  </code></pre>
  
  <p>And the code which calls the kernel initializes the texture with some constant non-zero values:</p>
  
  <pre><code>float* d_data_grad = NULL;
  
  cudaMalloc((void**) &amp;d_data_grad, gradientSize * sizeof(float));
  CHECK_CUDA_ERROR;
  
  texGrad.addressMode[0] = cudaAddressModeClamp;
  texGrad.addressMode[1] = cudaAddressModeClamp;
  texGrad.filterMode = cudaFilterModeLinear;
  texGrad.normalized = false;
  
  cudaMemset(d_data_grad, 50, gradientSize * sizeof(float));
  CHECK_CUDA_ERROR;
  
  cudaBindTexture(NULL, texGrad, d_data_grad, cudaCreateChannelDesc&lt;float2&gt;(), gradientSize * sizeof(float));
  
  float* d_data_barture = NULL;
  cudaMalloc((void**) &amp;d_data_barture, outputSize * sizeof(float));
  CHECK_CUDA_ERROR;
  
  dim3 dimBlock(8, 8, 1);
  dim3 dimGrid( ((width-1) / dimBlock.x)+1, ((height-1) / dimBlock.y)+1, 1);
  
  bartureKernel&lt;&lt;&lt; dimGrid, dimBlock, 0 &gt;&gt;&gt;( d_data_barture, width, height);
  </code></pre>
  
  <p>I know, setting the texture bytes to all "50" doesn't make much sense in the context of floats, but it should at least give me some non-zero values to read. </p>
  
  <p>I can only read zeros though...</p>