title: Benefit of splitting a big CUDA kernel and using dynamic parallelism
body: <p>I have a big kernel in which an initial state is evolved using different techniques. That is, I have a loop in the kernel, in this loop a certain predicate is evaluated on the current state and on the result of this predicate, a certain action is taken.</p>
  
  <p>The kernel needs a bit of temporary data and shared memory, but since it is big it uses 63 registers and the occupancy is very very low.</p>
  
  <p>I would like to split the kernel in many little kernels, but every block is totally independent from the others and I (think I) can't use a single thread on the host code to launch multiple small kernels.</p>
  
  <p>I am not sure if streams are adequate for this kind of work, I never used them, but since I have the option to use the dynamic parallelism, I would like if that is a good option to implement this kind of job.
  Is it fast to launch a kernel from a kernel?
  Do I need to copy data in global memory to make them available to a sub-kernel?</p>
  
  <p>If I split my big kernel in many little ones, and leave the first kernel with a main loop which calls the required kernel when necessary (which allows me to move temporary variables in every sub-kernel), will help me increase the occupancy?</p>
  
  <p>I know it is a bit generic question, but I do not know this technology and I would like if it fits my case or if streams are better.</p>
  
  <p>EDIT:
  To provide some other details, you can imagine my kernel to have this kind of structure:</p>
  
  <pre><code>__global__ void kernel(int *sampleData, int *initialData) {
      __shared__ int systemState[N];
      __shared__ int someTemp[N * 3];
      __shared__ int time;
      int tid = ...;
      systemState[tid] = initialData[tid];
  
      while (time &lt; TIME_END) {
          bool c = calc_something(systemState);
          if (c)
              break;
          someTemp[tid] = do_something(systemState);
          c = do_check(someTemp);
          if (__syncthreads_or(c))
              break;
          sample(sampleData, systemState);
          if (__syncthreads_and(...)) {
              do_something(systemState);
              sync();
              time += some_increment(systemState);
          }
          else {
              calcNewTemp(someTemp, systemState);
              sync();
              do_something_else(someTemp, systemState);
              time += some_other_increment(someTemp, systemState);
          }
      }
      do_some_stats();
  }
  </code></pre>
  
  <p>this is to show you that there is a main loop, that there are temporary data which are used somewhere and not in other points, that there are shared data, synchronization points, etc.</p>
  
  <p>Threads are used to compute vectorial data, while there is, ideally, one single loop in each block (well, of course it is not true, but logically it is)... One "big flow" for each block.</p>
  
  <p>Now, I am not sure about how to use streams in this case... Where is the "big loop"? On the host I guess... But how do I coordinate, from a single loop, all the blocks? This is what leaves me most dubious. May I use streams from different host threads (One thread per block)?</p>
  
  <p>I am less dubious about dynamic parallelism, because I could easily keep the big loop running, but I am not sure if I could have advantages here.</p>