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

[via cfe-commits]

https://github.com/jofrn updated https://github.com/llvm/llvm-project/pull/107458

>From b21b6c38f6a4cfb3103bb60b3122be4b5253b57f Mon Sep 17 00:00:00 2001
From: jofernau <Joe.Fernau at 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 -
 clang/test/Driver/hip-partial-link.hip |  8 --------
 3 files changed, 10 insertions(+), 17 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
diff --git a/clang/test/Driver/hip-partial-link.hip b/clang/test/Driver/hip-partial-link.hip
index c8451ec81ed37e..3deea52086c874 100644
--- a/clang/test/Driver/hip-partial-link.hip
+++ b/clang/test/Driver/hip-partial-link.hip
@@ -15,12 +15,10 @@
 // RUN: llvm-nm  %t.1.o | FileCheck -check-prefix=OBJ1 %s
 // OBJ1:  B __hip_cuid_[[ID:[0-9a-f]+]]
 // OBJ1:  U __hip_fatbin_[[ID]]
-// OBJ1:  U __hip_gpubin_handle_[[ID]]
 
 // RUN: llvm-nm  %t.2.o | FileCheck -check-prefix=OBJ2 %s
 // OBJ2:  B __hip_cuid_[[ID:[0-9a-f]+]]
 // OBJ2:  U __hip_fatbin_[[ID]]
-// OBJ2:  U __hip_gpubin_handle_[[ID]]
 
 // Link %t.1.o and %t.2.o by -r and then link with %t.main.o
 
@@ -30,8 +28,6 @@
 // RUN:   2>&1 | FileCheck -check-prefix=LD-R %s
 // LD-R: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID1:[0-9a-f]+]]
 // LD-R: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID2:[0-9a-f]+]]
-// LD-R: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID1]]
-// LD-R: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID2]]
 // LD-R: "{{.*}}/clang-offload-bundler" {{.*}}-unbundle
 // LD-R: "{{.*}}/lld" -flavor gnu -m elf64_amdgpu
 // LD-R: "{{.*}}/clang-offload-bundler"
@@ -43,8 +39,6 @@
 // OBJ:  B __hip_cuid_[[ID2:[0-9a-f]+]]
 // OBJ:  R __hip_fatbin_[[ID1]]
 // OBJ:  R __hip_fatbin_[[ID2]]
-// OBJ:  D __hip_gpubin_handle_[[ID1]]
-// OBJ:  D __hip_gpubin_handle_[[ID2]]
 
 // RUN: %clang -v --target=x86_64-unknown-linux-gnu --no-offload-new-driver \
 // RUN:   --hip-link -fgpu-rdc --offload-arch=gfx906 \
@@ -60,8 +54,6 @@
 // RUN:   2>&1 | FileCheck -check-prefix=STATIC %s
 // STATIC: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID1:[0-9a-f]+]]
 // STATIC: Found undefined HIP fatbin symbol: __hip_fatbin_[[ID2:[0-9a-f]+]]
-// STATIC: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID1]]
-// STATIC: Found undefined HIP gpubin handle symbol: __hip_gpubin_handle_[[ID2]]
 // STATIC: "{{.*}}/clang-offload-bundler" {{.*}}-unbundle
 // STATIC: "{{.*}}/lld" -flavor gnu -m elf64_amdgpu
 // STATIC: "{{.*}}/clang-offload-bundler"