<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=""><br class=""><div><blockquote type="cite" class=""><div class="">On Aug 31, 2015, at 10:11 PM, Mehdi Amini <<a href="mailto:mehdi.amini@apple.com" class="">mehdi.amini@apple.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><meta http-equiv="Content-Type" content="text/html charset=utf-8" class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><br class=""><div class=""><blockquote type="cite" class=""><div class="">On Aug 31, 2015, at 10:01 PM, Owen Anderson <<a href="mailto:resistor@mac.com" class="">resistor@mac.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><meta http-equiv="Content-Type" content="text/html charset=utf-8" class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><br class=""><div class=""><blockquote type="cite" class=""><div class="">On Aug 31, 2015, at 9:55 PM, Mehdi Amini <<a href="mailto:mehdi.amini@apple.com" class="">mehdi.amini@apple.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div class="" style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;"><div class=""><blockquote type="cite" class="" style="font-family: Helvetica; font-size: 12px; font-style: normal; font-variant: normal; font-weight: normal; letter-spacing: normal; line-height: normal; orphans: auto; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: auto; word-spacing: 0px; -webkit-text-stroke-width: 0px;"><div class=""><div class="" style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;"><div class=""><div class="">Can you define what you mean by “predicates under which the original executed”?</div></div></div></div></blockquote><div style="font-family: Helvetica; font-size: 12px; font-style: normal; font-variant: normal; font-weight: normal; letter-spacing: normal; line-height: normal; orphans: auto; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: auto; word-spacing: 0px; -webkit-text-stroke-width: 0px;" class=""><br class=""></div></div>I think I see what you mean: the syncthreads() call is made control dependent on b while it was not before. Your sentence made me think about something quite different.<div class=""></div></div></div></blockquote></div><br class=""><div class="">There are multiple ways to view this problem, none of which are particularly native to LLVM. Control dependence is one way. Execution predicates (as seen in IRs like Sea of Nodes) is another, equivalent, way.</div></div></div></blockquote></div><br class=""><div class="">Sure, I was not trying to correct anything you wrote, just expressing that I finally understood what you meant :)</div><div class="">I don’t know about Sea of Nodes and I associated the formulation you used to a totally different compiler concept than “execution predicates”. </div><div class=""><br class=""></div><div class="">What about full unrolling?</div><div class=""><br class=""></div><div class="">for(int i = 0; i < 2; i++) {</div><div class=""> __syncthreads();</div><div class="">}</div><div class=""><br class=""></div><div class="">to</div><div class=""><br class=""></div><div class=""><div class="">__syncthreads();</div></div><div class=""><div class="">__syncthreads();</div></div><div class=""><br class=""></div><div class="">This does not seem to preserve the "set of execution predicates”.</div><div class="">(I’ll read about the "Sea of Nodes" model)</div></div></div></blockquote><br class=""></div><div>Yeah, this is where it gets murkier, and possibly where some variation of control dependence becomes more informative. In some sense the dependence on i (or in the predicates world, the gating on the values of i) is trivial in that there’s no conditionality involved, so unrolling the loop does not introduce or remove conditionality.</div><div><br class=""></div><div>—Owen</div><br class=""></body></html>