Index: clang/include/clang/AST/ASTContext.h
===================================================================
--- clang/include/clang/AST/ASTContext.h
+++ clang/include/clang/AST/ASTContext.h
@@ -677,6 +677,22 @@
     ~CUDAConstantEvalContextRAII() { Ctx.CUDAConstantEvalCtx = SavedCtx; }
   };
 
+  struct CUDANameMangleContext {
+    /// Current name mangling is for device name in host compilation.
+    bool MangleDeviceNameInHostCompilation = false;
+  } CUDANameMangleCtx;
+  struct CUDANameMangleContextRAII {
+    ASTContext &Ctx;
+    CUDANameMangleContext SavedCtx;
+    CUDANameMangleContextRAII(ASTContext &Ctx_,
+                              bool MangleDeviceNameInHostCompilation)
+        : Ctx(Ctx_), SavedCtx(Ctx_.CUDANameMangleCtx) {
+      Ctx_.CUDANameMangleCtx.MangleDeviceNameInHostCompilation =
+          MangleDeviceNameInHostCompilation;
+    }
+    ~CUDANameMangleContextRAII() { Ctx.CUDANameMangleCtx = SavedCtx; }
+  };
+
   /// Returns the dynamic AST node parent map context.
   ParentMapContext &getParentMapContext();
 
Index: clang/lib/AST/ASTContext.cpp
===================================================================
--- clang/lib/AST/ASTContext.cpp
+++ clang/lib/AST/ASTContext.cpp
@@ -11753,7 +11753,14 @@
 
 unsigned ASTContext::getManglingNumber(const NamedDecl *ND) const {
   auto I = MangleNumbers.find(ND);
-  return I != MangleNumbers.end() ? I->second : 1;
+  unsigned Res = I != MangleNumbers.end() ? I->second : 1;
+  if (!LangOpts.CUDA || LangOpts.CUDAIsDevice)
+    return Res;
+
+  auto Cutoff = [](unsigned V) { return V > 1 ? V : 1; };
+  if (CUDANameMangleCtx.MangleDeviceNameInHostCompilation)
+    return Cutoff(Res >> 16);
+  return Cutoff(Res & 0xffff);
 }
 
 void ASTContext::setStaticLocalNumber(const VarDecl *VD, unsigned Number) {
Index: clang/lib/AST/MicrosoftCXXABI.cpp
===================================================================
--- clang/lib/AST/MicrosoftCXXABI.cpp
+++ clang/lib/AST/MicrosoftCXXABI.cpp
@@ -76,6 +76,15 @@
   unsigned getDeviceManglingNumber(const CXXMethodDecl *CallOperator) override {
     return DeviceCtx->getManglingNumber(CallOperator);
   }
+
+  unsigned getManglingNumber(const TagDecl *TD,
+                             unsigned MSLocalManglingNumber) override {
+    unsigned DeviceN = DeviceCtx->getManglingNumber(TD, MSLocalManglingNumber);
+    unsigned HostN =
+        MicrosoftNumberingContext::getManglingNumber(TD, MSLocalManglingNumber);
+    assert(DeviceN <= 0xffff && HostN <= 0xffff);
+    return (DeviceN << 16) | HostN;
+  }
 };
 
 class MSSYCLNumberingContext : public MicrosoftNumberingContext {
Index: clang/lib/CodeGen/CGCUDANV.cpp
===================================================================
--- clang/lib/CodeGen/CGCUDANV.cpp
+++ clang/lib/CodeGen/CGCUDANV.cpp
@@ -260,6 +260,8 @@
 }
 
 std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
+  ASTContext::CUDANameMangleContextRAII X(
+      CGM.getContext(), /*MangleDeviceNameInHostCompilation=*/true);
   GlobalDecl GD;
   // D could be either a kernel or a variable.
   if (auto *FD = dyn_cast<FunctionDecl>(ND))
Index: clang/test/CodeGenCUDA/struct-mangling-number.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/struct-mangling-number.cu
@@ -0,0 +1,63 @@
+// RUN: %clang_cc1 -emit-llvm -o - -aux-triple x86_64-pc-windows-msvc \
+// RUN:   -o %t.dev -fms-extensions -triple amdgcn-amd-amdhsa \
+// RUN:   -target-cpu gfx1030 -fcuda-is-device -x hip %s
+
+// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \
+// RUN:   -o %t.host -fms-extensions -aux-triple amdgcn-amd-amdhsa \
+// RUN:   -aux-target-cpu gfx1030 -x hip %s
+
+// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \
+// RUN:   -o %t.as_cpp -fms-extensions -x c++ %s
+
+// RUN: cat %t.dev %t.host | FileCheck %s
+
+// RUN: cat %t.host %t.as_cpp | FileCheck -check-prefix=CPP %s
+
+#if __HIP__
+#include "Inputs/cuda.h"
+#endif
+
+// Check local struct 'Op' uses Itanium mangling number instead of MSVC mangling
+// number in device side name mangling. It is the same in device and host
+// compilation.
+
+// CHECK: define amdgpu_kernel void @[[KERN:_Z6kernelIZN4TestIiE3runEvE2OpEvv]](
+// CHECK: @{{.*}} = {{.*}}c"[[KERN]]\00"
+
+// CHECK-NOT: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2Op_1Evv\00"
+#if __HIP__
+template<typename T>
+__attribute__((global)) void kernel()
+{
+}
+#endif
+
+// Check local struct 'Op' uses MSVC mangling number in host function name mangling.
+// It is the same when compiled as HIP or C++ program.
+
+// CPP: call void @[[FUN:"\?\?\$fun@UOp@\?2\?\?run@\?\$Test@H@@QEAAXXZ@@@YAXXZ"]]()
+// CPP: call void @[[FUN]]()
+template<typename T>
+void fun()
+{
+}
+
+template <typename T>
+class Test {
+public:
+  void run()
+  {
+    struct Op
+    {
+    };
+#if __HIP__
+    kernel<Op><<<1, 1>>>();
+#endif
+    fun<Op>();
+  }
+};
+
+int main() {
+  Test<int> A;
+  A.run();
+}