Index: clang/lib/AST/MicrosoftCXXABI.cpp
===================================================================
--- clang/lib/AST/MicrosoftCXXABI.cpp
+++ clang/lib/AST/MicrosoftCXXABI.cpp
@@ -76,6 +76,11 @@
   unsigned getDeviceManglingNumber(const CXXMethodDecl *CallOperator) override {
     return DeviceCtx->getManglingNumber(CallOperator);
   }
+
+  unsigned getManglingNumber(const TagDecl *TD,
+                             unsigned MSLocalManglingNumber) override {
+    return DeviceCtx->getManglingNumber(TD, MSLocalManglingNumber);
+  }
 };
 
 class MSSYCLNumberingContext : public MicrosoftNumberingContext {
Index: clang/test/CodeGenCUDA/struct-mangling-number.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/struct-mangling-number.cu
@@ -0,0 +1,32 @@
+// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \
+// RUN:   -fms-extensions -aux-triple amdgcn-amd-amdhsa \
+// RUN:   -aux-target-cpu gfx1030 -x hip %s | FileCheck %s
+ 
+#include "Inputs/cuda.h"
+
+// Check local struct 'Op' uses Itanium mangling number instead of MSVC mangling
+// number.
+
+// CHECK: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2OpEvv\00"
+// CHECK-NOT: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2Op_1Evv\00"
+template<typename T>
+__attribute__((global)) void kernel()
+{
+}
+
+template <typename T>
+class Test {
+public:
+  void run()
+  {
+    struct Op 
+    {
+    };
+    kernel<Op><<<1, 1>>>();
+  }
+};
+
+int main() {
+  Test<int> A;
+  A.run();
+}