Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -3724,6 +3724,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. Index: clang/lib/Sema/Sema.cpp =================================================================== --- clang/lib/Sema/Sema.cpp +++ clang/lib/Sema/Sema.cpp @@ -1279,6 +1279,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/lib/Sema/SemaDeclCXX.cpp =================================================================== --- clang/lib/Sema/SemaDeclCXX.cpp +++ clang/lib/Sema/SemaDeclCXX.cpp @@ -11129,10 +11129,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); @@ -12036,8 +12035,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) @@ -12361,8 +12361,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); @@ -12739,8 +12740,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); @@ -12871,8 +12873,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/SemaLambda.cpp =================================================================== --- clang/lib/Sema/SemaLambda.cpp +++ clang/lib/Sema/SemaLambda.cpp @@ -887,6 +887,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/SemaType.cpp =================================================================== --- clang/lib/Sema/SemaType.cpp +++ clang/lib/Sema/SemaType.cpp @@ -4916,7 +4916,9 @@ .getScopeRep() ->getKind() == NestedNameSpecifier::TypeSpec) || state.getDeclarator().getContext() == - DeclaratorContext::MemberContext; + DeclaratorContext::MemberContext || + state.getDeclarator().getContext() == + DeclaratorContext::LambdaExprContext; }; if (state.getSema().getLangOpts().OpenCLCPlusPlus && IsClassMember()) { @@ -4935,7 +4937,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/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)(); +}