caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[HIP][Clang][CodeGen] Handle hip bin symbols properly. (#107458)

Browse Source 
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.
This commit is contained in:
jofrn
2024-09-11 18:46:46 -04:00
committed by GitHub
parent 2a9208b0c9
commit b5fd9463a3
2 changed files with 12 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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 different shared libraries being merged.
if (Linkage != llvm::GlobalValue::InternalLinkage)

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

@@ -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