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 @@ -584,7 +584,10 @@ unsigned NumCaptures : 15; /// The number of explicit captures in this lambda. - unsigned NumExplicitCaptures : 13; + unsigned NumExplicitCaptures : 12; + + /// Has known `internal` linkage. + unsigned HasKnownInternalLinkage : 1; /// The number used to indicate this lambda expression for name /// mangling in the Itanium C++ ABI. @@ -603,12 +606,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), + MethodTyInfo(Info) { IsLambda = true; // C++1z [expr.prim.lambda]p4: @@ -1902,6 +1905,13 @@ return getLambdaData().ManglingNumber; } + /// The lambda is known to has internal linkage no matter whether it has name + /// mangling number. + bool hasKnownInternalLinkage() 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. /// @@ -1913,11 +1923,18 @@ /// the declaration context suffices. Decl *getLambdaContextDecl() const; + void setLambdaContextDecl(Decl *ContextDecl) { + assert(isLambda() && "Not a lambda closure type!"); + getLambdaData().ContextDecl = ContextDecl; + } + /// Set the mangling number and context declaration for a lambda /// class. - void setLambdaMangling(unsigned ManglingNumber, Decl *ContextDecl) { + void setLambdaManglingNumber(unsigned ManglingNumber, + 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 @@ -5853,12 +5853,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 @@ -2755,7 +2755,9 @@ ExpectedDecl CDeclOrErr = import(DCXX->getLambdaContextDecl()); if (!CDeclOrErr) return CDeclOrErr.takeError(); - D2CXX->setLambdaMangling(DCXX->getLambdaManglingNumber(), *CDeclOrErr); + D2CXX->setLambdaManglingNumber(DCXX->getLambdaManglingNumber(), + DCXX->hasKnownInternalLinkage()); + D2CXX->setLambdaContextDecl(*CDeclOrErr); } 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->hasKnownInternalLinkage() || + !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->hasKnownInternalLinkage() || + !OuterMostLambda->getLambdaManglingNumber()) return getInternalLinkageFor(D); return getLVForClosure( diff --git a/clang/lib/AST/DeclCXX.cpp b/clang/lib/AST/DeclCXX.cpp --- a/clang/lib/AST/DeclCXX.cpp +++ b/clang/lib/AST/DeclCXX.cpp @@ -1487,6 +1487,9 @@ Decl *CXXRecordDecl::getLambdaContextDecl() const { assert(isLambda() && "Not a lambda closure type!"); + // Skip if this lambda is not numbered. + if (!getLambdaData().ManglingNumber) + return nullptr; ExternalASTSource *Source = getParentASTContext().getExternalSource(); return getLambdaData().ContextDecl.get(Source); } 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,72 @@ 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->setLambdaManglingNumber(ManglingNumber, HasKnownInternalLinkage); + Class->setLambdaContextDecl(ManglingContextDecl); + return; } - return Method; + auto getMangleNumberingContext = + [this](CXXRecordDecl *Class) -> MangleNumberingContext * { + // Get mangle numbering context if there's any extra decl context. + if (auto ManglingContextDecl = Class->getLambdaContextDecl()) + 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()); + // Remember the extra mangle context decl even this lambda is not numbered so + // far. A lambda may need numbering later once an inner nested lambda is an + // extended one. + Class->setLambdaContextDecl(ManglingContextDecl); + if (MCtx) { + unsigned ManglingNumber = MCtx->getManglingNumber(Method); + Class->setLambdaManglingNumber(ManglingNumber); + } + // For an extended lambda, i.e. a lambda annotated with `__device__`, it must + // have a mangling number so that it could be named consistently between + // host- and device-side compilations following ODR rule. Furthermore, if + // that extended lambda is nested within other lambdas, all its parent + // lambdas must have mangling number as well in order to have a consistently + // mangled name. But, the mangler needs informing those re-numbered lambdas + // still have `internal` linkage. + if (getLangOpts().CUDA && Method->hasAttr()) { + for (DeclContext *DC = Method; DC && !DC->isTranslationUnit(); + DC = getLambdaAwareParentOfDeclContext(DC)) { + // Skip if it's not a lambda. + if (!isLambdaCallOperator(DC)) + continue; + Method = cast(DC); + Class = Method->getParent(); + // Skip if it already has mangling number. + if (Class->getLambdaManglingNumber()) + continue; + // Allocate/get mangle numbering context. + MCtx = getMangleNumberingContext(Class); + Class->setLambdaManglingNumber(MCtx->getManglingNumber(Method), + /*HasKnownInternalLinkage=*/true); + } + } } void Sema::buildLambdaScope(LambdaScopeInfo *LSI, @@ -951,6 +1004,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 @@ -11373,17 +11373,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->hasKnownInternalLinkage(), + 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; @@ -11403,6 +11404,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 @@ -1732,6 +1732,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 @@ -6222,6 +6222,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 -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