[clang] 980133a - [OpenCL] Use generic addr space for lambda call operator
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Tue Dec 3 08:08:05 PST 2019
Author: Anastasia Stulova
Date: 2019-12-03T16:07:18Z
New Revision: 980133a2098cf6159785b8ac0cbe4d8fbf99bfea
URL: https://github.com/llvm/llvm-project/commit/980133a2098cf6159785b8ac0cbe4d8fbf99bfea
DIFF: https://github.com/llvm/llvm-project/commit/980133a2098cf6159785b8ac0cbe4d8fbf99bfea.diff
LOG: [OpenCL] Use generic addr space for lambda call operator
Since lambdas are represented by callable objects, we add
generic addr space for implicit object parameter in call
operator.
Any lambda variable declared in __constant addr space
(which is not convertible to generic) fails to compile with
a diagnostic. To support constant addr space we need to
add a way to qualify the lambda call operators.
Tags: #clang
Differential Revision: https://reviews.llvm.org/D69938
Added:
clang/test/SemaOpenCLCXX/address-space-lambda.cl
Modified:
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
Removed:
################################################################################
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 59e8f3439669..bb8fd4c9c857 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -3907,6 +3907,9 @@ class Sema final {
/// 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.
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 035cb2e15ca0..5ade6bbea074 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -1290,6 +1290,12 @@ NamedDecl *Sema::getCurFunctionOrMethodDecl() {
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
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index f469580dd866..36f7bb320e08 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -11417,10 +11417,9 @@ void Sema::setupImplicitSpecialMemberType(CXXMethodDecl *SpecialMem,
// 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 @@ CXXMethodDecl *Sema::DeclareImplicitCopyAssignment(CXXRecordDecl *ClassDecl) {
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 @@ CXXMethodDecl *Sema::DeclareImplicitMoveAssignment(CXXRecordDecl *ClassDecl) {
// 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 @@ CXXConstructorDecl *Sema::DeclareImplicitCopyConstructor(
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 @@ CXXConstructorDecl *Sema::DeclareImplicitMoveConstructor(
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,
diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp
index c6b19a0b195c..14b443e9dac0 100644
--- a/clang/lib/Sema/SemaLambda.cpp
+++ b/clang/lib/Sema/SemaLambda.cpp
@@ -917,6 +917,10 @@ void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
/*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'
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index 4cbf041533b2..52a0581643bc 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -4950,7 +4950,9 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
.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 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state,
// 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);
diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl
new file mode 100644
index 000000000000..8f7a839da448
--- /dev/null
+++ b/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)();
+}
More information about the cfe-commits
mailing list