<div dir="ltr"><div class="gmail_extra"><div class="gmail_quote">On Fri, Aug 28, 2015 at 5:15 PM, escha <span dir="ltr"><<a href="mailto:escha@apple.com" target="_blank">escha@apple.com</a>></span> wrote:<blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex"><span>[...] I'm still concerned about loop unrolling in a case such as this:</span> </blockquote><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left-width:1px;border-left-color:rgb(204,204,204);border-left-style:solid;padding-left:1ex"><span>><br>
>  for (int i = 0; i < *bound; ++i) {<br>
>    if (i == 0)<br>
>      __syncthreads();<br>
>  }<br>
><br>
> 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.<br>
<br>
</span>This would be a violation of convergent semantics, IMO. Anything that causes a convergent instruction that was uniform to become not-uniform is, I think, duplication or not.<br>
<span><font color="#888888"><br>
—escha</font></span></blockquote><div><br></div><div>I agree. What my new example shows is that loop unrolling should take the convergent attribute into account, but it doesn't currently do that AFAICT. Sinking::SinkInstruction is the only place that I've found where the convergent attributed is considered for IR right now.</div><div><br></div><div>Loop unswitching also needs to take the convergent attribute into account. Take this example:</div><div><br></div><div>  bool b = *ptr;</div><div>  for (int i = 0; i < 100; ++i) {</div><div>    if (b) {</div><div>      // do something</div><div>    }</div><div>    __syncthreads();</div><div>  }</div><div><br></div><div>Loop unswitching could decide to unswitch on b even if b is divergent, which would duplicate the `__syncthreads()`. That's normally a performance problem, but with convergent/__syncthreads() it's a correctness problem.</div><div><br></div><div>The langref doesn't give semantics for when instructions marked convergent can be duplicated, only for when they can be moved. Owen, do you have some semantics for this in mind?<br></div><div><br></div><div>Bjarke<br></div></div></div></div>