Author: jofrn Date: 2024-09-11T18:46:46-04:00 New Revision: b5fd9463a3b9aecfc132828510f7e2a47b581b14
URL: https://github.com/llvm/llvm-project/commit/b5fd9463a3b9aecfc132828510f7e2a47b581b14 DIFF: https://github.com/llvm/llvm-project/commit/b5fd9463a3b9aecfc132828510f7e2a47b581b14.diff LOG: [HIP][Clang][CodeGen] Handle hip bin symbols properly. (#107458) Remove '_' in fatbin and gpubin symbol suffixes when missing TU hash ID. Internalize gpubin symbol so that it is not unresolved at link-time when symbol is not relocatable. 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 59c5927717933d..ae14d74f2d9151 100644 --- a/clang/lib/CodeGen/CGCUDANV.cpp +++ b/clang/lib/CodeGen/CGCUDANV.cpp @@ -840,8 +840,10 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() { FatBinStr = new llvm::GlobalVariable( CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr, - "__hip_fatbin_" + CGM.getContext().getCUIDHash(), nullptr, - llvm::GlobalVariable::NotThreadLocal); + "__hip_fatbin" + (CGM.getLangOpts().CUID.empty() + ? "" + : "_" + CGM.getContext().getCUIDHash()), + nullptr, llvm::GlobalVariable::NotThreadLocal); cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName); } @@ -894,8 +896,8 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() { // thread safety of the loaded program. Therefore we can assume sequential // execution of constructor functions here. if (IsHIP) { - auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage - : llvm::GlobalValue::ExternalLinkage; + auto Linkage = RelocatableDeviceCode ? llvm::GlobalValue::ExternalLinkage + : llvm::GlobalValue::InternalLinkage; llvm::BasicBlock *IfBlock = llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc); llvm::BasicBlock *ExitBlock = @@ -905,10 +907,11 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() { GpuBinaryHandle = new llvm::GlobalVariable( TheModule, PtrTy, /*isConstant=*/false, Linkage, /*Initializer=*/ - CudaGpuBinary ? llvm::ConstantPointerNull::get(PtrTy) : nullptr, - CudaGpuBinary - ? "__hip_gpubin_handle" - : "__hip_gpubin_handle_" + CGM.getContext().getCUIDHash()); + !RelocatableDeviceCode ? llvm::ConstantPointerNull::get(PtrTy) + : nullptr, + "__hip_gpubin_handle" + (CGM.getLangOpts().CUID.empty() + ? "" + : "_" + CGM.getContext().getCUIDHash())); GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign()); // Prevent the weak symbol in diff erent shared libraries being merged. if (Linkage != llvm::GlobalValue::InternalLinkage) diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu index 60304647bd4c54..8695433f6df10c 100644 --- a/clang/test/CodeGenCUDA/device-stub.cu +++ b/clang/test/CodeGenCUDA/device-stub.cu @@ -175,7 +175,7 @@ __device__ void device_use() { // HIP-SAME: section ".hipFatBinSegment" // * variable to save GPU binary handle after initialization // CUDANORDC: @__[[PREFIX]]_gpubin_handle = internal global ptr null -// HIPNEF: @__[[PREFIX]]_gpubin_handle_{{[0-9a-f]+}} = external hidden global ptr, align 8 +// HIPNEF: @__[[PREFIX]]_gpubin_handle_{{[0-9a-f]+}} = internal global ptr null, align 8 // * constant unnamed string with NVModuleID // CUDARDC: [[MODULE_ID_GLOBAL:@.*]] = private constant // CUDARDC-SAME: c"[[MODULE_ID:.+]]\00", section "__nv_module_id", align 32 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
