title: Implementing a critical section in CUDA
body: <p>I'm trying to implement a <em>critical section</em> in CUDA using atomic instructions, but I ran into some trouble. I have created the test program to show the problem:</p>
  
  <pre><code>#include &lt;cuda_runtime.h&gt;
  #include &lt;cutil_inline.h&gt;
  #include &lt;stdio.h&gt;
  
  __global__ void k_testLocking(unsigned int* locks, int n) {
      int id = threadIdx.x % n;
      while (atomicExch(&amp;(locks[id]), 1u) != 0u) {} //lock
      //critical section would go here
      atomicExch(&amp;(locks[id]),0u); //unlock
  }
  
  int main(int argc, char** argv) {
      //initialize the locks array on the GPU to (0...0)
      unsigned int* locks;
      unsigned int zeros[10]; for (int i = 0; i &lt; 10; i++) {zeros[i] = 0u;}
      cutilSafeCall(cudaMalloc((void**)&amp;locks, sizeof(unsigned int)*10));
      cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice));
  
      //Run the kernel:
      k_testLocking&lt;&lt;&lt;dim3(1), dim3(256)&gt;&gt;&gt;(locks, 10);
  
      //Check the error messages:
      cudaError_t error = cudaGetLastError();
      cutilSafeCall(cudaFree(locks));
      if (cudaSuccess != error) {
          printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
          exit(-1);
      }
      return 0;
  }
  </code></pre>
  
  <p>This code, unfortunately, hard freezes my machine for several seconds and finally exits, printing out the message:</p>
  
  <pre><code>fcudaSafeCall() Runtime API error in file &lt;XXX.cu&gt;, line XXX : the launch timed out and was terminated.
  </code></pre>
  
  <p>which means that one of those while loops is not returning, but it seems like this should work. </p>
  
  <p>As a reminder <code>atomicExch(unsigned int* address, unsigned int val)</code> atomically sets the value of the memory location stored in address to <code>val</code> and returns the <code>old</code> value. So the idea behind my locking mechanism is that it is initially <code>0u</code>, so one thread should get past the <code>while</code> loop and all other threads should wait on the <code>while</code> loop since they will read <code>locks[id]</code> as <code>1u</code>. Then when the thread is done with the <em>critical section</em>, it resets the lock to <code>0u</code> so another thread can enter.</p>
  
  <p>What am I missing?</p>
  
  <p>By the way, I am compiling with:</p>
  
  <pre><code>nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu
  </code></pre>