<html><head><meta http-equiv="Content-Type" content="text/html charset=utf-8"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class="">I don’t think the example code here is legal under any SPMD models I am aware of. It’s generally not legal to have barrier operations under divergent control flow, such as divergent trip-count loops.<div class=""><br class=""></div><div class="">From the CUDA docs:</div><div class=""><br class=""></div><div class=""><blockquote type="cite" class=""><samp class="ph codeph" style="font-family: Consolas, Courier, 'Courier New', monospace; color: rgb(34, 68, 0); background-color: rgb(244, 247, 240); padding: 0px 0.2em; margin-bottom: 1em; font-size: 14px; border: 0px !important;">__syncthreads()</samp><span style="font-family: 'Trebuchet MS', 'DIN Pro', sans-serif; font-size: 14px; background-color: rgb(255, 255, 255);" class=""> is allowed in conditional code but only if the conditional evaluates identically across the entire thread block, otherwise the code execution is likely to hang or produce unintended side effects.</span></blockquote></div><div class=""><br class=""></div><div class="">—Owen</div><div class=""><br class=""><div><blockquote type="cite" class=""><div class="">On Aug 21, 2015, at 12:29 PM, Bjarke Hammersholt Roune <<a href="mailto:broune@google.com" class="">broune@google.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class="">broune added a subscriber: broune.<br class="">broune added a comment.<br class=""><br class="">I think we'd need a change in loop unrolling for this. Here's an example, where the trip count is divergent:<br class=""><br class=""> for (int j = 0; j <= 31 - threadIdx.x; ++j) {<br class=""> for (int i = 0; i <= threadIdx.x; ++i) {<br class=""> // do something<br class=""> __syncthreads();<br class=""> }<br class=""> }<br class=""><br class="">We can't allow unrolling of the inner loop here, since then threads that were previously able to meet up at the single syncthreads will instead be distributed among the unrolled syncthreads copies.<br class=""><br class="">I think that loop unrolling will be OK if we change it so that it only unrolls loops that contain syncthreads if the trip count is known to be not divergent (i.e. convergent, but in the CUDA sense, not in the LLVM sense). Jingyue's divergence analysis pass can prove non-divergence.<br class=""><br class=""><br class=""><a href="http://reviews.llvm.org/D12246" class="">http://reviews.llvm.org/D12246</a><br class=""><br class=""><br class=""><br class=""></div></blockquote></div><br class=""></div></body></html>