[Clang] Convergent Attribute

[David Majnemer via cfe-commits]

I think it could be useful for CUDA too.

On Friday, May 6, 2016, Anastasia Stulova via cfe-commits <
cfe-commits at lists.llvm.org> wrote:

> Hi Ettore,
>
> LGTM generally!
>
> 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.
>
> Thanks!
> Anastasia
>
> -----Original Message-----
> From: Ettore Speziale [mailto:speziale.ettore at gmail.com <javascript:;>]
> Sent: 05 May 2016 01:48
> To: Aaron Ballman
> Cc: Ettore Speziale; Anastasia Stulova; Clang Commits
> Subject: Re: [Clang] Convergent Attribute
>
> 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:
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org <javascript:;>
> 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/de2ac4db/attachment.html>