[clang] [CIR][HIP] Emit module ctor without a GPU binary (PR #234855)
Arseniy Obolenskiy via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 22 23:21:57 PDT 2026
https://github.com/aobolensk created https://github.com/llvm/llvm-project/pull/234855
Mirror classic codegen, reference an external __hip_fatbin and add the CUID hash to the fatbin and handle names
>From c167baaa8e0d27e2d4a79d3decbec41605918ddd Mon Sep 17 00:00:00 2001
From: Arseniy Obolenskiy <arseniy.obolenskiy at amd.com>
Date: Fri, 23 Oct 2026 08:21:18 +0200
Subject: [PATCH] [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]]
More information about the cfe-commits
mailing list