<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 see thanks.<div class="">For the use case Owen is talking about I don’t think it would be allowed, because (if I understood correctly the use case) if the original was:</div><div class=""><br class=""></div><div class="">for (int i = 0; i < 100; ++i) {</div><div class=""> k0 = 0;</div><div class=""> if (b)</div><div class=""> k0 = 15;</div><div class=""> gradient_operation(k0);</div><div class="">}</div><div class=""><br class=""></div><div class="">transforming it to</div><div class=""><br class=""></div><div class="">if (b) {</div><div class=""> k0 = 15;</div><div class=""><div class=""> for (int i = 0; i < 100; ++i) {</div><div class=""> gradient_operation(k0);</div><div class=""> }</div><div class="">} else {</div></div><div class=""> k0 = 0;</div><div class=""><div class=""> for (int i = 0; i < 100; ++i) {</div><div class=""> gradient_operation(k0);</div><div class=""> }</div></div><div class="">}</div><div class=""><br class=""></div><div class="">In the first case when the gradient_operation is executed all the registers the operation depends on are set to a correct value, while in the second case in the “true” part of the if the values for the gradient operation wouldn’t be correctly set as they wouldn’t have the value for the “else” part set correctly.</div><div class=""><br class=""></div><div class="">Marcello</div><div class=""><br class=""></div><div class=""><div><blockquote type="cite" class=""><div class="">On 31 Aug 2015, at 19:51, 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 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="Apple-interchange-newline"><br class="">On Aug 31, 2015, at 7:29 PM, Marcello Maggioni <<a href="mailto:mmaggioni@apple.com" class="">mmaggioni@apple.com</a>> wrote:<br class=""><br class=""></div><blockquote type="cite" 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=""><div class=""><br class=""><div class=""><blockquote type="cite" class=""><div class="">On 31 Aug 2015, at 17:57, Mehdi Amini via llvm-commits <<a href="mailto:llvm-commits@lists.llvm.org" class="">llvm-commits@lists.llvm.org</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div 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;"><blockquote type="cite" class=""><div class=""><br class="Apple-interchange-newline">On Aug 31, 2015, at 4:55 PM, Owen Anderson via llvm-commits <<a href="mailto:llvm-commits@lists.llvm.org" class="">llvm-commits@lists.llvm.org</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div 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;"><blockquote type="cite" class=""><div class=""><br class="Apple-interchange-newline">On Aug 31, 2015, at 4:08 PM, Bjarke Roune <<a href="mailto:broune@google.com" class="">broune@google.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div dir="ltr" class=""><div class="gmail_extra"><br class=""><div class="gmail_quote">On Mon, Aug 31, 2015 at 10:42 AM, Owen Anderson<span class="Apple-converted-space"> </span><span dir="ltr" class=""><<a href="mailto:resistor@mac.com" target="_blank" class="">resistor@mac.com</a>></span><span class="Apple-converted-space"> </span>wrote:<br class=""><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;"><div class="" style="word-wrap: break-word;"><span class=""><br class=""><div class=""><blockquote type="cite" class=""><div class="">On Aug 31, 2015, at 10:01 AM, Bjarke Roune <<a href="mailto:broune@google.com" target="_blank" class="">broune@google.com</a>> wrote:</div><br class=""><div class=""><span class="" style="font-family: Helvetica; font-size: 12px; font-style: normal; font-variant: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; float: none; display: inline !important;">I'm not sure if you're saying that there should be no limitations as the use cases you had mind for convergent are always OK with such duplication, or that we should be figuring this stuff out on a case-by-case basis, or something else?</span></div></blockquote></div><br class=""></span><div class="">I designed the semantics of convergent to meet the needs of texture and gradient operations, without a lot of consideration for barriers. IMO, It would be a nice end result if it turns out be just what barriers need as well, but I don’t want to accidentally hose the use cases I intended it for in pursuit of making it work for barriers.</div><div class=""><br class=""></div><div class="">Everything you’ve identified re: full vs partial unrolling, unstitching, etc. seems fine to me so far.</div><span class="HOEnZb"><font color="#888888" class=""><div class=""><br class=""></div></font></span></div></blockquote></div>So texture and gradient operations also requires these restrictions for unswitching and partial unrolling, or is it just that imposing such restrictions even if unnecessary is fine?</div><div class="gmail_extra"><br class=""></div><div class="gmail_extra">Would you say that we need the langref to say something about when convergent operations can be duplicated? I think it's hard to reason about whether unrolling, loop unswitching and other passes are doing the right thing without that.</div></div></div></blockquote><br class=""></div><div 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;">I really want to avoid talking about duplication as part of the semantics of convergent, because I fundamentally believe the focus on duplication is a red herring. It’s not the duplication that’s the problem, it’s the introduction of conditionals.</div><div 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;"><br class=""></div><div 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;">Within a SPMD model, the constraint that textures and gradients, and to some extent barriers, must preserve is that the set of executing threads must be the same at the location where they end up as the place where the user wrote them. </div></div></blockquote><div class=""><br class=""></div><div class="">What is the last “them” of your sentence referring to? I read it as referring to “the set of executing threads”, but then I can’t make sense of the sentence.</div><div class=""><br class=""></div><div class="">The definition in the Lang.Ref. (i.e. " A dominates B and B post-dominates A”) makes sense for moving such an instruction, but I’m having hard time to apply it for the duplication.</div><div class=""><br class=""></div><blockquote type="cite" class=""><div class=""><div 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;">To phrase that without talking about SPMD models or threads, it suffices to say that they must be executed in under the same set of execution predicates as where the user wrote them, aka in a control-equivalent location.</div></div></blockquote><div class=""><br class=""></div><div class="">I’m not sure if you mean to allow or disallow the loop unswitching transformation mentioned by Bjarke before, i.e.:</div><div class=""><br class=""></div><div class="">if (b) {<br class=""> for (int i = 0; i < 100; ++i) {<br class=""> // do something<br class=""> __syncthreads();<br class=""> }<br class="">} else {<br class=""> for (int i = 0; i < 100; ++i) {<br class=""> __syncthreads();<br class=""> }<br class="">}<br class=""><br class=""></div><div class="">I think it would match your constraint, i.e. "they are be executed in under the same set of execution predicates as where the user wrote them”. I think that because you exclude the SPMD/threads from your definition you lose the fact that b can be divergent in the thread block. </div><div class=""><br class=""></div></div></div></blockquote><div class=""><br class=""></div>What do you mean by “B can be divergent in the thread block” exactly?<br class=""></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 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="">Maybe not the right word, apologize.</div><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="">I mean b can be a loop invariant but not the same for all threads in a block. Non-uniform condition is maybe a more correct term?</div><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 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 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="">-- </div><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="">Mehdi </div><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><br 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=""><blockquote type="cite" 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=""><div class=""><div class=""><blockquote type="cite" class=""><div class=""><div 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="">Looks like defining what we need for the “convergent” attribute to work without introducing threads is uneasy.</div><div class=""><br class=""></div><div class=""><div class="">— </div><div class="">Mehdi</div><div class=""><br class=""></div></div></div><span 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; float: none; display: inline !important;">_______________________________________________</span><br 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;"><span 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; float: none; display: inline !important;">llvm-commits mailing list</span><br 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;"><a href="mailto:llvm-commits@lists.llvm.org" 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;">llvm-commits@lists.llvm.org</a><br 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;"><a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits" 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;">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits</a></div></blockquote></div></div></blockquote></div></blockquote></div><br class=""></div></body></html>