diff --git a/clang/include/clang/AST/DeclCXX.h b/clang/include/clang/AST/DeclCXX.h --- a/clang/include/clang/AST/DeclCXX.h +++ b/clang/include/clang/AST/DeclCXX.h @@ -389,9 +389,12 @@ /// The number of explicit captures in this lambda. unsigned NumExplicitCaptures : 13; + /// Has known `internal` linkage. + unsigned HasKnownInternalLinkage : 1; + /// The number used to indicate this lambda expression for name /// mangling in the Itanium C++ ABI. - unsigned ManglingNumber = 0; + unsigned ManglingNumber : 31; /// The declaration that provides context for this lambda, if the /// actual DeclContext does not suffice. This is used for lambdas that @@ -406,12 +409,12 @@ /// The type of the call method. TypeSourceInfo *MethodTyInfo; - LambdaDefinitionData(CXXRecordDecl *D, TypeSourceInfo *Info, - bool Dependent, bool IsGeneric, - LambdaCaptureDefault CaptureDefault) - : DefinitionData(D), Dependent(Dependent), IsGenericLambda(IsGeneric), - CaptureDefault(CaptureDefault), NumCaptures(0), NumExplicitCaptures(0), - MethodTyInfo(Info) { + LambdaDefinitionData(CXXRecordDecl *D, TypeSourceInfo *Info, bool Dependent, + bool IsGeneric, LambdaCaptureDefault CaptureDefault) + : DefinitionData(D), Dependent(Dependent), IsGenericLambda(IsGeneric), + CaptureDefault(CaptureDefault), NumCaptures(0), + NumExplicitCaptures(0), HasKnownInternalLinkage(0), ManglingNumber(0), + MethodTyInfo(Info) { IsLambda = true; // C++1z [expr.prim.lambda]p4: @@ -1705,6 +1708,13 @@ return getLambdaData().ManglingNumber; } + /// The lambda is known to has internal linkage no matter whether it has name + /// mangling number. + bool hasKnownLambdaInternalLinkage() const { + assert(isLambda() && "Not a lambda closure type!"); + return getLambdaData().HasKnownInternalLinkage; + } + /// Retrieve the declaration that provides additional context for a /// lambda, when the normal declaration context is not specific enough. /// @@ -1718,9 +1728,12 @@ /// Set the mangling number and context declaration for a lambda /// class. - void setLambdaMangling(unsigned ManglingNumber, Decl *ContextDecl) { + void setLambdaMangling(unsigned ManglingNumber, Decl *ContextDecl, + bool HasKnownInternalLinkage = false) { + assert(isLambda() && "Not a lambda closure type!"); getLambdaData().ManglingNumber = ManglingNumber; getLambdaData().ContextDecl = ContextDecl; + getLambdaData().HasKnownInternalLinkage = HasKnownInternalLinkage; } /// Returns the inheritance model used for this record. 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 @@ -5926,12 +5926,17 @@ LambdaCaptureDefault CaptureDefault); /// Start the definition of a lambda expression. - CXXMethodDecl * - startLambdaDefinition(CXXRecordDecl *Class, SourceRange IntroducerRange, - TypeSourceInfo *MethodType, SourceLocation EndLoc, - ArrayRef Params, - ConstexprSpecKind ConstexprKind, - Optional> Mangling = None); + CXXMethodDecl *startLambdaDefinition(CXXRecordDecl *Class, + SourceRange IntroducerRange, + TypeSourceInfo *MethodType, + SourceLocation EndLoc, + ArrayRef Params, + ConstexprSpecKind ConstexprKind); + + /// Number lambda for linkage purposes if necessary. + void handleLambdaNumbering( + CXXRecordDecl *Class, CXXMethodDecl *Method, + Optional> Mangling = None); /// Endow the lambda scope info with the relevant properties. void buildLambdaScope(sema::LambdaScopeInfo *LSI, diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp --- a/clang/lib/AST/ASTImporter.cpp +++ b/clang/lib/AST/ASTImporter.cpp @@ -2694,7 +2694,8 @@ ExpectedDecl CDeclOrErr = import(DCXX->getLambdaContextDecl()); if (!CDeclOrErr) return CDeclOrErr.takeError(); - D2CXX->setLambdaMangling(DCXX->getLambdaManglingNumber(), *CDeclOrErr); + D2CXX->setLambdaMangling(DCXX->getLambdaManglingNumber(), *CDeclOrErr, + DCXX->hasKnownLambdaInternalLinkage()); } else if (DCXX->isInjectedClassName()) { // We have to be careful to do a similar dance to the one in // Sema::ActOnStartCXXMemberDeclarations diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp --- a/clang/lib/AST/Decl.cpp +++ b/clang/lib/AST/Decl.cpp @@ -1385,7 +1385,8 @@ case Decl::CXXRecord: { const auto *Record = cast(D); if (Record->isLambda()) { - if (!Record->getLambdaManglingNumber()) { + if (Record->hasKnownLambdaInternalLinkage() || + !Record->getLambdaManglingNumber()) { // This lambda has no mangling number, so it's internal. return getInternalLinkageFor(D); } @@ -1402,7 +1403,8 @@ // }; const CXXRecordDecl *OuterMostLambda = getOutermostEnclosingLambda(Record); - if (!OuterMostLambda->getLambdaManglingNumber()) + if (OuterMostLambda->hasKnownLambdaInternalLinkage() || + !OuterMostLambda->getLambdaManglingNumber()) return getInternalLinkageFor(D); return getLVForClosure( 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 @@ -335,7 +335,7 @@ case StaticDataMember: // -- the initializers of nonspecialized static members of template classes if (!IsInNonspecializedTemplate) - return std::make_tuple(nullptr, nullptr); + return std::make_tuple(nullptr, ManglingContextDecl); // Fall through to get the current context. LLVM_FALLTHROUGH; @@ -356,14 +356,15 @@ llvm_unreachable("unexpected context"); } -CXXMethodDecl *Sema::startLambdaDefinition( - CXXRecordDecl *Class, SourceRange IntroducerRange, - TypeSourceInfo *MethodTypeInfo, SourceLocation EndLoc, - ArrayRef Params, ConstexprSpecKind ConstexprKind, - Optional> Mangling) { +CXXMethodDecl *Sema::startLambdaDefinition(CXXRecordDecl *Class, + SourceRange IntroducerRange, + TypeSourceInfo *MethodTypeInfo, + SourceLocation EndLoc, + ArrayRef Params, + ConstexprSpecKind ConstexprKind) { QualType MethodType = MethodTypeInfo->getType(); TemplateParameterList *TemplateParams = - getGenericLambdaTemplateParameterList(getCurLambda(), *this); + getGenericLambdaTemplateParameterList(getCurLambda(), *this); // If a lambda appears in a dependent context or is a generic lambda (has // template parameters) and has an 'auto' return type, deduce it to a // dependent type. @@ -425,20 +426,55 @@ P->setOwningFunction(Method); } + return Method; +} + +void Sema::handleLambdaNumbering( + CXXRecordDecl *Class, CXXMethodDecl *Method, + Optional> Mangling) { if (Mangling) { - Class->setLambdaMangling(Mangling->first, Mangling->second); - } else { - MangleNumberingContext *MCtx; + unsigned ManglingNumber; + bool HasKnownInternalLinkage; Decl *ManglingContextDecl; - std::tie(MCtx, ManglingContextDecl) = - getCurrentMangleNumberContext(Class->getDeclContext()); - if (MCtx) { - unsigned ManglingNumber = MCtx->getManglingNumber(Method); - Class->setLambdaMangling(ManglingNumber, ManglingContextDecl); - } + std::tie(ManglingNumber, HasKnownInternalLinkage, ManglingContextDecl) = + Mangling.getValue(); + Class->setLambdaMangling(ManglingNumber, ManglingContextDecl, + HasKnownInternalLinkage); + return; } - return Method; + auto getMangleNumberingContext = + [this](CXXRecordDecl *Class, Decl *ManglingContextDecl) -> MangleNumberingContext * { + // Get mangle numbering context if there's any extra decl context. + if (ManglingContextDecl) + return &Context.getManglingNumberContext( + ASTContext::NeedExtraManglingDecl, ManglingContextDecl); + // Otherwise, from that lambda's decl context. + auto DC = Class->getDeclContext(); + while (auto *CD = dyn_cast(DC)) + DC = CD->getParent(); + return &Context.getManglingNumberContext(DC); + }; + + MangleNumberingContext *MCtx; + Decl *ManglingContextDecl; + std::tie(MCtx, ManglingContextDecl) = + getCurrentMangleNumberContext(Class->getDeclContext()); + bool HasKnownInternalLinkage = false; + if (!MCtx && getLangOpts().CUDA) { + // Force lambda numbering in CUDA/HIP as we need to name lambdas following + // ODR. Both device- and host-compilation need to have a consistent naming + // on kernel functions. As lambdas are potential part of these `__global__` + // function names, they needs numbering following ODR. + MCtx = getMangleNumberingContext(Class, ManglingContextDecl); + assert(MCtx && "Retrieving mangle numbering context failed!"); + HasKnownInternalLinkage = true; + } + if (MCtx) { + unsigned ManglingNumber = MCtx->getManglingNumber(Method); + Class->setLambdaMangling(ManglingNumber, ManglingContextDecl, + HasKnownInternalLinkage); + } } void Sema::buildLambdaScope(LambdaScopeInfo *LSI, @@ -951,6 +987,9 @@ if (getLangOpts().CUDA) CUDASetLambdaAttrs(Method); + // Number the lambda for linkage purposes if necessary. + handleLambdaNumbering(Class, Method); + // Introduce the function call operator as the current declaration context. PushDeclContext(CurScope, Method); diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -11497,17 +11497,18 @@ E->getCaptureDefault()); getDerived().transformedLocalDecl(OldClass, {Class}); - Optional> Mangling; + Optional> Mangling; if (getDerived().ReplacingOriginal()) - Mangling = std::make_pair(OldClass->getLambdaManglingNumber(), - OldClass->getLambdaContextDecl()); + Mangling = std::make_tuple(OldClass->getLambdaManglingNumber(), + OldClass->hasKnownLambdaInternalLinkage(), + OldClass->getLambdaContextDecl()); // Build the call operator. CXXMethodDecl *NewCallOperator = getSema().startLambdaDefinition( Class, E->getIntroducerRange(), NewCallOpTSI, E->getCallOperator()->getEndLoc(), NewCallOpTSI->getTypeLoc().castAs().getParams(), - E->getCallOperator()->getConstexprKind(), Mangling); + E->getCallOperator()->getConstexprKind()); LSI->CallOperator = NewCallOperator; @@ -11527,6 +11528,9 @@ getDerived().transformAttrs(E->getCallOperator(), NewCallOperator); getDerived().transformedLocalDecl(E->getCallOperator(), {NewCallOperator}); + // Number the lambda for linkage purposes if necessary. + getSema().handleLambdaNumbering(Class, NewCallOperator, Mangling); + // Introduce the context of the call operator. Sema::ContextRAII SavedContext(getSema(), NewCallOperator, /*NewThisContext*/false); diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp --- a/clang/lib/Serialization/ASTReaderDecl.cpp +++ b/clang/lib/Serialization/ASTReaderDecl.cpp @@ -1690,6 +1690,7 @@ Lambda.CaptureDefault = Record.readInt(); Lambda.NumCaptures = Record.readInt(); Lambda.NumExplicitCaptures = Record.readInt(); + Lambda.HasKnownInternalLinkage = Record.readInt(); Lambda.ManglingNumber = Record.readInt(); Lambda.ContextDecl = ReadDeclID(); Lambda.Captures = (Capture *)Reader.getContext().Allocate( diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp --- a/clang/lib/Serialization/ASTWriter.cpp +++ b/clang/lib/Serialization/ASTWriter.cpp @@ -6224,6 +6224,7 @@ Record->push_back(Lambda.CaptureDefault); Record->push_back(Lambda.NumCaptures); Record->push_back(Lambda.NumExplicitCaptures); + Record->push_back(Lambda.HasKnownInternalLinkage); Record->push_back(Lambda.ManglingNumber); AddDeclRef(D->getLambdaContextDecl()); AddTypeSourceInfo(Lambda.MethodTyInfo); diff --git a/clang/test/CodeGenCUDA/unnamed-types.cu b/clang/test/CodeGenCUDA/unnamed-types.cu new file mode 100644 --- /dev/null +++ b/clang/test/CodeGenCUDA/unnamed-types.cu @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -std=c++11 -x hip -triple x86_64-linux-gnu -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=HOST +// RUN: %clang_cc1 -std=c++11 -x hip -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm %s -o - | FileCheck %s --check-prefix=DEVICE + +#include "Inputs/cuda.h" + +// HOST: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1 + +__device__ float d0(float x) { + return [](float x) { return x + 2.f; }(x); +} + +__device__ float d1(float x) { + return [](float x) { return x * 2.f; }(x); +} + +// DEVICE: amdgpu_kernel void @_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_( +template +__global__ void k0(float *p, F f) { + p[0] = f(p[0]) + d0(p[1]) + d1(p[2]); +} + +void f0(float *p) { + [](float *p) { + *p = 1.f; + }(p); +} + +// The inner/outer lambdas are required to be mangled following ODR but their +// linkages are still required to keep the original `internal` linkage. + +// HOST: define internal void @_ZZ2f1PfENKUlS_E_clES_( +// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf( +void f1(float *p) { + [](float *p) { + k0<<<1,1>>>(p, [] __device__ (float x) { return x + 1.f; }); + }(p); +} +// HOST: @__hip_register_globals +// HOST: __hipRegisterFunction{{.*}}@_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_{{.*}}@0