title: Cuda unified memory between gpu and host
body: <p>I'm writing a cuda-based program that needs to periodically transfer a set of items from the GPU to the Host memory.  In order to keep the process asynchronous, I was hoping to use cuda's UMA to have a memory buffer and flag in the host memory (so both the GPU and the CPU can access it).  The GPU would make sure the flag is clear, add its items to the buffer, and set the flag.  The CPU waits for the flag to be set, copies things out of the buffer, and clears the flag.  As far as I can see, this doesn't produce any race condition because it forces the GPU and CPU to take turns, always reading and writing to the flag opposite each other.</p>
  
  <p>So far I haven't been able to get this to work because there does seem to be some sort of race condition.  I came up with a simpler example that has a similar issue:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  
  __global__
  void uva_counting_test(int n, int *h_i);
  
  int main() {
      int *h_i;
      int n;
  
      cudaMallocHost(&amp;h_i, sizeof(int));
  
      *h_i = 0;
      n = 2;
  
      uva_counting_test&lt;&lt;&lt;1, 1&gt;&gt;&gt;(n, h_i);
  
      //even numbers
      for(int i = 1; i &lt;= n; ++i) {
          //wait for a change to odd from gpu
          while(*h_i == (2*(i - 1)));
  
          printf("host h_i: %d\n", *h_i);
          *h_i = 2*i;
      }
  
      return 0;
  }
  
  __global__
  void uva_counting_test(int n, int *h_i) {
      //odd numbers
      for(int i = 0; i &lt; n; ++i) {
          //wait for a change to even from host
          while(*h_i == (2*(i - 1) + 1));
  
          *h_i = 2*i + 1;
      }
  }
  </code></pre>
  
  <p>For me, this case always hangs after the first print statement from the CPU (<code>host h_i: 1</code>).  The really unusual thing (which may be a clue) is that I can get it to work in cuda-gdb.  If I run it in cuda-gdb, it will hang as before.  If I press ctrl+C, it will bring me to the while() loop line in the kernel.  From there, surprisingly, I can tell it to continue and it will finish.  For n > 2, it will freeze on the while() loop in the kernel again after each kernel, but I can keep pushing it forward with ctrl+C and continue.</p>
  
  <p>If there's a better way to accomplish what I'm trying to do, that would also be helpful.</p>