[clang] [CIR][HIP] Handle HIP module constructor and destructor emission (PR #195391)
via cfe-commits
cfe-commits at lists.llvm.org
Sun May 3 15:16:03 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: David Rivera (RiverDave)
<details>
<summary>Changes</summary>
Related: https://github.com/llvm/llvm-project/issues/179278, https://github.com/llvm/llvm-project/issues/175871
Similar to https://github.com/llvm/llvm-project/pull/188673, This adds the HIP host-side module registration path in CIR lowering for the non-RDC, included-fatbin case.
Generated sequence for HIP, non-RDC, with `-fcuda-include-gpubinary`:
```c
void **__hip_gpubin_handle = nullptr;
void __hip_module_ctor() {
if (__hip_gpubin_handle == nullptr)
__hip_gpubin_handle = __hipRegisterFatBinary(&__hip_fatbin_wrapper);
__hip_register_globals(__hip_gpubin_handle); // we only register kernels so far.
atexit(__hip_module_dtor);
}
void __hip_module_dtor() {
if (__hip_gpubin_handle != nullptr) {
__hipUnregisterFatBinary(__hip_gpubin_handle);
__hip_gpubin_handle = nullptr;
}
}
```
Another divergence I added (but not depicted) is the per-kernel `__hipRegisterFunction` calls using the HIP kernel-handle global, not the device-stub function symbol.
---
Full diff: https://github.com/llvm/llvm-project/pull/195391.diff
2 Files Affected:
- (modified) clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp (+148-11)
- (modified) clang/test/CIR/CodeGenCUDA/device-stub.cu (+121)
``````````diff
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index bd3c8bc0aa8d1..578cc32ae4ed1 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -146,6 +146,7 @@ struct LoweringPreparePass
/// with the CUDA runtime.
void buildCUDAModuleCtor();
std::optional<FuncOp> buildCUDAModuleDtor();
+ std::optional<FuncOp> buildHIPModuleDtor();
std::optional<FuncOp> buildCUDARegisterGlobals();
void buildCUDARegisterGlobalFunctions(cir::CIRBaseBuilderTy &builder,
FuncOp regGlobalFunc);
@@ -1890,8 +1891,6 @@ static std::string addUnderscoredPrefix(llvm::StringRef prefix,
void LoweringPreparePass::buildCUDAModuleCtor() {
bool isHIP = astCtx->getLangOpts().HIP;
- if (isHIP)
- assert(!cir::MissingFeatures::hipModuleCtor());
if (astCtx->getLangOpts().GPURelocatableDeviceCode)
llvm_unreachable("GPU RDC NYI");
@@ -2018,8 +2017,67 @@ void LoweringPreparePass::buildCUDAModuleCtor() {
builder.setInsertionPointToStart(moduleCtor.addEntryBlock());
assert(!cir::MissingFeatures::opGlobalCtorPriority());
if (isHIP) {
- llvm_unreachable("HIP Module Constructor Support");
- } else if (!astCtx->getLangOpts().GPURelocatableDeviceCode) {
+ // --- Create HIP CTOR ---
+ // if (__hip_gpubin_handle == nullptr)
+ // __hip_gpubin_handle = __hipRegisterFatBinary(&fatbinWrapper);
+ // __hip_register_globals(__hip_gpubin_handle);
+ // atexit(__hip_module_dtor);
+ mlir::Block *entryBlock = builder.getInsertionBlock();
+ mlir::Region *parent = entryBlock->getParent();
+ mlir::Block *ifBlock = builder.createBlock(parent);
+ mlir::Block *exitBlock = builder.createBlock(parent);
+ {
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(entryBlock);
+ mlir::Value handle =
+ builder.createLoad(loc, builder.createGetGlobal(gpuBinHandle));
+ auto handlePtrTy = mlir::cast<cir::PointerType>(handle.getType());
+ mlir::Value nullPtr = builder.getNullPtr(handlePtrTy, loc);
+ mlir::Value isNull =
+ builder.createCompare(loc, cir::CmpOpKind::eq, handle, nullPtr);
+ cir::BrCondOp::create(builder, loc, isNull, ifBlock, exitBlock);
+ }
+ {
+ // Handle is null: load the fatbin and register it.
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(ifBlock);
+ mlir::Value wrapper = builder.createGetGlobal(fatbinWrapper);
+ mlir::Value fatbinVoidPtr = builder.createBitcast(wrapper, voidPtrTy);
+ cir::CallOp gpuBinaryHandleCall =
+ builder.createCallOp(loc, regFunc, fatbinVoidPtr);
+ mlir::Value gpuBinaryHandle = gpuBinaryHandleCall.getResult();
+ // Store the value back to the global `__hip_gpubin_handle`.
+ mlir::Value gpuBinaryHandleGlobal = builder.createGetGlobal(gpuBinHandle);
+ builder.createStore(loc, gpuBinaryHandle, gpuBinaryHandleGlobal);
+ cir::BrOp::create(builder, loc, exitBlock);
+ }
+ {
+ // Exit block: load the (possibly newly-registered) handle, call
+ // __hip_register_globals, and register the module dtor with atexit().
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToStart(exitBlock);
+ mlir::Value gHandle =
+ builder.createLoad(loc, builder.createGetGlobal(gpuBinHandle));
+
+ if (std::optional<FuncOp> regGlobal = buildCUDARegisterGlobals())
+ builder.createCallOp(loc, *regGlobal, gHandle);
+
+ if (std::optional<FuncOp> dtor = buildHIPModuleDtor()) {
+ cir::CIRBaseBuilderTy globalBuilder(getContext());
+ globalBuilder.setInsertionPointToStart(mlirModule.getBody());
+ FuncOp atexit = buildRuntimeFunction(
+ globalBuilder, "atexit", loc,
+ FuncType::get(PointerType::get(dtor->getFunctionType()), intTy));
+ mlir::Value dtorFunc = GetGlobalOp::create(
+ builder, loc, PointerType::get(dtor->getFunctionType()),
+ mlir::FlatSymbolRefAttr::get(dtor->getSymNameAttr()));
+ builder.createCallOp(loc, atexit, dtorFunc);
+ }
+ cir::ReturnOp::create(builder, loc);
+ }
+ return;
+ }
+ if (!astCtx->getLangOpts().GPURelocatableDeviceCode) {
// --- Create CUDA CTOR-DTOR ---
// Register binary with CUDA runtime. This is substantially different in
@@ -2120,6 +2178,77 @@ std::optional<FuncOp> LoweringPreparePass::buildCUDAModuleDtor() {
return dtor;
}
+/// Build the HIP module dtor:
+///
+/// void __hip_module_dtor() {
+/// if (__hip_gpubin_handle != nullptr) {
+/// __hipUnregisterFatBinary(__hip_gpubin_handle);
+/// __hip_gpubin_handle = nullptr;
+/// }
+/// }
+///
+/// Despite the name, OG doesn't treat this as a real destructor: putting it on
+/// 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::getCUDABinaryHandleAttrName()))
+ return {};
+
+ llvm::StringRef prefix = getCUDAPrefix(astCtx);
+
+ VoidType voidTy = VoidType::get(&getContext());
+ PointerType voidPtrPtrTy = PointerType::get(PointerType::get(voidTy));
+
+ mlir::Location loc = mlirModule.getLoc();
+
+ cir::CIRBaseBuilderTy builder(getContext());
+ builder.setInsertionPointToStart(mlirModule.getBody());
+
+ // void __hipUnregisterFatBinary(void ** handle);
+ std::string unregisterFuncName =
+ addUnderscoredPrefix(prefix, "UnregisterFatBinary");
+ FuncOp unregisterFunc = buildRuntimeFunction(
+ builder, unregisterFuncName, loc, FuncType::get({voidPtrPtrTy}, voidTy));
+
+ std::string dtorName = addUnderscoredPrefix(prefix, "_module_dtor");
+ FuncOp dtor =
+ buildRuntimeFunction(builder, dtorName, loc, FuncType::get({}, voidTy),
+ GlobalLinkageKind::InternalLinkage);
+
+ std::string gpubinName = addUnderscoredPrefix(prefix, "_gpubin_handle");
+ GlobalOp gpuBinGlobal = cast<GlobalOp>(mlirModule.lookupSymbol(gpubinName));
+
+ mlir::Block *entryBlock = dtor.addEntryBlock();
+ mlir::Block *ifBlock = builder.createBlock(&dtor.getBody());
+ mlir::Block *exitBlock = builder.createBlock(&dtor.getBody());
+
+ mlir::OpBuilder::InsertionGuard guard(builder);
+ builder.setInsertionPointToEnd(entryBlock);
+ mlir::Value handle =
+ builder.createLoad(loc, builder.createGetGlobal(gpuBinGlobal));
+ auto handlePtrTy = mlir::cast<cir::PointerType>(handle.getType());
+ mlir::Value nullPtr = builder.getNullPtr(handlePtrTy, loc);
+ mlir::Value isNotNull =
+ builder.createCompare(loc, cir::CmpOpKind::ne, handle, nullPtr);
+ cir::BrCondOp::create(builder, loc, isNotNull, ifBlock, exitBlock);
+
+ {
+ // Handle is non-null: unregister and clear it.
+ mlir::OpBuilder::InsertionGuard ifGuard(builder);
+ builder.setInsertionPointToStart(ifBlock);
+ builder.createCallOp(loc, unregisterFunc, handle);
+ builder.createStore(loc, nullPtr, builder.createGetGlobal(gpuBinGlobal));
+ cir::BrOp::create(builder, loc, exitBlock);
+ }
+ {
+ mlir::OpBuilder::InsertionGuard exitGuard(builder);
+ builder.setInsertionPointToStart(exitBlock);
+ cir::ReturnOp::create(builder, loc);
+ }
+
+ return dtor;
+}
+
std::optional<FuncOp> LoweringPreparePass::buildCUDARegisterGlobals() {
// There is nothing to register.
if (cudaKernelMap.empty())
@@ -2212,20 +2341,28 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
mlir::Value deviceFunc = builder.createBitcast(
builder.createGetGlobal(deviceFuncStr), voidPtrTy);
+ mlir::Value hostFunc;
if (isHIP) {
- llvm_unreachable("HIP kernel registration NYI");
+ // Under HIP, the kernel-handle is a GlobalOp shadow created by CIR
+ // codegen and named with the kernel-reference mangled name (e.g.
+ // `@_Z2fnv` pointing at the device-stub function
+ // `_Z17__device_stub__fnv`). The CUDAKernelNameAttr on the device-stub
+ // uses the same name, so we can resolve the shadow by symbol lookup.
+ auto funcHandle = cast<GlobalOp>(mlirModule.lookupSymbol(kernelName));
+ hostFunc =
+ builder.createBitcast(builder.createGetGlobal(funcHandle), voidPtrTy);
} else {
- mlir::Value hostFunc = builder.createBitcast(
+ hostFunc = builder.createBitcast(
GetGlobalOp::create(
builder, loc, PointerType::get(deviceStub.getFunctionType()),
mlir::FlatSymbolRefAttr::get(deviceStub.getSymNameAttr())),
voidPtrTy);
- builder.createCallOp(
- loc, cudaRegisterFunction,
- {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
- ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)),
- cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
}
+ builder.createCallOp(
+ loc, cudaRegisterFunction,
+ {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
+ ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)), cirNullPtr,
+ cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
}
}
diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu
index 0f9d4d68d67ff..385a359eddc2f 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -15,6 +15,18 @@
// RUN: -target-sdk-version=12.3 -o %t.nogpu.cir
// RUN: FileCheck --input-file=%t.nogpu.cir %s --check-prefix=NOGPUBIN
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-cir %s -x hip \
+// RUN: -fhip-new-launch-api -fcuda-include-gpubinary %t -o %t.hip.cir
+// RUN: FileCheck --input-file=%t.hip.cir %s --check-prefix=HIP-CIR
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-llvm %s -x hip \
+// RUN: -fhip-new-launch-api -fcuda-include-gpubinary %t -o %t.hip.ll
+// RUN: FileCheck --input-file=%t.hip.ll %s --check-prefix=HIP-OGCG
+
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -emit-cir %s -x hip \
+// RUN: -fhip-new-launch-api -o %t.nogpu.hip.cir
+// RUN: FileCheck --input-file=%t.nogpu.hip.cir %s --check-prefix=HIP-NOGPUBIN
+
#include "Inputs/cuda.h"
__global__ void kernelfunc(int i, int j, int k) {}
@@ -111,3 +123,112 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// NOGPUBIN-NOT: __cuda_register_globals
// NOGPUBIN-NOT: __cuda_module_ctor
// NOGPUBIN-NOT: __cuda_module_dtor
+
+// =============================================================================
+// HIP host-side registration (`buildCUDAModuleCtor` / `buildHIPModuleDtor` /
+// `buildCUDARegisterGlobalFunctions` HIP arms in CIR LoweringPrepare).
+// =============================================================================
+
+// HIP module ctor is registered with the default global-ctor priority.
+// HIP-CIR: cir.global_ctors = [#cir.global_ctor<"__hip_module_ctor", 65535>]
+
+// Runtime function decls.
+// HIP-CIR: cir.func private @atexit(!cir.ptr<!cir.func<()>>) -> !s32i
+// HIP-CIR: cir.func private @__hipUnregisterFatBinary(!cir.ptr<!cir.ptr<!void>>)
+
+// Module dtor: only unregister when the handle is non-null, then null it out.
+// Reuses the SSA value loaded in the entry block for the unregister call.
+// HIP-CIR: cir.func internal private @__hip_module_dtor()
+// HIP-CIR: %[[DH0:.*]] = cir.get_global @__hip_gpubin_handle
+// HIP-CIR: %[[H0:.*]] = cir.load %[[DH0]]
+// HIP-CIR: %[[NULL0:.*]] = cir.const #cir.ptr<null>
+// HIP-CIR: %[[NE:.*]] = cir.cmp ne %[[H0]], %[[NULL0]]
+// HIP-CIR: cir.brcond %[[NE]] ^bb1, ^bb2
+// HIP-CIR: ^bb1:
+// HIP-CIR: cir.call @__hipUnregisterFatBinary(%[[H0]])
+// HIP-CIR: %[[DH1:.*]] = cir.get_global @__hip_gpubin_handle
+// HIP-CIR: cir.store %[[NULL0]], %[[DH1]]
+// HIP-CIR: cir.br ^bb2
+// HIP-CIR: ^bb2:
+// HIP-CIR: cir.return
+
+// __hipRegisterFunction runtime declaration.
+// HIP-CIR: cir.func private @__hipRegisterFunction(!cir.ptr<!cir.ptr<!void>>, !cir.ptr<!void>, !cir.ptr<!void>, !cir.ptr<!void>, !s32i, !cir.ptr<!void>, !cir.ptr<!void>, !cir.ptr<!void>, !cir.ptr<!void>, !cir.ptr<!void>) -> !s32i
+
+// __hip_register_globals: under -fhip-new-launch-api the host-side argument is
+// the kernel-handle GlobalOp shadow (e.g. @_Z10kernelfunciii) — not the
+// device-stub function pointer that the CUDA arm uses.
+// HIP-CIR: cir.global "private" constant cir_private @".str_Z10kernelfunciii" = #cir.const_array<"_Z10kernelfunciii", trailing_zeros>
+// HIP-CIR: cir.func internal private @__hip_register_globals(%[[FATBIN:.*]]: !cir.ptr<!cir.ptr<!void>>
+// HIP-CIR: %[[NULL1:.*]] = cir.const #cir.ptr<null> : !cir.ptr<!void>
+// HIP-CIR: %[[STR_ADDR:.*]] = cir.get_global @".str_Z10kernelfunciii"
+// HIP-CIR: %[[DEVICE_FUNC:.*]] = cir.cast bitcast %[[STR_ADDR]]
+// HIP-CIR: %[[KH:.*]] = cir.get_global @_Z10kernelfunciii
+// HIP-CIR: %[[HOST_FUNC:.*]] = cir.cast bitcast %[[KH]]
+// HIP-CIR: %[[MINUS_ONE:.*]] = cir.const #cir.int<-1> : !s32i
+// HIP-CIR: cir.call @__hipRegisterFunction(%[[FATBIN]], %[[HOST_FUNC]], %[[DEVICE_FUNC]], %[[DEVICE_FUNC]], %[[MINUS_ONE]], %[[NULL1]], %[[NULL1]], %[[NULL1]], %[[NULL1]], %[[NULL1]])
+// HIP-CIR: cir.return
+
+// 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> {alignment = 8 : 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,
+// HIP-CIR-SAME: #cir.global_view<@__hip_fatbin_str> : !cir.ptr<!void>,
+// HIP-CIR-SAME: #cir.ptr<null> : !cir.ptr<!void>
+// HIP-CIR-SAME: }> : !rec_anon_struct {section = ".hipFatBinSegment"}
+
+// HIP-CIR: cir.global "private" internal @__hip_gpubin_handle = #cir.ptr<null> : !cir.ptr<!cir.ptr<!void>>
+// HIP-CIR: cir.func private @__hipRegisterFatBinary(!cir.ptr<!void>) -> !cir.ptr<!cir.ptr<!void>>
+
+// Module ctor: guard registration on a null handle, register globals from the
+// (possibly newly-stored) handle, then atexit(__hip_module_dtor).
+// HIP-CIR: cir.func internal private @__hip_module_ctor()
+// HIP-CIR: %[[GHA:.*]] = cir.get_global @__hip_gpubin_handle
+// HIP-CIR: %[[H:.*]] = cir.load %[[GHA]]
+// HIP-CIR: %[[NULLPTR:.*]] = cir.const #cir.ptr<null>
+// HIP-CIR: %[[EQ:.*]] = cir.cmp eq %[[H]], %[[NULLPTR]]
+// HIP-CIR: cir.brcond %[[EQ]] ^bb1, ^bb2
+// HIP-CIR: ^bb1:
+// HIP-CIR: %[[WRAPPER:.*]] = cir.get_global @__hip_fatbin_wrapper
+// HIP-CIR: %[[VOID_PTR:.*]] = cir.cast bitcast %[[WRAPPER]]
+// HIP-CIR: %[[REG:.*]] = cir.call @__hipRegisterFatBinary(%[[VOID_PTR]])
+// HIP-CIR: %[[GHA2:.*]] = cir.get_global @__hip_gpubin_handle
+// HIP-CIR: cir.store %[[REG]], %[[GHA2]]
+// HIP-CIR: cir.br ^bb2
+// HIP-CIR: ^bb2:
+// HIP-CIR: %[[GHA3:.*]] = cir.get_global @__hip_gpubin_handle
+// HIP-CIR: %[[H2:.*]] = cir.load %[[GHA3]]
+// HIP-CIR: cir.call @__hip_register_globals(%[[H2]])
+// HIP-CIR: %[[DTOR_PTR:.*]] = cir.get_global @__hip_module_dtor
+// HIP-CIR: {{.*}} = cir.call @atexit(%[[DTOR_PTR]])
+// HIP-CIR: cir.return
+
+// 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: @__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
+
+// HIP-OGCG: define internal void @__hip_module_ctor()
+// HIP-OGCG: load ptr, ptr @__hip_gpubin_handle
+// HIP-OGCG: icmp eq ptr {{.*}}, null
+// HIP-OGCG: call ptr @__hipRegisterFatBinary(ptr @__hip_fatbin_wrapper)
+// HIP-OGCG: store ptr {{.*}}, ptr @__hip_gpubin_handle
+// HIP-OGCG: call void @__hip_register_globals(
+// HIP-OGCG: call i32 @atexit(ptr @__hip_module_dtor)
+// HIP-OGCG: ret void
+
+// HIP-OGCG: define internal void @__hip_module_dtor()
+// HIP-OGCG: load ptr, ptr @__hip_gpubin_handle
+// HIP-OGCG: icmp ne ptr {{.*}}, null
+// HIP-OGCG: call void @__hipUnregisterFatBinary
+// HIP-OGCG: store ptr null, ptr @__hip_gpubin_handle
+
+// No GPU binary: no fatbin, no handle, no registration scaffolding.
+// 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
``````````
</details>
https://github.com/llvm/llvm-project/pull/195391
More information about the cfe-commits
mailing list