Author: David Rivera Date: 2026-08-12T05:47:03-04:00 New Revision: 742031c8aaf9c6509139f2242475250d6548a657
URL: https://github.com/llvm/llvm-project/commit/742031c8aaf9c6509139f2242475250d6548a657 DIFF: https://github.com/llvm/llvm-project/commit/742031c8aaf9c6509139f2242475250d6548a657.diff LOG: [CIR][HIP] Adjust HIP fatbin alignment to 4096 (#215314) Added: Modified: clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp clang/test/CIR/CodeGenCUDA/device-stub.cu Removed: ################################################################################ diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index 12d544fecbcc9..bab06ba2856ca 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -2423,7 +2423,13 @@ 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 67846c1fa572e..1de4e5575ce1d 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
