I think it could<span></span> be useful for CUDA too.<br><br>On Friday, May 6, 2016, Anastasia Stulova via cfe-commits <<a href="mailto:cfe-commits@lists.llvm.org">cfe-commits@lists.llvm.org</a>> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Hi Ettore,<br>
<br>
LGTM generally!<br>
<br>
I was just wondering whether it would make sense to restrict the usage of the attribute to OpenCL language i.e. to add  "let LangOpts = [OpenCL];" in the attribute definition.<br>
<br>
Thanks!<br>
Anastasia<br>
<br>
-----Original Message-----<br>
From: Ettore Speziale [mailto:<a href="javascript:;" onclick="_e(event, 'cvml', 'speziale.ettore@gmail.com')">speziale.ettore@gmail.com</a>]<br>
Sent: 05 May 2016 01:48<br>
To: Aaron Ballman<br>
Cc: Ettore Speziale; Anastasia Stulova; Clang Commits<br>
Subject: Re: [Clang] Convergent Attribute<br>
<br>
Hello,<br>
<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>
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>
<br>
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>
<br>
>> diff --git a/include/clang/Basic/Attr.td<br>
>> b/include/clang/Basic/Attr.td 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>
>> +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>
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>
<br>
>> +  let Subjects = SubjectList<[Function]>;  let Documentation =<br>
>> + [Undocumented];<br>
><br>
> Please, no new undocumented attributes.<br>
<br>
Fixed, here is updated patch:<br>
<br>
_______________________________________________<br>
cfe-commits mailing list<br>
<a href="javascript:;" onclick="_e(event, 'cvml', 'cfe-commits@lists.llvm.org')">cfe-commits@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits</a><br>
</blockquote>