[clang] e6522a9 - [OpenCL] Allow addr space qualifiers on lambda call expressions

[Anastasia Stulova via cfe-commits]

Author: Anastasia Stulova
Date: 2019-12-04T12:25:20Z
New Revision: e6522a96f56ce0257ab8cc6fca77bf9ea4462fa6

URL: https://github.com/llvm/llvm-project/commit/e6522a96f56ce0257ab8cc6fca77bf9ea4462fa6
DIFF: https://github.com/llvm/llvm-project/commit/e6522a96f56ce0257ab8cc6fca77bf9ea4462fa6.diff

LOG: [OpenCL] Allow addr space qualifiers on lambda call expressions

The addr space qualifier can be added optionally for lambdas after
the attributes. They will alter the default addr space of lambda
call operator that is in generic address space by default for OpenCL.

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 addr space conversion rules
and there will be no difference to how addr spaces work in method
qualifiers.

Tags: #clang

Differential Revision: https://reviews.llvm.org/D70242

Added: 
    

Modified: 
    clang/lib/Parse/ParseExprCXX.cpp
    clang/test/SemaOpenCLCXX/address-space-lambda.cl

Removed: 
    


################################################################################
diff  --git a/clang/lib/Parse/ParseExprCXX.cpp b/clang/lib/Parse/ParseExprCXX.cpp
index 77eed5437609..7dfe71fb9ebc 100644
--- a/clang/lib/Parse/ParseExprCXX.cpp
+++ b/clang/lib/Parse/ParseExprCXX.cpp
@@ -1352,6 +1352,13 @@ ExprResult Parser::ParseLambdaExpressionAfterIntroducer(
     // 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].
@@ -1380,10 +1387,12 @@ ExprResult Parser::ParseLambdaExpressionAfterIntroducer(
                       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.
@@ -1392,6 +1401,11 @@ ExprResult Parser::ParseLambdaExpressionAfterIntroducer(
     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;

diff  --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl
index 8f7a839da448..cf87bfaeede2 100644
--- a/clang/test/SemaOpenCLCXX/address-space-lambda.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-lambda.cl
@@ -3,7 +3,7 @@
 //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++;};
@@ -23,3 +23,31 @@ __kernel void foo() {
   (*(__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}}
+}
+