title: How to read successfully from a 2D texture
body: <p>How can I:</p>
  
  <ol>
  <li>Bind cudaMallocPitch float memory to a 2D texture reference</li>
  <li>Copy some host data to the 2D array on the device</li>
  <li>Add one to the texture reference and write to either a.) the Pitch 2D array OR b.) write to a linear memory array </li>
  <li>Read the answer back and display it.</li>
  </ol>
  
  <p>Below is a code that should accomplish this. Note that for NxN array sizes, my code works. For NxM where N!=M, my code bites the dust (not the correct result). If you can solve this problem I will award you 1 internets (supply limited). Maybe I'm crazy, but according to the documentation this should work (and it does work for square arrays!). The attached code should run with 'nvcc whateveryoucallit.cu -o runit'.</p>
  
  <p>Help is appreciated!</p>
  
  <pre><code>#include&lt;stdio.h&gt;
  #include&lt;cuda.h&gt;
  #include&lt;iostream&gt;
  #define height 16
  #define width 11
  #define BLOCKSIZE 16
  
  using namespace std;
  
  // Device Kernels
  
  //Texture reference Declaration
  texture&lt;float,2&gt; texRefEx;
  
  
  __global__ void kernel_w_textures(float* devMPPtr, float * devMPtr, int pitch)
  {
   // Thread indexes
          unsigned int idx = blockIdx.x*blockDim.x + threadIdx.x;
          unsigned int idy = blockIdx.y*blockDim.y + threadIdx.y;
  
   // Texutre Coordinates
   float u=(idx)/float(width);
   float v=(idy)/float(height);
   devMPtr[idy*width+idx]=devMPPtr[idy*pitch/sizeof(float)+idx];
   // Write Texture Contents to malloc array +1
   devMPtr[idy*width+idx]= tex2D(texRefEx,u,v);//+1.0f;
  }
  int main()
  {
   // memory size
   size_t memsize=height*width;
   size_t offset;
   float * data,  // input from host
    *h_out,  // host space for output
    *devMPPtr, // malloc Pitch ptr
    *devMPtr; // malloc ptr
  
   size_t pitch;
  
   // Allocate space on the host
   data=(float *)malloc(sizeof(float)*memsize);
   h_out=(float *)malloc(sizeof(float)*memsize);
  
  
  // Define data
  for (int i = 0; i &lt;  height; i++)
   for (int j=0; j &lt; width; j++)
    data[i*width+j]=float(j);
  
  // Define the grid
  dim3 grid((int)(width/BLOCKSIZE)+1,(int)(height/BLOCKSIZE)+1), threads(BLOCKSIZE,BLOCKSIZE);
  
  // allocate Malloc Pitch
  cudaMallocPitch((void**)&amp;devMPPtr,&amp;pitch, width * sizeof(float), height);
  
  // Print the pitch
  printf("The pitch is %d \n",pitch/sizeof(float));
  
  // Texture Channel Description
  //cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc&lt;float&gt;();
  cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindFloat);
  
  // Bind texture to pitch mem:
  cudaBindTexture2D(&amp;offset,&amp;texRefEx,devMPPtr,&amp;channelDesc,width,height,pitch);
  cout &lt;&lt; "My Description x is " &lt;&lt; channelDesc.x &lt;&lt; endl;
  cout &lt;&lt; "My Description y is " &lt;&lt; channelDesc.y &lt;&lt; endl;
  cout &lt;&lt; "My Description z is " &lt;&lt; channelDesc.z &lt;&lt; endl;
  cout &lt;&lt; "My Description w is " &lt;&lt; channelDesc.w &lt;&lt; endl;
  cout &lt;&lt; "My Description kind is " &lt;&lt; channelDesc.f &lt;&lt; endl;
  cout &lt;&lt; "Offset is " &lt;&lt; offset &lt;&lt; endl;
  
  // Set mutable properties:
  texRefEx.normalized=true;
  texRefEx.addressMode[0]=cudaAddressModeWrap;
  texRefEx.addressMode[1]=cudaAddressModeWrap;
  texRefEx.filterMode= cudaFilterModePoint;
  
  // Allocate cudaMalloc memory
  cudaMalloc((void**)&amp;devMPtr,memsize*sizeof(float));
  
  // Read data from host to device
  cudaMemcpy2D((void*)devMPPtr,pitch,(void*)data,sizeof(float)*width,
    sizeof(float)*width,height,cudaMemcpyHostToDevice);
  
  //Read back and check this memory
  cudaMemcpy2D((void*)h_out,width*sizeof(float),(void*)devMPPtr,pitch,
    sizeof(float)*width,height,cudaMemcpyDeviceToHost);
  
  // Print the memory
   for (int i=0; i&lt;height; i++){
    for (int j=0; j&lt;width; j++){
     printf("%2.2f ",h_out[i*width+j]);
    }
   cout &lt;&lt; endl;
   }
  
   cout &lt;&lt; "Done" &lt;&lt; endl;
  // Memory is fine... 
  
  kernel_w_textures&lt;&lt;&lt;grid,threads&gt;&gt;&gt;(devMPPtr, devMPtr, pitch);
  
  // Copy back data to host
  cudaMemcpy((void*)h_out,(void*)devMPtr,width*height*sizeof(float),cudaMemcpyDeviceToHost);
  
  
  // Print the Result
   cout &lt;&lt; endl;
   for (int i=0; i&lt;height; i++){
    for (int j=0; j&lt;width; j++){
     printf("%2.2f ",h_out[i*width+j]);
    }
   cout &lt;&lt; endl;
   }
   cout &lt;&lt; "Done" &lt;&lt; endl;
  
  return(0);
  }
  </code></pre>
  
  <p>Edit October 17: So I still haven't found a solution to this issue. Nvidia is pretty silent on this seems that the world is too. I found a workaround using shared mem but if anyone has a texture solution I would be very please.</p>
  
  <p><strong>Edit</strong> Octoboer 26: Still no soltuion, but still interested in one if anyone knows.</p>
  
  <p><strong>Edit</strong> July 26: Wow it has been 9 months - and I had overlooked the correct answer the whole time. The trick was:</p>
  
  <pre><code>if ( idx &lt; width  &amp;&amp; idy &lt; height){//.... code }
  </code></pre>
  
  <p>As had been pointed out before. Thanks to all of those who contributed!</p>