yaxunl updated this revision to Diff 422001.
yaxunl retitled this revision from "[HIP] Fix mangling number for local struct"
to "[CUDA][HIP] Fix mangling number for local struct".
yaxunl edited the summary of this revision.
yaxunl added a comment.
Herald added a subscriber: carlosgalvezp.
Use Itaninium mangling number for mangling device side name
in host compilation only. Keep host name mangling consistent
with C++ programs.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D122734/new/
https://reviews.llvm.org/D122734
Files:
clang/include/clang/AST/ASTContext.h
clang/lib/AST/ASTContext.cpp
clang/lib/AST/MicrosoftCXXABI.cpp
clang/lib/CodeGen/CGCUDANV.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,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();
+}
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/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/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/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();
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
