[clang] [CIR][HIP] Use the kernel handle as a kernel's address on the host (PR #228377)
via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 2 03:23:53 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clangir
Author: Steffen Larsen (steffenlarsen)
<details>
<summary>Changes</summary>
For HIP, a __global__ function referenced from host code is represented by its kernel handle. CIR used the address of the device stub instead, so APIs that take a kernel pointer failed.
Match classic codegen:
- emitFunctionDeclLValue gives the address of the kernel handle.
- Constant initializers, such as tables of kernel pointers, refer to the kernel handle.
- A launch through a kernel pointer (f<<<...>>>) loads the device stub from the handle and calls it.
CUDA has no separate kernel handle, instead the address of a kernel is the device stub. The new test covers both HIP and CUDA.
Assisted-by: Claude Opus 5.5
---
Full diff: https://github.com/llvm/llvm-project/pull/228377.diff
3 Files Affected:
- (modified) clang/lib/CIR/CodeGen/CIRGenExpr.cpp (+38-3)
- (modified) clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp (+10-2)
- (added) clang/test/CIR/CodeGenCUDA/kernel-address.cu (+75)
``````````diff
diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
index 08901b30084bd..5bb0b7e4907dc 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp
@@ -983,8 +983,26 @@ static LValue emitFunctionDeclLValue(CIRGenFunction &cgf, const Expr *e,
mlir::Type fnTy = funcOp.getFunctionType();
mlir::Type ptrTy = cir::PointerType::get(fnTy);
- mlir::Value addr = cir::GetGlobalOp::create(cgf.getBuilder(), loc, ptrTy,
- funcOp.getSymName());
+ mlir::Value addr;
+
+ // On the HIP host, a reference to a __global__ kernel must resolve to the
+ // address of the kernel handle registered with the offload runtime, not
+ // the device stub's own address. CUDA uses the device stub itself as the
+ // kernel handle.
+ if (cgf.cgm.getLangOpts().HIP && !cgf.cgm.getLangOpts().CUDAIsDevice &&
+ fd->hasAttr<CUDAGlobalAttr>()) {
+ auto handle = mlir::cast<cir::GlobalOp>(
+ cgf.cgm.getCUDARuntime().getKernelHandle(funcOp, gd));
+ cir::PointerType handlePtrTy = cir::PointerType::get(handle.getSymType());
+ mlir::Value handleAddr = cir::GetGlobalOp::create(
+ cgf.getBuilder(), loc, handlePtrTy, handle.getSymName());
+ addr = cir::CastOp::create(cgf.getBuilder(), loc, ptrTy,
+ cir::CastKind::bitcast, handleAddr);
+ }
+
+ if (!addr)
+ addr = cir::GetGlobalOp::create(cgf.getBuilder(), loc, ptrTy,
+ funcOp.getSymName());
if (funcOp.getFunctionType() != cgf.convertType(fd->getType())) {
fnTy = cgf.convertType(fd->getType());
@@ -2431,7 +2449,24 @@ RValue CIRGenFunction::emitCall(clang::QualType calleeTy,
}
assert(!cir::MissingFeatures::opCallFnInfoOpts());
- assert(!cir::MissingFeatures::hip());
+
+ // HIP function pointer contains kernel handle when it is used in triple
+ // chevron. The kernel stub needs to be loaded from kernel handle and used
+ // as callee.
+ const clang::Decl *targetDecl =
+ origCallee.getAbstractInfo().getCalleeDecl().getDecl();
+ if (getLangOpts().HIP && !getLangOpts().CUDAIsDevice &&
+ isa<CUDAKernelCallExpr>(e) &&
+ (!targetDecl || !isa<FunctionDecl>(targetDecl))) {
+ mlir::Value handleAddr = callee.getFunctionPointer()->getResult(0);
+ mlir::Location loc = getLoc(e->getSourceRange());
+ auto handlePtrTy = mlir::cast<cir::PointerType>(handleAddr.getType());
+ mlir::Value handleAddrAddr =
+ builder.createBitcast(handleAddr, cir::PointerType::get(handlePtrTy));
+ cir::LoadOp stub = builder.createLoad(
+ loc, Address(handleAddrAddr, handlePtrTy, getPointerAlign()));
+ callee.setFunctionPointer(stub.getOperation());
+ }
cir::CIRCallOpInterface callOp;
RValue callResult = emitCall(funcInfo, callee, returnValue, args, &callOp,
diff --git a/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp b/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp
index 46aabba5b567a..6b422e76b04ab 100644
--- a/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp
@@ -918,9 +918,17 @@ ConstantLValueEmitter::tryEmitBase(const APValue::LValueBase &base) {
mlir::isa<cir::PointerType>(destTy)
? mlir::cast<cir::PointerType>(destTy)
: cir::PointerType::get(fop.getFunctionType());
+ mlir::StringAttr symName = fop.getSymNameAttr();
+ // On the HIP host, the address of a kernel is the address of its kernel
+ // handle, not of its device stub. CUDA uses the device stub itself as
+ // the kernel handle.
+ if (cgm.getLangOpts().HIP && !cgm.getLangOpts().CUDAIsDevice &&
+ fd->hasAttr<CUDAGlobalAttr>())
+ symName = mlir::cast<cir::GlobalOp>(
+ cgm.getCUDARuntime().getKernelHandle(fop, fd))
+ .getSymNameAttr();
return cir::GlobalViewAttr::get(
- ptrTy,
- mlir::FlatSymbolRefAttr::get(mlirContext, fop.getSymNameAttr()));
+ ptrTy, mlir::FlatSymbolRefAttr::get(mlirContext, symName));
}
if (auto *vd = dyn_cast<VarDecl>(d)) {
diff --git a/clang/test/CIR/CodeGenCUDA/kernel-address.cu b/clang/test/CIR/CodeGenCUDA/kernel-address.cu
new file mode 100644
index 0000000000000..39ce76aba653d
--- /dev/null
+++ b/clang/test/CIR/CodeGenCUDA/kernel-address.cu
@@ -0,0 +1,75 @@
+#include "Inputs/cuda.h"
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefixes=CIR,CIR-HIP
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-HIP
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -fhip-new-launch-api -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-HIP
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda -target-sdk-version=12.0 -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck %s --check-prefixes=CIR,CIR-CUDA
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda -target-sdk-version=12.0 -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-CUDA
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x cuda -target-sdk-version=12.0 -emit-llvm %s -o - \
+// RUN: | FileCheck %s --check-prefixes=LLVM,LLVM-CUDA
+
+// On the host, the address of a kernel is the pointer the offload runtime
+// registers for it. For HIP that is the kernel handle, a global holding the
+// address of the device stub, so launching through a kernel pointer loads the
+// stub from the handle. For CUDA it is the device stub itself.
+
+__global__ void kern(int *p) {}
+template <class T> __global__ void tkern(T *p) {}
+
+const void *table[] = {(const void *)kern, (const void *)tkern<float>};
+
+// CIR-HIP: cir.global external @table = #cir.const_array<[#cir.global_view<@_Z4kernPi> : !cir.ptr<!void>, #cir.global_view<@_Z5tkernIfEvPT_> : !cir.ptr<!void>]>
+// CIR-CUDA: cir.global external @table = #cir.const_array<[#cir.global_view<@_Z19__device_stub__kernPi> : !cir.ptr<!void>, #cir.global_view<@_Z20__device_stub__tkernIfEvPT_> : !cir.ptr<!void>]>
+
+// LLVM-HIP: @table = global [2 x ptr] [ptr @_Z4kernPi, ptr @_Z5tkernIfEvPT_]
+// LLVM-CUDA: @table = global [2 x ptr] [ptr @_Z19__device_stub__kernPi, ptr @_Z20__device_stub__tkernIfEvPT_]
+
+const void *addr() { return (const void *)kern; }
+
+// CIR-LABEL: cir.func {{.*}} @_Z4addrv(
+// CIR-HIP: %[[HANDLE:.+]] = cir.get_global @_Z4kernPi : !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>>
+// CIR-HIP: cir.cast bitcast %[[HANDLE]] : !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>> -> !cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>
+// CIR-CUDA: cir.get_global @_Z19__device_stub__kernPi : !cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>
+
+// LLVM-LABEL: define {{.*}} ptr @_Z4addrv(
+// LLVM-HIP: {{(ret|store)}} ptr @_Z4kernPi
+// LLVM-CUDA: {{(ret|store)}} ptr @_Z19__device_stub__kernPi
+
+const void *tmpl_addr() { return (const void *)tkern<float>; }
+
+// CIR-LABEL: cir.func {{.*}} @_Z9tmpl_addrv(
+// CIR-HIP: cir.get_global @_Z5tkernIfEvPT_ : !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!cir.float>)>>>
+// CIR-CUDA: cir.get_global @_Z20__device_stub__tkernIfEvPT_ : !cir.ptr<!cir.func<(!cir.ptr<!cir.float>)>>
+
+// LLVM-LABEL: define {{.*}} ptr @_Z9tmpl_addrv(
+// LLVM-HIP: {{(ret|store)}} ptr @_Z5tkernIfEvPT_
+// LLVM-CUDA: {{(ret|store)}} ptr @_Z20__device_stub__tkernIfEvPT_
+
+void indirect(void (*f)(int *), int *p) { f<<<1, 1>>>(p); }
+
+// CIR-LABEL: cir.func {{.*}} @_Z8indirectPFvPiES_(
+// CIR: %[[F:.+]] = cir.load {{.*}} : !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>>, !cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>
+// CIR-HIP: %[[HANDLE:.+]] = cir.cast bitcast %[[F]] : !cir.ptr<!cir.func<(!cir.ptr<!s32i>)>> -> !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>>
+// CIR-HIP: %[[STUB:.+]] = cir.load {{.*}} %[[HANDLE]] : !cir.ptr<!cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>>, !cir.ptr<!cir.func<(!cir.ptr<!s32i>)>>
+// CIR-HIP: cir.call %[[STUB]](
+// CIR-CUDA: cir.call %[[F]](
+
+// LLVM-LABEL: define {{.*}} void @_Z8indirectPFvPiES_(
+// LLVM: [[F:%.+]] = load ptr, ptr %{{.+}}, align 8
+// LLVM-HIP: [[STUB:%.+]] = load ptr, ptr [[F]], align 8
+// LLVM-HIP: call void [[STUB]](
+// LLVM-CUDA: call void [[F]](
+
+void direct(int *p) { kern<<<1, 1>>>(p); }
+
+// CIR-LABEL: cir.func {{.*}} @_Z6directPi(
+// CIR: cir.call @_Z19__device_stub__kernPi(
+
+// LLVM-LABEL: define {{.*}} void @_Z6directPi(
+// LLVM: call void @_Z19__device_stub__kernPi(
``````````
</details>
https://github.com/llvm/llvm-project/pull/228377
More information about the cfe-commits
mailing list