[clang] 11d3e31 - [CUDA][HIP] Fix mangling number for local struct
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 28 16:55:24 PDT 2022
Author: Yaxun (Sam) Liu
Date: 2022-04-28T19:54:43-04:00
New Revision: 11d3e31c60bdc9e491c51b97a964b6289575edfa
URL: https://github.com/llvm/llvm-project/commit/11d3e31c60bdc9e491c51b97a964b6289575edfa
DIFF: https://github.com/llvm/llvm-project/commit/11d3e31c60bdc9e491c51b97a964b6289575edfa.diff
LOG: [CUDA][HIP] Fix mangling number for local struct
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 for mangling device side names in CUDA/HIP host
compilation on Windows to fix this issue.
A state is added to ASTContext to indicate whether the
current name mangling is for device side names in host
compilation. Device and host mangling number
are encoded/decoded as upper and lower half of 32 bit
unsigned integer to fit into the original mangling number
field for AST. Diagnostic will be emitted if a manglining
number exceeds limit.
Reviewed by: Artem Belevich, Reid Kleckner
Differential Revision: https://reviews.llvm.org/D122734
Fixes: SWDEV-328515
Added:
clang/test/CodeGenCUDA/struct-mangling-number.cu
Modified:
clang/include/clang/AST/ASTContext.h
clang/lib/AST/ASTContext.cpp
clang/lib/AST/MicrosoftCXXABI.cpp
clang/lib/CodeGen/CGCUDANV.cpp
Removed:
################################################################################
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index 03d6a0fbe6eea..e04766f93d69e 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -677,6 +677,9 @@ class ASTContext : public RefCountedBase<ASTContext> {
~CUDAConstantEvalContextRAII() { Ctx.CUDAConstantEvalCtx = SavedCtx; }
};
+ /// Current CUDA name mangling is for device name in host compilation.
+ bool CUDAMangleDeviceNameInHostCompilation = false;
+
/// Returns the dynamic AST node parent map context.
ParentMapContext &getParentMapContext();
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 85d2bcf268f3a..c91839c50e95c 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -11762,7 +11762,14 @@ void ASTContext::setManglingNumber(const NamedDecl *ND, unsigned Number) {
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;
+
+ // CUDA/HIP host compilation encodes host and device mangling numbers
+ // as lower and upper half of 32 bit integer.
+ Res = CUDAMangleDeviceNameInHostCompilation ? Res >> 16 : Res & 0xFFFF;
+ return Res > 1 ? Res : 1;
}
void ASTContext::setStaticLocalNumber(const VarDecl *VD, unsigned Number) {
diff --git a/clang/lib/AST/MicrosoftCXXABI.cpp b/clang/lib/AST/MicrosoftCXXABI.cpp
index b7dc0e62e66af..b3e26bc889c1a 100644
--- a/clang/lib/AST/MicrosoftCXXABI.cpp
+++ b/clang/lib/AST/MicrosoftCXXABI.cpp
@@ -76,6 +76,20 @@ class MSHIPNumberingContext : public MicrosoftNumberingContext {
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);
+ if (DeviceN > 0xFFFF || HostN > 0xFFFF) {
+ DiagnosticsEngine &Diags = TD->getASTContext().getDiagnostics();
+ unsigned DiagID = Diags.getCustomDiagID(
+ DiagnosticsEngine::Error, "Mangling number exceeds limit (65535)");
+ Diags.Report(TD->getLocation(), DiagID);
+ }
+ return (DeviceN << 16) | HostN;
+ }
};
class MSSYCLNumberingContext : public MicrosoftNumberingContext {
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 4390228297d0e..961f7e39c8f6a 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -24,6 +24,7 @@
#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/ReplaceConstant.h"
#include "llvm/Support/Format.h"
+#include "llvm/Support/SaveAndRestore.h"
using namespace clang;
using namespace CodeGen;
@@ -260,6 +261,8 @@ llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
}
std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
+ llvm::SaveAndRestore<bool> MangleAsDevice(
+ CGM.getContext().CUDAMangleDeviceNameInHostCompilation, true);
GlobalDecl GD;
// D could be either a kernel or a variable.
if (auto *FD = dyn_cast<FunctionDecl>(ND))
diff --git a/clang/test/CodeGenCUDA/struct-mangling-number.cu b/clang/test/CodeGenCUDA/struct-mangling-number.cu
new file mode 100644
index 0000000000000..e9e80671c28c0
--- /dev/null
+++ b/clang/test/CodeGenCUDA/struct-mangling-number.cu
@@ -0,0 +1,68 @@
+// RUN: %clang_cc1 -emit-llvm -o - -aux-triple x86_64-pc-windows-msvc \
+// RUN: -fms-extensions -triple amdgcn-amd-amdhsa \
+// RUN: -target-cpu gfx1030 -fcuda-is-device -x hip %s \
+// RUN: | FileCheck -check-prefix=DEV %s
+
+// 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 \
+// RUN: | FileCheck -check-prefix=HOST %s
+
+// 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 \
+// RUN: | FileCheck -check-prefix=HOST-NEG %s
+
+// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \
+// RUN: -fms-extensions -x c++ %s \
+// RUN: | 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.
+
+// DEV: define amdgpu_kernel void @_Z6kernelIZN4TestIiE3runEvE2OpEvv(
+
+// HOST-DAG: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2OpEvv\00"
+
+// HOST-NEG-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.
+
+// HOST-DAG: call void @"??$fun at UOp@?2??run@?$Test at H@@QEAAXXZ@@@YAXXZ"()
+// CPP: call void @"??$fun at UOp@?2??run@?$Test at H@@QEAAXXZ@@@YAXXZ"()
+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();
+}
More information about the cfe-commits
mailing list