maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA/HIP] Remove argument from module ctor/dtor signatures 

In theory, constructors can take arguments when called via .init_array
where at least glibc passes in (argc, argv, envp). This isn't used in
the generated code and if it was, the first argument should be an
integer, not a pointer. For destructors registered via atexit, the
function should never take an argument.

Differential Revision: https://reviews.llvm.org/D123370
This commit is contained in:
Jonas Hahnfeld 2022-04-08 10:30:30 +02:00
parent f49a763c4d
commit e4903d8be3
2 changed files with 8 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
clang/lib/CodeGen/CGCUDANV.cpp
View File

@ -659,7 +659,7 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
///
/// For CUDA:
/// \code
/// void __cuda_module_ctor(void*) {
/// void __cuda_module_ctor() {
/// Handle = __cudaRegisterFatBinary(GpuBinaryBlob);
/// __cuda_register_globals(Handle);
/// }
@ -667,7 +667,7 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
///
/// For HIP:
/// \code
/// void __hip_module_ctor(void*) {
/// void __hip_module_ctor() {
/// if (__hip_gpubin_handle == 0) {
/// __hip_gpubin_handle = __hipRegisterFatBinary(GpuBinaryBlob);
/// __hip_register_globals(__hip_gpubin_handle);
@ -717,7 +717,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
}
llvm::Function *ModuleCtorFunc = llvm::Function::Create(
llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
llvm::FunctionType::get(VoidTy, false),
llvm::GlobalValue::InternalLinkage,
addUnderscoredPrefixToName("_module_ctor"), &TheModule);
llvm::BasicBlock *CtorEntryBB =
@ -931,14 +931,14 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
///
/// For CUDA:
/// \code
/// void __cuda_module_dtor(void*) {
/// void __cuda_module_dtor() {
/// __cudaUnregisterFatBinary(Handle);
/// }
/// \endcode
///
/// For HIP:
/// \code
/// void __hip_module_dtor(void*) {
/// void __hip_module_dtor() {
/// if (__hip_gpubin_handle) {
/// __hipUnregisterFatBinary(__hip_gpubin_handle);
/// __hip_gpubin_handle = 0;
@ -956,7 +956,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
addUnderscoredPrefixToName("UnregisterFatBinary"));
llvm::Function *ModuleDtorFunc = llvm::Function::Create(
llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
llvm::FunctionType::get(VoidTy, false),
llvm::GlobalValue::InternalLinkage,
addUnderscoredPrefixToName("_module_dtor"), &TheModule);

4
clang/test/CodeGenCUDA/device-stub.cu
View File

@ -257,8 +257,8 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// CUDANORDC-NEXT: call void @__[[PREFIX]]_register_globals
// HIP-NEXT: call void @__[[PREFIX]]_register_globals
// * In separate mode we also register a destructor.
// CUDANORDC-NEXT: call i32 @atexit(void (i8*)* @__[[PREFIX]]_module_dtor)
// HIP-NEXT: call i32 @atexit(void (i8*)* @__[[PREFIX]]_module_dtor)
// CUDANORDC-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
// HIP-NEXT: call i32 @atexit(void ()* @__[[PREFIX]]_module_dtor)
// With relocatable device code we call __[[PREFIX]]RegisterLinkedBinary%NVModuleID%
// CUDARDC: call{{.*}}__[[PREFIX]]RegisterLinkedBinary[[MODULE_ID]](