title: OpenCL barrier usage within conditional
body: <p>According to the <a href="https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/barrier.html" rel="nofollow noreferrer">specs</a>:</p>
  <blockquote>
  <p>If barrier is inside a conditional statement, then all work-items must enter the conditional if any work-item enters the conditional statement and executes the barrier.</p>
  <p>If barrier is inside a loop, all work-items must execute the barrier for each iteration of the loop before any are allowed to continue execution beyond the barrier.</p>
  </blockquote>
  <p>In my understanding this means, that in any Kernel:</p>
  <pre class="lang-c prettyprint-override"><code>if(0 == (get_local_id(0)%2)){
    //block a
    work_group_barrier(CLK_GLOBAL_MEM_FENCE);
    //block a part 2
  }else{
    //block b
    work_group_barrier(CLK_GLOBAL_MEM_FENCE);
    //block b part 2
  }
  //common operations
  </code></pre>
  <p>Should one worker reach <code>//block a</code>, every other worker needs to reach it too.</p>
  <p>By this logic it is not possible to correctly synchronize every odd local worker with every even one ( blocks a and b ) to be run at the same time.</p>
  <p>Is this understanding correct?</p>
  <p>What would be a good synchronisation strategy for a situation like this ?</p>
  <ul>
  <li>where every other worker would need to do a different logic, but by <code>//block a part 2</code> and <code>//block b part 2</code> the workers inside one working group be synced up.</li>
  <li>in the requested usecase there are more phases, than 2, and I'd like to keep every phase to be synchronised.</li>
  </ul>
  <p>Would a logic like this be an acceptable solution?</p>
  <pre class="lang-c prettyprint-override"><code>__local number_finished = 0;
  if(0 == (get_local_id(0)%2)){
    //block a
    atomic_add(&amp;number_finished, 1);
    while(number_finished &lt; get_local_size(0));
    //block a part 2
  }else{
    //block b
    atomic_add(&amp;number_finished, 1);
    while(number_finished &lt; get_local_size(0));
    //block b part 2
  }
  work_group_barrier(CLK_GLOBAL_MEM_FENCE);
  //common operations
  </code></pre>