title: Is 1D texture memory access faster than 1D global memory access?
body: <p>I am measuring the difference between the standard and 1Dtexture  access to memory. To do so I have created two kernels</p>
  
  <pre><code>__global__ void texture1D(float* doarray,int size)
  {
    int index;
    //calculate each thread global index
    index=blockIdx.x*blockDim.x+threadIdx.x;
    //fetch global memory through texture reference
    doarray[index]=tex1Dfetch(texreference,index);
    return;
  }
  __global__ void standard1D(float* diarray, float* doarray, int size)
  {
    int index;
    //calculate each thread global index
    index=blockIdx.x*blockDim.x+threadIdx.x;
    //fetch global memory through texture reference
    doarray[index]= diarray[index];
    return;
  }
  </code></pre>
  
  <p>Then, I call eache kernel measuring the time it takes:</p>
  
  <pre><code>//copy array from host to device memory
    cudaMemcpy(diarray,harray,sizeof(float)*size,cudaMemcpyHostToDevice);
  
    checkCuda( cudaEventCreate(&amp;startEvent) );
    checkCuda( cudaEventCreate(&amp;stopEvent) );
    checkCuda( cudaEventRecord(startEvent, 0) );
  
    //bind texture reference with linear memory
    cudaBindTexture(0,texreference,diarray,sizeof(float)*size);
  
    //execute device kernel
    texture1D&lt;&lt;&lt;(int)ceil((float)size/threadSize),threadSize&gt;&gt;&gt;(doarray,size);
  
    //unbind texture reference to free resource
    cudaUnbindTexture(texreference);
  
    checkCuda( cudaEventRecord(stopEvent, 0) );
    checkCuda( cudaEventSynchronize(stopEvent) );
  
    //copy result array from device to host memory
    cudaMemcpy(horray,doarray,sizeof(float)*size,cudaMemcpyDeviceToHost);
  
    //check result
    checkResutl(horray, harray, size);
  
    cudaEvent_t startEvent2, stopEvent2;
    checkCuda( cudaEventCreate(&amp;startEvent2) );
    checkCuda( cudaEventCreate(&amp;stopEvent2) );
    checkCuda( cudaEventRecord(startEvent2, 0) );
    standard1D&lt;&lt;&lt;(int)ceil((float)size/threadSize),threadSize&gt;&gt;&gt;(diarray,doarray,size);
    checkCuda( cudaEventRecord(stopEvent2, 0) );
    checkCuda( cudaEventSynchronize(stopEvent2) );
  
    //copy back to CPU
    cudaMemcpy(horray,doarray,sizeof(float)*size,cudaMemcpyDeviceToHost);
  </code></pre>
  
  <p>and print results:</p>
  
  <pre><code>  float time,time2;
    checkCuda( cudaEventElapsedTime(&amp;time, startEvent, stopEvent) );
    checkCuda( cudaEventElapsedTime(&amp;time2, startEvent2, stopEvent2) );
    printf("Texture  bandwidth (GB/s): %f\n",bytes * 1e-6 / time);
    printf("Standard bandwidth (GB/s): %f\n",bytes * 1e-6 / time2);
  </code></pre>
  
  <p>It turns out that, no matters the size of the array I am allocating (<code>size</code>), the standard bandwidth is always much higher. 
  Is that how it suppose to be or am I screwing it up at some point? 
  My understanding of Texture memory access was that it can speed up global memory access.</p>