[OpenCL] Allow addr space qualifiers on lambda call expressions

[Anastasia]

Summary

The addr space qualifier can be added optionally for lambdas after the attributes. They will alter the default address space of lambda call operator that is in __generic address space by default for OpenCL (see https://reviews.llvm.org/D69938).

Syntax:

[ captures ] ( params ) specifiers exception attr opencl_addrspace -> ret { body }

Example:

[&] (int i) mutable __global { ... };

On the call into lambda a compatibility check will be performed to determine whether address space of lambda object and its call operator are compatible. This will follow regular address space conversion rules and there will be no difference to the behavior of address spaces in method qualifiers (see: https://reviews.llvm.org/D55850)

[rjmccall]

Well that was easy.

Do we accept the address-space attribute in this position, or is that TBD?

[Anastasia]

In D70242#1748313, @rjmccall wrote:

Well that was easy.

Do we accept the address-space attribute in this position, or is that TBD?

We accept the OpenCL one right at the end. I might need to test more in C++ though...

clang/lib/Parse/ParseExprCXX.cpp

ESpecType = tryParseExceptionSpecification(/*Delayed=*/false,
                                                ExceptionSpecTokens);
 
     if (ESpecType != EST_None)
       DeclEndLoc = ESpecRange.getEnd();
 
     // Parse attribute-specifier[opt].
     MaybeParseCXX11Attributes(Attr, &DeclEndLoc);
 
+    // Parse OpenCL addr space attribute.
+    if (Tok.isOneOf(tok::kw___private, tok::kw___global, tok::kw___local,
+                    tok::kw___constant, tok::kw___generic)) {
+      ParseOpenCLQualifiers(DS.getAttributes());
+      ConsumeToken();
+    }
+
     SourceLocation FunLocalRangeEnd = DeclEndLoc;
 
     // Parse trailing-return-type[opt].
     if (Tok.is(tok::arrow)) {
       FunLocalRangeEnd = Tok.getLocation();
       SourceRange Range;
       TrailingReturnType =
           ParseTrailingReturnType(Range, /*MayBeFollowedByDirectInit*/ false);
D.AddTypeInfo(DeclaratorChunk::getFunction(
                       ParamInfo.size(), EllipsisLoc, RParenLoc,
                       /*RefQualifierIsLvalueRef=*/true,
                       /*RefQualifierLoc=*/NoLoc, MutableLoc, ESpecType,
                       ESpecRange, DynamicExceptions.data(),
                       DynamicExceptionRanges.data(), DynamicExceptions.size(),
                       NoexceptExpr.isUsable() ? NoexceptExpr.get() : nullptr,
                       /*ExceptionSpecTokens*/ nullptr,
                       /*DeclsInPrototype=*/None, LParenLoc, FunLocalRangeEnd, D,
-                      TrailingReturnType),
+                      TrailingReturnType, &DS),
                   std::move(Attr), DeclEndLoc);
   } else if (Tok.isOneOf(tok::kw_mutable, tok::arrow, tok::kw___attribute,
-                         tok::kw_constexpr, tok::kw_consteval) ||
+                         tok::kw_constexpr, tok::kw_consteval,
+                         tok::kw___private, tok::kw___global, tok::kw___local,
+                         tok::kw___constant, tok::kw___generic) ||
              (Tok.is(tok::l_square) && NextToken().is(tok::l_square))) {
     // It's common to forget that one needs '()' before 'mutable', an attribute
     // specifier, or the result type. Deal with this.
     unsigned TokKind = 0;
     switch (Tok.getKind()) {
     case tok::kw_mutable: TokKind = 0; break;
     case tok::arrow: TokKind = 1; break;
     case tok::kw___attribute:
+    case tok::kw___private:
+    case tok::kw___global:
+    case tok::kw___local:
+    case tok::kw___constant:
+    case tok::kw___generic:
     case tok::l_square: TokKind = 2; break;
     case tok::kw_constexpr: TokKind = 3; break;
     case tok::kw_consteval: TokKind = 4; break;
     default: llvm_unreachable("Unknown token kind");
     }
 
     Diag(Tok, diag::err_lambda_missing_parens)
       << TokKind

clang/test/SemaOpenCLCXX/address-space-lambda.cl

 //RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
 
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic'
 auto glambda = [](auto a) { return a; };
 
-__kernel void foo() {
+__kernel void test() {
   int i;
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   auto  llambda = [&]() {i++;};
   llambda();
   glambda(1);
   // Test lambda with default parameters
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
   [&] {i++;} ();
   __constant auto err = [&]() {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}})'}}
   err();                          //expected-error-re{{no matching function for call to object of type '__constant (lambda at {{.*}})'}}
   // FIXME: There is very limited addr space functionality
   // we can test when taking lambda type from the object.
   // The limitation is due to addr spaces being added to all
   // objects in OpenCL. Once we add metaprogramming utility
   // for removing address spaces from a type we can enhance
   // testing here.
   (*(__constant decltype(llambda) *)nullptr)(); //expected-error{{multiple address spaces specified for type}}
   (*(decltype(llambda) *)nullptr)();
 }
+
+__kernel void test_qual() {
+//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const'
+  auto priv1 = []() __private {};
+  priv1();
+//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
+  auto priv2 = []() __generic {};
+  priv2();
+  auto priv3 = []() __global {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const __global (lambda at {{.*}})'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  priv3(); //expected-error{{no matching function for call to object of type}}
+
+  __constant auto const1 = []() __private{}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const (lambda at {{.*}}'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
+  __constant auto const2 = []() __generic{}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}}'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
+//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __constant'
+  __constant auto const3 = []() __constant{};
+  const3();
+
+  [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const __global (lambda at {{.*}})'}}
+  [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('(lambda at {{.*}})'), parameter type must be 'const (lambda at {{.*}})'}}
+
+  [&] __private {} (); //expected-error{{lambda requires '()' before attribute specifier}} expected-error{{expected body of lambda expression}}
+
+  [&] () mutable __private {} ();
+  [&] () __private mutable {} (); //expected-error{{expected body of lambda expression}}
+}
+