<div dir="ltr">In Bjarke's example, all threads call __syncthreads in and only in their first iteration (assuming *bounds > 0). </div><div class="gmail_extra"><br><div class="gmail_quote">On Sun, Aug 30, 2015 at 10:12 PM, Xuetian Weng <span dir="ltr"><<a href="mailto:wengxt@gmail.com" target="_blank">wengxt@gmail.com</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"><span class="">wengxt marked an inline comment as done.<br>
</span>wengxt added a comment.<br>
<span class=""><br>
In <a href="http://reviews.llvm.org/D12246#235675" rel="noreferrer" target="_blank">http://reviews.llvm.org/D12246#235675</a>, @broune wrote:<br>
<br>
>   for (int i = 0; i < *bound; ++i) {<br>
>     if (i == 0)<br>
>       __syncthreads();<br>
>   }<br>
><br>
><br>
</span><span class="">> This input program is valid as long as *bound > 0 has the same value across the block. Here loop-unrolling by a factor of 2 will separate off the first iteration of the loop into a duplicate body for the case where *bound is odd. I checked with an example loop that's similar but that doesn't use __syncthreads() and LLVM does do unrolling by a factor of 2 in this way. If whether *bound is odd is divergent, then only part of the warp would execute the __syncthreads() in the duplicate odd-case unrolled loop body. So I think that unrolling does have to be careful with divergent trip counts for loops that include __syncthreads() in cases such as this.<br>
<br>
<br>
</span>IMHO, for loop, it is not possible for two thread running on different iteration to sync together, that is actually divergent. In that sense, if two function call instruction converge in the original code, they will also converge after unrolling.<br>
<br>
Thus I don't think there is any problem in this example.<br>
<br>
<br>
<a href="http://reviews.llvm.org/D12246" rel="noreferrer" target="_blank">http://reviews.llvm.org/D12246</a><br>
<br>
<br>
<br>
</blockquote></div><br></div>