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 {
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 {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
