<div dir="ltr"><div class="gmail_extra"><div class="gmail_quote">On Wed, May 4, 2016 at 5:47 PM, Ettore Speziale via cfe-commits <span dir="ltr"><<a href="mailto:cfe-commits@lists.llvm.org" target="_blank">cfe-commits@lists.llvm.org</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Hello,<br>
<span class=""><br>
> I would appreciate a bit more background on this attribute's<br>
> semantics. How would a user know when to add this attribute to their<br>
> function definition? Are there other attributes that cannot be used in<br>
> conjunction with this one? Should this apply to member functions? What<br>
> about Objective-C methods?<br>
<br>
</span>The convergent attribute is meant to be used with languages supporting the SIMT execution model, like OpenCL.<br>
<br>
I put the following example in the documentation:<br>
<br>
__attribute__((convergent)) __attribute__((pure)) int foo(void) {<br>
int x;<br>
...<br>
barrier(CLK_GLOBAL_MEM_FENCE);<br>
...<br>
return x;<br>
}<br>
<br>
kernel void bar(global int *y) {<br>
int z = foo();<br>
*y = get_global_id() == 0 ? z : 0;<br>
}<br>
<br>
The call to barrier must be either executed by all work-items in a work-group, or by none of them.<br>
This is a requirement of OpenCL, and is left to the programmer to ensure that happens.<br>
<br>
In the case of foo, there could be a problem.<br>
If you do not mark it convergent, the LLVM sink pass push the call to foo to the then branch of the ternary operator, hence the program has been incorrectly optimized.<br></blockquote><div><br></div><div>Really? It looks like the problem is that you lied to the compiler by marking the function as 'pure'. The barrier is a side-effect that cannot be removed or duplicated, so it's not correct to mark this function as pure.</div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
The LLVM convergent attribute has been introduced in order to solve this problem for intrinsic functions.<br>
The goal of this patch is to expose that attribute at the CLANG level, so it can be used on all functions.<br>
<br>
The user is supposed to add such attribute when the function requires convergent execution, like in the example above.<br>
<br>
I’m not aware of any attribute that would conflict with convergent.<br>
<br>
The convergent attribute can be applied as well to member functions.<br>
<br>
The convergent attribute cannot be applied to Objective-C methods right now — it will be ignored:<br>
<br>
test.c:14:27: warning: 'convergent' attribute only applies to functions [-Wignored-attributes]<br>
- (void) x __attribute__((convergent));<br>
<br>
Since convergent is meant for languages supporting the SIMT execution model, and to the best of my knowledge I’m not aware of any language based on Objective-C supporting that, I would guess there is no benefit in supporting convergent on ObjectiveC methods.<br>
<span class=""><br>
>> diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td<br>
>> index df41aeb..eafafc6 100644<br>
>> --- a/include/clang/Basic/Attr.td<br>
>> +++ b/include/clang/Basic/Attr.td<br>
>> @@ -580,6 +580,12 @@ def Constructor : InheritableAttr {<br>
>> let Documentation = [Undocumented];<br>
>> }<br>
>><br>
>> +def Convergent : InheritableAttr {<br>
>> + let Spellings = [GNU<"convergent">];<br>
><br>
> Is there a reason to not support this under CXX11<"clang",<br>
> "convergent"> as well?<br>
<br>
</span>I’ve just used the most basic spelling, which fit the OpenCL case.<br>
I can add support for the CXX11 spelling if you find it valuable.<br>
<span class=""><br>
>> + let Subjects = SubjectList<[Function]>;<br>
>> + let Documentation = [Undocumented];<br>
><br>
> Please, no new undocumented attributes.<br>
<br>
</span>Fixed, here is updated patch:<br>
<br>
<br><br>
<br>
Thanks!<br>
<br>
--------------------------------------------------<br>
Ettore Speziale — Compiler Engineer<br>
<a href="mailto:speziale.ettore@gmail.com">speziale.ettore@gmail.com</a><br>
<a href="mailto:espeziale@apple.com">espeziale@apple.com</a><br>
--------------------------------------------------<br>
<br>
<br>_______________________________________________<br>
cfe-commits mailing list<br>
<a href="mailto:cfe-commits@lists.llvm.org">cfe-commits@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits</a><br>
<br></blockquote></div><br></div></div>