https://github.com/RiverDave updated https://github.com/llvm/llvm-project/pull/179823
>From 2a625851b643eb9a4fffa496605e010172cee80e Mon Sep 17 00:00:00 2001 From: David Rivera <[email protected]> Date: Wed, 4 Feb 2026 19:15:32 -0500 Subject: [PATCH] [CIR][HIP] Add Stub body emission test coverage and Fix kernelHandle storage --- clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp | 13 +++++-------- clang/test/CIR/CodeGenCUDA/kernel-call.cu | 15 ++++++++++++--- 2 files changed, 17 insertions(+), 11 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp index ad5da0d11ff02..451c28c3cccc1 100644 --- a/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenCUDANV.cpp @@ -121,9 +121,6 @@ void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf, if (cgm.getLangOpts().OffloadViaLLVM) cgm.errorNYI("CIRGenNVCUDARuntime: Offload via LLVM"); - if (cgm.getLangOpts().HIP) - cgm.errorNYI("CIRGenNVCUDARuntime: HIP Support"); - CIRGenBuilderTy &builder = cgm.getBuilder(); mlir::Location loc = fn.getLoc(); @@ -213,7 +210,8 @@ void CIRGenNVCUDARuntime::emitDeviceStubBodyNew(CIRGenFunction &cgf, cir::PointerType kernelTy = cir::PointerType::get(globalOp.getSymType()); mlir::Value kernelVal = cir::GetGlobalOp::create(builder, loc, kernelTy, globalOp.getSymName()); - return kernelVal; + mlir::Value func = builder.createBitcast(kernelVal, cgm.voidPtrTy); + return func; } if (cir::FuncOp funcOp = llvm::dyn_cast_or_null<cir::FuncOp>( kernelHandles[fn.getSymName()])) { @@ -325,10 +323,9 @@ mlir::Operation *CIRGenNVCUDARuntime::getKernelHandle(cir::FuncOp fn, CIRGenBuilderTy &builder = cgm.getBuilder(); StringRef globalName = cgm.getMangledName( gd.getWithKernelReferenceKind(KernelReferenceKind::Kernel)); - const VarDecl *varDecl = llvm::dyn_cast_or_null<VarDecl>(gd.getDecl()); - cir::GlobalOp globalOp = - cgm.getOrCreateCIRGlobal(globalName, fn.getFunctionType().getReturnType(), - LangAS::Default, varDecl, NotForDefinition); + cir::GlobalOp globalOp = CIRGenModule::createGlobalOp( + cgm, fn.getLoc(), globalName, fn.getFunctionType(), + /*isConstant=*/true); globalOp->setAttr("alignment", builder.getI64IntegerAttr( cgm.getPointerAlign().getQuantity())); diff --git a/clang/test/CIR/CodeGenCUDA/kernel-call.cu b/clang/test/CIR/CodeGenCUDA/kernel-call.cu index ccc33461567bf..05c48625335a6 100644 --- a/clang/test/CIR/CodeGenCUDA/kernel-call.cu +++ b/clang/test/CIR/CodeGenCUDA/kernel-call.cu @@ -1,16 +1,19 @@ // Based on clang/test/CodeGenCUDA/kernel-call.cu. -// Tests device stub body emission for CUDA kernels. +// Tests device stub body emission for CUDA and HIP kernels. // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -target-sdk-version=9.2 \ // RUN: -emit-cir %s -x cuda -o %t.cir // RUN: FileCheck --input-file=%t.cir %s --check-prefix=CUDA-NEW +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fhip-new-launch-api \ +// RUN: -x hip -emit-cir %s -o %t.hip.cir +// RUN: FileCheck --input-file=%t.hip.cir %s --check-prefix=HIP-NEW + #include "Inputs/cuda.h" // TODO: Test CUDA legacy (< 9.0) when legacy stub body is implemented -// TODO: Test HIP when HIP stub body support is complete // Check that the stub function is generated with the correct name // CUDA-NEW-LABEL: cir.func {{.*}} @_Z21__device_stub__kernelif @@ -48,5 +51,11 @@ // Check cudaLaunchKernel is called with all 6 arguments: // func ptr, gridDim, blockDim, args, sharedMem, stream // CUDA-NEW: cir.call @cudaLaunchKernel({{.*}}) : (!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_cudaStream>) -> (!u32i {llvm.noundef}) - +// +// HIP-NEW: cir.global constant external @_Z6kernelif = #cir.global_view<@_Z21__device_stub__kernelif> : !cir.func<(!s32i, !cir.float)> +// 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.call @hipLaunchKernel({{.*}}) : (!cir.ptr<!void>, !rec_dim3, !rec_dim3, !cir.ptr<!cir.ptr<!void>>, !u64i, !cir.ptr<!rec_hipStream>) -> (!u32i {llvm.noundef}) __global__ void kernel(int x, float y) {} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
