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 <[email protected]> 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
