https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/226435
>From fae8370258f54af17a63d1fd333a0a28fc8290bb Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 1 Sep 2026 09:00:32 -0500 Subject: [PATCH 1/3] [CIR][AMDGPU] Route printf in device code through the OpenCL runtime This commit implements CIR support for device-side printf on AMDGPU targets. Previously, any call to printf fell through to a library call, which does not exist on the device, so any kernel calling printf failed to build. For OGCG, the AMDGPU target routes printf calls through the OpenCL printf runtime, which takes the format string and a buffer of appended arguments. This implementation leverages this lowering by first emitting a call to a marker function `__cir_amdgpu_printf` during CIRGen, which is then lowered to the appropriate OpenCL printf runtime call during the LLVM IR lowering, leveraging the existing `llvm::emitAMDGPUPrintfCall` function. --- clang/include/clang/CIR/LowerToLLVM.h | 6 +++ clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 2 +- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 39 ++++++++++++++++ clang/lib/CIR/CodeGen/CIRGenFunction.h | 3 ++ clang/lib/CIR/FrontendAction/CIRGenAction.cpp | 8 +++- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 46 +++++++++++++++++++ clang/test/CIR/CodeGenHIP/device-printf.hip | 29 ++++++++++++ 7 files changed, 130 insertions(+), 3 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/device-printf.hip diff --git a/clang/include/clang/CIR/LowerToLLVM.h b/clang/include/clang/CIR/LowerToLLVM.h index d97ac5265c282b..68a9460281baf5 100644 --- a/clang/include/clang/CIR/LowerToLLVM.h +++ b/clang/include/clang/CIR/LowerToLLVM.h @@ -33,6 +33,12 @@ lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule, llvm::LLVMContext &llvmCtx, bool enableOpenMP, llvm::StringRef mlirSaveTempsOutFile = {}, llvm::vfs::FileSystem *fs = nullptr); + +// Expand calls to the internal __cir_amdgpu_printf marker CIRGen emits for a +// device-side printf into the real AMDGPU sequence. Must run after the +// module has been translated to LLVM IR and before device-library bitcode +// linking. +void expandAMDGPUDevicePrintf(llvm::Module &module); } // namespace direct } // namespace cir diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 2b223c8ae19392..d4d31a7f2ae650 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -3105,7 +3105,7 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, if ((getTarget().getTriple().isAMDGCN() || getTarget().getTriple().isSPIRV()) && getLangOpts().HIP) - return errorBuiltinNYI(*this, e, builtinID); + return RValue::get(emitAMDGPUDevicePrintfCallExpr(e)); } break; case Builtin::BI__builtin_canonicalize: diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 58ef4b1bd4b61d..a0daed0938a335 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1096,3 +1096,42 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return std::nullopt; } } + +// Emit AMDGPU printf CIR stand-in function call. This stand-in function call is +// lowered to the appropriate call structure during LLVM IR lowering. +mlir::Value +CIRGenFunction::emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr) { + assert(cgm.getTriple().isAMDGCN() || + (cgm.getTriple().isSPIRV() && + cgm.getTriple().getVendor() == llvm::Triple::AMD)); + assert(expr->getBuiltinCallee() == Builtin::BIprintf || + expr->getBuiltinCallee() == Builtin::BI__builtin_printf); + assert(expr->getNumArgs() >= 1); + + const FunctionProtoType *funcPrototype = + expr->getDirectCallee()->getType()->getAs<FunctionProtoType>(); + CallArgList args; + emitCallArgs(args, funcPrototype, expr->arguments(), expr->getDirectCallee()); + + mlir::Location loc = getLoc(expr->getBeginLoc()); + + // We don't know how to emit non-scalar varargs. + bool hasNonScalar = llvm::any_of(args, [&](const CallArg &a) { + return a.hasLValue() || !a.getKnownRValue().isScalar(); + }); + if (hasNonScalar) { + cgm.errorUnsupported(expr, "non-scalar args to printf"); + return builder.getConstInt(loc, builder.getSInt32Ty(), 0); + } + + llvm::SmallVector<mlir::Value, 8> callArgs; + for (const CallArg &a : args) + callArgs.push_back(a.getKnownRValue().getValue()); + + // int __cir_amdgpu_printf(char *format, ...); + auto fnTy = cir::FuncType::get({cir::PointerType::get(builder.getSInt8Ty())}, + builder.getSInt32Ty(), + /*isVarArg=*/true); + cir::FuncOp fn = cgm.createRuntimeFunction(fnTy, "__cir_amdgpu_printf"); + return builder.createCallOp(loc, fn, callArgs).getResult(); +} diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index b232abf5d9299d..0f657533748432 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -2363,6 +2363,9 @@ class CIRGenFunction : public CIRGenTypeCache { /// Emit a device-side printf call for NVPTX targets. mlir::Value emitNVPTXDevicePrintfCallExpr(const CallExpr *expr); + /// Emit a device-side printf call for AMDGPU targets. + mlir::Value emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr); + LValue emitOpaqueValueLValue(const OpaqueValueExpr *e); LValue emitConditionalOperatorLValue(const AbstractConditionalOperator *expr); diff --git a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp index 240601f9834e5d..e25a251150cf9a 100644 --- a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp +++ b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp @@ -65,8 +65,12 @@ lowerFromCIRToLLVMIR(mlir::ModuleOp MLIRModule, llvm::LLVMContext &LLVMCtx, bool EnableOpenMP, llvm::StringRef mlirSaveTempsOutFile = {}, llvm::vfs::FileSystem *fs = nullptr) { - return direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx, EnableOpenMP, - mlirSaveTempsOutFile, fs); + std::unique_ptr<llvm::Module> LLVMModule = + direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx, EnableOpenMP, + mlirSaveTempsOutFile, fs); + if (LLVMModule) + direct::expandAMDGPUDevicePrintf(*LLVMModule); + return LLVMModule; } class CIRGenConsumer : public clang::ASTConsumer { diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index bd529f71b38edf..c55de0d19fc5de 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -49,12 +49,15 @@ #include "llvm/ADT/MapVector.h" #include "llvm/ADT/StringMap.h" #include "llvm/ADT/TypeSwitch.h" +#include "llvm/IR/IRBuilder.h" #include "llvm/IR/Module.h" #include "llvm/Support/Casting.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/TimeProfiler.h" #include "llvm/Support/VirtualFileSystem.h" #include "llvm/Support/raw_ostream.h" +#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h" +#include "llvm/Transforms/Utils/Local.h" using namespace cir; using namespace llvm; @@ -5914,6 +5917,49 @@ void populateCIRToLLVMPasses(mlir::OpPassManager &pm, bool enableOpenMP) { pm.addPass(mlir::omp::createHostOpFilteringPass()); } +// Expand calls to the internal __cir_amdgpu_printf marker CIRGen emits for a +// device-side printf into the real AMDGPU sequence. +void expandAMDGPUDevicePrintf(llvm::Module &module) { + llvm::Function *marker = module.getFunction("__cir_amdgpu_printf"); + if (!marker) + return; + + // CIR records the requested lowering as a module flag. + bool isBuffered = false; + if (llvm::Metadata *md = + module.getModuleFlag(cir::CIRDialect::getAMDGPUPrintfKindAttrName())) + if (auto *mdStr = llvm::dyn_cast<llvm::MDString>(md)) + isBuffered = mdStr->getString() == "buffered"; + + // Snapshot marker's users before mutating anything. + llvm::SmallVector<llvm::User *, 8> users(marker->user_begin(), + marker->user_end()); + for (llvm::User *u : users) { + auto *cb = llvm::cast<llvm::CallBase>(u); + + // CIR emits an invoke rather than a call when the marker call site is + // inside a region that requires unwinding, even though the device printf + // sequence emitted below can never throw. Normalize those invokes to + // calls. + llvm::CallInst *ci = llvm::dyn_cast<llvm::CallInst>(cb); + if (!ci) { + assert(llvm::isa<llvm::InvokeInst>(cb) && + "unexpected non-call user of printf marker"); + ci = llvm::changeToCall(llvm::cast<llvm::InvokeInst>(cb)); + } + + llvm::IRBuilder<> irb(ci); + llvm::SmallVector<llvm::Value *, 8> args(ci->args()); + llvm::Value *res = llvm::emitAMDGPUPrintfCall(irb, args, isBuffered); + if (res && !ci->use_empty()) + ci->replaceAllUsesWith(res); + ci->eraseFromParent(); + } + + assert(marker->use_empty() && "printf marker should have no remaining uses"); + marker->eraseFromParent(); +} + std::unique_ptr<llvm::Module> lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule, LLVMContext &llvmCtx, bool enableOpenMP, StringRef mlirSaveTempsOutFile, diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip new file mode 100644 index 00000000000000..5e7b3039cccf40 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -0,0 +1,29 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ +// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ +// RUN: -std=c++11 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// printf fell through to a library call, which does not exist on the device: +// the reference survived to the device link and failed there as an undefined +// symbol. AMDGPU implements it through the OpenCL printf runtime. +// +// CIRGen cannot build the sequence itself -- llvm::emitAMDGPUPrintfCall works +// through an IRBuilder -- so it emits a marker call that is expanded once the +// module has been translated to LLVM IR. It has to happen there rather than +// later, because the __ockl_* calls are what make the device-library linker +// pull in ockl. +// +// The checks are shared with the classic CodeGen run line. + +#define __device__ __attribute__((device)) +extern "C" __device__ int printf(const char *, ...); + +// LLVM-LABEL: @_Z1pi +// LLVM: call i64 @__ockl_printf_begin(i64 +// LLVM: call i64 @__ockl_printf_append_string_n(i64 +// LLVM: call i64 @__ockl_printf_append_args(i64 +// LLVM-NOT: call{{.*}}@__cir_amdgpu_printf +__device__ void p(int v) { printf("%d\n", v); } >From 11652782b7f435a11921d549c5d18a000e5ba1bc Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Fri, 25 Sep 2026 06:43:03 -0500 Subject: [PATCH 2/3] Remove unnecessary over-verbose test comment Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/CIR/CodeGenHIP/device-printf.hip | 8 -------- 1 file changed, 8 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 5e7b3039cccf40..539106d7200584 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -9,14 +9,6 @@ // printf fell through to a library call, which does not exist on the device: // the reference survived to the device link and failed there as an undefined // symbol. AMDGPU implements it through the OpenCL printf runtime. -// -// CIRGen cannot build the sequence itself -- llvm::emitAMDGPUPrintfCall works -// through an IRBuilder -- so it emits a marker call that is expanded once the -// module has been translated to LLVM IR. It has to happen there rather than -// later, because the __ockl_* calls are what make the device-library linker -// pull in ockl. -// -// The checks are shared with the classic CodeGen run line. #define __device__ __attribute__((device)) extern "C" __device__ int printf(const char *, ...); >From 6115a12e18e3441bca633474739f7174aa0de3e5 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Fri, 25 Sep 2026 06:46:21 -0500 Subject: [PATCH 3/3] Use new triple Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/CIR/CodeGenHIP/device-printf.hip | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 539106d7200584..8f08a5982dcfd6 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -1,9 +1,9 @@ // REQUIRES: amdgpu-registered-target -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ -// RUN: -std=c++11 -fclangir -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fclangir -fcuda-is-device -emit-llvm %s -o %t-cir.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s -// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \ -// RUN: -std=c++11 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s // printf fell through to a library call, which does not exist on the device: _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
