summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
-rw-r--r--clang/lib/Parse/ParseExprCXX.cpp18
-rw-r--r--clang/test/SemaOpenCLCXX/address-space-lambda.cl30
2 files changed, 45 insertions, 3 deletions
diff --git a/clang/lib/Parse/ParseExprCXX.cpp b/clang/lib/Parse/ParseExprCXX.cpp
index 77eed543760..7dfe71fb9eb 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 8f7a839da44..cf87bfaeede 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}}
+}
+
OpenPOWER on IntegriCloud