Author: Steffen Larsen Date: 2026-10-05T07:57:35Z New Revision: 8ddf7e4a2ae80cc90d97a1f30d23967dcb4870ec
URL: https://github.com/llvm/llvm-project/commit/8ddf7e4a2ae80cc90d97a1f30d23967dcb4870ec DIFF: https://github.com/llvm/llvm-project/commit/8ddf7e4a2ae80cc90d97a1f30d23967dcb4870ec.diff LOG: [CIR][HIP] Use the kernel handle as a kernel's address on the host (#228377) 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 Signed-off-by: Steffen Holst Larsen <[email protected]> Added: clang/test/CIR/CodeGenCUDA/kernel-address.cu Modified: clang/lib/CIR/CodeGen/CIRGenExpr.cpp clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp index f5eafaf393c02..1f738f9737d31 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExpr.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExpr.cpp @@ -985,8 +985,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()); @@ -2448,7 +2466,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 22b2e79f771ac..afc6b8058066d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenExprConstant.cpp @@ -968,9 +968,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( _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
