[PATCH] D63335: [HIP] Change kernel stub name again
Michael Liao via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Fri Jun 14 08:25:33 PDT 2019
hliao created this revision.
hliao added reviewers: yaxunl, tra.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
- Prefix kernel stub with `__device_stub__` to avoid potential symbol name conflicts in debugger.
- Revise the interface to derive the stub name and simplify the assertion of it.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D63335
Files:
clang/lib/CodeGen/CGCUDANV.cpp
clang/lib/CodeGen/CGCUDARuntime.h
clang/lib/CodeGen/CodeGenModule.cpp
clang/test/CodeGenCUDA/kernel-stub-name.cu
Index: clang/test/CodeGenCUDA/kernel-stub-name.cu
===================================================================
--- clang/test/CodeGenCUDA/kernel-stub-name.cu
+++ clang/test/CodeGenCUDA/kernel-stub-name.cu
@@ -10,7 +10,7 @@
__global__ void kernelfunc() {}
// CHECK-LABEL: define{{.*}}@_Z8hostfuncv()
-// CHECK: call void @[[STUB:_Z10kernelfuncIiEvv.stub]]()
+// CHECK: call void @[[STUB:__device_stub___Z10kernelfuncIiEvv]]()
void hostfunc(void) { kernelfunc<int><<<1, 1>>>(); }
// CHECK: define{{.*}}@[[STUB]]
Index: clang/lib/CodeGen/CodeGenModule.cpp
===================================================================
--- clang/lib/CodeGen/CodeGenModule.cpp
+++ clang/lib/CodeGen/CodeGenModule.cpp
@@ -1088,13 +1088,10 @@
const auto *ND = cast<NamedDecl>(GD.getDecl());
std::string MangledName = getMangledNameImpl(*this, GD, ND);
- // Postfix kernel stub names with .stub to differentiate them from kernel
- // names in device binaries. This is to facilitate the debugger to find
- // the correct symbols for kernels in the device binary.
+ // Derive the kernel stub from CUDA runtime.
if (auto *FD = dyn_cast<FunctionDecl>(GD.getDecl()))
- if (getLangOpts().HIP && !getLangOpts().CUDAIsDevice &&
- FD->hasAttr<CUDAGlobalAttr>())
- MangledName = MangledName + ".stub";
+ if (!getLangOpts().CUDAIsDevice && FD->hasAttr<CUDAGlobalAttr>())
+ MangledName = getCUDARuntime().getDeviceStubName(MangledName);
auto Result = Manglings.insert(std::make_pair(MangledName, GD));
return MangledDeclNames[CanonicalGD] = Result.first->first();
Index: clang/lib/CodeGen/CGCUDARuntime.h
===================================================================
--- clang/lib/CodeGen/CGCUDARuntime.h
+++ clang/lib/CodeGen/CGCUDARuntime.h
@@ -15,6 +15,8 @@
#ifndef LLVM_CLANG_LIB_CODEGEN_CGCUDARUNTIME_H
#define LLVM_CLANG_LIB_CODEGEN_CGCUDARUNTIME_H
+#include "llvm/ADT/StringRef.h"
+
namespace llvm {
class Function;
class GlobalVariable;
@@ -63,6 +65,9 @@
/// Returns a module cleanup function or nullptr if it's not needed.
/// Must be called after ModuleCtorFunction
virtual llvm::Function *makeModuleDtorFunction() = 0;
+
+ /// Construct and return the stub name of a kernel.
+ virtual std::string getDeviceStubName(llvm::StringRef Name) const = 0;
};
/// Creates an instance of a CUDA runtime class.
Index: clang/lib/CodeGen/CGCUDANV.cpp
===================================================================
--- clang/lib/CodeGen/CGCUDANV.cpp
+++ clang/lib/CodeGen/CGCUDANV.cpp
@@ -132,6 +132,8 @@
llvm::Function *makeModuleCtorFunction() override;
/// Creates module destructor function
llvm::Function *makeModuleDtorFunction() override;
+ /// Construct and return the stub name of a kernel.
+ std::string getDeviceStubName(llvm::StringRef Name) const override;
};
}
@@ -217,10 +219,11 @@
void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
FunctionArgList &Args) {
- assert(getDeviceSideName(CGF.CurFuncDecl) == CGF.CurFn->getName() ||
- getDeviceSideName(CGF.CurFuncDecl) + ".stub" == CGF.CurFn->getName() ||
- CGF.CGM.getContext().getTargetInfo().getCXXABI() !=
- CGF.CGM.getContext().getAuxTargetInfo()->getCXXABI());
+ assert((CGF.CGM.getContext().getAuxTargetInfo() &&
+ (CGF.CGM.getContext().getAuxTargetInfo()->getCXXABI() !=
+ CGF.CGM.getContext().getTargetInfo().getCXXABI())) ||
+ getDeviceStubName(getDeviceSideName(CGF.CurFuncDecl)) ==
+ CGF.CurFn->getName());
EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
@@ -780,6 +783,12 @@
return ModuleDtorFunc;
}
+std::string CGNVCUDARuntime::getDeviceStubName(llvm::StringRef Name) const {
+ if (!CGM.getLangOpts().HIP)
+ return Name;
+ return std::move(("__device_stub__" + Name).str());
+}
+
CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
return new CGNVCUDARuntime(CGM);
}
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D63335.204771.patch
Type: text/x-patch
Size: 4039 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20190614/7ca9fa60/attachment.bin>
More information about the cfe-commits
mailing list