[hip][cuda] Fix the extended lambda name mangling issue.

[hliao]

Summary

• HIP/CUDA host side needs to use device kernel symbol name to match the device side binaries. Without a consistent naming between host- and device-side compilations, it's risky that wrong device binaries are executed. Consistent naming is usually not an issue until unnamed types are used, especially the lambda. In this patch, the consistent name mangling is addressed for the extended lambdas, i.e. the lambdas annotated with __device__.
• In [Itanium C++ ABI][1], the mangling of the lambda is generally unspecified unless, in certain cases, ODR rule is required to ensure consisent naming cross TUs. The extended lambda is such a case as its name may be part of a device kernel function, e.g., the extended lambda is used as a template argument and etc. Thus, we need to force ODR for extended lambdas as they are referenced in both device- and host-side TUs. Furthermore, if a extended lambda is nested in other (extended or not) lambdas, those lambdas are required to follow ODR naming as well. This patch revises the current lambda mangle numbering to force ODR from an extended lambda to all its parent lambdas.
• On the other side, the aforementioned ODR naming should not change those lambdas' original linkages, i.e., we cannot replace the original internal with linkonce_odr; otherwise, we may violate ODR in general. This patch introduces a new field HasKnownInternalLinkage in lambda data to decouple the current linkage calculation based on mangling number assigned.

[1]: https://itanium-cxx-abi.github.io/cxx-abi/abi.html

[hliao]

this's a patch address the same issue previously proposed to be worked around in https://reviews.llvm.org/D63164

[hliao]

minor comment to help review.

BTW, as the AST ser/deser is changed as well, not sure we have compatibility issue and, if there is, how to handle that. please advice me on that concern. Thanks.

clang/include/clang/AST/DeclCXX.h

390–394

Try to avoid inflating memory footprint by borrowing one bit from NumExplicitCaptures. It seems to me that we won't have that large number (8192) of *explicit* capture.

clang/lib/Sema/SemaLambda.cpp

432

extracted from startLambdaDefinition so that we could have a single place to handle mangling numbers based on CUDA attributes.

[hliao]

PING for review, thanks

[hliao]

PING for review

[tra]

@rsmith Richard, could you take a look, please? Lambdas, mangling, ODR rules & ABI scare me. :-)

[hliao]

In D68818#1709688, @tra wrote:

@rsmith Richard, could you take a look, please? Lambdas, mangling, ODR rules & ABI scare me. :-)

@tra thanks for promoting the review. This patch is quite critical to support extended lambda in clang. We have several workloads that have this mangle numbering issue.

[hliao]

PING for review

[rsmith]

Broadly, I think it's reasonable to number additional lambda expressions in CUDA compilations. However:

• This is (in theory) an ABI break on the host side, as it changes the lambda numbering in inline functions and function templates and the like. That could be mitigated by using a different numbering sequence for the lambdas that are only numbered for this reason.
• Depending on whether the call operator is a device function is unstable. If I understand the CUDA rules correctly, then in practice, because constexpr functions are implicitly host device, all lambdas will get numbered in CUDA on C++14 onwards but not in CUDA on C++11, and we generally want those modes to be ABI-compatible. I'd suggest you simplify and stabilize this by simply numbering all lambdas in CUDA mode.

clang/lib/Sema/SemaLambda.cpp

477

What happens if there are other enclosing constructs that would give the lambda internal linkage (eg, an anonymous namespace or static function that might collide with one in another translation unit, or a declaration involving a type with internal linkage)? Presumably you can still suffer from mangling collisions in those cases, at least if you link together multiple translation units containing device code. Do we need (something like) unique identifiers for device code TUs to use in manglings of ostensibly internal linkage entities?

[hliao]

In D68818#1711698, @rsmith wrote:

Broadly, I think it's reasonable to number additional lambda expressions in CUDA compilations. However:

• This is (in theory) an ABI break on the host side, as it changes the lambda numbering in inline functions and function templates and the like. That could be mitigated by using a different numbering sequence for the lambdas that are only numbered for this reason.
• Depending on whether the call operator is a device function is unstable. If I understand the CUDA rules correctly, then in practice, because constexpr functions are implicitly host device, all lambdas will get numbered in CUDA on C++14 onwards but not in CUDA on C++11, and we generally want those modes to be ABI-compatible. I'd suggest you simplify and stabilize this by simply numbering all lambdas in CUDA mode.

I vote for the numbering of all lambdas in the CUDA mode once addressed in (D63164 as long as we decouple the numbering from the linkage. As long as all lambdas are numbered, the first one is also mitigated. I will prepare the changes forcing all lambdas to be numbered under CUDA/HIP mode.

clang/lib/Sema/SemaLambda.cpp

477

Those unnamed ones will be addressed in late patches as, currently, we only have lambda numbering issues from real workloads. I am preparing changes to address other unnamed types, namespaces, and etc.

[rjmccall]

I agree with Richard's suggestion of just numbering all the lambdas in both modes if that's viable.

clang/include/clang/AST/DeclCXX.h

390–394

If this field is necessary, please borrow its space from ManglingNumber below rather than the number of explicit captures.

[hliao]

Force numbering on all lambdas in CUDA/HIP.

[hliao]

PING for review

[hliao]

minor test case revising.

[rsmith]

clang/include/clang/AST/DeclCXX.h

1713

This function name should mention that it's only applicable to lambdas.

clang/include/clang/AST/DeclCXX.h

struct LambdaDefinitionData : public DefinitionData {
 
     /// The Default Capture.
     unsigned CaptureDefault : 2;
 
     /// The number of captures in this lambda is limited 2^NumCaptures.
     unsigned NumCaptures : 15;
 
     /// 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
     /// occur within default arguments of function parameters within the class
     /// or within a data member initializer.
     LazyDeclPtr ContextDecl;
 
     /// The list of captures, both explicit and implicit, for this
     /// lambda.
     Capture *Captures = nullptr;
 
     /// The type of the call method.
     TypeSourceInfo *MethodTyInfo;
 
-    LambdaDefinitionData(CXXRecordDecl *D, TypeSourceInfo *Info,
-                         bool Dependent, bool IsGeneric,
-                         LambdaCaptureDefault CaptureDefault)
+    LambdaDefinitionData(CXXRecordDecl *D, TypeSourceInfo *Info, bool Dependent,
+                         bool IsGeneric, LambdaCaptureDefault CaptureDefault)
         : DefinitionData(D), Dependent(Dependent), IsGenericLambda(IsGeneric),
-        CaptureDefault(CaptureDefault), NumCaptures(0), NumExplicitCaptures(0),
+          CaptureDefault(CaptureDefault), NumCaptures(0),
+          NumExplicitCaptures(0), HasKnownInternalLinkage(0), ManglingNumber(0),
           MethodTyInfo(Info) {
       IsLambda = true;
 
       // C++1z [expr.prim.lambda]p4:
       //   This class type is not an aggregate type.
       Aggregate = false;
       PlainOldData = false;
     }
   };

   /// Zero indicates that this closure type has internal linkage, so the
   /// mangling number does not matter, while a non-zero value indicates which
   /// lambda expression this is in this particular context.
   unsigned getLambdaManglingNumber() const {
     assert(isLambda() && "Not a lambda closure type!");
     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.
   ///
   /// Certain contexts (default arguments of in-class function parameters and
   /// the initializers of data members) have separate name mangling rules for
   /// lambdas within the Itanium C++ ABI. For these cases, this routine provides
   /// the declaration in which the lambda occurs, e.g., the function parameter
   /// or the non-static data member. Otherwise, it returns NULL to imply that
   /// the declaration context suffices.
   Decl *getLambdaContextDecl() const;
 
   /// 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.
   MSInheritanceAttr::Spelling getMSInheritanceModel() const;
 
   /// Calculate what the inheritance model would be for this class.
   MSInheritanceAttr::Spelling calculateInheritanceModel() const;
 

clang/include/clang/Sema/Sema.h

public:
 
   /// Create a new lambda closure type.
   CXXRecordDecl *createLambdaClosureType(SourceRange IntroducerRange,
                                          TypeSourceInfo *Info,
                                          bool KnownDependent,
                                          LambdaCaptureDefault CaptureDefault);
 
   /// Start the definition of a lambda expression.
-  CXXMethodDecl *
-  startLambdaDefinition(CXXRecordDecl *Class, SourceRange IntroducerRange,
-                        TypeSourceInfo *MethodType, SourceLocation EndLoc,
+  CXXMethodDecl *startLambdaDefinition(CXXRecordDecl *Class,
+                                       SourceRange IntroducerRange,
+                                       TypeSourceInfo *MethodType,
+                                       SourceLocation EndLoc,
                                        ArrayRef<ParmVarDecl *> Params,
-                        ConstexprSpecKind ConstexprKind,
-                        Optional<std::pair<unsigned, Decl *>> Mangling = None);
+                                       ConstexprSpecKind ConstexprKind);
+
+  /// Number lambda for linkage purposes if necessary.
+  void handleLambdaNumbering(
+      CXXRecordDecl *Class, CXXMethodDecl *Method,
+      Optional<std::tuple<unsigned, bool, Decl *>> Mangling = None);
 
   /// Endow the lambda scope info with the relevant properties.
   void buildLambdaScope(sema::LambdaScopeInfo *LSI,
                         CXXMethodDecl *CallOperator,
                         SourceRange IntroducerRange,
                         LambdaCaptureDefault CaptureDefault,
                         SourceLocation CaptureDefaultLoc,
                         bool ExplicitParams,

clang/lib/AST/ASTImporter.cpp

if (DCXX->isLambda()) {
       if (GetImportedOrCreateSpecialDecl(
               D2CXX, CXXRecordDecl::CreateLambda, D, Importer.getToContext(),
               DC, *TInfoOrErr, Loc, DCXX->isDependentLambda(),
               DCXX->isGenericLambda(), DCXX->getLambdaCaptureDefault()))
         return D2CXX;
       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
       const bool DelayTypeCreation = true;
       if (GetImportedOrCreateDecl(
               D2CXX, D, Importer.getToContext(), D->getTagKind(), DC,
               *BeginLocOrErr, Loc, Name.getAsIdentifierInfo(),
               cast_or_null<CXXRecordDecl>(PrevDecl), DelayTypeCreation))

clang/lib/AST/Decl.cpp

switch (D->getKind()) {
     case Decl::ObjCProperty:
     case Decl::ObjCPropertyImpl:
     case Decl::ObjCProtocol:
       return getExternalLinkageFor(D);
 
     case Decl::CXXRecord: {
       const auto *Record = cast<CXXRecordDecl>(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);
         }
 
         // This lambda has its linkage/visibility determined:
         //  - either by the outermost lambda if that lambda has no mangling
         //    number.
         //  - or by the parent of the outer most lambda
         // This prevents infinite recursion in settings such as nested lambdas
         // used in NSDMI's, for e.g.
         //  struct L {
         //    int t{};
         //    int t2 = ([](int a) { return [](int b) { return b; };})(t)(t);
         //  };
         const CXXRecordDecl *OuterMostLambda =
             getOutermostEnclosingLambda(Record);
-        if (!OuterMostLambda->getLambdaManglingNumber())
+        if (OuterMostLambda->hasKnownLambdaInternalLinkage() ||
+            !OuterMostLambda->getLambdaManglingNumber())
           return getInternalLinkageFor(D);
 
         return getLVForClosure(
                   OuterMostLambda->getDeclContext()->getRedeclContext(),
                   OuterMostLambda->getLambdaContextDecl(), computation);
       }
 
       break;

clang/lib/Sema/SemaLambda.cpp

case Normal: {
     }
 
     return std::make_tuple(nullptr, nullptr);
   }
 
   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;
 
   case DataMember:
     //  -- the in-class initializers of class members
   case DefaultArgument:
     //  -- default arguments appearing in class definitions
   case InlineVariable:
     //  -- the initializers of inline variables
   case VariableTemplate:
     //  -- the initializers of templated variables
     return std::make_tuple(
         &Context.getManglingNumberContext(ASTContext::NeedExtraManglingDecl,
                                           ManglingContextDecl),
         ManglingContextDecl);
   }
 
   llvm_unreachable("unexpected context");
 }
 
-CXXMethodDecl *Sema::startLambdaDefinition(
-    CXXRecordDecl *Class, SourceRange IntroducerRange,
-    TypeSourceInfo *MethodTypeInfo, SourceLocation EndLoc,
-    ArrayRef<ParmVarDecl *> Params, ConstexprSpecKind ConstexprKind,
-    Optional<std::pair<unsigned, Decl *>> Mangling) {
+CXXMethodDecl *Sema::startLambdaDefinition(CXXRecordDecl *Class,
+                                           SourceRange IntroducerRange,
+                                           TypeSourceInfo *MethodTypeInfo,
+                                           SourceLocation EndLoc,
+                                           ArrayRef<ParmVarDecl *> Params,
+                                           ConstexprSpecKind ConstexprKind) {
   QualType MethodType = MethodTypeInfo->getType();
   TemplateParameterList *TemplateParams =
-            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.
   if (Class->isDependentContext() || TemplateParams) {
     const FunctionProtoType *FPT = MethodType->castAs<FunctionProtoType>();
     QualType Result = FPT->getReturnType();
     if (Result->isUndeducedType()) {
       Result = SubstAutoType(Result, Context.DependentTy);
if (!Params.empty()) {
     Method->setParams(Params);
     CheckParmsForFunctionDef(Params,
                              /*CheckParameterNames=*/false);
 
     for (auto P : Method->parameters())
       P->setOwningFunction(Method);
   }
 
+  return Method;
+}
+
+void Sema::handleLambdaNumbering(
+    CXXRecordDecl *Class, CXXMethodDecl *Method,
+    Optional<std::tuple<unsigned, bool, Decl *>> Mangling) {
   if (Mangling) {
-    Class->setLambdaMangling(Mangling->first, Mangling->second);
-  } else {
-    MangleNumberingContext *MCtx;
-    Decl *ManglingContextDecl;
-    std::tie(MCtx, ManglingContextDecl) =
-        getCurrentMangleNumberContext(Class->getDeclContext());
-    if (MCtx) {
-      unsigned ManglingNumber = MCtx->getManglingNumber(Method);
-      Class->setLambdaMangling(ManglingNumber, ManglingContextDecl);
-    }
-  }
+    unsigned ManglingNumber;
+    bool HasKnownInternalLinkage;
+    Decl *ManglingContextDecl;
+    std::tie(ManglingNumber, HasKnownInternalLinkage, ManglingContextDecl) =
+        Mangling.getValue();
+    Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
+                             HasKnownInternalLinkage);
+    return;
+  }
+
+  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<CapturedDecl>(DC))
+      DC = CD->getParent();
+    return &Context.getManglingNumberContext(DC);
+  };
+
+  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;
+  }
+    Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
+                             HasKnownInternalLinkage);
 
-  return Method;
-}
-
 void Sema::buildLambdaScope(LambdaScopeInfo *LSI,
                                         CXXMethodDecl *CallOperator,
                                         SourceRange IntroducerRange,
                                         LambdaCaptureDefault CaptureDefault,
                                         SourceLocation CaptureDefaultLoc,
                                         bool ExplicitParams,
                                         bool ExplicitResultType,
                                         bool Mutable) {
void Sema::ActOnStartOfLambdaDefinition(LambdaIntroducer &Intro,
   // Attributes on the lambda apply to the method.
   ProcessDeclAttributes(CurScope, Method, ParamInfo);
 
   // CUDA lambdas get implicit attributes based on the scope in which they're
   // declared.
   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);
 
   // Build the lambda scope.
   buildLambdaScope(LSI, Method, Intro.Range, Intro.Default, Intro.DefaultLoc,
                    ExplicitParams, ExplicitResultType, !Method->isConst());
 
   // C++11 [expr.prim.lambda]p9:

clang/lib/Sema/TreeTransform.h

TreeTransform<Derived>::TransformLambdaExpr(LambdaExpr *E) {
   CXXRecordDecl *OldClass = E->getLambdaClass();
   CXXRecordDecl *Class
     = getSema().createLambdaClosureType(E->getIntroducerRange(),
                                         NewCallOpTSI,
                                         /*KnownDependent=*/false,
                                         E->getCaptureDefault());
   getDerived().transformedLocalDecl(OldClass, {Class});
 
-  Optional<std::pair<unsigned, Decl*>> Mangling;
+  Optional<std::tuple<unsigned, bool, Decl *>> Mangling;
   if (getDerived().ReplacingOriginal())
-    Mangling = std::make_pair(OldClass->getLambdaManglingNumber(),
+    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<FunctionProtoTypeLoc>().getParams(),
-      E->getCallOperator()->getConstexprKind(), Mangling);
+      E->getCallOperator()->getConstexprKind());
 
   LSI->CallOperator = NewCallOperator;
 
   for (unsigned I = 0, NumParams = NewCallOperator->getNumParams();
        I != NumParams; ++I) {
     auto *P = NewCallOperator->getParamDecl(I);
     if (P->hasUninstantiatedDefaultArg()) {
       EnterExpressionEvaluationContext Eval(
           getSema(),
           Sema::ExpressionEvaluationContext::PotentiallyEvaluatedIfUsed, P);
       ExprResult R = getDerived().TransformExpr(
           E->getCallOperator()->getParamDecl(I)->getDefaultArg());
       P->setDefaultArg(R.get());
     }
   }
 
   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);
 
   // Enter the scope of the lambda.
   getSema().buildLambdaScope(LSI, NewCallOperator,
                              E->getIntroducerRange(),
                              E->getCaptureDefault(),

clang/lib/Serialization/ASTReaderDecl.cpp

if (Data.IsLambda) {
     using Capture = LambdaCapture;
 
     auto &Lambda = static_cast<CXXRecordDecl::LambdaDefinitionData &>(Data);
     Lambda.Dependent = Record.readInt();
     Lambda.IsGenericLambda = Record.readInt();
     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(
         sizeof(Capture) * Lambda.NumCaptures);
     Capture *ToCapture = Lambda.Captures;
     Lambda.MethodTyInfo = GetTypeSourceInfo();
     for (unsigned I = 0, N = Lambda.NumCaptures; I != N; ++I) {
       SourceLocation Loc = ReadSourceLocation();

clang/lib/Serialization/ASTWriter.cpp

void ASTRecordWriter::AddCXXDefinitionData(const CXXRecordDecl *D) {
   // Add lambda-specific data.
   if (Data.IsLambda) {
     auto &Lambda = D->getLambdaData();
     Record->push_back(Lambda.Dependent);
     Record->push_back(Lambda.IsGenericLambda);
     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);
     for (unsigned I = 0, N = Lambda.NumCaptures; I != N; ++I) {
       const LambdaCapture &Capture = Lambda.Captures[I];
       AddSourceLocation(Capture.getLocation());
       Record->push_back(Capture.isImplicit());
       Record->push_back(Capture.getCaptureKind());

clang/test/CodeGenCUDA/unnamed-types.cu

+// 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 <typename F>
+__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