title: CUDA Atomic Operations thread or block level?
body: <p>When occurs an atomic operation serialized at block level? 
  If I have the following code : </p>
  
  <pre><code>__global__ void sum (int *input){
  
     if ( threadIdx.x == 0) 
       __shared__ int result = 0;
     __syncthreads(); 
     atomicAdd(result,input[threadIdx.x+blockDim.x*blockId.x]);
  }
  </code></pre>
  
  <p>do this serialization occurs at block level? 
  I don't understand in general what "block level serialization" means because from what I know operations are always performed by threads.</p>