I think it could be useful for CUDA too. On Friday, May 6, 2016, Anastasia Stulova via cfe-commits < cfe-commits@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.ett...@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@lists.llvm.org <javascript:;> > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits >
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
