[clang] [CIR][HIP] Handle HIP module constructor and destructor emission (PR #195391)
David Rivera via cfe-commits
cfe-commits at lists.llvm.org
Wed May 6 01:30:01 PDT 2026
https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/195391
>From 4248db26a4828a589812d51ad3cba7a2043b05a8 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Fri, 1 May 2026 20:33:59 -0400
Subject: [PATCH 1/4] [CIR][HIP] Handle HIP module constructor and destructor
emission
---
.../Dialect/Transforms/LoweringPrepare.cpp | 152 +++++++++++++++++-
clang/test/CIR/CodeGenCUDA/device-stub.cu | 121 ++++++++++++++
2 files changed, 268 insertions(+), 5 deletions(-)
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index dc56a2f21aced..e60dfd778fc83 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())
@@ -2213,7 +2342,20 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
builder.createGetGlobal(deviceFuncStr), voidPtrTy);
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));
+ mlir::Value hostFunc =
+ builder.createBitcast(builder.createGetGlobal(funcHandle), voidPtrTy);
+ builder.createCallOp(
+ loc, cudaRegisterFunction,
+ {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
+ ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)),
+ cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
} else {
mlir::Value hostFunc = builder.createBitcast(
GetGlobalOp::create(
diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu
index 0e95c74324592..e9cf3b23e08f0 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -19,6 +19,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) {}
@@ -134,3 +146,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
>From b2dd24894f89d7d4fd6484fcd034a2c197e0c6ec Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Fri, 1 May 2026 20:36:26 -0400
Subject: [PATCH 2/4] fix fmt
---
clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp | 9 ++++-----
1 file changed, 4 insertions(+), 5 deletions(-)
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index e60dfd778fc83..6e3f800d1556c 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -2344,11 +2344,10 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
if (isHIP) {
// 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));
+ // `@_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));
mlir::Value hostFunc =
builder.createBitcast(builder.createGetGlobal(funcHandle), voidPtrTy);
builder.createCallOp(
>From 6ab5d93ae324a581945ac29a210161cea80066be Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Sun, 3 May 2026 17:55:00 -0400
Subject: [PATCH 3/4] Simplify `RegisterFunction` call on target divergance
---
.../Dialect/Transforms/LoweringPrepare.cpp | 20 ++++++++-----------
1 file changed, 8 insertions(+), 12 deletions(-)
diff --git a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
index 6e3f800d1556c..63f961b0c0647 100644
--- a/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
+++ b/clang/lib/CIR/Dialect/Transforms/LoweringPrepare.cpp
@@ -2341,6 +2341,7 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
mlir::Value deviceFunc = builder.createBitcast(
builder.createGetGlobal(deviceFuncStr), voidPtrTy);
+ mlir::Value hostFunc;
if (isHIP) {
// Under HIP, the kernel-handle is a GlobalOp shadow created by CIR
// codegen and named with the kernel-reference mangled name (e.g.
@@ -2348,25 +2349,20 @@ void LoweringPreparePass::buildCUDARegisterGlobalFunctions(
// `_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));
- mlir::Value hostFunc =
+ hostFunc =
builder.createBitcast(builder.createGetGlobal(funcHandle), voidPtrTy);
- builder.createCallOp(
- loc, cudaRegisterFunction,
- {fatbinHandle, hostFunc, deviceFunc, deviceFunc,
- ConstantOp::create(builder, loc, IntAttr::get(intTy, -1)),
- cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr, cirNullPtr});
} 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});
}
}
>From 8079fea07a8e34b0ba9472d177e6309b24a227b7 Mon Sep 17 00:00:00 2001
From: David Rivera <davidriverg at gmail.com>
Date: Wed, 6 May 2026 04:29:42 -0400
Subject: [PATCH 4/4] Add HIP LLVM lowering and proper HIP host kernel handle
emission.
---
clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 4 +--
clang/test/CIR/CodeGenCUDA/device-stub.cu | 42 +++++++++++++++++++++--
clang/test/CIR/CodeGenCUDA/kernel-call.cu | 4 +--
clang/test/CIR/CodeGenHIP/simple.cpp | 2 +-
4 files changed, 44 insertions(+), 8 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
index 1bae5e470aadd..0466c212fe165 100644
--- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp
@@ -333,9 +333,9 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn,
CIRGenBuilderTy &builder = cgm.getBuilder();
StringRef globalName = cgm.getMangledName(
gd.getWithKernelReferenceKind(KernelReferenceKind::Kernel));
+ cir::PointerType fnPtrTy = builder.getPointerTo(fn.getFunctionType());
cir::GlobalOp globalOp = CIRGenModule::createGlobalOp(
- cgm, fn.getLoc(), globalName, fn.getFunctionType(),
- /*isConstant=*/true);
+ cgm, fn.getLoc(), globalName, fnPtrTy, /*isConstant=*/true);
globalOp->setAttr("alignment", builder.getI64IntegerAttr(
cgm.getPointerAlign().getQuantity()));
diff --git a/clang/test/CIR/CodeGenCUDA/device-stub.cu b/clang/test/CIR/CodeGenCUDA/device-stub.cu
index e9cf3b23e08f0..b635f5931df77 100644
--- a/clang/test/CIR/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CIR/CodeGenCUDA/device-stub.cu
@@ -23,6 +23,10 @@
// 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 -fclangir -emit-llvm %s -x hip \
+// RUN: -fhip-new-launch-api -fcuda-include-gpubinary %t -o %t.hip-cir.ll
+// RUN: FileCheck --input-file=%t.hip-cir.ll %s --check-prefix=HIP-LLVM
+//
// 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
@@ -181,12 +185,12 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// __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.global "private" constant cir_private @".str_Z10kernelfunciii" = #cir.const_array<"_Z10kernelfunciii" : !cir.array<!u8i x 18>, trailing_zeros> : !cir.array<!u8i x 18>
// 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: %[[KH:.*]] = cir.get_global @_Z10kernelfunciii : !cir.ptr<!cir.ptr<!cir.func<(!s32i, !s32i, !s32i)>>>
// 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]])
@@ -194,7 +198,7 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// 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 "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 constant cir_private @__hip_fatbin_wrapper = #cir.const_record<{
// HIP-CIR-SAME: #cir.int<1212764230> : !s32i,
// HIP-CIR-SAME: #cir.int<1> : !s32i,
@@ -228,6 +232,8 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// HIP-CIR: {{.*}} = cir.call @atexit(%[[DTOR_PTR]])
// HIP-CIR: cir.return
+// 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: @__hip_fatbin_wrapper = internal constant { i32, i32, ptr, ptr } { i32 1212764230, i32 1, ptr @{{.*}}, ptr null }, section ".hipFatBinSegment"
@@ -249,6 +255,36 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
// HIP-OGCG: call void @__hipUnregisterFatBinary
// 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: @__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
+// HIP-LLVM: @llvm.global_ctors = appending global {{.*}}@__hip_module_ctor
+
+// HIP-LLVM: define internal void @__hip_module_dtor()
+// HIP-LLVM: load ptr, ptr @__hip_gpubin_handle
+// HIP-LLVM: icmp ne ptr {{.*}}, null
+// HIP-LLVM: br i1 {{.*}}, label %{{.*}}, label %{{.*}}
+// HIP-LLVM: call void @__hipUnregisterFatBinary(ptr {{.*}})
+// HIP-LLVM: store ptr null, ptr @__hip_gpubin_handle
+// HIP-LLVM: ret void
+
+// HIP-LLVM: define internal void @__hip_register_globals(ptr %[[FATBIN:.*]])
+// HIP-LLVM: call{{.*}}@__hipRegisterFunction(ptr %[[FATBIN]], ptr @_Z10kernelfunciii, ptr @{{.*}}, ptr @{{.*}}, i32 -1, ptr null, ptr null, ptr null, ptr null, ptr null)
+// HIP-LLVM: ret void
+
+// HIP-LLVM: define internal void @__hip_module_ctor()
+// HIP-LLVM: load ptr, ptr @__hip_gpubin_handle
+// HIP-LLVM: icmp eq ptr {{.*}}, null
+// HIP-LLVM: br i1 {{.*}}, label %{{.*}}, label %{{.*}}
+// HIP-LLVM: call ptr @__hipRegisterFatBinary(ptr @__hip_fatbin_wrapper)
+// HIP-LLVM: store ptr {{.*}}, ptr @__hip_gpubin_handle
+// HIP-LLVM: load ptr, ptr @__hip_gpubin_handle
+// HIP-LLVM: call void @__hip_register_globals(ptr {{.*}})
+// HIP-LLVM: call i32 @atexit(ptr @__hip_module_dtor)
+// HIP-LLVM: ret void
+
// No GPU binary: no fatbin, no handle, no registration scaffolding.
// HIP-NOGPUBIN-NOT: __hip_fatbin
// HIP-NOGPUBIN-NOT: __hip_gpubin_handle
diff --git a/clang/test/CIR/CodeGenCUDA/kernel-call.cu b/clang/test/CIR/CodeGenCUDA/kernel-call.cu
index 34719a2d3acb1..924f26c6dafbc 100644
--- a/clang/test/CIR/CodeGenCUDA/kernel-call.cu
+++ b/clang/test/CIR/CodeGenCUDA/kernel-call.cu
@@ -66,11 +66,11 @@
// CUDA-PTH: cir.call @cudaLaunchKernel_ptsz
//
-// HIP-NEW: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)>
+// HIP-NEW: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.ptr<!cir.func<(!s32i, !cir.float)>> {alignment = 8 : i64}
// HIP-NEW-LABEL: cir.func {{.*}} @_Z21__device_stub__kernelif
// HIP-NEW: cir.alloca !cir.ptr<!rec_hipStream>, {{.*}} ["stream"]
// HIP-NEW: cir.call @__hipPopCallConfiguration({{.*}}) : (!cir.ptr<!rec_dim3>, !cir.ptr<!rec_dim3>, !cir.ptr<!u64i>, !cir.ptr<!cir.ptr<!rec_hipStream>>) -> !s32i
-// HIP-NEW: cir.get_global @_Z6kernelif : !cir.ptr<!cir.func<(!s32i, !cir.float)>>
+// HIP-NEW: cir.get_global @_Z6kernelif : !cir.ptr<!cir.ptr<!cir.func<(!s32i, !cir.float)>>>
// HIP-NEW: cir.call @hipLaunchKernel({{.*}}) : (!cir.ptr<!void> {{.*}}, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>{{.*}}, !u64i{{.*}}, !cir.ptr<!rec_hipStream>{{.*}}) -> (!u32i {llvm.noundef})
// HIP-PTH: cir.call @hipLaunchKernel_spt
diff --git a/clang/test/CIR/CodeGenHIP/simple.cpp b/clang/test/CIR/CodeGenHIP/simple.cpp
index f4d2645c7803a..e684fc2104f4f 100644
--- a/clang/test/CIR/CodeGenHIP/simple.cpp
+++ b/clang/test/CIR/CodeGenHIP/simple.cpp
@@ -46,7 +46,7 @@ __global__ void global_fn(int a) {}
// CIR-HOST: %[[#CIRKernelArgs:]] = cir.alloca {{.*}}"kernel_args"
// CIR-HOST: %[[#Decayed:]] = cir.cast array_to_ptrdecay %[[#CIRKernelArgs]]
// CIR-HOST: cir.call @__hipPopCallConfiguration
-// CIR-HOST: cir.get_global @_Z9global_fni : !cir.ptr<!cir.func<(!s32i)>>
+// CIR-HOST: cir.get_global @_Z9global_fni : !cir.ptr<!cir.ptr<!cir.func<(!s32i)>>>
// CIR-HOST: cir.call @hipLaunchKernel
// OGCG-HOST: define dso_local void @_Z24__device_stub__global_fni
More information about the cfe-commits
mailing list