diff --git a/clang/lib/Parse/ParseExprCXX.cpp b/clang/lib/Parse/ParseExprCXX.cpp --- a/clang/lib/Parse/ParseExprCXX.cpp +++ b/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; diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl --- 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 @@ (*(__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}} +} +