[Clang] Convergent Attribute

[Ettore Speziale via cfe-commits]

Hello,

> I would appreciate a bit more background on this attribute's
> semantics. How would a user know when to add this attribute to their
> function definition? Are there other attributes that cannot be used in
> conjunction with this one? Should this apply to member functions? What
> about Objective-C methods?

The convergent attribute is meant to be used with languages supporting the SIMT execution model, like OpenCL.

I put the following example in the documentation:

  __attribute__((convergent)) __attribute__((pure)) int foo(void) {
    int x;
    ...
    barrier(CLK_GLOBAL_MEM_FENCE);
    ...
    return x;
  }

  kernel void bar(global int *y) {
    int z = foo();
    *y = get_global_id() == 0 ? z : 0;
  }

The call to barrier must be either executed by all work-items in a work-group, or by none of them.
This is a requirement of OpenCL, and is left to the programmer to ensure that happens.

In the case of foo, there could be a problem.
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.

The LLVM convergent attribute has been introduced in order to solve this problem for intrinsic functions.
The goal of this patch is to expose that attribute at the CLANG level, so it can be used on all functions.

The user is supposed to add such attribute when the function requires convergent execution, like in the example above.

I’m not aware of any attribute that would conflict with convergent.

The convergent attribute can be applied as well to member functions.

The convergent attribute cannot be applied to Objective-C methods right now — it will be ignored:

test.c:14:27: warning: 'convergent' attribute only applies to functions [-Wignored-attributes]
- (void) x __attribute__((convergent));

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.

>> diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
>> index df41aeb..eafafc6 100644
>> --- a/include/clang/Basic/Attr.td
>> +++ b/include/clang/Basic/Attr.td
>> @@ -580,6 +580,12 @@ def Constructor : InheritableAttr {
>>   let Documentation = [Undocumented];
>> }
>> 
>> +def Convergent : InheritableAttr {
>> +  let Spellings = [GNU<"convergent">];
> 
> Is there a reason to not support this under CXX11<"clang",
> "convergent"> as well?

I’ve just used the most basic spelling, which fit the OpenCL case.
I can add support for the CXX11 spelling if you find it valuable.

>> +  let Subjects = SubjectList<[Function]>;
>> +  let Documentation = [Undocumented];
> 
> Please, no new undocumented attributes.

Fixed, here is updated patch:

-------------- next part --------------
A non-text attachment was scrubbed...
Name: convergent.diff
Type: application/octet-stream
Size: 4241 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160504/3a20a932/attachment.obj>
-------------- next part --------------


Thanks!

--------------------------------------------------
Ettore Speziale — Compiler Engineer
speziale.ettore at gmail.com
espeziale at apple.com
--------------------------------------------------