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

[Anastasia Stulova via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGe6522a96f56c: [OpenCL] Allow addr space qualifiers on lambda call expressions (authored by Anastasia).
Herald added a project: clang.

Changed prior to commit:
  https://reviews.llvm.org/D70242?vs=229305&id=232086#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D70242/new/

https://reviews.llvm.org/D70242

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


Index: clang/test/SemaOpenCLCXX/address-space-lambda.cl
===================================================================
--- clang/test/SemaOpenCLCXX/address-space-lambda.cl
+++ 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 @@
   (*(__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}}
+}
+
Index: clang/lib/Parse/ParseExprCXX.cpp
===================================================================
--- clang/lib/Parse/ParseExprCXX.cpp
+++ clang/lib/Parse/ParseExprCXX.cpp
@@ -1352,6 +1352,13 @@
     // 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 @@
                       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 @@
     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;


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D70242.232086.patch
Type: text/x-patch
Size: 5140 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20191204/e9b44af6/attachment.bin>