title: Race Condition in CUDA programs
body: <p>I have two pieces of code. One written in C and the corresponding operation written in CUDA.
  Please help me understand how <code>__syncthreads()</code> works in context of the following programs. As per my understanding, <code>__syncthreads()</code> ensures synchronization of threads limited to one block.</p>
  
  <p>C program :</p>
  
  <pre><code>{
      for(i=1;i&lt;10000;i++)
      {
          t=a[i]+b[i];
          a[i-1]=t;
      }
  }
  </code></pre>
  
  <p>`</p>
  
  <p>The equivalent CUDA program :
  `    </p>
  
  <pre><code>__global__ void kernel0(int *b, int *a, int *t, int N)
  {
      int b0=blockIdx.x;
      int t0=threadIdx.x;
      int tid=b0*blockDim.x+t0;
      int private_t;
      if(tid&lt;10000)
      {
          private_t=a[tid]+b[tid];
          if(tid&gt;1)
              a[tid-1]=private_t;
          __syncthreads();
          if(tid==9999)
          *t=private_t;
      }
  }
  </code></pre>
  
  <p>Kernel Dimensions:</p>
  
  <pre><code>dim3 k0_dimBlock(32);
  dim3 k0_dimGrid(313);
  kernel0 &lt;&lt;&lt;k0_dimGrid, k0_dimBlock&gt;&gt;&gt;
  </code></pre>
  
  <p>The surprising fact is output from C and CUDA program are identical. Given the nature of problem, which has dependency of a[] onto itself, a[i] is loaded by thrad-ID i and written to a[i-1] by the same thread. Now the same happens for thread-ID i-1. Had the problem size been lesser than 32, the output is obvious. But for a problem of size 10000 with 313 blocks and blocks, how does the dependency gets respected ?</p>