Author: Jonas Hahnfeld Date: 2022-04-09T12:34:41+02:00 New Revision: e4903d8be399864cc978236fc4a28087f91c20fe
URL: https://github.com/llvm/llvm-project/commit/e4903d8be399864cc978236fc4a28087f91c20fe DIFF: https://github.com/llvm/llvm-project/commit/e4903d8be399864cc978236fc4a28087f91c20fe.diff LOG: [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 Added: Modified: clang/lib/CodeGen/CGCUDANV.cpp clang/test/CodeGenCUDA/device-stub.cu Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp index 3ae152d743206..187817d0e5059 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 aa7211aeaf8e7..0f925a29c215d 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]]( _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
