title: Reduction in CUDA
body: <p>I'm just starting to learn CUDA programming, and I have some confusion about reduction. </p>
  
  <p>I know that global memory has much visiting delay as compared to shared memory, but can I use global memory to (at least) simulate a behavior similar to shared memory?</p>
  
  <p>For example, I want to sum the elements of a large array whose length is exactly <code>BLOCK_SIZE * THREAD_SIZE</code> (both the dimensions of grid and block are power of <code>2</code>), and I've tried to use the code below:</p>
  
  <pre><code>    __global__ void parallelSum(unsigned int* array) {
  
      unsigned int totalThreadsNum = gridDim.x * blockDim.x;
      unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x;
  
      int i = totalThreadsNum / 2;
      while (i != 0) {
              if (idx &lt; i) {
                  array[idx] += array[idx + i];
          }
          __syncthreads();
          i /= 2;
      }
  }
  </code></pre>
  
  <p>I compared the result of this code and the result generated serially on the host, and the weird thing is: sometimes the results are the same, but sometimes they are apparently different. Is there any reason related to using global memory here?</p>