title: CUDA: Shared memory allocation with overlapping borders
body: <p>Is there an easy way (google hasn't delivered...) to allocate per-block shared memory regions from a single input array such that there can be an overlap?</p>
  
  <p>The simple example is string searching; Saw I want to dice up the input text, have each thread in each block search for a pattern starting from text[thread_id], but want the data assigned to each block to overlap by the pattern length so matching cases that fall across the border are still found.</p>
  
  <p>I.e the total memory size allocated to shared memory on each block is </p>
  
  <pre><code>(blocksize+patternlength)*sizeof(char)
  </code></pre>
  
  <p>I'm probably missing something simple and am currently diving through the CUDA guide, but would appreciate some guidance.</p>
  
  <p>UPDATE: I suspect some people have misunderstood my question (or I mis-explained it).</p>
  
  <p>Say I have a dataset QWERTYUIOP, and i want to search for a 3 character match, and i dice up the dataset (arbitrarily) into 4's for each thread block; QWER  TYUI  OPxx</p>
  
  <p>This is simple enough to accomplish but the algorithm fails if the 3 character match is actually looking for IOP.</p>
  
  <p>In this case, what I want is for each block to have in shared memory:</p>
  
  <p>QWERTY TYUIOP OPxxxx</p>
  
  <p>ie each block gets assigned the blocksize+patternlength-1 characters so no memory border issues occur. </p>
  
  <p>Hope that explains things better.</p>
  
  <p>Since @jmilloy is being persistent... :P </p>
  
  <pre><code>//VERSION 1: Simple
  __global__ void gpuSearchSimple(char *T, int lenT, char *P, int lenP, int *pFound)
  {
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
      int fMatch = 1;
      for (int i=0; i &lt; lenP; i++)
      {
        if (T[startIndex+i] != P[i]) fMatch = 0;
      }
      if (fMatch) atomicMin(pFound, startIndex);
  }
  //VERSION 2: Texture
  __global__ void gpuSearchTexture(int lenT, int lenP, int *pFound)
  {
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
      int fMatch = 1;
      for (int i=0; i &lt; lenP; i++)
      {
        if (tex1Dfetch(texT,startIndex+i) != tex1Dfetch(texP,i)) fMatch = 0;
      }
      if (fMatch) atomicMin(pFound, startIndex);
  }
  //Version 3: Shared
  __global__ void gpuSearchTexSha(int lenT, int lenP, int *pFound)
  {
    extern __shared__ char shaP[];
    for (int i=0;threadIdx.x+i&lt;lenP; i+=blockDim.x){
      shaP[threadIdx.x+i]= tex1Dfetch(texP,threadIdx.x+i);
    }
    __syncthreads();
  
    //At this point shaP is populated with the pattern
    int startIndex = blockDim.x*blockIdx.x + threadIdx.x;
      // only continue if an earlier instance hasn't already been found
      int fMatch = 1;
      for (int i=0; i &lt; lenP; i++)
      {
        if (tex1Dfetch(texT,startIndex+i) != shaP[i]) fMatch = 0;
      }
      if (fMatch) atomicMin(pFound, startIndex);
  }
  </code></pre>
  
  <p>What I would like to have done is to put the text into shared memory chunks, as described in the rest of the question, instead of keeping the text in texture memory for the later versions.</p>