[hip][cuda] Enable extended lambda support on Windows.

[hliao]

Summary

• On Windows, extended lambda has extra issues due to the numbering schemes are different between the host compilation (Microsoft C++ ABI) and the device compilation (Itanium C++ ABI. Additional device side lambda number is required per lambda for the host compilation to correctly mangle the device-side lambda name.
• A hybrid numbering context MSHIPNumberingContext is introduced to number a lambda for both host- and device-compilations.

[hliao]

PING for review

[rjmccall]

Please don't ping for review after a single day.

[hliao]

PING for review

[tra]

@rnk Reid, would you be the right person to look at the change on the Windows' side?

[rnk]

clang/include/clang/AST/MangleNumberingContext.h

56–57

It would be nicer if there were a single entry point that does the right thing for all mangling contexts.

clang/lib/AST/MicrosoftCXXABI.cpp

68–69

I know it's inheritance instead of composition, but maybe the simple thing to do here would be to inherit from the MicrosoftMangleNumberingContext, and then implement an override for getDeviceManglingNumber that delegates to the Itanium one.

90

This is actually the device's mangling context, and it's an Itanium mangling context. I think it needs comments and a more descriptive variable name.

[hliao]

revise MSHIPNumberingContext

[hliao]

clang/include/clang/AST/MangleNumberingContext.h

56–57

could you elaborate in more details?

[rnk]

clang/include/clang/AST/DeclCXX.h

394–398

It seems a shame to grow LambdaDefinitionData by a pointer for all users of C++ that do not use CUDA. Optimizing bitfields may be worth the time, but I'll leave it to @rjmccall or @rsmith to give guidance on whether that's worth it.

An alternative would be to store the device numbers in the mangling context and look them up when needed, since they are so rarely needed.

clang/include/clang/AST/MangleNumberingContext.h

56–57

I'll add comments at the call site.

clang/lib/Sema/SemaLambda.cpp

477–479

Rather than having a virtual method that returns bool, call a virtual method that does the entire thing you are trying to do. For example, MSHIP* could be responsible for setting the mangling number on the class. Although, maybe it should be storing the number in some side table now that I think about it.

[hliao]

simplify again following suggestion.

[hliao]

clang/include/clang/AST/DeclCXX.h

394–398

I like the alternative way by storing all numbering into the mangle/numbering context instead of AST itself. But, it needs additional numbering post-processing after AST importing. Sound to me a major refactoring work likely to be addressed later.

[hliao]

kindly PING for review

[hliao]

Looks like holidays are approaching, :). Anyway, it's really appreciated that you could review this patch to enable CUDA/HIP applications on Windows.

Thanks.

[hliao]

PING for review

[hliao]

PING for review

[tra]

@hliao -- Can you take a look at https://bugs.llvm.org/show_bug.cgi?id=48866. This patch may be relevant there.

[tra]

In D69322#2523058, @tra wrote:

@hliao -- Can you take a look at https://bugs.llvm.org/show_bug.cgi?id=48866. This patch may be relevant there.

@rnk Reid, looks like this patch does fix a lambda mangling issue in CUDA on Windows. Could you take another look at it?

[rnk]

clang/include/clang/AST/DeclCXX.h

394–398

Generally, I don't think we can count on contributors to come back later and optimize memory usage, so it seems reasonable to ask to avoid the regression in the first place. Above I see the bitfield usage optimizing memory usage of the number of captures, and then here we spent lots of memory storing device mangling numbers that are only used for CUDA. I think the memory usage concern still stands. I don't think it's unreasonable to maintain these numbers on the side in a DenseMap.

@rsmith or @rjmccall, do you agree with that reasoning?

[rjmccall]

I agree that we should aim to avoid broad memory impact for uncommon features, and this seems to apply.

clang/include/clang/AST/DeclCXX.h

394–398

I think that's the right idea in general. I don't know how much it really matters for LambdaData, because there aren't a huge number of lambdas in typical translation units. But if you want to optimize this, a hashtable in the ASTContext or ManglingContext seems reasonable.

[hliao]

Rebase to the main branch.

[hliao]

• Remove DeviceManglingNumber to reduce memory usage.
• Add a table in ASTContext for device lambda mangling numbers if that lamba requires.

[hliao]

Revise comments.

[hliao]

PING for review

[hliao]

@rsmith, @rnk and @rjmccall, could you review the latest change addressing the memory usage issue?

[rnk]

Looks good with one suggestion.

clang/include/clang/AST/ASTContext.h

540

This is a nit, but I'd group the device mangling number map here with the other mangling number side tables.

[thakis]

Hi, a whole bunch of cuda-related tests failed on windows with this in (http://45.33.8.238/win/32620/summary.html). Did you run check-clang before committing? I reverted this for now in 4874ff02417916cc9ff994b34abcb5e563056546 . After the revert, all tests pass again.

clang/include/clang/AST/ASTContext.h

private:
   llvm::DenseMap<const DeclContext *, std::unique_ptr<MangleNumberingContext>>
       MangleNumberingContexts;
   llvm::DenseMap<const Decl *, std::unique_ptr<MangleNumberingContext>>
       ExtraMangleNumberingContexts;
 
   /// Side-table of mangling numbers for declarations which rarely
   /// need them (like static local vars).
   llvm::MapVector<const NamedDecl *, unsigned> MangleNumbers;
   llvm::MapVector<const VarDecl *, unsigned> StaticLocalNumbers;
+  /// Mapping the associated device lambda mangling number if present.
+  mutable llvm::DenseMap<const CXXRecordDecl *, unsigned>
+      DeviceLambdaManglingNumbers;
 
   /// Mapping that stores parameterIndex values for ParmVarDecls when
   /// that value exceeds the bitfield size of ParmVarDeclBits.ParameterIndex.
   using ParameterIndexTable = llvm::DenseMap<const VarDecl *, unsigned>;
   ParameterIndexTable ParamIndices;
 
   ImportDecl *FirstLocalImport = nullptr;
   ImportDecl *LastLocalImport = nullptr;

clang/include/clang/AST/DeclCXX.h

struct LambdaDefinitionData : public DefinitionData {
     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 : 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;

   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;
   }
 
+  /// Set the device side mangling number.
+  void setDeviceLambdaManglingNumber(unsigned Num) const;
+
+  /// Retrieve the device side mangling number.
+  unsigned getDeviceLambdaManglingNumber() const;
+
   /// Returns the inheritance model used for this record.
   MSInheritanceModel getMSInheritanceModel() const;
 
   /// Calculate what the inheritance model would be for this class.
   MSInheritanceModel calculateInheritanceModel() const;
 
   /// In the Microsoft C++ ABI, use zero for the field offset of a null data
   /// member pointer if we can guarantee that zero is not a valid field offset,

clang/include/clang/AST/Mangle.h

public:
 
   /// @name Mangler Entry Points
   /// @{
 
   bool shouldMangleDeclName(const NamedDecl *D);
   virtual bool shouldMangleCXXName(const NamedDecl *D) = 0;
   virtual bool shouldMangleStringLiteral(const StringLiteral *SL) = 0;
 
+  virtual bool isDeviceMangleContext() const { return false; }
+  virtual void setDeviceMangleContext(bool) {}
+
   // FIXME: consider replacing raw_ostream & with something like SmallString &.
   void mangleName(GlobalDecl GD, raw_ostream &);
   virtual void mangleCXXName(GlobalDecl GD, raw_ostream &) = 0;
   virtual void mangleThunk(const CXXMethodDecl *MD,
                           const ThunkInfo &Thunk,
                           raw_ostream &) = 0;
   virtual void mangleCXXDtorThunk(const CXXDestructorDecl *DD, CXXDtorType Type,
                                   const ThisAdjustment &ThisAdjustment,

clang/include/clang/AST/MangleNumberingContext.h

public:
   /// this context.
   virtual unsigned getManglingNumber(const VarDecl *VD,
                                      unsigned MSLocalManglingNumber) = 0;
 
   /// Retrieve the mangling number of a static local variable within
   /// this context.
   virtual unsigned getManglingNumber(const TagDecl *TD,
                                      unsigned MSLocalManglingNumber) = 0;
+
+  /// Retrieve the mangling number of a new lambda expression with the
+  /// given call operator within the device context. No device number is
+  /// assigned if there's no device numbering context is associated.
+  virtual unsigned getDeviceManglingNumber(const CXXMethodDecl *) { return 0; }
 };
 
 } // end namespace clang
 #endif

clang/include/clang/Sema/Sema.h

CXXMethodDecl *startLambdaDefinition(CXXRecordDecl *Class,
                                        SourceLocation EndLoc,
                                        ArrayRef<ParmVarDecl *> Params,
                                        ConstexprSpecKind ConstexprKind,
                                        Expr *TrailingRequiresClause);
 
   /// Number lambda for linkage purposes if necessary.
   void handleLambdaNumbering(
       CXXRecordDecl *Class, CXXMethodDecl *Method,
-      Optional<std::tuple<unsigned, bool, Decl *>> Mangling = None);
+      Optional<std::tuple<bool, unsigned, unsigned, 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()) {
               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,
                                DCXX->hasKnownLambdaInternalLinkage());
+      D2CXX->setDeviceLambdaManglingNumber(
+          DCXX->getDeviceLambdaManglingNumber());
    } 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/CXXABI.h

 #include "clang/AST/Type.h"
 
 namespace clang {
 
 class ASTContext;
 class CXXConstructorDecl;
 class DeclaratorDecl;
 class Expr;
-class MemberPointerType;
+class MangleContext;
 class MangleNumberingContext;
+class MemberPointerType;
 
 /// Implements C++ ABI-specific semantic analysis functions.
 class CXXABI {
 public:
   virtual ~CXXABI();
 
   struct MemberPointerInfo {
     uint64_t Width;
virtual void addDeclaratorForUnnamedTagDecl(TagDecl *TD,
                                               DeclaratorDecl *DD) = 0;
 
   virtual DeclaratorDecl *getDeclaratorForUnnamedTagDecl(const TagDecl *TD) = 0;
 };
 
 /// Creates an instance of a C++ ABI class.
 CXXABI *CreateItaniumCXXABI(ASTContext &Ctx);
 CXXABI *CreateMicrosoftCXXABI(ASTContext &Ctx);
+std::unique_ptr<MangleNumberingContext>
+createItaniumNumberingContext(MangleContext *);
 }
 
 #endif

clang/lib/AST/DeclCXX.cpp

 }
 
 Decl *CXXRecordDecl::getLambdaContextDecl() const {
   assert(isLambda() && "Not a lambda closure type!");
   ExternalASTSource *Source = getParentASTContext().getExternalSource();
   return getLambdaData().ContextDecl.get(Source);
 }
 
+void CXXRecordDecl::setDeviceLambdaManglingNumber(unsigned Num) const {
+  assert(isLambda() && "Not a lambda closure type!");
+  if (Num)
+    getASTContext().DeviceLambdaManglingNumbers[this] = Num;
+}
+
+unsigned CXXRecordDecl::getDeviceLambdaManglingNumber() const {
+  assert(isLambda() && "Not a lambda closure type!");
+  auto I = getASTContext().DeviceLambdaManglingNumbers.find(this);
+  if (I != getASTContext().DeviceLambdaManglingNumbers.end())
+    return I->second;
+  return 0;
+}
+
 static CanQualType GetConversionType(ASTContext &Context, NamedDecl *Conv) {
   QualType T =
       cast<CXXConversionDecl>(Conv->getUnderlyingDecl()->getAsFunction())
           ->getConversionType();
   return Context.getCanonicalType(T);
 }
 
 /// Collect the visible conversions of a base class.

clang/lib/AST/ItaniumCXXABI.cpp

return std::make_unique<ItaniumNumberingContext>(
         cast<ItaniumMangleContext>(Mangler.get()));
   }
 };
 }
 
 CXXABI *clang::CreateItaniumCXXABI(ASTContext &Ctx) {
   return new ItaniumCXXABI(Ctx);
 }
+
+std::unique_ptr<MangleNumberingContext>
+clang::createItaniumNumberingContext(MangleContext *Mangler) {
+  return std::make_unique<ItaniumNumberingContext>(
+      cast<ItaniumMangleContext>(Mangler));
+}

clang/lib/AST/ItaniumMangle.cpp

 
 static const unsigned UnknownArity = ~0U;
 
 class ItaniumMangleContextImpl : public ItaniumMangleContext {
   typedef std::pair<const DeclContext*, IdentifierInfo*> DiscriminatorKeyTy;
   llvm::DenseMap<DiscriminatorKeyTy, unsigned> Discriminator;
   llvm::DenseMap<const NamedDecl*, unsigned> Uniquifier;
 
+  bool IsDevCtx = false;
+
 public:
   explicit ItaniumMangleContextImpl(ASTContext &Context,
                                     DiagnosticsEngine &Diags)
       : ItaniumMangleContext(Context, Diags) {}
 
   /// @name Mangler Entry Points
   /// @{
 
   bool shouldMangleCXXName(const NamedDecl *D) override;
   bool shouldMangleStringLiteral(const StringLiteral *) override {
     return false;
   }
+
+  bool isDeviceMangleContext() const override { return IsDevCtx; }
+  void setDeviceMangleContext(bool IsDev) override { IsDevCtx = IsDev; }
+
   void mangleCXXName(GlobalDecl GD, raw_ostream &) override;
   void mangleThunk(const CXXMethodDecl *MD, const ThunkInfo &Thunk,
                    raw_ostream &) override;
   void mangleCXXDtorThunk(const CXXDestructorDecl *DD, CXXDtorType Type,
                           const ThisAdjustment &ThisAdjustment,
                           raw_ostream &) override;
   void mangleReferenceTemporary(const VarDecl *D, unsigned ManglingNumber,
                                 raw_ostream &) override;
void CXXNameMangler::mangleLambda(const CXXRecordDecl *Lambda) {
   mangleLambdaSig(Lambda);
   Out << "E";
 
   // The number is omitted for the first closure type with a given
   // <lambda-sig> in a given context; it is n-2 for the nth closure type
   // (in lexical order) with that same <lambda-sig> and context.
   //
   // The AST keeps track of the number for us.
-  unsigned Number = Lambda->getLambdaManglingNumber();
+  //
+  // In CUDA/HIP, to ensure the consistent lamba numbering between the device-
+  // and host-side compilations, an extra device mangle context may be created
+  // if the host-side CXX ABI has different numbering for lambda. In such case,
+  // if the mangle context is that device-side one, use the device-side lambda
+  // mangling number for this lambda.
+  unsigned Number = Context.isDeviceMangleContext()
+                        ? Lambda->getDeviceLambdaManglingNumber()
+                        : Lambda->getLambdaManglingNumber();
   assert(Number > 0 && "Lambda should be mangled as an unnamed class");
   if (Number > 1)
     mangleNumber(Number - 2);
   Out << '_';
 }
 
 void CXXNameMangler::mangleLambdaSig(const CXXRecordDecl *Lambda) {
   for (auto *D : Lambda->getLambdaExplicitTemplateParameters())

clang/lib/AST/MicrosoftCXXABI.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "CXXABI.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/Attr.h"
 #include "clang/AST/CXXInheritance.h"
 #include "clang/AST/DeclCXX.h"
+#include "clang/AST/Mangle.h"
 #include "clang/AST/MangleNumberingContext.h"
 #include "clang/AST/RecordLayout.h"
 #include "clang/AST/Type.h"
 #include "clang/Basic/TargetInfo.h"
 
 using namespace clang;
 
 namespace {
public:
   }
 
   unsigned getManglingNumber(const TagDecl *TD,
                              unsigned MSLocalManglingNumber) override {
     return MSLocalManglingNumber;
   }
 };
 
+class MSHIPNumberingContext : public MicrosoftNumberingContext {
+  std::unique_ptr<MangleNumberingContext> DeviceCtx;
+
+public:
+  MSHIPNumberingContext(MangleContext *Mangler) {
+    DeviceCtx = createItaniumNumberingContext(Mangler);
+  }
+
+  unsigned getDeviceManglingNumber(const CXXMethodDecl *CallOperator) override {
+    return DeviceCtx->getManglingNumber(CallOperator);
+  }
+};
+
 class MicrosoftCXXABI : public CXXABI {
   ASTContext &Context;
   llvm::SmallDenseMap<CXXRecordDecl *, CXXConstructorDecl *> RecordToCopyCtor;
 
   llvm::SmallDenseMap<TagDecl *, DeclaratorDecl *>
       UnnamedTagDeclToDeclaratorDecl;
   llvm::SmallDenseMap<TagDecl *, TypedefNameDecl *>
       UnnamedTagDeclToTypedefNameDecl;
 
+  // MangleContext for device numbering context, which is based on Itanium C++
+  // ABI.
+  std::unique_ptr<MangleContext> Mangler;
+
 public:
-  MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) { }
+  MicrosoftCXXABI(ASTContext &Ctx) : Context(Ctx) {
+    if (Context.getLangOpts().CUDA) {
+      assert(Context.getTargetInfo().getCXXABI().isMicrosoft() &&
+             Context.getAuxTargetInfo()->getCXXABI().isItaniumFamily() &&
+             "Unexpected combination of C++ ABIs.");
+      Mangler.reset(Context.createMangleContext(Context.getAuxTargetInfo()));
+    }
+  }
 
   MemberPointerInfo
   getMemberPointerInfo(const MemberPointerType *MPT) const override;
 
   CallingConv getDefaultMethodCallConv(bool isVariadic) const override {
     if (!isVariadic &&
         Context.getTargetInfo().getTriple().getArch() == llvm::Triple::x86)
       return CC_X86ThisCall;
public:
 
   DeclaratorDecl *getDeclaratorForUnnamedTagDecl(const TagDecl *TD) override {
     return UnnamedTagDeclToDeclaratorDecl.lookup(
         const_cast<TagDecl *>(TD->getCanonicalDecl()));
   }
 
   std::unique_ptr<MangleNumberingContext>
   createMangleNumberingContext() const override {
+    if (Context.getLangOpts().CUDA)
+      return std::make_unique<MSHIPNumberingContext>(Mangler.get());
     return std::make_unique<MicrosoftNumberingContext>();
   }
 };
 }
 
 // getNumBases() seems to only give us the number of direct bases, and not the
 // total.  This function tells us if we inherit from anybody that uses MI, or if
 // we have a non-primary base class, which uses the multiple inheritance model.
if (Target.getTriple().isArch64Bit()) {
     MPI.HasPadding = MPI.Width != (Ptrs * PtrSize + Ints * IntSize);
   }
   return MPI;
 }
 
 CXXABI *clang::CreateMicrosoftCXXABI(ASTContext &Ctx) {
   return new MicrosoftCXXABI(Ctx);
 }
-

clang/lib/CodeGen/CGCUDANV.cpp

CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
 
   IntTy = CGM.IntTy;
   SizeTy = CGM.SizeTy;
   VoidTy = CGM.VoidTy;
 
   CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
   VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
   VoidPtrPtrTy = VoidPtrTy->getPointerTo();
+  // If the host and device have different C++ ABIs, mark it as the device
+  // mangle context so that the mangling needs to retrieve the additonal device
+  // lambda mangling number instead of the regular host one.
+  DeviceMC->setDeviceMangleContext(
+      CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
+      CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily());
 }
 
 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
   // cudaError_t cudaSetupArgument(void *, size_t, size_t)
   llvm::Type *Params[] = {VoidPtrTy, SizeTy, SizeTy};
   return CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, Params, false),
       addPrefixToName("SetupArgument"));

clang/lib/Sema/SemaLambda.cpp

for (auto P : Method->parameters())
       P->setOwningFunction(Method);
   }
 
   return Method;
 }
 
 void Sema::handleLambdaNumbering(
     CXXRecordDecl *Class, CXXMethodDecl *Method,
-    Optional<std::tuple<unsigned, bool, Decl *>> Mangling) {
+    Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling) {
   if (Mangling) {
-    unsigned ManglingNumber;
     bool HasKnownInternalLinkage;
+    unsigned ManglingNumber, DeviceManglingNumber;
     Decl *ManglingContextDecl;
-    std::tie(ManglingNumber, HasKnownInternalLinkage, ManglingContextDecl) =
-        Mangling.getValue();
+    std::tie(HasKnownInternalLinkage, ManglingNumber, DeviceManglingNumber,
+             ManglingContextDecl) = Mangling.getValue();
     Class->setLambdaMangling(ManglingNumber, ManglingContextDecl,
                              HasKnownInternalLinkage);
+    Class->setDeviceLambdaManglingNumber(DeviceManglingNumber);
     return;
   }
 
   auto getMangleNumberingContext =
       [this](CXXRecordDecl *Class,
              Decl *ManglingContextDecl) -> MangleNumberingContext * {
     // Get mangle numbering context if there's any extra decl context.
     if (ManglingContextDecl)
if (!MCtx && getLangOpts().CUDA) {
     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);
+    Class->setDeviceLambdaManglingNumber(MCtx->getDeviceManglingNumber(Method));
   }
 }
 
 void Sema::buildLambdaScope(LambdaScopeInfo *LSI,
                                         CXXMethodDecl *CallOperator,
                                         SourceRange IntroducerRange,
                                         LambdaCaptureDefault CaptureDefault,
                                         SourceLocation CaptureDefaultLoc,
                                         bool ExplicitParams,
                                         bool ExplicitResultType,

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::tuple<unsigned, bool, Decl *>> Mangling;
+  Optional<std::tuple<bool, unsigned, unsigned, Decl *>> Mangling;
   if (getDerived().ReplacingOriginal())
-    Mangling = std::make_tuple(OldClass->getLambdaManglingNumber(),
-                               OldClass->hasKnownLambdaInternalLinkage(),
+    Mangling = std::make_tuple(OldClass->hasKnownLambdaInternalLinkage(),
+                               OldClass->getLambdaManglingNumber(),
+                               OldClass->getDeviceLambdaManglingNumber(),
                                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(),

clang/lib/Serialization/ASTReaderDecl.cpp

if (Data.IsLambda) {
     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();
+    D->setDeviceLambdaManglingNumber(Record.readInt());
     Lambda.ContextDecl = readDeclID();
     Lambda.Captures = (Capture *)Reader.getContext().Allocate(
         sizeof(Capture) * Lambda.NumCaptures);
     Capture *ToCapture = Lambda.Captures;
     Lambda.MethodTyInfo = readTypeSourceInfo();
     for (unsigned I = 0, N = Lambda.NumCaptures; I != N; ++I) {
       SourceLocation Loc = readSourceLocation();
       bool IsImplicit = Record.readInt();

clang/lib/Serialization/ASTWriter.cpp

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);
+    Record->push_back(D->getDeviceLambdaManglingNumber());
     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());
       switch (Capture.getCaptureKind()) {

clang/test/CodeGenCUDA/ms-linker-options.cu

 // RUN: %clang_cc1 -emit-llvm -o - -fcuda-is-device -fms-extensions -x hip %s \
 // RUN:   -fno-autolink -triple amdgcn-amd-amdhsa \
 // RUN:   | FileCheck -check-prefix=DEV %s
 // RUN: %clang_cc1 -emit-llvm -o - -fms-extensions -x hip %s -triple \
-// RUN:    x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s
+// RUN:    x86_64-pc-windows-msvc -aux-triple amdgcn | FileCheck -check-prefix=HOST %s
 // RUN: %clang_cc1 -emit-llvm -o - -fcuda-is-device -fms-extensions %s \
 // RUN:   -fno-autolink -triple amdgcn-amd-amdhsa \
 // RUN:   | FileCheck -check-prefix=DEV %s
 // RUN: %clang_cc1 -emit-llvm -o - -fms-extensions %s -triple \
-// RUN:    x86_64-pc-windows-msvc | FileCheck -check-prefix=HOST %s
+// RUN:    x86_64-pc-windows-msvc -aux-triple amdgcn | FileCheck -check-prefix=HOST %s
 
 // DEV-NOT: llvm.linker.options
 // DEV-NOT: llvm.dependent-libraries
 // HOST: lvm.linker.options
 // HOST: "/DEFAULTLIB:libcpmt.lib"
 // HOST: "/FAILIFMISMATCH:\22myLib_version=9\22"
 
 #pragma comment(lib, "libcpmt")
 #pragma detect_mismatch("myLib_version", "9")

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 x86_64-pc-windows-msvc -aux-triple amdgcn-amd-amdhsa -emit-llvm %s -o - | FileCheck %s --check-prefix=MSVC
 // 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
+// HOST: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
+// Check that, on MSVC, the same device kernel mangling name is generated.
+// MSVC: @0 = private unnamed_addr constant [43 x i8] c"_Z2k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_\00", align 1
+// MSVC: @1 = private unnamed_addr constant [60 x i8] c"_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_\00", align 1
 
 __device__ float d0(float x) {
-  return [](float x) { return x + 2.f; }(x);
+  return [](float x) { return x + 1.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_(
+// DEVICE: define internal float @_ZZZ2f1PfENKUlS_E_clES_ENKUlfE_clEf(
 template <typename F>
 __global__ void k0(float *p, F f) {
   p[0] = f(p[0]) + d0(p[1]) + d1(p[2]);
 }
 
+// DEVICE: amdgpu_kernel void @_Z2k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_(
+// DEVICE: define internal float @_ZZ2f1PfENKUlfE_clEf(
+// DEVICE: define internal float @_ZZ2f1PfENKUlffE_clEff(
+// DEVICE: define internal float @_ZZ2f1PfENKUlfE0_clEf(
+template <typename F0, typename F1, typename F2>
+__global__ void k1(float *p, F0 f0, F1 f1, F2 f2) {
+  p[0] = f0(p[0]) + f1(p[1], p[2]) + f2(p[3]);
+}
+
 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; });
+    k0<<<1,1>>>(p, [] __device__ (float x) { return x + 3.f; });
   }(p);
+  k1<<<1,1>>>(p,
+              [] __device__ (float x) { return x + 4.f; },
+              [] __device__ (float x, float y) { return x * y; },
+              [] __device__ (float x) { return x + 5.f; });
 }
 // HOST: @__hip_register_globals
 // HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k0IZZ2f1PfENKUlS0_E_clES0_EUlfE_EvS0_T_{{.*}}@0
+// HOST: __hipRegisterFunction{{.*}}@_Z17__device_stub__k1IZ2f1PfEUlfE_Z2f1S0_EUlffE_Z2f1S0_EUlfE0_EvS0_T_T0_T1_{{.*}}@1
+// MSVC: __hipRegisterFunction{{.*}}@"??$k0@V<lambda_1>@?0???R1?0??f1@@YAXPEAM@Z@QEBA@0@Z@@@YAXPEAMV<lambda_1>@?0???R0?0??f1@@YAX0@Z@QEBA@0@Z@@Z{{.*}}@0
+// MSVC: __hipRegisterFunction{{.*}}@"??$k1@V<lambda_2>@?0??f1@@YAXPEAM@Z@V<lambda_3>@?0??2@YAX0@Z@V<lambda_4>@?0??2@YAX0@Z@@@YAXPEAMV<lambda_2>@?0??f1@@YAX0@Z@V<lambda_3>@?0??1@YAX0@Z@V<lambda_4>@?0??1@YAX0@Z@@Z{{.*}}@1