[clang] [CIR][HIP] Adjust HIP fatbin alignment to 4096 (PR #215314)

[David Rivera via cfe-commits]

https://github.com/RiverDave created https://github.com/llvm/llvm-project/pull/215314

None

>From b4497a77b353d2cd6bd16c601eee90929984fa02 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Mon, 10 Aug 2026 11:53:19 -0400
Subject: [PATCH] [CIR][HIP] Adjust HIP fatbin alignment to 4096

---
 clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp | 6 +++++-
 clang/test/CIR/CodeGenCUDA/device-stub.cu            | 6 +++---
 2 files changed, 8 insertions(+), 4 deletions(-)

diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index 12d544fecbcc9..06a44c025c0ec 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -2423,7 +2423,11 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
   GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType,
                                         /*isConstant=*/true, {},
                                         GlobalLinkageKind::PrivateLinkage);
-  fatbinStr.setAlignment(8);
+  if (isHIP) {
+    const unsigned HIPCodeObjectAlign = 4096;
+    fatbinStr.setAlignment(HIPCodeObjectAlign);
+  } else
+    fatbinStr.setAlignment(8);
   fatbinStr.setInitialValueAttr(cir::ConstArrayAttr::get(
       fatbinType, StringAttr::get(gpuBinary->getBuffer(), fatbinType)));
   fatbinStr.setSection(fatbinConstName);
diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu
index ca5e185add5fc..2c6bbd77e3333 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -249,7 +249,7 @@ __device__ _BitInt(36) c;
 
 // Fatbin string + wrapper live in the HIP-specific sections; magic
 // 0x48495046 = 1212764230.
-// HIP-CIR: cir.global "private" constant cir_private @__hip_fatbin_str = #cir.const_array<"GPU binary would be here." : !cir.array<!u8i x 25>> : !cir.array<!u8i x 25> {alignment = 8 : i64, section = ".hip_fatbin"}
+// HIP-CIR: cir.global "private" constant cir_private @__hip_fatbin_str = #cir.const_array<"GPU binary would be here." : !cir.array<!u8i x 25>> : !cir.array<!u8i x 25> {alignment = 4096 : i64, section = ".hip_fatbin"}
 // HIP-CIR: cir.global constant cir_private @__hip_fatbin_wrapper = #cir.const_record<{
 // HIP-CIR-SAME: #cir.int<1212764230> : !s32i,
 // HIP-CIR-SAME: #cir.int<1> : !s32i,
@@ -286,7 +286,7 @@ __device__ _BitInt(36) c;
 // HIP-CIR: cir.global constant external @_Z10kernelfunciii = #cir.global_view<@_Z25__device_stub__kernelfunciii> : !cir.ptr<!cir.func<(!s32i, !s32i, !s32i)>> {alignment = 8 : i64}
 
 // HIP OGCG cross-check (LLVM IR matches what OG codegen emits for HIP).
-// HIP-OGCG: @{{.*}} = private constant [25 x i8] c"GPU binary would be here.", section ".hip_fatbin"
+// HIP-OGCG: @{{.*}} = private constant [25 x i8] c"GPU binary would be here.", section ".hip_fatbin", align 4096
 // HIP-OGCG: @__hip_fatbin_wrapper = internal constant { i32, i32, ptr, ptr } { i32 1212764230, i32 1, ptr @{{.*}}, ptr null }, section ".hipFatBinSegment"
 // HIP-OGCG: @__hip_gpubin_handle = internal global ptr null
 // HIP-OGCG: @llvm.global_ctors = appending global {{.*}}@__hip_module_ctor
@@ -307,7 +307,7 @@ __device__ _BitInt(36) c;
 // HIP-OGCG:   store ptr null, ptr @__hip_gpubin_handle
 
 // HIP LLVM lowering cross-check.
-// HIP-LLVM: @{{.*}} = private constant [25 x i8] c"GPU binary would be here.", section ".hip_fatbin", align 8
+// HIP-LLVM: @{{.*}} = private constant [25 x i8] c"GPU binary would be here.", section ".hip_fatbin", align 4096
 // HIP-LLVM: @__hip_fatbin_wrapper = {{.*}}constant { i32, i32, ptr, ptr } { i32 1212764230, i32 1, ptr @{{.*}}, ptr null }, section ".hipFatBinSegment"
 // HIP-LLVM: @__hip_gpubin_handle = internal global ptr null
 // HIP-LLVM: @_Z10kernelfunciii = constant ptr @_Z25__device_stub__kernelfunciii, align 8