https://github.com/jofrn updated
https://github.com/llvm/llvm-project/pull/107458
>From d385e0810bc24dcaf8296501922b50fd88e18415 Mon Sep 17 00:00:00 2001
From: jofernau <joe.fer...@amd.com>
Date: Thu, 5 Sep 2024 23:31:55 -0400
Subject: [PATCH] [HIP][Clang][CodeGen] Handle hip bin symbols properly.
Remove '_' in fatbin symbol suffix when missing TU hash ID. Internalize gpubin
symbol so that it is not unresolved at link-time.
---
clang/lib/CodeGen/CGCUDANV.cpp | 18 ++++++++++--------
clang/test/CodeGenCUDA/device-stub.cu | 1 -
2 files changed, 10 insertions(+), 9 deletions(-)
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 59c5927717933d..7988d4ce462caf 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,10 @@ 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());
+ llvm::ConstantPointerNull::get(PtrTy),
+ "__hip_gpubin_handle" + (CGM.getLangOpts().CUID.empty()
+ ? ""
+ : "_" + CGM.getContext().getCUIDHash()));
GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
// Prevent the weak symbol in different 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..11d1f5a867b09a 100644
--- a/clang/test/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CodeGenCUDA/device-stub.cu
@@ -175,7 +175,6 @@ __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
// * 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
