https://github.com/aobolensk updated https://github.com/llvm/llvm-project/pull/234855
>From c167baaa8e0d27e2d4a79d3decbec41605918ddd Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 23 Oct 2026 08:21:18 +0200 Subject: [PATCH 1/3] [CIR][HIP] Emit module ctor without a GPU binary Mirror classic codegen, reference an external __hip_fatbin and add the CUID hash to the fatbin and handle names --- .../clang/CIR/Dialect/IR/CIRDialect.td | 3 + clang/include/clang/CIR/MissingFeatures.h | 1 - clang/lib/CIR/CodeGen/CIRGenModule.cpp | 3 + .../Dialect/Transforms/LoweringPrepare.cpp | 65 ++++++++++++------- clang/test/CIR/CodeGenCUDA/device-stub.cu | 37 +++++++++-- 5 files changed, 78 insertions(+), 31 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index cd36d618ca14b3..ebc5ab7ac8d17a 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -126,6 +126,9 @@ def CIR_Dialect : Dialect { // Platform SDK version from -target-sdk-version, as printed by // llvm::VersionTuple (e.g. "12.3"). Absent when none was given. static llvm::StringRef getSDKVersionAttrName() { return "cir.sdk_version"; } + // Hash of the -cuid compilation unit ID, as returned by + // ASTContext::getCUIDHash. Absent when no CUID was given. + static llvm::StringRef getCUIDHashAttrName() { return "cir.cuid_hash"; } static llvm::StringRef getMustTailAttrName() { return "musttail"; } static llvm::StringRef getCatchCopyThunkAttrName() { return "cir.eh.catch_copy_thunk"; } diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index cc826f7816ebbc..68afad25887e1b 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -216,7 +216,6 @@ struct MissingFeatures { static bool ctorMemcpyizer() { return false; } static bool cudaSupport() { return false; } static bool cudaTextureType() { return false; } - static bool hipModuleCtor() { return false; } static bool dataLayoutTypeAllocSize() { return false; } static bool dataLayoutPtrHandlingBasedOnLangAS() { return false; } static bool deadOnReturnAttr() { return false; } diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index 8801a730e8d678..f91074afe6f18b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -179,6 +179,9 @@ CIRGenModule::CIRGenModule(mlir::MLIRContext &mlirContext, !sdkVersion.empty()) getModule()->setAttr(cir::CIRDialect::getSDKVersionAttrName(), builder.getStringAttr(sdkVersion.getAsString())); + if (langOpts.CUDA && !langOpts.CUID.empty()) + getModule()->setAttr(cir::CIRDialect::getCUIDHashAttrName(), + builder.getStringAttr(astContext.getCUIDHash())); // TODO(CIR): These attributes should eventually be replaced by // TypeSizeInfoAttr once it is upstreamed. getModule()->setAttr(cir::CIRDialect::getSizeTypeWidthAttrName(), diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index 495109e0dc4350..220a5eea9b0873 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -350,6 +350,16 @@ struct LoweringPreparePass return sdkVersion; } + /// Appends the CUID hash recorded on the module by CIRGen, if any, to + /// \p name. + std::string addCUIDSuffix(llvm::StringRef name) { + std::string result = name.str(); + if (auto cuidHash = mlirModule->getAttrOfType<mlir::StringAttr>( + CIRDialect::getCUIDHashAttrName())) + result += ("_" + cuidHash.getValue()).str(); + return result; + } + /// Tracks current module. mlir::ModuleOp mlirModule; @@ -2849,14 +2859,11 @@ void LoweringPreparePass::buildCUDAModuleCtor() { return; // There's no device-side binary, so no need to proceed for CUDA. - // HIP has to create an external symbol in this case, which is NYI. + // HIP creates an external symbol in this case. auto deviceBinaryAttr = mlirModule->getAttrOfType<mlir::StringAttr>( CIRDialect::getCUDADeviceBinaryAttrName()); - if (!deviceBinaryAttr) { - if (isHIP) - assert(!cir::MissingFeatures::hipModuleCtor()); + if (!deviceBinaryAttr && !isHIP) return; - } // Set up common types and builder. llvm::StringRef cudaPrefix = getCUDAPrefix(getLangOpts()); @@ -2878,22 +2885,32 @@ void LoweringPreparePass::buildCUDAModuleCtor() { llvm::StringRef fatbinSectionName = getLangOpts().HIP ? ".hipFatBinSegment" : ".nvFatBinSegment"; - // Create the fatbin string constant with GPU binary contents. - // The dialect verifier guarantees the attribute is typed as the array. - auto fatbinType = mlir::cast<ArrayType>(deviceBinaryAttr.getType()); - std::string fatbinStrName = addUnderscoredPrefix(cudaPrefix, "_fatbin_str"); - GlobalOp fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType, - /*isConstant=*/true, {}, - GlobalLinkageKind::PrivateLinkage); - if (isHIP) { - const unsigned HIPCodeObjectAlign = 4096; - fatbinStr.setAlignment(HIPCodeObjectAlign); + GlobalOp fatbinStr; + if (deviceBinaryAttr) { + // Create the fatbin string constant with GPU binary contents. + // The dialect verifier guarantees the attribute is typed as the array. + auto fatbinType = mlir::cast<ArrayType>(deviceBinaryAttr.getType()); + std::string fatbinStrName = addUnderscoredPrefix(cudaPrefix, "_fatbin_str"); + fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType, + /*isConstant=*/true, {}, + GlobalLinkageKind::PrivateLinkage); + if (isHIP) { + const unsigned HIPCodeObjectAlign = 4096; + fatbinStr.setAlignment(HIPCodeObjectAlign); + } else { + fatbinStr.setAlignment(8); + } + + fatbinStr.setInitialValueAttr( + cir::ConstArrayAttr::get(fatbinType, deviceBinaryAttr)); } else { - fatbinStr.setAlignment(8); + // If fatbin is not available, create an external symbol __hip_fatbin in + // section .hip_fatbin. It is populated elsewhere, e.g. by lld through a + // linker script. + fatbinStr = GlobalOp::create( + builder, loc, addCUIDSuffix("__hip_fatbin"), builder.getUIntNTy(8), + /*isConstant=*/true, {}, GlobalLinkageKind::ExternalLinkage); } - - fatbinStr.setInitialValueAttr( - cir::ConstArrayAttr::get(fatbinType, deviceBinaryAttr)); fatbinStr.setSection(fatbinConstName); fatbinStr.setPrivate(); @@ -2928,6 +2945,8 @@ void LoweringPreparePass::buildCUDAModuleCtor() { // Create the GPU binary handle global variable. std::string gpubinHandleName = addUnderscoredPrefix(cudaPrefix, "_gpubin_handle"); + if (isHIP) + gpubinHandleName = addCUIDSuffix(gpubinHandleName); GlobalOp gpuBinHandle = GlobalOp::create( builder, loc, gpubinHandleName, voidPtrPtrTy, @@ -3130,9 +3149,6 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() { /// the dtor list would cause a double-free. It is meant to be registered via /// atexit() at the end of the module ctor. std::optional<FuncOp> LoweringPreparePass::buildHIPModuleDtor() { - if (!mlirModule->getAttr(CIRDialect::getCUDADeviceBinaryAttrName())) - return {}; - llvm::StringRef prefix = getCUDAPrefix(getLangOpts()); VoidType voidTy = VoidType::get(&getContext()); @@ -3154,7 +3170,8 @@ std::optional<FuncOp> LoweringPreparePass::buildHIPModuleDtor() { buildRuntimeFunction(builder, dtorName, loc, FuncType::get({}, voidTy), GlobalLinkageKind::InternalLinkage); - std::string gpubinName = addUnderscoredPrefix(prefix, "_gpubin_handle"); + std::string gpubinName = + addCUIDSuffix(addUnderscoredPrefix(prefix, "_gpubin_handle")); GlobalOp gpuBinGlobal = cast<GlobalOp>(mlirModule.lookupSymbol(gpubinName)); mlir::Block *entryBlock = dtor.addEntryBlock(); @@ -3451,7 +3468,7 @@ void LoweringPreparePass::runOnOperation() { buildCUDAModuleCtor(); // The fatbin global now references the same attribute; drop the module's // reference so an emitted .cir doesn't print the bytes twice. This has to - // happen out here because the ctor and both dtor builders test the + // happen out here because the ctor and the CUDA dtor builder test the // attribute to decide whether a device-side binary exists at all. mlirModule->removeAttr(CIRDialect::getCUDADeviceBinaryAttrName()); } diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu index aa513f35374545..dde05578b1cb9b 100644 --- a/clang/test/CIR/CodeGenCUDA/device-stub.cu +++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu @@ -42,6 +42,14 @@ // RUN: -fhip-new-launch-api -o %t.nogpu.hip.cir // RUN: FileCheck --input-file=%t.nogpu.hip.cir %s --check-prefix=HIP-NOGPUBIN +// RUN: %clang_cc1 -triple x86_64-linux-gnu -fclangir -emit-llvm %s -x hip \ +// RUN: -fhip-new-launch-api -cuid=123 -o %t.cuid.hip-cir.ll +// RUN: FileCheck --input-file=%t.cuid.hip-cir.ll %s --check-prefix=HIP-CUID + +// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x hip \ +// RUN: -fhip-new-launch-api -cuid=123 -o %t.cuid.hip.ll +// RUN: FileCheck --input-file=%t.cuid.hip.ll %s --check-prefix=HIP-CUID + #include "Inputs/cuda.h" __global__ void kernelfunc(int i, int j, int k) {} @@ -358,10 +366,27 @@ __device__ _BitInt(36) c; // HIP-LLVM: call i32 @atexit(ptr @__hip_module_dtor) // HIP-LLVM: ret void -// No GPU binary: no fatbin, no handle, no registration scaffolding. +// No GPU binary: the wrapper points to an external __hip_fatbin, populated +// at link time. // HIP-NOGPUBIN-NOT: cir.cu.device_binary -// HIP-NOGPUBIN-NOT: __hip_fatbin -// HIP-NOGPUBIN-NOT: __hip_gpubin_handle -// HIP-NOGPUBIN-NOT: __hip_register_globals -// HIP-NOGPUBIN-NOT: __hip_module_ctor -// HIP-NOGPUBIN-NOT: __hip_module_dtor +// HIP-NOGPUBIN: cir.global_ctors = [#cir.global_ctor<"__hip_module_ctor", 65535>] +// HIP-NOGPUBIN: cir.func internal private @__hip_module_dtor() +// HIP-NOGPUBIN: cir.func internal private @__hip_register_globals( +// HIP-NOGPUBIN: cir.global "private" constant external @__hip_fatbin : !u8i section(".hip_fatbin") +// HIP-NOGPUBIN: cir.global constant cir_private @__hip_fatbin_wrapper = #cir.const_record<{ +// HIP-NOGPUBIN-SAME: #cir.global_view<@__hip_fatbin> : !cir.ptr<!void> +// HIP-NOGPUBIN: cir.global "private" internal @__hip_gpubin_handle = #cir.ptr<null> +// HIP-NOGPUBIN: cir.func internal private @__hip_module_ctor() +// HIP-NOGPUBIN: cir.call @__hipRegisterFatBinary( +// HIP-NOGPUBIN: cir.call @__hip_register_globals( +// HIP-NOGPUBIN: cir.call @atexit( + +// With a CUID, the external fatbin and the handle get its hash as a suffix. +// HIP-CUID: @[[FATBIN:__hip_fatbin_[0-9a-f]+]] = external constant i8, section ".hip_fatbin" +// HIP-CUID: @__hip_fatbin_wrapper = {{.*}}constant { i32, i32, ptr, ptr } { i32 1212764230, i32 1, ptr @[[FATBIN]], ptr null }, section ".hipFatBinSegment" +// HIP-CUID: @[[HANDLE:__hip_gpubin_handle_[0-9a-f]+]] = internal global ptr null +// HIP-CUID-DAG: define internal void @__hip_module_ctor() +// HIP-CUID-DAG: store ptr %{{.*}}, ptr @[[HANDLE]] +// HIP-CUID-DAG: call i32 @atexit(ptr @__hip_module_dtor) +// HIP-CUID-DAG: define internal void @__hip_module_dtor() +// HIP-CUID-DAG: store ptr null, ptr @[[HANDLE]] >From 5cfb0cd12b9d0d725ae5f5bb794fb9cd1363b4ce Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 23 Oct 2026 09:52:07 +0200 Subject: [PATCH 2/3] Update clang/include/clang/CIR/Dialect/IR/CIRDialect.td Co-authored-by: Steffen Larsen <[email protected]> --- clang/include/clang/CIR/Dialect/IR/CIRDialect.td | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td index ebc5ab7ac8d17a..4cccbb745b9774 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRDialect.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRDialect.td @@ -126,8 +126,7 @@ def CIR_Dialect : Dialect { // Platform SDK version from -target-sdk-version, as printed by // llvm::VersionTuple (e.g. "12.3"). Absent when none was given. static llvm::StringRef getSDKVersionAttrName() { return "cir.sdk_version"; } - // Hash of the -cuid compilation unit ID, as returned by - // ASTContext::getCUIDHash. Absent when no CUID was given. + // Hash of the -cuid compilation unit ID. Absent when no CUID was given. static llvm::StringRef getCUIDHashAttrName() { return "cir.cuid_hash"; } static llvm::StringRef getMustTailAttrName() { return "musttail"; } static llvm::StringRef getCatchCopyThunkAttrName() { return "cir.eh.catch_copy_thunk"; } >From 7c357266896142a7353756d9d0efe5b4c12f1006 Mon Sep 17 00:00:00 2001 From: Arseniy Obolenskiy <[email protected]> Date: Fri, 23 Oct 2026 09:52:33 +0200 Subject: [PATCH 3/3] Update clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp Co-authored-by: Steffen Larsen <[email protected]> --- clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp index 220a5eea9b0873..f39a2e499fa576 100644 --- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp +++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp @@ -2894,12 +2894,8 @@ void LoweringPreparePass::buildCUDAModuleCtor() { fatbinStr = GlobalOp::create(builder, loc, fatbinStrName, fatbinType, /*isConstant=*/true, {}, GlobalLinkageKind::PrivateLinkage); - if (isHIP) { - const unsigned HIPCodeObjectAlign = 4096; - fatbinStr.setAlignment(HIPCodeObjectAlign); - } else { - fatbinStr.setAlignment(8); - } + constexpr unsigned HIPCodeObjectAlign = 4096; + fatbinStr.setAlignment(isHIP ? HIPCodeObjectAlign : 8); fatbinStr.setInitialValueAttr( cir::ConstArrayAttr::get(fatbinType, deviceBinaryAttr)); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
