title: Cuda: Kernel launch queue
body: <p>I'm not finding much info on the mechanics of a kernel launch operation.  The <a href="http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_Toolkit_Reference_Manual.pdf" rel="nofollow">API</a> say to see the <a href="http://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Programming_Guide.pdf" rel="nofollow">CudaProgGuide</a>.  And I'm not finding much there either.<br>
  Being that kernel execution is asynch, and some machines support concurrent execution, I'm lead to believe there is a queue for the kernels.  </p>
  
  <pre><code>    Host code:      
      1. malloc(hostArry, ......);  
      2. cudaMalloc(deviceArry, .....);  
      3. cudaMemcpy(deviceArry, hostArry, ... hostToDevice);
      4. kernelA&lt;&lt;&lt;1,300&gt;&gt;&gt;(int, int);  
      5. kernelB&lt;&lt;&lt;10,2&gt;&gt;&gt;(float, int));  
      6. cudaMemcpy(hostArry, deviceArry, ... deviceToHost);  
      7. cudaFree(deviceArry);
  </code></pre>
  
  <p>Line 3 is synchronous.  Line 4 &amp; 5 are asynchronous, and the machine supports concurrent execution.  So at some point, both of these kernels are running on the GPU.  (There is the possibility that kernelB starts and finishes, before kernelA finishes.)  While this is happening, the host is executing line 6.  Line 6 is synchronous with respect to the copy operation, but there is nothing preventing it from executing before kernelA or kernelB has finished.</p>
  
  <p>1) Is there a kernel queue in the GPU?  (Does the GPU block/stall the host?)<br>
  2) How does the host know that the kernel has finished, and it is "safe" to Xfer the results from the device to the host?</p>