title: Why is texture memory version of below program slower than global memory  version
body: <p>i am confused why my texture version is slower than my global memory version because the texture version should exploit spatial locality. I am trying to compute the dot product in the below case. Thus, if one thread accesses index i, its neighbour should access i+1. Thus, we see spatial locality.</p>
  
  <p>Below is the texture memory version:</p>
  
  <pre><code>#include&lt;cuda_runtime.h&gt;
  #include&lt;cuda.h&gt;
  #include&lt;stdio.h&gt;
  #include&lt;stdlib.h&gt;
  #define intMin(a,b) ((a&lt;b)?a:b)
  //Threads per block
  #define TPB 128
  //blocks per grid
  #define BPG intMin(128, ((n+TPB-1)/TPB))
  
  texture&lt;float&gt; arr1;
  texture&lt;float&gt; arr2;
  
  
  const int n = 4;
  
  __global__ void addVal( float *c){
      int tid = blockIdx.x * blockDim.x + threadIdx.x;
  
      //Using shared memory to temporary store results
      __shared__ float cache[TPB];
      float temp = 0;
      while(tid &lt; n){
          temp += tex1Dfetch(arr1,tid) * tex1Dfetch(arr2,tid);
          tid += gridDim.x * blockDim.x;
  
  
      }
      cache[threadIdx.x] = temp;
      __syncthreads();
      int i = blockDim.x/2;
      while( i !=0){
          if(threadIdx.x &lt; i){
              cache[threadIdx.x] = cache[threadIdx.x] +cache[threadIdx.x + i] ;
  
          }
      __syncthreads();
      i = i/2;
  
      }
      if(threadIdx.x == 1){
          c[blockIdx.x ] = cache[0];
      }
  
  
  
  }
  
  int main(){
  
  float a[n] , b[n] , c[BPG];
  float *deva, *devb, *devc;
  int i;
  //Filling with random values to test
  for( i =0; i&lt; n; i++){
      a[i] = i;
      b[i] = i*2;
  }
  printf("Not using constant memory\n");
  cudaMalloc((void**)&amp;deva, n * sizeof(float));
  cudaMalloc((void**)&amp;devb, n * sizeof(float));
  cudaMalloc((void**)&amp;devc, BPG * sizeof(float));
  
  
  cudaMemcpy(deva, a, n *sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(devb, b, n *sizeof(float), cudaMemcpyHostToDevice);
  cudaBindTexture(NULL,arr1, deva,sizeof(float) * n); // note: deva shd be in gpu
  cudaBindTexture(NULL,arr2, devb,sizeof(float) * n); // note: deva shd be in gpu
  cudaEvent_t start, stop;
  cudaEventCreate(&amp;start);
  cudaEventCreate(&amp;stop);
  cudaEventRecord(start, 0);
  
  //Call function to do dot product
  addVal&lt;&lt;&lt;BPG, TPB&gt;&gt;&gt;(devc);
  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  float time;
  cudaEventElapsedTime(&amp;time,start, stop);
  printf("The elapsed time is: %f\n", time);
  
  
  //copy result back
  cudaMemcpy(c, devc, BPG * sizeof(float), cudaMemcpyDeviceToHost);
  float sum =0 ;
  for ( i = 0 ; i&lt; BPG; i++){
      sum+=c[i];
  
  }
  //display answer
  printf("%f\n",sum);
  cudaUnbindTexture(arr1);
  cudaUnbindTexture(arr2);
  cudaFree(devc);
  
  getchar();
  
  return 0;
  }
  </code></pre>
  
  <p>Global Memory version:</p>
  
  <pre><code>#include&lt;cuda_runtime.h&gt;
  #include&lt;cuda.h&gt;
  #include&lt;stdio.h&gt;
  #include&lt;stdlib.h&gt;
  #define intMin(a,b) ((a&lt;b)?a:b)
  //Threads per block
  #define TPB 128
  //blocks per grid
  #define BPG intMin(128, ((n+TPB-1)/TPB))
  
  const int n = 4;
  
  __global__ void addVal(float *a, float *b, float *c){
      int tid = blockIdx.x * blockDim.x + threadIdx.x;
  
      //Using shared memory to temporary store results
      __shared__ float cache[TPB];
      float temp = 0;
      while(tid &lt; n){
          temp += a[tid] * b[tid];
          tid += gridDim.x * blockDim.x;
  
  
      }
      cache[threadIdx.x] = temp;
      __syncthreads();
      int i = blockDim.x/2;
      while( i !=0){
          if(threadIdx.x &lt; i){
              cache[threadIdx.x] = cache[threadIdx.x] +cache[threadIdx.x + i] ;
  
          }
      __syncthreads();
      i = i/2;
  
      }
      if(threadIdx.x == 1){
          c[blockIdx.x ] = cache[0];
      }
  
  
  
  }
  
  int main(){
  
  float a[n] , b[n] , c[BPG];
  float *deva, *devb, *devc;
  int i;
  //Filling with random values to test
  for( i =0; i&lt; n; i++){
      a[i] = i;
      b[i] = i*2;
  }
  printf("Not using constant memory\n");
  cudaMalloc((void**)&amp;deva, n * sizeof(float));
  cudaMalloc((void**)&amp;devb, n * sizeof(float));
  cudaMalloc((void**)&amp;devc, BPG * sizeof(float));
  cudaMemcpy(deva, a, n *sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(devb, b, n *sizeof(float), cudaMemcpyHostToDevice);
  
  cudaEvent_t start, stop;
  cudaEventCreate(&amp;start);
  cudaEventCreate(&amp;stop);
  cudaEventRecord(start, 0);
  
  //Call function to do dot product
  addVal&lt;&lt;&lt;BPG, TPB&gt;&gt;&gt;(deva, devb, devc);
  cudaEventRecord(stop, 0);
  cudaEventSynchronize(stop);
  float time;
  cudaEventElapsedTime(&amp;time,start, stop);
  printf("The elapsed time is: %f\n", time);
  
  
  //copy result back
  cudaMemcpy(c, devc, BPG * sizeof(float), cudaMemcpyDeviceToHost);
  float sum =0 ;
  for ( i = 0 ; i&lt; BPG; i++){
      sum+=c[i];
  
  }
  //display answer
  printf("%f\n",sum);
  
  
  getchar();
  
  return 0;
  }
  </code></pre>