title: CUDA shared memory atomics error
body: <p>I am using a Tesla C1060 with 1.3 compute capability and nvcc compiler driver 4.0. I am trying to do some computation local to thread block. Each thread block is provided with a shared array which is first initialized to zero values. For synchronizing concurrent updates (addition) to shared data by threads of the thread block, I use CUDA atomicAdd primitive. </p>
  
  <p>Once each thread block is ready with the results in its shared data array, each entry in shared data array is iteratively merged (using atomicAdd) to corresponding entries in global data array.</p>
  
  <p>Below is a code very similar to what I basically trying to do. </p>
  
  <pre><code>#define DATA_SZ 16
  typedef unsigned long long int ULLInt;
  
  __global__ void kernel( ULLInt* data, ULLInt ThreadCount )
  {
    ULLInt thid = threadIdx.x + blockIdx.x * blockDim.x;
    __shared__ ULLInt sharedData[DATA_SZ];
  
    // Initialize the shared data
    if( threadIdx.x == 0 )
    {
      for( int i = 0; i &lt; DATA_SZ; i++ ) { sharedData[i] = 0; }
    }
    __syncthreads();
  
    //..some code here
  
    if( thid &lt; ThreadCount )
    {
      //..some code here
  
      atomicAdd( &amp;sharedData[getIndex(thid), thid );
  
      //..some code here        
  
      for(..a loop...)
      { 
        //..some code here
  
        if(thid % 2 == 0)
        {           
          // getIndex() returns a value in [0, DATA_SZ )
          atomicAdd( &amp;sharedData[getIndex(thid)], thid * thid );
        }
      }
    }
    __syncthreads();
  
    if( threadIdx.x == 0 )
    {
      // ...
      for( int i = 0; i &lt; DATA_SZ; i++ ) { atomicAdd( &amp;Data[i], sharedData[i] ); }
      //...
    }
  }
  </code></pre>
  
  <p>If I compile with -arch=sm_20 I dont get any errors. However when I compile the kernel using -arch=sm_13 option I get the following errors:</p>
  
  <pre><code>ptxas /tmp/tmpxft_00004dcf_00000000-2_mycode.ptx, line error   : Global state space expected for instruction 'atom'
  ptxas /tmp/tmpxft_00004dcf_00000000-2_mycode.ptx, line error   : Global state space expected for instruction 'atom'
  ptxas fatal   : Ptx assembly aborted due to errors
  </code></pre>
  
  <p>If I comment the following two lines I dont get any errors with -arch=sm_13:</p>
  
  <pre><code>atomicAdd( &amp;sharedData[getIndex(thid), thid );
  atomicAdd( &amp;sharedData[getIndex(thid)], thid * thid );
  </code></pre>
  
  <p>Can  someone suggest what I might be doing wrong ? </p>