title: CUDA: Shift arrays on shared memory
body: <p>I am trying to load a flattened 2D matrix into shared memory, shift the data along x, write back to global memory shifting also along y. The input data is therefore shifted along x and y. What I have:  </p>
  
  <pre><code>__global__ void test_shift(float *data_old, float *data_new)
  
  {
  
  uint glob_index = threadIdx.x + blockIdx.y*blockDim.x;
  
  __shared__ float VAR;
  __shared__ float VAR2[NUM_THREADS];
  
  // load from global to shared
  
  VAR = data_old[glob_index];
  
  // do some stuff on VAR 
  
  if (threadIdx.x &lt; NUM_THREADS - 1)
  {
  VAR2[threadIdx.x + 1] = VAR; // shift (+1) along x
  }
  
  __syncthreads();
  
  // write to global memory
  
  if (threadIdx.y &lt; ny - 1)
  {
  glob_index = threadIdx.x + (blockIdx.y + 1)*blockDim.x; // redefine glob_index to shift along y (+1)
  data_new[glob_index] = VAR2[threadIdx.x];
  }
  </code></pre>
  
  <p>The call to the kernel:</p>
  
  <pre><code>test_shift &lt;&lt;&lt; grid, block &gt;&gt;&gt; (data_old, data_new);
  </code></pre>
  
  <p>and grid and blocks (blockDim.x is equal to the matrix width, i.e. 64):</p>
  
  <pre><code>dim3 block(NUM_THREADS, 1);
  dim3 grid(1, ny); 
  </code></pre>
  
  <p>I am not able to achieve it. Could someone please point out what's wrong with this? Should I use a strided index or an offset?</p>