diff options
| author | Jonas Hahnfeld <jonas.hahnfeld@cern.ch> | 2022-04-08 10:30:30 +0200 |
|---|---|---|
| committer | Jonas Hahnfeld <jonas.hahnfeld@cern.ch> | 2022-04-09 12:34:41 +0200 |
| commit | e4903d8be399864cc978236fc4a28087f91c20fe (patch) | |
| tree | 9fb2fd6489eb48549158514d46ef0a2af4a458da | |
| parent | f49a763c4da9a01fcd263f31e4dca72e0d930f11 (diff) | |
| download | llvm-e4903d8be399864cc978236fc4a28087f91c20fe.zip llvm-e4903d8be399864cc978236fc4a28087f91c20fe.tar.gz llvm-e4903d8be399864cc978236fc4a28087f91c20fe.tar.bz2 | |
[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
| -rw-r--r-- | clang/lib/CodeGen/CGCUDANV.cpp | 12 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/device-stub.cu | 4 |
2 files changed, 8 insertions, 8 deletions
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 3ae152d..187817d 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -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); diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu index aa7211a..0f925a2 100644 --- a/clang/test/CodeGenCUDA/device-stub.cu +++ b/clang/test/CodeGenCUDA/device-stub.cu @@ -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]]( |