This revision was automatically updated to reflect the committed changes.
Closed by commit rG980133a2098c: [OpenCL] Use generic addr space for lambda
call operator (authored by Anastasia).
Herald added a project: clang.
Changed prior to commit:
https://reviews.llvm.org/D69938?vs=231409&id=231917#toc
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D69938/new/
https://reviews.llvm.org/D69938
Files:
clang/include/clang/Sema/Sema.h
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaDeclCXX.cpp
clang/lib/Sema/SemaLambda.cpp
clang/lib/Sema/SemaType.cpp
clang/test/SemaOpenCLCXX/address-space-lambda.cl
Index: clang/test/SemaOpenCLCXX/address-space-lambda.cl
===================================================================
--- /dev/null
+++ clang/test/SemaOpenCLCXX/address-space-lambda.cl
@@ -0,0 +1,25 @@
+//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic'
+auto glambda = [](auto a) { return a; };
+
+__kernel void foo() {
+ int i;
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
+ auto llambda = [&]() {i++;};
+ llambda();
+ glambda(1);
+ // Test lambda with default parameters
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
+ [&] {i++;} ();
+ __constant auto err = [&]() {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}})'}}
+ err(); //expected-error-re{{no matching function for call to object of type '__constant (lambda at {{.*}})'}}
+ // FIXME: There is very limited addr space functionality
+ // we can test when taking lambda type from the object.
+ // The limitation is due to addr spaces being added to all
+ // objects in OpenCL. Once we add metaprogramming utility
+ // for removing address spaces from a type we can enhance
+ // testing here.
+ (*(__constant decltype(llambda) *)nullptr)(); //expected-error{{multiple address spaces specified for type}}
+ (*(decltype(llambda) *)nullptr)();
+}
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -4950,7 +4950,9 @@
.getScopeRep()
->getKind() == NestedNameSpecifier::TypeSpec) ||
state.getDeclarator().getContext() ==
- DeclaratorContext::MemberContext;
+ DeclaratorContext::MemberContext ||
+ state.getDeclarator().getContext() ==
+ DeclaratorContext::LambdaExprContext;
};
if (state.getSema().getLangOpts().OpenCLCPlusPlus && IsClassMember()) {
@@ -4969,7 +4971,8 @@
// If a class member function's address space is not set, set it to
// __generic.
LangAS AS =
- (ASIdx == LangAS::Default ? LangAS::opencl_generic : ASIdx);
+ (ASIdx == LangAS::Default ? S.getDefaultCXXMethodAddrSpace()
+ : ASIdx);
EPI.TypeQuals.addAddressSpace(AS);
}
T = Context.getFunctionType(T, ParamTys, EPI);
Index: clang/lib/Sema/SemaLambda.cpp
===================================================================
--- clang/lib/Sema/SemaLambda.cpp
+++ clang/lib/Sema/SemaLambda.cpp
@@ -917,6 +917,10 @@
/*IsVariadic=*/false, /*IsCXXMethod=*/true));
EPI.HasTrailingReturn = true;
EPI.TypeQuals.addConst();
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default)
+ EPI.TypeQuals.addAddressSpace(AS);
+
// C++1y [expr.prim.lambda]:
// The lambda return type is 'auto', which is replaced by the
// trailing-return type if provided and/or deduced from 'return'
Index: clang/lib/Sema/SemaDeclCXX.cpp
===================================================================
--- clang/lib/Sema/SemaDeclCXX.cpp
+++ clang/lib/Sema/SemaDeclCXX.cpp
@@ -11417,10 +11417,9 @@
// Build an exception specification pointing back at this constructor.
FunctionProtoType::ExtProtoInfo EPI = getImplicitMethodEPI(*this, SpecialMem);
- if (getLangOpts().OpenCLCPlusPlus) {
- // OpenCL: Implicitly defaulted special member are of the generic address
- // space.
- EPI.TypeQuals.addAddressSpace(LangAS::opencl_generic);
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default) {
+ EPI.TypeQuals.addAddressSpace(AS);
}
auto QT = Context.getFunctionType(ResultTy, Args, EPI);
@@ -12330,8 +12329,9 @@
return nullptr;
QualType ArgType = Context.getTypeDeclType(ClassDecl);
- if (Context.getLangOpts().OpenCLCPlusPlus)
- ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default)
+ ArgType = Context.getAddrSpaceQualType(ArgType, AS);
QualType RetType = Context.getLValueReferenceType(ArgType);
bool Const = ClassDecl->implicitCopyAssignmentHasConstParam();
if (Const)
@@ -12656,8 +12656,9 @@
// constructor rules.
QualType ArgType = Context.getTypeDeclType(ClassDecl);
- if (Context.getLangOpts().OpenCLCPlusPlus)
- ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default)
+ ArgType = Context.getAddrSpaceQualType(ArgType, AS);
QualType RetType = Context.getLValueReferenceType(ArgType);
ArgType = Context.getRValueReferenceType(ArgType);
@@ -13034,8 +13035,9 @@
if (Const)
ArgType = ArgType.withConst();
- if (Context.getLangOpts().OpenCLCPlusPlus)
- ArgType = Context.getAddrSpaceQualType(ArgType, LangAS::opencl_generic);
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default)
+ ArgType = Context.getAddrSpaceQualType(ArgType, AS);
ArgType = Context.getLValueReferenceType(ArgType);
@@ -13166,8 +13168,9 @@
QualType ClassType = Context.getTypeDeclType(ClassDecl);
QualType ArgType = ClassType;
- if (Context.getLangOpts().OpenCLCPlusPlus)
- ArgType = Context.getAddrSpaceQualType(ClassType, LangAS::opencl_generic);
+ LangAS AS = getDefaultCXXMethodAddrSpace();
+ if (AS != LangAS::Default)
+ ArgType = Context.getAddrSpaceQualType(ClassType, AS);
ArgType = Context.getRValueReferenceType(ArgType);
bool Constexpr = defaultedSpecialMemberIsConstexpr(*this, ClassDecl,
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -1290,6 +1290,12 @@
return nullptr;
}
+LangAS Sema::getDefaultCXXMethodAddrSpace() const {
+ if (getLangOpts().OpenCL)
+ return LangAS::opencl_generic;
+ return LangAS::Default;
+}
+
void Sema::EmitCurrentDiagnostic(unsigned DiagID) {
// FIXME: It doesn't make sense to me that DiagID is an incoming argument here
// and yet we also use the current diag ID on the DiagnosticsEngine. This has
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -3907,6 +3907,9 @@
/// Add the given method to the list of globally-known methods.
void addMethodToGlobalList(ObjCMethodList *List, ObjCMethodDecl *Method);
+ /// Returns default addr space for method qualifiers.
+ LangAS getDefaultCXXMethodAddrSpace() const;
+
private:
/// AddMethodToGlobalPool - Add an instance or factory method to the global
/// pool. See descriptoin of AddInstanceMethodToGlobalPool.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
