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/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(getDefaultCXXMethodAddrSpace()); + // 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()) { 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)(); +}