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/7] [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 d97ac5265c282..68a9460281baf 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 2b223c8ae1939..d4d31a7f2ae65 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 58ef4b1bd4b61..a0daed0938a33 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 b232abf5d9299..0f65753374843 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 240601f9834e5..e25a251150cf9 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 bd529f71b38ed..c55de0d19fc5d 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 0000000000000..5e7b3039cccf4 --- /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/7] 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 5e7b3039cccf4..539106d720058 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/7] 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 539106d720058..8f08a5982dcfd 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: >From 44b51b829bf56a93a71b21a706e04bd8d9447210 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 01:14:16 -0500 Subject: [PATCH 4/7] Fix buffered printf Signed-off-by: Steffen Holst Larsen <[email protected]> --- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 24 +++++++++++++++---- clang/test/CIR/CodeGenHIP/device-printf.hip | 11 +++++++++ 2 files changed, 31 insertions(+), 4 deletions(-) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index c55de0d19fc5d..e0a6df59c8ab7 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5924,10 +5924,11 @@ void expandAMDGPUDevicePrintf(llvm::Module &module) { if (!marker) return; - // CIR records the requested lowering as a module flag. + // CIR records the requested lowering as a module flag. The flag is stored + // under the LLVM-side name (see amendModule in LowerToLLVMIR.cpp), not the + // CIR attribute name. bool isBuffered = false; - if (llvm::Metadata *md = - module.getModuleFlag(cir::CIRDialect::getAMDGPUPrintfKindAttrName())) + if (llvm::Metadata *md = module.getModuleFlag("amdgpu_printf_kind")) if (auto *mdStr = llvm::dyn_cast<llvm::MDString>(md)) isBuffered = mdStr->getString() == "buffered"; @@ -5948,9 +5949,24 @@ void expandAMDGPUDevicePrintf(llvm::Module &module) { ci = llvm::changeToCall(llvm::cast<llvm::InvokeInst>(cb)); } - llvm::IRBuilder<> irb(ci); + // Buffered lowering splits the call site's block and build new control flow + // of its own. It expects to be the one driving codegen for the rest of the + // block, as it would if called from normal frontend codegen. Since we're + // expanding a marker call after the fact, split off everything that + // was already emitted after it into its own block first, then + // reconnect to that block once the real printf sequence has been + // built. + llvm::BasicBlock *originalBB = ci->getParent(); llvm::SmallVector<llvm::Value *, 8> args(ci->args()); + llvm::BasicBlock *continuation = + originalBB->splitBasicBlock(ci->getNextNode()); + originalBB->getTerminator()->eraseFromParent(); + + llvm::IRBuilder<> irb(originalBB); + irb.SetCurrentDebugLocation(ci->getDebugLoc()); llvm::Value *res = llvm::emitAMDGPUPrintfCall(irb, args, isBuffered); + irb.CreateBr(continuation); + if (res && !ci->use_empty()) ci->replaceAllUsesWith(res); ci->eraseFromParent(); diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 8f08a5982dcfd..815ea2f732a39 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -5,6 +5,12 @@ // 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 +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fclangir -fcuda-is-device -mprintf-kind=buffered -emit-llvm %s -o %t-cir-buffered.ll +// RUN: FileCheck --check-prefix=BUFFERED --input-file=%t-cir-buffered.ll %s +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcuda-is-device -mprintf-kind=buffered -emit-llvm %s -o %t-buffered.ll +// RUN: FileCheck --check-prefix=BUFFERED --input-file=%t-buffered.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 @@ -18,4 +24,9 @@ extern "C" __device__ int printf(const char *, ...); // LLVM: call i64 @__ockl_printf_append_string_n(i64 // LLVM: call i64 @__ockl_printf_append_args(i64 // LLVM-NOT: call{{.*}}@__cir_amdgpu_printf +// BUFFERED-LABEL: @_Z1pi +// BUFFERED: call ptr addrspace(1) @__printf_alloc(i32 +// BUFFERED: end.block: +// BUFFERED: argpush.block: +// BUFFERED-NOT: call{{.*}}@__cir_amdgpu_printf __device__ void p(int v) { printf("%d\n", v); } >From 8acd4ab4b7621241537a3bea5eb1ea748da43e1f Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 01:26:46 -0500 Subject: [PATCH 5/7] cir-translate and nothrow Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 ++- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 15 ++------- clang/test/CIR/CodeGenHIP/device-printf.hip | 33 +++++++++++++++++++ clang/test/CIR/Tools/amdgpu-printf.cir | 21 ++++++++++++ clang/tools/cir-translate/cir-translate.cpp | 1 + 5 files changed, 61 insertions(+), 14 deletions(-) create mode 100644 clang/test/CIR/Tools/amdgpu-printf.cir diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index a0daed0938a33..8ba724bf758d5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -1133,5 +1133,8 @@ CIRGenFunction::emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr) { builder.getSInt32Ty(), /*isVarArg=*/true); cir::FuncOp fn = cgm.createRuntimeFunction(fnTy, "__cir_amdgpu_printf"); - return builder.createCallOp(loc, fn, callArgs).getResult(); + cir::CallOp call = builder.createCallOp(loc, fn, callArgs); + // The sequence the marker expands into never unwinds. + call.setNothrowAttr(builder.getUnitAttr()); + return call.getResult(); } diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index e0a6df59c8ab7..3a98c2a1e5789 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -57,7 +57,6 @@ #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; @@ -5936,18 +5935,8 @@ void expandAMDGPUDevicePrintf(llvm::Module &module) { 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)); - } + // CIRGen marks the marker call nothrow, so it is never an invoke. + auto *ci = llvm::cast<llvm::CallInst>(u); // Buffered lowering splits the call site's block and build new control flow // of its own. It expects to be the one driving codegen for the rest of the diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 815ea2f732a39..13e36580e8ca1 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -1,14 +1,18 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ // 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=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ // RUN: -fcuda-is-device -emit-llvm %s -o %t.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s // RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ // RUN: -fclangir -fcuda-is-device -mprintf-kind=buffered -emit-llvm %s -o %t-cir-buffered.ll // RUN: FileCheck --check-prefix=BUFFERED --input-file=%t-cir-buffered.ll %s // RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ // RUN: -fcuda-is-device -mprintf-kind=buffered -emit-llvm %s -o %t-buffered.ll // RUN: FileCheck --check-prefix=BUFFERED --input-file=%t-buffered.ll %s @@ -30,3 +34,32 @@ extern "C" __device__ int printf(const char *, ...); // BUFFERED: argpush.block: // BUFFERED-NOT: call{{.*}}@__cir_amdgpu_printf __device__ void p(int v) { printf("%d\n", v); } + +// The result of printf must be forwarded to its users. CIR places the +// expansion blocks after the block that uses the result, so the use is matched +// in any order. +// LLVM-LABEL: @_Z1ri +// LLVM-DAG: [[RES:%.*]] = call i64 @__ockl_printf_append_args(i64 +// LLVM-DAG: [[RET:%.*]] = trunc i64 [[RES]] to i32 +// LLVM-DAG: {{store|ret}} i32 [[RET]] +// LLVM-NOT: call{{.*}}@__cir_amdgpu_printf +// BUFFERED-LABEL: @_Z1ri +// BUFFERED-DAG: %printf_result = sext i1 %{{.*}} to i32 +// BUFFERED-DAG: {{store|ret}} i32 %printf_result +// BUFFERED-NOT: call{{.*}}@__cir_amdgpu_printf +__device__ int r(int v) { return printf("%d\n", v); } + +// Device printf cannot throw, so a pending cleanup must not turn it into an +// invoke with an exception-handling path. +struct D { __device__ ~D(); }; +// LLVM-LABEL: @_Z2ehi( +// LLVM-NOT: personality +// LLVM: call i64 @__ockl_printf_begin(i64 +// LLVM-NOT: {{invoke|landingpad|resume}} +// LLVM-NOT: call{{.*}}@__cir_amdgpu_printf +// BUFFERED-LABEL: @_Z2ehi( +// BUFFERED-NOT: personality +// BUFFERED: call ptr addrspace(1) @__printf_alloc(i32 +// BUFFERED-NOT: {{invoke|landingpad|resume}} +// BUFFERED-NOT: call{{.*}}@__cir_amdgpu_printf +__device__ void eh(int v) { D d; printf("%d\n", v); } diff --git a/clang/test/CIR/Tools/amdgpu-printf.cir b/clang/test/CIR/Tools/amdgpu-printf.cir new file mode 100644 index 0000000000000..8383b9ed4c3ee --- /dev/null +++ b/clang/test/CIR/Tools/amdgpu-printf.cir @@ -0,0 +1,21 @@ +// RUN: cir-translate --cir-to-llvmir --target amdgcn-amd-amdhsa --disable-cc-lowering %s -o %t.ll +// RUN: FileCheck %s -input-file %t.ll -check-prefix=LLVM + +!s32i = !cir.int<s, 32> +!s8i = !cir.int<s, 8> + +module { + cir.func @p(%fmt: !cir.ptr<!s8i>, %v: !s32i) -> !s32i { + %0 = cir.call @__cir_amdgpu_printf(%fmt, %v) nothrow : (!cir.ptr<!s8i>, !s32i) -> !s32i + cir.return %0 : !s32i + } + cir.func private @__cir_amdgpu_printf(!cir.ptr<!s8i>, ...) -> !s32i +} + +// LLVM-LABEL: define{{.*}} i32 @p( +// LLVM: call i64 @__ockl_printf_begin(i64 +// LLVM-DAG: call i64 @__ockl_printf_append_string_n(i64 +// LLVM-DAG: [[RES:%.*]] = call i64 @__ockl_printf_append_args(i64 +// LLVM-DAG: [[RET:%.*]] = trunc i64 [[RES]] to i32 +// LLVM-DAG: ret i32 [[RET]] +// LLVM-NOT: __cir_amdgpu_printf diff --git a/clang/tools/cir-translate/cir-translate.cpp b/clang/tools/cir-translate/cir-translate.cpp index cefa7d2996f30..aa9cd06094665 100644 --- a/clang/tools/cir-translate/cir-translate.cpp +++ b/clang/tools/cir-translate/cir-translate.cpp @@ -173,6 +173,7 @@ void registerToLLVMTranslation() { enableOpenMP); if (!llvmModule) return mlir::failure(); + cir::direct::expandAMDGPUDevicePrintf(*llvmModule); llvmModule->renumberMetadataForAssembly(); llvmModule->print(output, nullptr); return mlir::success(); >From beaf34b82b7103a9a93c9aa6e4e098bd4b4c115d Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 01:48:11 -0500 Subject: [PATCH 6/7] Link TransformUtils Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt index d3d0805afb4d8..2cc003d4af1c4 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt +++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt @@ -1,6 +1,7 @@ set(LLVM_LINK_COMPONENTS Core Support + TransformUtils ) add_clang_library(clangCIRLoweringDirectToLLVM >From 90402a3fe2358a2fd3b634e232d8686257cebb25 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 28 Sep 2026 07:21:57 -0500 Subject: [PATCH 7/7] Move emission back into direct and add CIR test cases Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/CIR/LowerToLLVM.h | 6 ----- clang/lib/CIR/FrontendAction/CIRGenAction.cpp | 8 ++----- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 4 +++- clang/test/CIR/CodeGenHIP/device-printf.hip | 24 +++++++++++++++++++ clang/tools/cir-translate/cir-translate.cpp | 1 - 5 files changed, 29 insertions(+), 14 deletions(-) diff --git a/clang/include/clang/CIR/LowerToLLVM.h b/clang/include/clang/CIR/LowerToLLVM.h index 68a9460281baf..d97ac5265c282 100644 --- a/clang/include/clang/CIR/LowerToLLVM.h +++ b/clang/include/clang/CIR/LowerToLLVM.h @@ -33,12 +33,6 @@ 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/FrontendAction/CIRGenAction.cpp b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp index e25a251150cf9..240601f9834e5 100644 --- a/clang/lib/CIR/FrontendAction/CIRGenAction.cpp +++ b/clang/lib/CIR/FrontendAction/CIRGenAction.cpp @@ -65,12 +65,8 @@ lowerFromCIRToLLVMIR(mlir::ModuleOp MLIRModule, llvm::LLVMContext &LLVMCtx, bool EnableOpenMP, llvm::StringRef mlirSaveTempsOutFile = {}, llvm::vfs::FileSystem *fs = nullptr) { - std::unique_ptr<llvm::Module> LLVMModule = - direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx, EnableOpenMP, - mlirSaveTempsOutFile, fs); - if (LLVMModule) - direct::expandAMDGPUDevicePrintf(*LLVMModule); - return LLVMModule; + return direct::lowerDirectlyFromCIRToLLVMIR(MLIRModule, LLVMCtx, EnableOpenMP, + mlirSaveTempsOutFile, fs); } class CIRGenConsumer : public clang::ASTConsumer { diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 3a98c2a1e5789..d1e41448d27c4 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -5918,7 +5918,7 @@ void populateCIRToLLVMPasses(mlir::OpPassManager &pm, bool enableOpenMP) { // 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) { +static void expandAMDGPUDevicePrintf(llvm::Module &module) { llvm::Function *marker = module.getFunction("__cir_amdgpu_printf"); if (!marker) return; @@ -6007,6 +6007,8 @@ lowerDirectlyFromCIRToLLVMIR(mlir::ModuleOp mlirModule, LLVMContext &llvmCtx, report_fatal_error("Lowering from LLVMIR dialect to llvm IR failed!"); } + expandAMDGPUDevicePrintf(*llvmModule); + return llvmModule; } } // namespace direct diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 13e36580e8ca1..a60540442afac 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -1,6 +1,14 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ // RUN: -fcxx-exceptions -fexceptions \ +// RUN: -fclangir -fcuda-is-device -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefixes=CIR,CIR-HOSTCALL --input-file=%t.cir %s +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ +// RUN: -fclangir -fcuda-is-device -mprintf-kind=buffered -emit-cir %s -o %t-buffered.cir +// RUN: FileCheck --check-prefixes=CIR,CIR-BUFFERED --input-file=%t-buffered.cir %s +// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ +// RUN: -fcxx-exceptions -fexceptions \ // 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=amdgpu9.0a-amd-amdhsa -x hip -std=c++11 \ @@ -23,6 +31,14 @@ #define __device__ __attribute__((device)) extern "C" __device__ int printf(const char *, ...); +// CIR-HOSTCALL: module {{.*}} attributes {{.*}}cir.amdgpu_printf_kind = "hostcall" +// CIR-BUFFERED: module {{.*}} attributes {{.*}}cir.amdgpu_printf_kind = "buffered" + +// CIR-LABEL: cir.func {{.*}} @_Z1pi( +// CIR: %[[FMT:.*]] = cir.cast array_to_ptrdecay +// CIR: %[[V:.*]] = cir.load {{.*}} : !cir.ptr<!s32i>, !s32i +// CIR: cir.call @__cir_amdgpu_printf(%[[FMT]], %[[V]]) nothrow : (!cir.ptr<!s8i>, !s32i) -> !s32i +// CIR: cir.func private {{.*}} @__cir_amdgpu_printf(!cir.ptr<!s8i>, ...) -> !s32i // LLVM-LABEL: @_Z1pi // LLVM: call i64 @__ockl_printf_begin(i64 // LLVM: call i64 @__ockl_printf_append_string_n(i64 @@ -38,6 +54,10 @@ __device__ void p(int v) { printf("%d\n", v); } // The result of printf must be forwarded to its users. CIR places the // expansion blocks after the block that uses the result, so the use is matched // in any order. +// CIR-LABEL: cir.func {{.*}} @_Z1ri( +// CIR: %[[RETVAL:.*]] = cir.alloca "__retval" +// CIR: %[[RES:.*]] = cir.call @__cir_amdgpu_printf({{.*}}) nothrow +// CIR: cir.store %[[RES]], %[[RETVAL]] // LLVM-LABEL: @_Z1ri // LLVM-DAG: [[RES:%.*]] = call i64 @__ockl_printf_append_args(i64 // LLVM-DAG: [[RET:%.*]] = trunc i64 [[RES]] to i32 @@ -52,6 +72,10 @@ __device__ int r(int v) { return printf("%d\n", v); } // Device printf cannot throw, so a pending cleanup must not turn it into an // invoke with an exception-handling path. struct D { __device__ ~D(); }; +// CIR-LABEL: cir.func {{.*}} @_Z2ehi( +// CIR: cir.cleanup.scope { +// CIR: cir.call @__cir_amdgpu_printf({{.*}}) nothrow +// CIR: } cleanup all { // LLVM-LABEL: @_Z2ehi( // LLVM-NOT: personality // LLVM: call i64 @__ockl_printf_begin(i64 diff --git a/clang/tools/cir-translate/cir-translate.cpp b/clang/tools/cir-translate/cir-translate.cpp index aa9cd06094665..cefa7d2996f30 100644 --- a/clang/tools/cir-translate/cir-translate.cpp +++ b/clang/tools/cir-translate/cir-translate.cpp @@ -173,7 +173,6 @@ void registerToLLVMTranslation() { enableOpenMP); if (!llvmModule) return mlir::failure(); - cir::direct::expandAMDGPUDevicePrintf(*llvmModule); llvmModule->renumberMetadataForAssembly(); llvmModule->print(output, nullptr); return mlir::success(); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
