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;
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;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
