diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/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. diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp --- a/clang/lib/Sema/Sema.cpp +++ b/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 diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/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, diff --git a/clang/lib/Sema/SemaLambda.cpp b/clang/lib/Sema/SemaLambda.cpp --- a/clang/lib/Sema/SemaLambda.cpp +++ b/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' diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp --- a/clang/lib/Sema/SemaType.cpp +++ b/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); diff --git a/clang/test/SemaOpenCLCXX/address-space-lambda.cl b/clang/test/SemaOpenCLCXX/address-space-lambda.cl new file mode 100644 --- /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)(); +}