title: CUDA: are access times for texture memory similar to coalesced global memory?
body: <p>My kernel threads access a linear character array in a coalesced fashion. If I map
  the array to texture I don't see any speedup. The running times are 
  almost the same. I'm working on a Tesla C2050 with compute capability 2.0 and read
  somewhere that global accesses are cached. Is that true? Perhaps that is why I
  am not seeing a difference in the running time. </p>
  
  <p>The array in the main program is </p>
  
  <pre><code>char *dev_database = NULL;
  cudaMalloc( (void**) &amp;dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );
  </code></pre>
  
  <p>and I bind it to texture <code>texture&lt;char&gt; texdatabase</code> with</p>
  
  <pre><code>cudaBindTexture(NULL, texdatabase, dev_database, JOBS * FRAGMENTSIZE * sizeof(char) );
  </code></pre>
  
  <p>Each thread then reads a character <code>ch = tex1Dfetch(texdatabase, p + id)</code> where id
  is <code>threadIdx.x + blockIdx.x * blockDim.x</code> and <code>p</code> is an offset.</p>
  
  <p>I'm binding only once and <code>dev_database</code> is a large array. Actually I found that
  if the size is too large the bind fails. Is there a limit on the size of the array
  to bind? Thanks very much.</p>