title: Understanding Dynamic Parallelism in CUDA
body: <p>Example of dynamic parallelism:</p>
  
  <pre><code>__global__ void nestedHelloWorld(int const iSize,int iDepth) {
      int tid = threadIdx.x;
      printf("Recursion=%d: Hello World from thread %d" "block %d\n",iDepth,tid,blockIdx.x);
      // condition to stop recursive execution
      if (iSize == 1) return;
      // reduce block size to half
      int nthreads = iSize&gt;&gt;1;
      // thread 0 launches child grid recursively
      if(tid == 0 &amp;&amp; nthreads &gt; 0) {
          nestedHelloWorld&lt;&lt;&lt;1, nthreads&gt;&gt;&gt;(nthreads,++iDepth);
          printf("-------&gt; nested execution depth: %d\n",iDepth);
      }
  }
  </code></pre>
  
  <p>Prints with one block, with two blocks the entire parent grid has finished:</p>
  
  <pre><code>./nestedHelloWorld Execution Configuration: grid 1 block 8
  Recursion=0: Hello World from thread 0 block 0
  Recursion=0: Hello World from thread 1 block 0
  Recursion=0: Hello World from thread 2 block 0
  Recursion=0: Hello World from thread 3 block 0
  Recursion=0: Hello World from thread 4 block 0
  Recursion=0: Hello World from thread 5 block 0
  Recursion=0: Hello World from thread 6 block 0
  Recursion=0: Hello World from thread 7 block 0
  -------&gt; nested execution depth: 1
  Recursion=1: Hello World from thread 0 block 0
  Recursion=1: Hello World from thread 1 block 0
  Recursion=1: Hello World from thread 2 block 0
  Recursion=1: Hello World from thread 3 block 0
  -------&gt; nested execution depth: 2
  Recursion=2: Hello World from thread 0 block 0
  Recursion=2: Hello World from thread 1 block 0
  -------&gt; nested execution depth: 3
  Recursion=3: Hello World from thread 0 block 0
  </code></pre>
  
  <p>Say I launch a child grid from one thread in a block at threadIdx.x==0. Can I assume that all other threads in the parent grid have finished executing up to the point I launched the child grid as well? </p>
  
  <p>If so, how does this work? All I'm reading is that a parent grid is not technically finished before a child grid. Nothing about guarantees of other parent threads that have not launched children.</p>