[PATCH] D122734: [HIP] Fix mangling number for local struct

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
Herald added a project: All.
yaxunl requested review of this revision.

MSVC and Itanium mangling use different mangling numbers
for function-scope structs, which causes inconsistent
mangled kernel names in device and host compilations.

This patch uses Itanium mangling number for structs
in HIP host compilation on Windows to fix this issue.


https://reviews.llvm.org/D122734

Files:
  clang/lib/AST/MicrosoftCXXABI.cpp
  clang/test/CodeGenCUDA/struct-mangling-number.cu


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();
+}
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 {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D122734.419152.patch
Type: text/x-patch
Size: 1531 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20220330/bcb1e961/attachment.bin>