title: Negative array indexing in shared memory based 1d stencil CUDA implementation
body: <p>I'm currently working with CUDA programming and I'm trying to learn off of slides from a workshop I found online, which can be found <a href="http://on-demand.gputechconf.com/gtc-express/2011/presentations/GTC_Express_Sarah_Tariq_June2011.pdf" rel="noreferrer">here</a>. The problem I am having is on slide <code>48</code>. The following code can be found there:</p>
  
  <pre><code>__global__ void stencil_1d(int *in, int *out) {
  
      __shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
  
      int gindex = threadIdx.x + blockIdx.x * blockDim.x;
      int lindex = threadIdx.x + RADIUS;
  
      // Read input elements into shared memory
      temp[lindex] = in[gindex];
      if (threadIdx.x &lt; RADIUS) {
          temp[lindex - RADIUS] = in[gindex - RADIUS];
          temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
      }
  
  ....
  </code></pre>
  
  <p>To add a bit of context. We have an array called <code>in</code> which as length say <code>N</code>. We then have another array <code>out</code> which has length <code>N+(2*RADIUS)</code>, where <code>RADIUS</code> has a value of <code>3</code> for this particular example. The idea is to copy array <code>in</code>, into array <code>out</code> but to place the array <code>in</code> in position <code>3</code> from the beginning of array <code>out</code> i.e <code>out = [RADIUS][in][RADIUS]</code>, see slide for graphical representation.</p>
  
  <p>The confusion comes in on the following line:</p>
  
  <pre><code> temp[lindex - RADIUS] = in[gindex - RADIUS];
  </code></pre>
  
  <p>If gindex is <code>0</code> then we have <code>in[-3]</code>. How can we read from a negative index in an array? Any help would really be appreciated.</p>