[Clang] Convergent Attribute

[Richard Smith via cfe-commits]

On Wed, May 4, 2016 at 5:47 PM, Ettore Speziale via cfe-commits <
cfe-commits at lists.llvm.org> wrote:

> 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.
>

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.


> 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:
>
>
>
>
> Thanks!
>
> --------------------------------------------------
> Ettore Speziale — Compiler Engineer
> speziale.ettore at gmail.com
> espeziale at apple.com
> --------------------------------------------------
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160506/36437546/attachment-0001.html>