title: OpenCL - Vectorization vs In-thread for loop
body: <p>I have a problem where I need to process a known number of threads in parallel (great), but for which each thread may have a vastly different number of internal iterations (not great). In my mind, this makes it better to do a kernel scheme like this:</p>
  
  <pre><code>__kernel something(whatever)
  {
     unsigned int glIDx = get_global_id(0);
  
     for(condition_from_whatever)
     {
  
     }//alternatively, do while
  
  }
  </code></pre>
  
  <p>where id(0) is known beforehand, rather than:</p>
  
  <pre><code>__kernel something(whatever)
  {
     unsigned int glIDx = get_global_id(0);
     unsigned int glIDy = get_global_id(1); // max "unroll dimension"
  
     if( glIDy_meets_condition)
        do_something();
     else
        dont_do_anything();
  
  }
  </code></pre>
  
  <p>which would necessarily execute for the FULL POSSIBLE RANGE of glIDy, with no way to terminate beforehand, as per this discussion:</p>
  
  <p><a href="https://stackoverflow.com/questions/16308667/killing-opencl-kernels">Killing OpenCL Kernels</a></p>
  
  <p>I can't seem to find any specific information about costs of dynamic-sized forloops / do-while statements within kernels, though I do see them everywhere in kernels in Nvidia's and AMD's SDK. I remember reading something about how the more aperiodic an intra-kernel condition branch is, the worse the performance. </p>
  
  <p>ACTUAL QUESTION: </p>
  
  <p>Is there a more efficient way to deal with this on a GPU architecture than the first scheme I proposed?</p>
  
  <p>I'm also open to general information about this topic.</p>
  
  <p>Thanks.</p>