<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 9:11 PM, Marcello Maggioni <<a href="mailto:mmaggioni@apple.com" class="">mmaggioni@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="">Ah, ok, you were referring to the fact that if “we can prove a condition uniform” actually the “control-equivalent locations” category broadens to more basic-blocks than you can actually infer from the shape of the CFG.</div></div></blockquote><div><br class=""></div><div>Not exactly, I was referring to the fact that it is not clear that the unswitching is disallowed by the current definition of the convergent attribute, since the duplication is not addressed. I am talking about the actual definition in the language reference, not what you have in your mind about what the attribute *should do* or what it *is actually trying to solve*.</div><div> </div><br class=""><blockquote type="cite" class=""><div class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><div class="">That is certainly true and probably without at least mentioning the concept of “threads” there is no way of define it in a clear way.</div><div class="">Exposing the SPMD though looks like an implementation detail.</div><div class="">Probably mentioning that it should be control equivalent WTR to the threads that would originally have executed the unmodified statement would be enough? Or do you think that it would still be overly restrictive?</div></div></div></blockquote><div><br class=""></div><div>You first have to define the programming model if you want to define it this way, and Owen was trying to formalize its definition without involving the GPU programming model, and the way he did is elegant but does not seems complete yet.</div><div><br class=""></div><div>— </div><div>Mehdi</div><div> </div><br class=""><blockquote type="cite" class=""><div class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><div class=""><br class=""></div><div class="">Marcello<br class=""><div class=""><blockquote type="cite" class=""><div class="">On 31 Aug 2015, at 20:55, 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 8:14 PM, Marcello Maggioni <<a href="mailto:mmaggioni@apple.com" class="">mmaggioni@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="">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></div></blockquote><div class=""><br class=""></div><div class="">I know quite well the problem ;) from the hardware and SPMD model point of view. This is why I said that it is invalid if the condition (b) is not uniform.</div><div class="">But the question here is can you specify the convergent attribute in an abstract manner at the IR level without introducing SPMD or threads at all.</div><div class=""><br class=""></div><div class="">The Language Reference only mentions *moving* a convergent instruction, maybe saying that in case of duplication the same constraint applies: i.e. "the duplicated basic block must be dominated by the original one, and post-dominate it”. </div><div class="">This would allow unrolling:</div><div class=""><br class=""></div><div class=""> for (int i = 0; i < 100; ++i) {</div><div class=""> // do stuff<br class=""> gradient_operation(k0);</div><div class=""> ++i;</div><div class=""> // do stuff<br class=""> gradient_operation(k0);<br class=""> }<br class=""><br class=""></div><div class="">Or more broadly it can probably extended to "the duplicated *CFG subset* is dominated by the original one and post-dominate it”, to allow:</div><div class=""><br class=""></div><div class=""><div class=""> for (int i = 0; i < 100; ++i) {</div><div class=""> // do stuff</div><div class=""> if (cond)<br class=""> gradient_operation(k0);</div><div class=""> else </div><div class=""> // other stuff</div><div class=""> ++i;</div><div class=""> // do stuff<br class=""><div class=""> if (cond)<br class=""> gradient_operation(k0);</div><div class=""> else </div><div class=""><div class=""> // other stuff</div></div> }<br class=""></div></div><div class=""><br class=""></div><div class=""><br class=""></div><div class="">This would still prevent the invalid unswitching. However this is overly conservative as we wouldn’t allow the loop unswitching even if the condition could be proven uniform.</div><div class="">This is where I don’t see how you can get away with introducing the SPMD/threading model.</div><div class=""><br class=""></div><div class=""><br class=""></div><div class="">— </div><div class="">Mehdi</div><div class=""><br class=""></div><div class=""><br class=""></div><div class=""><br class=""></div><div class=""><br class=""></div><br class=""><blockquote type="cite" class=""><div class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; -webkit-line-break: after-white-space;" class=""><div class=""><br class=""></div><div class="">Marcello</div><div class=""><br class=""></div><div class=""><div class=""><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></div></div></blockquote></div><br class=""></div></div></blockquote></div><br class=""></div></div></div></blockquote></div><br class=""></body></html>