[clang] [llvm] [CIR][AMDGPU] Route printf in device code through the OpenCL runtime (PR #226435)
Steffen Larsen via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 29 22:51:51 PDT 2026
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 <sholstla at amd.com>
Date: Tue, 1 Sep 2026 09:00:32 -0500
Subject: [PATCH 01/10] [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 <sholstla at amd.com>
Date: Fri, 25 Sep 2026 06:43:03 -0500
Subject: [PATCH 02/10] Remove unnecessary over-verbose test comment
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
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 <sholstla at amd.com>
Date: Fri, 25 Sep 2026 06:46:21 -0500
Subject: [PATCH 03/10] Use new triple
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
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 <sholstla at amd.com>
Date: Mon, 28 Sep 2026 01:14:16 -0500
Subject: [PATCH 04/10] Fix buffered printf
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
.../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 <sholstla at amd.com>
Date: Mon, 28 Sep 2026 01:26:46 -0500
Subject: [PATCH 05/10] cir-translate and nothrow
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
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 <sholstla at amd.com>
Date: Mon, 28 Sep 2026 01:48:11 -0500
Subject: [PATCH 06/10] Link TransformUtils
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
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 <sholstla at amd.com>
Date: Mon, 28 Sep 2026 07:21:57 -0500
Subject: [PATCH 07/10] Move emission back into direct and add CIR test cases
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
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();
>From c0bfd2dedfd2f9b21391f33329fdf752757e8535 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Mon, 28 Sep 2026 08:46:47 -0500
Subject: [PATCH 08/10] Refactor run lines and add SPIR-V cases
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/test/CIR/CodeGenHIP/device-printf.hip | 62 ++++++++++-----------
1 file changed, 31 insertions(+), 31 deletions(-)
diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip
index a60540442afac..7a8ffe297a2d9 100644
--- a/clang/test/CIR/CodeGenHIP/device-printf.hip
+++ b/clang/test/CIR/CodeGenHIP/device-printf.hip
@@ -1,28 +1,28 @@
// 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 \
-// 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
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck --check-prefixes=CIR,CIR-HOSTCALL %s
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -mprintf-kind=buffered -emit-cir %s -o - \
+// RUN: | FileCheck --check-prefixes=CIR,CIR-BUFFERED %s
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcxx-exceptions -fexceptions -fclangir -fcuda-is-device -emit-llvm %s -o - | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -mprintf-kind=buffered -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=BUFFERED %s
+// RUN: %clang_cc1 -triple=amdgpu9.0a-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -mprintf-kind=buffered -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=BUFFERED %s
+
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-cir %s -o - \
+// RUN: | FileCheck --check-prefixes=CIR,CIR-HOSTCALL %s
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -mprintf-kind=buffered -emit-cir %s -o - \
+// RUN: | FileCheck --check-prefixes=CIR,CIR-BUFFERED %s
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=LLVM %s
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir -mprintf-kind=buffered -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=BUFFERED %s
+// RUN: %clang_cc1 -triple=spirv64-amd-amdhsa -x hip -fcuda-is-device -fcxx-exceptions -fexceptions -mprintf-kind=buffered -emit-llvm %s -o - \
+// RUN: | FileCheck --check-prefix=BUFFERED %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
@@ -40,12 +40,12 @@ extern "C" __device__ int printf(const char *, ...);
// 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
-// LLVM: call i64 @__ockl_printf_append_args(i64
+// 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
// BUFFERED-LABEL: @_Z1pi
-// BUFFERED: call ptr addrspace(1) @__printf_alloc(i32
+// BUFFERED: call{{.*}} ptr addrspace(1) @__printf_alloc(i32
// BUFFERED: end.block:
// BUFFERED: argpush.block:
// BUFFERED-NOT: call{{.*}}@__cir_amdgpu_printf
@@ -59,7 +59,7 @@ __device__ void p(int v) { printf("%d\n", v); }
// 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: [[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
@@ -78,12 +78,12 @@ struct D { __device__ ~D(); };
// CIR: } cleanup all {
// LLVM-LABEL: @_Z2ehi(
// LLVM-NOT: personality
-// LLVM: call i64 @__ockl_printf_begin(i64
+// 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: 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); }
>From 7c2c40599b266e9eb7661fc6c874a3b154cb17a3 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 29 Sep 2026 03:56:03 -0500
Subject: [PATCH 09/10] Change to CIR offload Op
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/include/clang/CIR/Dialect/IR/CIROps.td | 34 +
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 25 +-
.../CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 +
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 58 +-
.../CIR/Lowering/DirectToLLVM/LowerToLLVM.h | 9 +
.../DirectToLLVM/LowerToLLVMOffloadPrintf.cpp | 674 ++++++++++++++++++
clang/test/CIR/CodeGenHIP/device-printf.hip | 192 ++++-
clang/test/CIR/IR/offload-printf.cir | 15 +
.../test/CIR/Lowering/offload-printf-nyi.cir | 13 +
clang/test/CIR/Tools/amdgpu-printf.cir | 21 -
clang/test/CIR/Tools/offload-printf.cir | 86 +++
.../llvm/Transforms/Utils/AMDGPUEmitPrintf.h | 28 +
.../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 68 +-
13 files changed, 1079 insertions(+), 145 deletions(-)
create mode 100644 clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp
create mode 100644 clang/test/CIR/IR/offload-printf.cir
create mode 100644 clang/test/CIR/Lowering/offload-printf-nyi.cir
delete mode 100644 clang/test/CIR/Tools/amdgpu-printf.cir
create mode 100644 clang/test/CIR/Tools/offload-printf.cir
diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td
index c1ba78eea2835..ee4c3faceb0ae 100644
--- a/clang/include/clang/CIR/Dialect/IR/CIROps.td
+++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td
@@ -9596,6 +9596,40 @@ def CIR_MemChrOp : CIR_Op<"libc.memchr",
let hasVerifier = 1;
}
+//===----------------------------------------------------------------------===//
+// OffloadPrintfOp
+//===----------------------------------------------------------------------===//
+
+def CIR_OffloadPrintfOp : CIR_Op<"offload.printf"> {
+ let summary = "Device-side printf in offload code";
+ let description = [{
+ Represents a call to `printf` in device code. The target has no libc
+ `printf` to call, so the op is expanded into the target's printf runtime
+ sequence when lowering to LLVM.
+
+ `format` is the format string and `args` are the remaining arguments,
+ after the default argument promotions. `result` is what `printf` returns
+ on the target.
+
+ Example:
+
+ ```
+ %r = cir.offload.printf(%fmt, %v) : (!cir.ptr<!s8i>, !s32i) -> !s32i
+ ```
+ }];
+
+ let arguments = (ins
+ CIR_PtrToChar8Type:$format,
+ Variadic<AnyTypeOf<[CIR_AnyIntType, CIR_AnyFloatType, CIR_AnyPtrType]>>:$args
+ );
+
+ let results = (outs CIR_SInt32:$result);
+
+ let assemblyFormat = [{
+ `(` $format (`,` $args^)? `)` attr-dict `:` functional-type(operands, results)
+ }];
+}
+
//===----------------------------------------------------------------------===//
// LaunderOp
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 8ba724bf758d5..d87c454b35048 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -1097,8 +1097,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
}
}
-// Emit AMDGPU printf CIR stand-in function call. This stand-in function call is
-// lowered to the appropriate call structure during LLVM IR lowering.
+// Emit an AMDGPU device printf as a cir.offload.printf, which is expanded into
+// the AMDGPU printf runtime sequence during LLVM lowering.
mlir::Value
CIRGenFunction::emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr) {
assert(cgm.getTriple().isAMDGCN() ||
@@ -1115,9 +1115,14 @@ CIRGenFunction::emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr) {
mlir::Location loc = getLoc(expr->getBeginLoc());
- // We don't know how to emit non-scalar varargs.
+ // We don't know how to emit non-scalar varargs, nor scalars the printf
+ // runtime has no encoding for, such as vectors.
bool hasNonScalar = llvm::any_of(args, [&](const CallArg &a) {
- return a.hasLValue() || !a.getKnownRValue().isScalar();
+ if (a.hasLValue() || !a.getKnownRValue().isScalar())
+ return true;
+ mlir::Type ty = a.getKnownRValue().getValue().getType();
+ return !mlir::isa<cir::IntType, cir::PointerType>(ty) &&
+ !cir::isAnyFloatingPointType(ty);
});
if (hasNonScalar) {
cgm.errorUnsupported(expr, "non-scalar args to printf");
@@ -1128,13 +1133,7 @@ CIRGenFunction::emitAMDGPUDevicePrintfCallExpr(const CallExpr *expr) {
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");
- cir::CallOp call = builder.createCallOp(loc, fn, callArgs);
- // The sequence the marker expands into never unwinds.
- call.setNothrowAttr(builder.getUnitAttr());
- return call.getResult();
+ return cir::OffloadPrintfOp::create(builder, loc, builder.getSInt32Ty(),
+ callArgs.front(),
+ llvm::drop_begin(callArgs));
}
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
index 2cc003d4af1c4..7dfca636ea882 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt
@@ -7,6 +7,7 @@ set(LLVM_LINK_COMPONENTS
add_clang_library(clangCIRLoweringDirectToLLVM
LowerToLLVM.cpp
LowerToLLVMIR.cpp
+ LowerToLLVMOffloadPrintf.cpp
LowerToLLVMOpenCLMetadata.cpp
DEPENDS
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
index d1e41448d27c4..5a97da3392308 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp
@@ -49,14 +49,12 @@
#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"
using namespace cir;
using namespace llvm;
@@ -4598,9 +4596,8 @@ mlir::LogicalResult CIRToLLVMInsertMemberOpLowering::matchAndRewrite(
void createLLVMFuncOpIfNotExist(mlir::ConversionPatternRewriter &rewriter,
mlir::SymbolTableCollection &symbolTables,
mlir::Operation *srcOp, llvm::StringRef fnName,
- mlir::Type fnTy,
- mlir::ArrayAttr argAttrs = nullptr,
- mlir::ArrayAttr resAttrs = nullptr) {
+ mlir::Type fnTy, mlir::ArrayAttr argAttrs,
+ mlir::ArrayAttr resAttrs) {
mlir::ModuleOp modOp = srcOp->getParentOfType<mlir::ModuleOp>();
mlir::Operation *sourceSymbol = symbolTables.lookupSymbolIn(
modOp, mlir::StringAttr::get(fnTy.getContext(), fnName));
@@ -5916,55 +5913,6 @@ 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.
-static 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. 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("amdgpu_printf_kind"))
- 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) {
- // 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
- // 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();
- }
-
- 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,
@@ -6007,8 +5955,6 @@ 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/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.h b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.h
index 56941b2aa51e7..6ca81904af559 100644
--- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.h
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.h
@@ -35,6 +35,15 @@ mlir::Value lowerCirAttrAsValue(mlir::Operation *parentOp, mlir::Attribute attr,
mlir::LLVM::Linkage convertLinkage(cir::GlobalLinkageKind linkage);
+/// Declare `fnName` with type `fnTy` before the function enclosing `srcOp`,
+/// unless the module already has a symbol of that name.
+void createLLVMFuncOpIfNotExist(mlir::ConversionPatternRewriter &rewriter,
+ mlir::SymbolTableCollection &symbolTables,
+ mlir::Operation *srcOp, llvm::StringRef fnName,
+ mlir::Type fnTy,
+ mlir::ArrayAttr argAttrs = nullptr,
+ mlir::ArrayAttr resAttrs = nullptr);
+
struct LLVMBlockAddressInfo {
// Get the next tag index
uint32_t getTagIndex() { return blockTagOpIndex++; }
diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp
new file mode 100644
index 0000000000000..fcfe6ec22b0bc
--- /dev/null
+++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp
@@ -0,0 +1,674 @@
+//===- LowerToLLVMOffloadPrintf.cpp - Offload printf lowering -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "LowerToLLVM.h"
+
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "clang/CIR/Dialect/IR/CIRDialect.h"
+#include "llvm/ADT/SparseBitVector.h"
+#include "llvm/Support/MathExtras.h"
+#include "llvm/TargetParser/Triple.h"
+#include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h"
+
+namespace cir {
+namespace direct {
+
+namespace {
+
+/// Get the contents of the constant string ptr points to, up to the first NUL.
+/// This is the counterpart of llvm::getConstantStringInfo, and looks through
+/// the CIR a string literal is addressed with.
+bool getConstantStringInfo(mlir::Value ptr, mlir::ModuleOp mod,
+ mlir::SymbolTableCollection &symbolTables,
+ std::string &str) {
+ auto cast = ptr.getDefiningOp<cir::CastOp>();
+ while (cast && (cast.getKind() == cir::CastKind::array_to_ptrdecay ||
+ cast.getKind() == cir::CastKind::bitcast ||
+ cast.getKind() == cir::CastKind::address_space)) {
+ ptr = cast.getSrc();
+ cast = ptr.getDefiningOp<cir::CastOp>();
+ }
+
+ auto getGlobal = ptr.getDefiningOp<cir::GetGlobalOp>();
+ if (!getGlobal)
+ return false;
+
+ // The global may or may not have been converted already.
+ mlir::Operation *symbol =
+ symbolTables.lookupSymbolIn(mod, getGlobal.getNameAttr());
+ llvm::StringRef data;
+ if (auto global = mlir::dyn_cast_if_present<cir::GlobalOp>(symbol)) {
+ std::optional<mlir::Attribute> init = global.getInitialValue();
+ if (!global.getConstant() || !init)
+ return false;
+ if (mlir::isa<cir::ZeroAttr>(*init)) {
+ str.clear();
+ return true;
+ }
+ auto constArr = mlir::dyn_cast<cir::ConstArrayAttr>(*init);
+ auto strAttr = constArr
+ ? mlir::dyn_cast<mlir::StringAttr>(constArr.getElts())
+ : nullptr;
+ if (!strAttr)
+ return false;
+ data = strAttr.getValue();
+ } else if (auto global =
+ mlir::dyn_cast_if_present<mlir::LLVM::GlobalOp>(symbol)) {
+ if (!global.getConstant())
+ return false;
+ mlir::Attribute value = global.getValueOrNull();
+ if (auto strAttr = mlir::dyn_cast_if_present<mlir::StringAttr>(value)) {
+ data = strAttr.getValue();
+ } else if (auto dense =
+ mlir::dyn_cast_if_present<mlir::DenseIntElementsAttr>(value);
+ dense && dense.getElementType().isInteger(8)) {
+ str.clear();
+ for (const llvm::APInt &c : dense.getValues<llvm::APInt>()) {
+ if (c.isZero())
+ break;
+ str.push_back(static_cast<char>(c.getZExtValue()));
+ }
+ return true;
+ } else {
+ return false;
+ }
+ } else {
+ return false;
+ }
+
+ str = data.substr(0, data.find('\0')).str();
+ return true;
+}
+
+// Helper struct to package the string related data.
+struct StringData {
+ std::string str;
+ mlir::Value realSize;
+ mlir::Value alignedSize;
+ bool isConst = true;
+};
+
+/// Expands one cir.offload.printf into the AMDGPU printf runtime sequence.
+///
+/// Everything after the op is split off into a continuation block first, so
+/// the expansion can build control flow of its own.
+class AMDGPUPrintfLowering {
+public:
+ AMDGPUPrintfLowering(cir::OffloadPrintfOp op,
+ mlir::ConversionPatternRewriter &rewriter,
+ const mlir::DataLayout &dataLayout,
+ mlir::SymbolTableCollection &symbolTables)
+ : op(op), rewriter(rewriter), dataLayout(dataLayout),
+ symbolTables(symbolTables), mod(op->getParentOfType<mlir::ModuleOp>()) {
+ }
+
+ mlir::LogicalResult lower(mlir::ValueRange args, bool isBuffered);
+
+private:
+ mlir::Value getI32(uint32_t value) {
+ return mlir::LLVM::ConstantOp::create(
+ rewriter, op.getLoc(), rewriter.getI32Type(), llvm::APInt(32, value));
+ }
+ mlir::Value getI64(uint64_t value) {
+ return mlir::LLVM::ConstantOp::create(
+ rewriter, op.getLoc(), rewriter.getI64Type(), llvm::APInt(64, value));
+ }
+
+ mlir::Value createCall(llvm::StringRef name, mlir::Type resTy,
+ mlir::ValueRange args);
+
+ mlir::Value fitArgInto64Bits(mlir::Value arg);
+ mlir::Value callAppendArgs(mlir::Value desc, int numArgs, mlir::Value arg0,
+ mlir::Value arg1, mlir::Value arg2,
+ mlir::Value arg3, mlir::Value arg4,
+ mlir::Value arg5, mlir::Value arg6, bool isLast);
+ mlir::Value appendArg(mlir::Value desc, mlir::Value arg, bool isLast);
+ mlir::Value getStrlenWithNull(mlir::Value str);
+ mlir::Value callAppendStringN(mlir::Value desc, mlir::Value str,
+ mlir::Value length, bool isLast);
+ mlir::Value appendString(mlir::Value desc, mlir::Value arg, bool isLast);
+ mlir::Value processArg(mlir::Value desc, mlir::Value arg, bool specIsCString,
+ bool isLast);
+ mlir::Value emitHostcall(mlir::ValueRange args);
+
+ bool isCStringArg(mlir::ValueRange args, size_t i) const {
+ // The format specifies a string but the argument is not a pointer. The
+ // frontend will have warned; send the argument as a scalar.
+ return specIsCString.test(i) &&
+ mlir::isa<mlir::LLVM::LLVMPointerType>(args[i].getType());
+ }
+ mlir::Value alignTo8(mlir::Value len);
+ mlir::Value
+ callBufferedPrintfStart(mlir::ValueRange args, bool isConstFmtStr,
+ llvm::SmallVectorImpl<StringData> &stringContents,
+ mlir::Value &argSize);
+ void
+ processConstantStringArg(const StringData &sd,
+ llvm::SmallVectorImpl<mlir::Value> &whatToStore);
+ mlir::Value processNonStringArg(mlir::Value arg);
+ void callBufferedPrintfArgPush(mlir::ValueRange args, mlir::Value ptrToStore,
+ llvm::ArrayRef<StringData> stringContents,
+ bool isConstFmtStr);
+ void addPrintfFormatMetadata(llvm::StringRef entry, bool onlyIfEmpty);
+ mlir::Value emitBuffered(mlir::ValueRange args);
+
+ cir::OffloadPrintfOp op;
+ mlir::ConversionPatternRewriter &rewriter;
+ const mlir::DataLayout &dataLayout;
+ mlir::SymbolTableCollection &symbolTables;
+ mlir::ModuleOp mod;
+
+ /// The block the rest of the original block was split off into.
+ mlir::Block *cont = nullptr;
+ std::string fmtStr;
+ llvm::SparseBitVector<8> specIsCString;
+};
+
+} // namespace
+
+mlir::Value AMDGPUPrintfLowering::createCall(llvm::StringRef name,
+ mlir::Type resTy,
+ mlir::ValueRange args) {
+ auto fnTy = mlir::LLVM::LLVMFunctionType::get(
+ resTy, llvm::to_vector(args.getTypes()), /*isVarArg=*/false);
+ createLLVMFuncOpIfNotExist(rewriter, symbolTables, op, name, fnTy);
+ return mlir::LLVM::CallOp::create(
+ rewriter, op.getLoc(), fnTy,
+ mlir::FlatSymbolRefAttr::get(rewriter.getContext(), name), args)
+ .getResult();
+}
+
+mlir::Value AMDGPUPrintfLowering::fitArgInto64Bits(mlir::Value arg) {
+ mlir::Type ty = arg.getType();
+
+ if (auto intTy = mlir::dyn_cast<mlir::IntegerType>(ty)) {
+ switch (intTy.getWidth()) {
+ case 32:
+ return mlir::LLVM::ZExtOp::create(rewriter, op.getLoc(),
+ rewriter.getI64Type(), arg);
+ case 64:
+ return arg;
+ }
+ }
+
+ if (ty.isF64())
+ return mlir::LLVM::BitcastOp::create(rewriter, op.getLoc(),
+ rewriter.getI64Type(), arg);
+
+ if (mlir::isa<mlir::LLVM::LLVMPointerType>(ty))
+ return mlir::LLVM::PtrToIntOp::create(rewriter, op.getLoc(),
+ rewriter.getI64Type(), arg);
+
+ llvm_unreachable("argument types are checked before lowering");
+}
+
+mlir::Value AMDGPUPrintfLowering::callAppendArgs(
+ mlir::Value desc, int numArgs, mlir::Value arg0, mlir::Value arg1,
+ mlir::Value arg2, mlir::Value arg3, mlir::Value arg4, mlir::Value arg5,
+ mlir::Value arg6, bool isLast) {
+ mlir::Value isLastValue = getI32(isLast);
+ mlir::Value numArgsValue = getI32(numArgs);
+ return createCall("__ockl_printf_append_args", rewriter.getI64Type(),
+ {desc, numArgsValue, arg0, arg1, arg2, arg3, arg4, arg5,
+ arg6, isLastValue});
+}
+
+mlir::Value AMDGPUPrintfLowering::appendArg(mlir::Value desc, mlir::Value arg,
+ bool isLast) {
+ mlir::Value arg0 = fitArgInto64Bits(arg);
+ mlir::Value zero = getI64(0);
+ return callAppendArgs(desc, 1, arg0, zero, zero, zero, zero, zero, zero,
+ isLast);
+}
+
+// The device library does not provide strlen, so we build our own loop
+// here. While we are at it, we also include the terminating null in the length.
+mlir::Value AMDGPUPrintfLowering::getStrlenWithNull(mlir::Value str) {
+ mlir::Block *prev = rewriter.getInsertionBlock();
+ mlir::Region *region = cont->getParent();
+ mlir::Type ptrTy = str.getType();
+
+ // The length is either zero for a null pointer, or the computed value for an
+ // actual string. The join block's argument represents the final value.
+ //
+ // Strictly speaking, the zero does not matter since
+ // __ockl_printf_append_string_n ignores the length if the pointer is null.
+ mlir::Block *whileBlock =
+ rewriter.createBlock(region, cont->getIterator(), {ptrTy}, {op.getLoc()});
+ mlir::Block *whileDone = rewriter.createBlock(region, cont->getIterator());
+ mlir::Block *join = rewriter.createBlock(
+ region, cont->getIterator(), {rewriter.getI64Type()}, {op.getLoc()});
+
+ // Emit an early return for when the pointer is null.
+ rewriter.setInsertionPointToEnd(prev);
+ mlir::Value zero = getI64(0);
+ mlir::Value cmpNull = mlir::LLVM::ICmpOp::create(
+ rewriter, op.getLoc(), mlir::LLVM::ICmpPredicate::eq, str,
+ mlir::LLVM::ZeroOp::create(rewriter, op.getLoc(), ptrTy));
+ mlir::LLVM::CondBrOp::create(rewriter, op.getLoc(), cmpNull, join, zero,
+ whileBlock, str);
+
+ // Entry to the while loop.
+ rewriter.setInsertionPointToEnd(whileBlock);
+ mlir::Value ptrPhi = whileBlock->getArgument(0);
+ mlir::Value ptrNext = mlir::LLVM::GEPOp::create(
+ rewriter, op.getLoc(), ptrTy, rewriter.getI8Type(), ptrPhi,
+ llvm::ArrayRef<mlir::LLVM::GEPArg>{getI64(1)});
+
+ // Condition for the while loop.
+ mlir::Value data = mlir::LLVM::LoadOp::create(rewriter, op.getLoc(),
+ rewriter.getI8Type(), ptrPhi);
+ mlir::Value cmp = mlir::LLVM::ICmpOp::create(
+ rewriter, op.getLoc(), mlir::LLVM::ICmpPredicate::eq, data,
+ mlir::LLVM::ConstantOp::create(rewriter, op.getLoc(),
+ rewriter.getI8Type(), llvm::APInt(8, 0)));
+ mlir::LLVM::CondBrOp::create(rewriter, op.getLoc(), cmp, whileDone,
+ mlir::ValueRange(), whileBlock, ptrNext);
+
+ // Add one to the computed length.
+ rewriter.setInsertionPointToEnd(whileDone);
+ auto addrTy = mlir::IntegerType::get(rewriter.getContext(),
+ *dataLayout.getTypeIndexBitwidth(ptrTy));
+ mlir::Value endAddr =
+ mlir::LLVM::PtrToAddrOp::create(rewriter, op.getLoc(), addrTy, ptrPhi);
+ mlir::Value beginAddr =
+ mlir::LLVM::PtrToAddrOp::create(rewriter, op.getLoc(), addrTy, str);
+ mlir::Value len =
+ mlir::LLVM::SubOp::create(rewriter, op.getLoc(), endAddr, beginAddr);
+ if (addrTy != rewriter.getI64Type())
+ len = mlir::LLVM::ZExtOp::create(rewriter, op.getLoc(),
+ rewriter.getI64Type(), len);
+ len = mlir::LLVM::AddOp::create(rewriter, op.getLoc(), len, getI64(1));
+
+ // Final join.
+ mlir::LLVM::BrOp::create(rewriter, op.getLoc(), len, join);
+ rewriter.setInsertionPointToEnd(join);
+ return join->getArgument(0);
+}
+
+mlir::Value AMDGPUPrintfLowering::callAppendStringN(mlir::Value desc,
+ mlir::Value str,
+ mlir::Value length,
+ bool isLast) {
+ mlir::Value isLastInt32 = getI32(isLast);
+ auto name = mlir::StringAttr::get(rewriter.getContext(),
+ "__ockl_printf_append_string_n");
+ // As in OGCG, the declaration takes the pointer type of the first string it
+ // is called with. Unlike an LLVM IR call, llvm.call must match its callee, so
+ // cast strings from another address space (CIR string literals on SPIR-V are
+ // not in the generic address space).
+ if (auto fn =
+ symbolTables.lookupSymbolIn<mlir::LLVM::LLVMFuncOp>(mod, name)) {
+ mlir::Type strTy = fn.getFunctionType().getParamType(1);
+ if (str.getType() != strTy)
+ str = mlir::LLVM::AddrSpaceCastOp::create(rewriter, op.getLoc(), strTy,
+ str);
+ }
+ return createCall(name, rewriter.getI64Type(),
+ {desc, str, length, isLastInt32});
+}
+
+mlir::Value AMDGPUPrintfLowering::appendString(mlir::Value desc,
+ mlir::Value arg, bool isLast) {
+ mlir::Value length = getStrlenWithNull(arg);
+ return callAppendStringN(desc, arg, length, isLast);
+}
+
+mlir::Value AMDGPUPrintfLowering::processArg(mlir::Value desc, mlir::Value arg,
+ bool specIsCString, bool isLast) {
+ if (specIsCString && mlir::isa<mlir::LLVM::LLVMPointerType>(arg.getType()))
+ return appendString(desc, arg, isLast);
+ // If the format specifies a string but the argument is not, the frontend will
+ // have printed a warning. We just rely on undefined behaviour and send the
+ // argument anyway.
+ return appendArg(desc, arg, isLast);
+}
+
+mlir::Value AMDGPUPrintfLowering::emitHostcall(mlir::ValueRange args) {
+ size_t numOps = args.size();
+ mlir::Value desc =
+ createCall("__ockl_printf_begin", rewriter.getI64Type(), getI64(0));
+ desc = appendString(desc, args[0], numOps == 1);
+
+ // FIXME: This invokes hostcall once for each argument. We can pack up to
+ // seven scalar printf arguments in a single hostcall. See the signature of
+ // callAppendArgs().
+ for (size_t i = 1; i != numOps; ++i) {
+ bool isLast = i == numOps - 1;
+ bool isCString = specIsCString.test(i);
+ desc = processArg(desc, args[i], isCString, isLast);
+ }
+
+ return mlir::LLVM::TruncOp::create(rewriter, op.getLoc(),
+ rewriter.getI32Type(), desc);
+}
+
+// Align the computed length to next 8 byte boundary.
+mlir::Value AMDGPUPrintfLowering::alignTo8(mlir::Value len) {
+ mlir::Value tempAdd =
+ mlir::LLVM::AddOp::create(rewriter, op.getLoc(), len, getI64(7));
+ // OGCG masks with the 32-bit ~7U zero-extended to i64; match it exactly.
+ return mlir::LLVM::AndOp::create(rewriter, op.getLoc(), tempAdd, getI64(~7U));
+}
+
+// Calculates frame size required for current printf expansion and allocates
+// space on printf buffer. Printf frame includes following contents
+// [ ControlDWord , format string/Hash , Arguments (each aligned to 8 byte) ]
+mlir::Value AMDGPUPrintfLowering::callBufferedPrintfStart(
+ mlir::ValueRange args, bool isConstFmtStr,
+ llvm::SmallVectorImpl<StringData> &stringContents, mlir::Value &argSize) {
+ mlir::Value nonConstStrLen;
+
+ // First 4 bytes to be reserved for control dword
+ uint64_t bufSize = 4;
+ if (isConstFmtStr) {
+ // First 8 bytes of MD5 hash
+ bufSize += 8;
+ } else {
+ mlir::Value lenWithNull = getStrlenWithNull(args[0]);
+ nonConstStrLen = alignTo8(lenWithNull);
+ stringContents.push_back({"", lenWithNull, nonConstStrLen, false});
+ }
+
+ mlir::OperandRange cirArgs = op->getOperands();
+ for (size_t i = 1; i < args.size(); i++) {
+ if (isCStringArg(args, i)) {
+ std::string argStr;
+ if (getConstantStringInfo(cirArgs[i], mod, symbolTables, argStr)) {
+ bufSize += llvm::alignTo(argStr.size() + 1, 8);
+ stringContents.push_back({argStr, {}, {}, true});
+ } else {
+ mlir::Value lenWithNull = getStrlenWithNull(args[i]);
+ mlir::Value lenWithNullAligned = alignTo8(lenWithNull);
+
+ if (nonConstStrLen)
+ nonConstStrLen = mlir::LLVM::AddOp::create(
+ rewriter, op.getLoc(), lenWithNullAligned, nonConstStrLen);
+ else
+ nonConstStrLen = lenWithNullAligned;
+
+ stringContents.push_back({"", lenWithNull, lenWithNullAligned, false});
+ }
+ } else {
+ uint64_t allocSize =
+ dataLayout.getTypeSize(args[i].getType()).getFixedValue();
+ // We end up expanding non string arguments to 8 bytes (args smaller than
+ // 8 bytes)
+ bufSize += std::max<uint64_t>(allocSize, 8);
+ }
+ }
+
+ // calculate final size value to be passed to printf_alloc
+ if (nonConstStrLen)
+ argSize = mlir::LLVM::TruncOp::create(
+ rewriter, op.getLoc(), rewriter.getI32Type(),
+ mlir::LLVM::AddOp::create(rewriter, op.getLoc(), nonConstStrLen,
+ getI64(bufSize)));
+ else
+ argSize = getI32(bufSize);
+
+ // call the printf_alloc function
+ unsigned globalAS = 0;
+ if (auto as = mlir::dyn_cast_if_present<mlir::IntegerAttr>(
+ dataLayout.getGlobalMemorySpace()))
+ globalAS = as.getValue().getZExtValue();
+ auto ptrTy =
+ mlir::LLVM::LLVMPointerType::get(rewriter.getContext(), globalAS);
+ auto allocName =
+ mlir::StringAttr::get(rewriter.getContext(), "__printf_alloc");
+ bool declared = symbolTables.lookupSymbolIn(mod, allocName);
+ mlir::Value ptr = createCall(allocName, ptrTy, argSize);
+ if (!declared)
+ mlir::cast<mlir::LLVM::LLVMFuncOp>(
+ symbolTables.lookupSymbolIn(mod, allocName))
+ .setNoUnwind(true);
+ return ptr;
+}
+
+// Prepare constant string argument to push onto the buffer
+void AMDGPUPrintfLowering::processConstantStringArg(
+ const StringData &sd, llvm::SmallVectorImpl<mlir::Value> &whatToStore) {
+ llvm::SmallVector<uint32_t, 16> words;
+ llvm::packAMDGPUPrintfConstantString(sd.str, words);
+ whatToStore.reserve(whatToStore.size() + words.size());
+ for (uint32_t word : words)
+ whatToStore.push_back(getI32(word));
+}
+
+mlir::Value AMDGPUPrintfLowering::processNonStringArg(mlir::Value arg) {
+ mlir::Type ty = arg.getType();
+
+ if (auto intTy = mlir::dyn_cast<mlir::IntegerType>(ty))
+ if (intTy.getWidth() < 64)
+ return mlir::LLVM::ZExtOp::create(rewriter, op.getLoc(),
+ rewriter.getI64Type(), arg);
+
+ if (mlir::isa<mlir::FloatType>(ty))
+ if (dataLayout.getTypeSize(ty).getFixedValue() < 8)
+ return mlir::LLVM::FPExtOp::create(rewriter, op.getLoc(),
+ rewriter.getF64Type(), arg);
+
+ return arg;
+}
+
+void AMDGPUPrintfLowering::callBufferedPrintfArgPush(
+ mlir::ValueRange args, mlir::Value ptrToStore,
+ llvm::ArrayRef<StringData> stringContents, bool isConstFmtStr) {
+ mlir::Type ptrTy = ptrToStore.getType();
+ auto gepInBounds = [&](mlir::Value ptr, mlir::LLVM::GEPArg offset) {
+ return mlir::LLVM::GEPOp::create(rewriter, op.getLoc(), ptrTy,
+ rewriter.getI8Type(), ptr,
+ llvm::ArrayRef<mlir::LLVM::GEPArg>{offset},
+ mlir::LLVM::GEPNoWrapFlags::inbounds);
+ };
+
+ const StringData *strIt = stringContents.begin();
+ size_t i = isConstFmtStr ? 1 : 0;
+ for (; i < args.size(); i++) {
+ llvm::SmallVector<mlir::Value, 32> whatToStore;
+ if ((i == 0) || isCStringArg(args, i)) {
+ if (strIt->isConst) {
+ processConstantStringArg(*strIt, whatToStore);
+ strIt++;
+ } else {
+ // This copies the contents of the string, however the next offset
+ // is at aligned length, the extra space that might be created due
+ // to alignment padding is not populated with any specific value
+ // here. This would be safe as long as runtime is sync with
+ // the offsets.
+ mlir::LLVM::MemcpyOp::create(rewriter, op.getLoc(), ptrToStore, args[i],
+ strIt->realSize, /*isVolatile=*/false);
+ ptrToStore = gepInBounds(ptrToStore, strIt->alignedSize);
+
+ // done with current argument, move to next
+ strIt++;
+ continue;
+ }
+ } else {
+ whatToStore.push_back(processNonStringArg(args[i]));
+ }
+
+ for (mlir::Value toStore : whatToStore) {
+ mlir::LLVM::StoreOp::create(rewriter, op.getLoc(), toStore, ptrToStore);
+ ptrToStore = gepInBounds(
+ ptrToStore,
+ dataLayout.getTypeSize(toStore.getType()).getFixedValue());
+ }
+ }
+}
+
+void AMDGPUPrintfLowering::addPrintfFormatMetadata(llvm::StringRef entry,
+ bool onlyIfEmpty) {
+ constexpr llvm::StringLiteral name = "llvm.printf.fmts";
+ mlir::Attribute node = mlir::LLVM::MDNodeAttr::get(
+ rewriter.getContext(),
+ {mlir::LLVM::MDStringAttr::get(
+ rewriter.getContext(),
+ mlir::StringAttr::get(rewriter.getContext(), entry))});
+
+ mlir::LLVM::NamedMetadataOp metaD;
+ for (auto md : mod.getOps<mlir::LLVM::NamedMetadataOp>()) {
+ if (md.getMetadataName() == name) {
+ metaD = md;
+ break;
+ }
+ }
+
+ if (!metaD) {
+ mlir::OpBuilder::InsertionGuard guard(rewriter);
+ rewriter.setInsertionPointToEnd(mod.getBody());
+ mlir::LLVM::NamedMetadataOp::create(rewriter, mod.getLoc(), name,
+ rewriter.getArrayAttr(node));
+ return;
+ }
+
+ if (onlyIfEmpty && !metaD.getNodes().empty())
+ return;
+ llvm::SmallVector<mlir::Attribute> nodes(metaD.getNodes().getValue());
+ nodes.push_back(node);
+ rewriter.modifyOpInPlace(
+ metaD, [&] { metaD.setNodesAttr(rewriter.getArrayAttr(nodes)); });
+}
+
+mlir::Value AMDGPUPrintfLowering::emitBuffered(mlir::ValueRange args) {
+ llvm::SmallVector<StringData, 8> stringContents;
+ bool isConstFmtStr = !fmtStr.empty();
+
+ mlir::Value argSize;
+ mlir::Value ptr =
+ callBufferedPrintfStart(args, isConstFmtStr, stringContents, argSize);
+
+ // The buffered version still follows OpenCL printf standards for
+ // printf return value, i.e 0 on success, -1 on failure.
+ mlir::Value cmp = mlir::LLVM::ICmpOp::create(
+ rewriter, op.getLoc(), mlir::LLVM::ICmpPredicate::ne, ptr,
+ mlir::LLVM::ZeroOp::create(rewriter, op.getLoc(), ptr.getType()));
+
+ // The continuation block doubles as the end block.
+ mlir::Block *prev = rewriter.getInsertionBlock();
+ mlir::Block *argPush =
+ rewriter.createBlock(cont->getParent(), cont->getIterator());
+ rewriter.setInsertionPointToEnd(prev);
+ mlir::LLVM::CondBrOp::create(rewriter, op.getLoc(), cmp, argPush, cont);
+ rewriter.setInsertionPointToEnd(argPush);
+
+ // Create controlDWord and store as the first entry, format as follows
+ // Bit 0 (LSB) -> stream (1 if stderr, 0 if stdout, printf always outputs to
+ // stdout) Bit 1 -> constant format string (1 if constant) Bits 2-31 -> size
+ // of printf data frame
+ mlir::Value controlDWord =
+ mlir::LLVM::ShlOp::create(rewriter, op.getLoc(), argSize, getI32(2));
+ if (isConstFmtStr)
+ controlDWord = mlir::LLVM::OrOp::create(rewriter, op.getLoc(), controlDWord,
+ getI32(2));
+
+ mlir::LLVM::StoreOp::create(rewriter, op.getLoc(), controlDWord, ptr);
+
+ ptr = mlir::LLVM::GEPOp::create(rewriter, op.getLoc(), ptr.getType(),
+ rewriter.getI8Type(), ptr,
+ llvm::ArrayRef<mlir::LLVM::GEPArg>{4},
+ mlir::LLVM::GEPNoWrapFlags::inbounds);
+
+ // Create MD5 hash for constant format string, push low 64 bits of the
+ // same onto buffer and metadata.
+ if (isConstFmtStr) {
+ addPrintfFormatMetadata(llvm::getAMDGPUPrintfFormatMetadata(fmtStr),
+ /*onlyIfEmpty=*/false);
+
+ mlir::LLVM::StoreOp::create(rewriter, op.getLoc(),
+ getI64(llvm::getAMDGPUPrintfFormatHash(fmtStr)),
+ ptr);
+ ptr = mlir::LLVM::GEPOp::create(rewriter, op.getLoc(), ptr.getType(),
+ rewriter.getI8Type(), ptr,
+ llvm::ArrayRef<mlir::LLVM::GEPArg>{8},
+ mlir::LLVM::GEPNoWrapFlags::inbounds);
+ } else {
+ // Include a dummy metadata instance in case of only non constant
+ // format string usage, This might be an absurd usecase but needs to
+ // be done for completeness
+ addPrintfFormatMetadata(llvm::AMDGPUPrintfNonConstFormatMetadata,
+ /*onlyIfEmpty=*/true);
+ }
+
+ // Push The printf arguments onto buffer
+ callBufferedPrintfArgPush(args, ptr, stringContents, isConstFmtStr);
+
+ // End block, returns -1 on failure
+ mlir::LLVM::BrOp::create(rewriter, op.getLoc(), cont);
+ rewriter.setInsertionPointToStart(cont);
+ mlir::Value notCmp = mlir::LLVM::XOrOp::create(
+ rewriter, op.getLoc(), cmp,
+ mlir::LLVM::ConstantOp::create(rewriter, op.getLoc(),
+ rewriter.getI1Type(), llvm::APInt(1, 1)));
+ return mlir::LLVM::SExtOp::create(rewriter, op.getLoc(),
+ rewriter.getI32Type(), notCmp);
+}
+
+mlir::LogicalResult AMDGPUPrintfLowering::lower(mlir::ValueRange args,
+ bool isBuffered) {
+ // Hostcall passes every argument as a 64-bit word, and only knows how to
+ // widen the types the default argument promotions produce.
+ if (!isBuffered) {
+ for (auto [cirArg, arg] : llvm::zip(op.getArgs(), args.drop_front())) {
+ mlir::Type ty = arg.getType();
+ if (ty.isInteger(32) || ty.isInteger(64) || ty.isF64() ||
+ mlir::isa<mlir::LLVM::LLVMPointerType>(ty))
+ continue;
+ return op.emitError() << "unsupported argument type " << cirArg.getType()
+ << " for AMDGPU printf";
+ }
+ }
+
+ if (getConstantStringInfo(op.getFormat(), mod, symbolTables, fmtStr))
+ llvm::locateAMDGPUPrintfCStrings(specIsCString, fmtStr);
+ else
+ fmtStr.clear();
+
+ mlir::Block *block = op->getBlock();
+ cont = rewriter.splitBlock(block, mlir::Block::iterator(op));
+ rewriter.setInsertionPointToEnd(block);
+
+ mlir::Value result;
+ if (isBuffered) {
+ result = emitBuffered(args);
+ } else {
+ result = emitHostcall(args);
+ mlir::LLVM::BrOp::create(rewriter, op.getLoc(), cont);
+ }
+
+ rewriter.replaceOp(op, result);
+ return mlir::success();
+}
+
+mlir::LogicalResult CIRToLLVMOffloadPrintfOpLowering::matchAndRewrite(
+ cir::OffloadPrintfOp op, OpAdaptor adaptor,
+ mlir::ConversionPatternRewriter &rewriter) const {
+ auto mod = op->getParentOfType<mlir::ModuleOp>();
+ llvm::StringRef tripleStr;
+ if (auto tripleAttr = mod->getAttrOfType<mlir::StringAttr>(
+ cir::CIRDialect::getTripleAttrName()))
+ tripleStr = tripleAttr.getValue();
+ llvm::Triple triple(tripleStr);
+
+ if (triple.isAMDGCN() ||
+ (triple.isSPIRV() && triple.getVendor() == llvm::Triple::AMD)) {
+ bool isBuffered = false;
+ if (auto kind = mod->getAttrOfType<mlir::StringAttr>(
+ cir::CIRDialect::getAMDGPUPrintfKindAttrName()))
+ isBuffered = kind.getValue() == "buffered";
+ return AMDGPUPrintfLowering(op, rewriter, dataLayout, symbolTables)
+ .lower(adaptor.getOperands(), isBuffered);
+ }
+
+ return op.emitError() << "offload printf lowering is NYI for target '"
+ << tripleStr << "'";
+}
+
+} // namespace direct
+} // namespace cir
diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip
index 7a8ffe297a2d9..c1c0bfb87cf3b 100644
--- a/clang/test/CIR/CodeGenHIP/device-printf.hip
+++ b/clang/test/CIR/CodeGenHIP/device-printf.hip
@@ -37,36 +37,41 @@ extern "C" __device__ int printf(const char *, ...);
// 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
-// 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
+// CIR: cir.offload.printf(%[[FMT]], %[[V]]) : (!cir.ptr<!s8i>, !s32i) -> !s32i
+// LLVM-LABEL: @_Z1pi(
+// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0)
+// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}@.str{{( to ptr.*\))?}}, i64 {{%.*}}, i32 0)
+// LLVM: [[A:%.*]] = zext i32 {{%.*}} to i64
+// LLVM: call{{.*}} i64 @__ockl_printf_append_args(i64 [[D1]], i32 1, i64 [[A]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// BUFFERED-LABEL: @_Z1pi(
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 20)
+// BUFFERED: [[OK:%.*]] = icmp ne ptr addrspace(1) [[BUF]], null
+// BUFFERED: br i1 [[OK]],
+// BUFFERED: store i32 82, ptr addrspace(1) [[BUF]], align 4
+// BUFFERED: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4
+// BUFFERED: store i64 -3402830678189845200, ptr addrspace(1) [[P0]], align 8
+// BUFFERED: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i32 8
+// BUFFERED: [[A:%.*]] = zext i32 {{%.*}} to i64
+// BUFFERED: store i64 [[A]], ptr addrspace(1) [[P1]], align 8
__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.
+// The result of printf must be forwarded to its users. The expansion's block
+// layout differs between CIR and classic codegen, 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: %[[RES:.*]] = cir.offload.printf({{.*}})
// CIR: cir.store %[[RES]], %[[RETVAL]]
-// LLVM-LABEL: @_Z1ri
+// 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
+// BUFFERED-LABEL: @_Z1ri(
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 20)
+// BUFFERED: [[OK:%.*]] = icmp ne ptr addrspace(1) [[BUF]], null
+// BUFFERED-DAG: [[FAIL:%.*]] = xor i1 [[OK]], true
+// BUFFERED-DAG: [[RES:%.*]] = sext i1 [[FAIL]] to i32
+// BUFFERED-DAG: {{store|ret}} i32 [[RES]]
__device__ int r(int v) { return printf("%d\n", v); }
// Device printf cannot throw, so a pending cleanup must not turn it into an
@@ -74,16 +79,153 @@ __device__ int r(int v) { return printf("%d\n", v); }
struct D { __device__ ~D(); };
// CIR-LABEL: cir.func {{.*}} @_Z2ehi(
// CIR: cir.cleanup.scope {
-// CIR: cir.call @__cir_amdgpu_printf({{.*}}) nothrow
+// CIR: cir.offload.printf({{.*}})
// CIR: } cleanup all {
// 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); }
+
+// %s arguments. A constant string is known at compile time; the length of a
+// runtime string is computed by a loop. The buffered kind stores a constant
+// string inline and copies a runtime string.
+// CIR-LABEL: cir.func {{.*}} @_Z1sPKc(
+// CIR: cir.offload.printf({{.*}}) : (!cir.ptr<!s8i>, !cir.ptr<!s8i>, !cir.ptr<!s8i{{.*}}>) -> !s32i
+// LLVM-LABEL: @_Z1sPKc(
+// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0)
+// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}@.str.1{{( to ptr.*\))?}}, i64 {{%.*}}, i32 0)
+// LLVM: [[D2:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D1]], ptr{{.*}}@.str.2{{( to ptr.*\))?}}, i64 {{%.*}}, i32 0)
+// LLVM: [[ISNULL:%.*]] = icmp eq ptr{{( addrspace\(4\))?}} [[STR:%[^,]+]], null
+// LLVM: br i1 [[ISNULL]], label %[[JOIN:.*]], label %[[WHILE:.*]]
+// LLVM: [[WHILE]]:
+// LLVM: [[CUR:%.*]] = phi ptr
+// LLVM: getelementptr i8, ptr{{( addrspace\(4\))?}} [[CUR]], i64 1
+// LLVM: [[CH:%.*]] = load i8, ptr{{( addrspace\(4\))?}} [[CUR]], align 1
+// LLVM: [[ISNUL:%.*]] = icmp eq i8 [[CH]], 0
+// LLVM: br i1 [[ISNUL]], label %[[DONE:.*]], label %[[WHILE]]
+// LLVM: [[DONE]]:
+// LLVM: [[END:%.*]] = ptrtoaddr ptr{{( addrspace\(4\))?}} [[CUR]] to i64
+// LLVM: [[BEGIN:%.*]] = ptrtoaddr ptr{{( addrspace\(4\))?}} [[STR]] to i64
+// LLVM: [[DIFF:%.*]] = sub i64 [[END]], [[BEGIN]]
+// LLVM: [[LEN:%.*]] = add i64 [[DIFF]], 1
+// LLVM: br label %[[JOIN]]
+// LLVM: [[JOIN]]:
+// LLVM: [[SIZE:%.*]] = phi i64 [ [[LEN]], %[[DONE]] ], [ 0, %{{.*}} ]
+// LLVM: call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D2]], ptr{{.*}}, i64 [[SIZE]], i32 1)
+// BUFFERED-LABEL: @_Z1sPKc(
+// BUFFERED: icmp eq ptr{{.*}}, null
+// BUFFERED: [[LEN:%.*]] = phi i64
+// BUFFERED: [[ADD:%.*]] = add i64 [[LEN]], 7
+// BUFFERED: [[ALIGNED:%.*]] = and i64 [[ADD]], 4294967288
+// BUFFERED: [[TOTAL:%.*]] = add i64 [[ALIGNED]], 20
+// BUFFERED: [[SIZE:%.*]] = trunc i64 [[TOTAL]] to i32
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 [[SIZE]])
+// BUFFERED: [[OK:%.*]] = icmp ne ptr addrspace(1) [[BUF]], null
+// BUFFERED: br i1 [[OK]],
+// BUFFERED: [[SHL:%.*]] = shl i32 [[SIZE]], 2
+// BUFFERED: [[CTRL:%.*]] = or i32 [[SHL]], 2
+// BUFFERED: store i32 [[CTRL]], ptr addrspace(1) [[BUF]], align 4
+// BUFFERED: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4
+// BUFFERED: store i64 5333508983605851890, ptr addrspace(1) [[P0]], align 8
+// BUFFERED: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i32 8
+// "cons" and "t\0" as little-endian i32 chunks.
+// BUFFERED: store i32 1936617315, ptr addrspace(1) [[P1]], align 4
+// BUFFERED: [[P2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P1]], i32 4
+// BUFFERED: store i32 116, ptr addrspace(1) [[P2]], align 4
+// BUFFERED: [[P3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P2]], i32 4
+// BUFFERED: call{{.*}} void @llvm.memcpy.p1.{{p0|p4}}.i64(ptr addrspace(1){{.*}} [[P3]], ptr{{.*}}, i64 [[LEN]], i1 false)
+__device__ void s(const char *str) { printf("%s|%s\n", "const", str); }
+
+// A runtime format string. The buffered kind copies it instead of storing its
+// hash, and does not set the constant-format bit in the control dword.
+// CIR-LABEL: cir.func {{.*}} @_Z2nfPKci(
+// CIR: cir.offload.printf({{.*}}) : (!cir.ptr<!s8i{{.*}}>, !s32i) -> !s32i
+// LLVM-LABEL: @_Z2nfPKci(
+// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0)
+// LLVM: icmp eq ptr{{.*}}, null
+// LLVM: [[LEN:%.*]] = phi i64
+// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}, i64 [[LEN]], i32 0)
+// LLVM: [[A:%.*]] = zext i32 {{%.*}} to i64
+// LLVM: call{{.*}} i64 @__ockl_printf_append_args(i64 [[D1]], i32 1, i64 [[A]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// BUFFERED-LABEL: @_Z2nfPKci(
+// BUFFERED: [[LEN:%.*]] = phi i64
+// BUFFERED: [[ADD:%.*]] = add i64 [[LEN]], 7
+// BUFFERED: [[ALIGNED:%.*]] = and i64 [[ADD]], 4294967288
+// BUFFERED: [[TOTAL:%.*]] = add i64 [[ALIGNED]], 12
+// BUFFERED: [[SIZE:%.*]] = trunc i64 [[TOTAL]] to i32
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 [[SIZE]])
+// BUFFERED: [[CTRL:%.*]] = shl i32 [[SIZE]], 2
+// BUFFERED-NEXT: store i32 [[CTRL]], ptr addrspace(1) [[BUF]], align 4
+// BUFFERED-NEXT: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4
+// BUFFERED-NEXT: call{{.*}} void @llvm.memcpy.p1.{{p0|p4}}.i64(ptr addrspace(1){{.*}} [[P0]], ptr{{.*}}, i64 [[LEN]], i1 false)
+// BUFFERED-NEXT: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i64 [[ALIGNED]]
+// BUFFERED-NEXT: [[A:%.*]] = zext i32 {{%.*}} to i64
+// BUFFERED-NEXT: store i64 [[A]], ptr addrspace(1) [[P1]], align 8
+__device__ void nf(const char *fmt, int v) { printf(fmt, v); }
+
+// No arguments: the format string is the last thing appended.
+// CIR-LABEL: cir.func {{.*}} @_Z1zv(
+// CIR: cir.offload.printf(%{{.*}}) : (!cir.ptr<!s8i>) -> !s32i
+// LLVM-LABEL: @_Z1zv(
+// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0)
+// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}@.str.3{{( to ptr.*\))?}}, i64 {{%.*}}, i32 1)
+// LLVM-NOT: @__ockl_printf_append_args
+// LLVM: trunc i64 [[D1]] to i32
+// BUFFERED-LABEL: @_Z1zv(
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 12)
+// BUFFERED: store i32 50, ptr addrspace(1) [[BUF]], align 4
+// BUFFERED: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4
+// BUFFERED: store i64 3806265321777173681, ptr addrspace(1) [[P0]], align 8
+__device__ void z() { printf("hello\n"); }
+
+// Every promoted scalar kind. Each argument is appended separately and widened
+// to 64 bits.
+// CIR-LABEL: cir.func {{.*}} @_Z4manyildPvfcys(
+// CIR: cir.offload.printf({{.*}}) : (!cir.ptr<!s8i>, !s32i, !s64i, !cir.double, !cir.ptr<!void{{.*}}>, !cir.double, !s32i, !u64i, !s32i) -> !s32i
+// LLVM-LABEL: @_Z4manyildPvfcys(
+// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0)
+// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}@.str.4{{( to ptr.*\))?}}, i64 {{%.*}}, i32 0)
+// LLVM: [[A1:%.*]] = zext i32 {{%.*}} to i64
+// LLVM: [[D2:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D1]], i32 1, i64 [[A1]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[D3:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D2]], i32 1, i64 {{%.*}}, i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[A3:%.*]] = bitcast double {{%.*}} to i64
+// LLVM: [[D4:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D3]], i32 1, i64 [[A3]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[A4:%.*]] = ptrtoint ptr{{( addrspace\(4\))?}} {{%.*}} to i64
+// LLVM: [[D5:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D4]], i32 1, i64 [[A4]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[A5:%.*]] = bitcast double {{%.*}} to i64
+// LLVM: [[D6:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D5]], i32 1, i64 [[A5]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[A6:%.*]] = zext i32 {{%.*}} to i64
+// LLVM: [[D7:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D6]], i32 1, i64 [[A6]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[D8:%.*]] = call{{.*}} i64 @__ockl_printf_append_args(i64 [[D7]], i32 1, i64 {{%.*}}, i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 0)
+// LLVM: [[A8:%.*]] = zext i32 {{%.*}} to i64
+// LLVM: call{{.*}} i64 @__ockl_printf_append_args(i64 [[D8]], i32 1, i64 [[A8]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// BUFFERED-LABEL: @_Z4manyildPvfcys(
+// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 76)
+// BUFFERED: store i32 306, ptr addrspace(1) [[BUF]], align 4
+// BUFFERED: store i64 -3000127804524445155, ptr addrspace(1)
+// BUFFERED: store i64 {{%.*}}, ptr addrspace(1)
+// BUFFERED: store i64 {{%.*}}, ptr addrspace(1)
+// BUFFERED: store double {{%.*}}, ptr addrspace(1)
+// BUFFERED: store ptr{{( addrspace\(4\))?}} {{%.*}}, ptr addrspace(1)
+// BUFFERED: store double {{%.*}}, ptr addrspace(1)
+// BUFFERED: store i64 {{%.*}}, ptr addrspace(1)
+// BUFFERED: store i64 {{%.*}}, ptr addrspace(1)
+// BUFFERED: store i64 {{%.*}}, ptr addrspace(1)
+__device__ void many(int a, long b, double c, void *d, float e, char f,
+ unsigned long long g, short h) {
+ printf("%d %ld %f %p %f %c %llu %hd\n", a, b, c, d, e, f, g, h);
+}
+
+// Constant formats are recorded with the low 64 bits of their MD5 hash, in
+// emission order. The runtime format in nf adds no entry.
+// LLVM-NOT: !llvm.printf.fmts
+// BUFFERED: !llvm.printf.fmts = !{[[F0:![0-9]+]], [[F0]], [[F0]], [[F1:![0-9]+]], [[F2:![0-9]+]], [[F3:![0-9]+]]}
+// BUFFERED-DAG: [[F0]] = !{!"0:0:d0c6b786f3b15530,%d\0A"}
+// BUFFERED-DAG: [[F1]] = !{!"0:0:4a046e29962e76f2,%s|%s\0A"}
+// BUFFERED-DAG: [[F2]] = !{!"0:0:34d29224c96a94b1,hello\0A"}
+// BUFFERED-DAG: [[F3]] = !{!"0:0:d65d67a83a8c421d,%d %ld %f %p %f %c %llu %hd\0A"}
diff --git a/clang/test/CIR/IR/offload-printf.cir b/clang/test/CIR/IR/offload-printf.cir
new file mode 100644
index 0000000000000..890c025f4af42
--- /dev/null
+++ b/clang/test/CIR/IR/offload-printf.cir
@@ -0,0 +1,15 @@
+// RUN: cir-opt %s --verify-roundtrip | FileCheck %s
+
+!s32i = !cir.int<s, 32>
+!s64i = !cir.int<s, 64>
+!s8i = !cir.int<s, 8>
+module {
+ cir.func @f(%fmt : !cir.ptr<!s8i>, %i : !s32i, %l : !s64i, %d : !cir.double,
+ %p : !cir.ptr<!s8i>) -> !s32i {
+ // CHECK: cir.offload.printf(%{{.*}}) : (!cir.ptr<!s8i>) -> !s32i
+ %0 = cir.offload.printf(%fmt) : (!cir.ptr<!s8i>) -> !s32i
+ // CHECK: cir.offload.printf(%{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}) : (!cir.ptr<!s8i>, !s32i, !s64i, !cir.double, !cir.ptr<!s8i>) -> !s32i
+ %1 = cir.offload.printf(%fmt, %i, %l, %d, %p) : (!cir.ptr<!s8i>, !s32i, !s64i, !cir.double, !cir.ptr<!s8i>) -> !s32i
+ cir.return %1 : !s32i
+ }
+}
diff --git a/clang/test/CIR/Lowering/offload-printf-nyi.cir b/clang/test/CIR/Lowering/offload-printf-nyi.cir
new file mode 100644
index 0000000000000..8b35ef52e7bc1
--- /dev/null
+++ b/clang/test/CIR/Lowering/offload-printf-nyi.cir
@@ -0,0 +1,13 @@
+// RUN: cir-opt %s --cir-to-llvm -verify-diagnostics
+
+!s32i = !cir.int<s, 32>
+!s8i = !cir.int<s, 8>
+
+module attributes {cir.triple = "nvptx64-nvidia-cuda"} {
+ cir.func @p(%fmt: !cir.ptr<!s8i>, %v: !s32i) -> !s32i {
+ // expected-error @below {{offload printf lowering is NYI for target 'nvptx64-nvidia-cuda'}}
+ // expected-error @below {{failed to legalize operation 'cir.offload.printf'}}
+ %0 = cir.offload.printf(%fmt, %v) : (!cir.ptr<!s8i>, !s32i) -> !s32i
+ cir.return %0 : !s32i
+ }
+}
diff --git a/clang/test/CIR/Tools/amdgpu-printf.cir b/clang/test/CIR/Tools/amdgpu-printf.cir
deleted file mode 100644
index 8383b9ed4c3ee..0000000000000
--- a/clang/test/CIR/Tools/amdgpu-printf.cir
+++ /dev/null
@@ -1,21 +0,0 @@
-// 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/test/CIR/Tools/offload-printf.cir b/clang/test/CIR/Tools/offload-printf.cir
new file mode 100644
index 0000000000000..bc1e96534c8b1
--- /dev/null
+++ b/clang/test/CIR/Tools/offload-printf.cir
@@ -0,0 +1,86 @@
+// RUN: cir-translate --cir-to-llvmir --target amdgcn-amd-amdhsa --disable-cc-lowering --split-input-file %s -o %t.ll
+// RUN: FileCheck %s -input-file %t.ll -check-prefixes=HOSTCALL,BUFFERED
+
+!s32i = !cir.int<s, 32>
+!s8i = !cir.int<s, 8>
+
+module attributes {cir.amdgpu_printf_kind = "hostcall"} {
+ cir.func @hostcall(%fmt: !cir.ptr<!s8i>, %v: !s32i) -> !s32i {
+ %0 = cir.offload.printf(%fmt, %v) : (!cir.ptr<!s8i>, !s32i) -> !s32i
+ cir.return %0 : !s32i
+ }
+}
+
+// The format is not a constant, so its length is computed at runtime.
+// HOSTCALL-LABEL: define{{.*}} i32 @hostcall(ptr %0, i32 %1)
+// HOSTCALL: [[DESC:%.*]] = call i64 @__ockl_printf_begin(i64 0)
+// HOSTCALL: [[ISNULL:%.*]] = icmp eq ptr %0, null
+// HOSTCALL: br i1 [[ISNULL]], label %[[JOIN:.*]], label %[[WHILE:.*]]
+// HOSTCALL: [[WHILE]]:
+// HOSTCALL: [[CUR:%.*]] = phi ptr [ [[NEXT:%.*]], %[[WHILE]] ], [ %0, %{{.*}} ]
+// HOSTCALL: [[NEXT]] = getelementptr i8, ptr [[CUR]], i64 1
+// HOSTCALL: [[CH:%.*]] = load i8, ptr [[CUR]], align 1
+// HOSTCALL: [[ISNUL:%.*]] = icmp eq i8 [[CH]], 0
+// HOSTCALL: br i1 [[ISNUL]], label %[[DONE:.*]], label %[[WHILE]]
+// HOSTCALL: [[DONE]]:
+// HOSTCALL: [[END:%.*]] = ptrtoaddr ptr [[CUR]] to i64
+// HOSTCALL: [[BEGIN:%.*]] = ptrtoaddr ptr %0 to i64
+// HOSTCALL: [[DIFF:%.*]] = sub i64 [[END]], [[BEGIN]]
+// HOSTCALL: [[LEN:%.*]] = add i64 [[DIFF]], 1
+// HOSTCALL: br label %[[JOIN]]
+// HOSTCALL: [[JOIN]]:
+// HOSTCALL: [[SIZE:%.*]] = phi i64 [ [[LEN]], %[[DONE]] ], [ 0, %{{.*}} ]
+// HOSTCALL: [[DESC1:%.*]] = call i64 @__ockl_printf_append_string_n(i64 [[DESC]], ptr %0, i64 [[SIZE]], i32 0)
+// HOSTCALL: [[ARG:%.*]] = zext i32 %1 to i64
+// HOSTCALL: [[DESC2:%.*]] = call i64 @__ockl_printf_append_args(i64 [[DESC1]], i32 1, i64 [[ARG]], i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1)
+// HOSTCALL: [[RES:%.*]] = trunc i64 [[DESC2]] to i32
+// HOSTCALL: ret i32 [[RES]]
+// HOSTCALL-NOT: !llvm.printf.fmts
+
+// -----
+
+!s32i = !cir.int<s, 32>
+!s8i = !cir.int<s, 8>
+
+module attributes {cir.amdgpu_printf_kind = "buffered"} {
+ cir.global "private" constant cir_private @".str" = #cir.const_array<"%d %s\0A" : !cir.array<!s8i x 6>, trailing_zeros> : !cir.array<!s8i x 7> {alignment = 1 : i64}
+ cir.func @buffered(%v: !s32i, %str: !cir.ptr<!s8i>) -> !s32i {
+ %0 = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 7>>
+ %1 = cir.cast array_to_ptrdecay %0 : !cir.ptr<!cir.array<!s8i x 7>> -> !cir.ptr<!s8i>
+ %2 = cir.offload.printf(%1, %v, %str) : (!cir.ptr<!s8i>, !s32i, !cir.ptr<!s8i>) -> !s32i
+ cir.return %2 : !s32i
+ }
+}
+
+// The format is a constant and only its hash (the low 64 bits of its MD5,
+// which also keys the !llvm.printf.fmts entry) is stored. The %s argument is not
+// a constant, so it is copied after its length is computed at runtime.
+// BUFFERED-LABEL: define{{.*}} i32 @buffered(i32 %0, ptr %1)
+// BUFFERED: icmp eq ptr %1, null
+// BUFFERED: getelementptr i8, ptr %{{.*}}, i64 1
+// BUFFERED: [[LEN:%.*]] = phi i64
+// BUFFERED: [[ADD:%.*]] = add i64 [[LEN]], 7
+// BUFFERED: [[ALIGNED:%.*]] = and i64 [[ADD]], 4294967288
+// BUFFERED: [[SUM:%.*]] = add i64 [[ALIGNED]], 20
+// BUFFERED: [[SIZE:%.*]] = trunc i64 [[SUM]] to i32
+// BUFFERED: [[BUF:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 [[SIZE]])
+// BUFFERED: [[OK:%.*]] = icmp ne ptr addrspace(1) [[BUF]], null
+// BUFFERED: br i1 [[OK]], label %[[PUSH:.*]], label %[[CONT:.*]]
+// BUFFERED: [[PUSH]]:
+// BUFFERED: [[SHL:%.*]] = shl i32 [[SIZE]], 2
+// BUFFERED: [[CTRL:%.*]] = or i32 [[SHL]], 2
+// BUFFERED: store i32 [[CTRL]], ptr addrspace(1) [[BUF]], align 4
+// BUFFERED: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4
+// BUFFERED: store i64 -3179066091271850508, ptr addrspace(1) [[P1]], align 8
+// BUFFERED: [[P2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P1]], i32 8
+// BUFFERED: [[ARG:%.*]] = zext i32 %0 to i64
+// BUFFERED: store i64 [[ARG]], ptr addrspace(1) [[P2]], align 8
+// BUFFERED: [[P3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P2]], i32 8
+// BUFFERED: call void @llvm.memcpy.p1.p0.i64(ptr addrspace(1) [[P3]], ptr %1, i64 [[LEN]], i1 false)
+// BUFFERED: br label %[[CONT]]
+// BUFFERED: [[CONT]]:
+// BUFFERED: [[FAIL:%.*]] = xor i1 [[OK]], true
+// BUFFERED: [[RES:%.*]] = sext i1 [[FAIL]] to i32
+// BUFFERED: ret i32 [[RES]]
+// BUFFERED: !llvm.printf.fmts = !{[[FMT:![0-9]+]]}
+// BUFFERED: [[FMT]] = !{!"0:0:d3e1b03bc04075f4,%d %s\0A"}
diff --git a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
index e7adf91b3f569..df4b414312a0c 100644
--- a/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
+++ b/llvm/include/llvm/Transforms/Utils/AMDGPUEmitPrintf.h
@@ -14,14 +14,42 @@
#ifndef LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H
#define LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/SparseBitVector.h"
+#include "llvm/ADT/StringRef.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/Support/Compiler.h"
+#include <string>
namespace llvm {
LLVM_ABI Value *emitAMDGPUPrintfCall(IRBuilder<> &Builder,
ArrayRef<Value *> Args, bool isBuffered);
+/// Marks the printf arguments that the format string \p Fmt specifies as
+/// strings, i.e. with a "%s" specifier. Argument 0 is the format string itself,
+/// and each '*' in a specifier consumes an argument of its own.
+LLVM_ABI void locateAMDGPUPrintfCStrings(SparseBitVector<8> &BV, StringRef Fmt);
+
+/// Returns the ID that buffered printf stores in place of the constant format
+/// string \p Fmt: the low 64 bits of its MD5 hash.
+LLVM_ABI uint64_t getAMDGPUPrintfFormatHash(StringRef Fmt);
+
+/// Returns the llvm.printf.fmts entry that maps the ID of the constant format
+/// string \p Fmt back to the string.
+LLVM_ABI std::string getAMDGPUPrintfFormatMetadata(StringRef Fmt);
+
+/// The llvm.printf.fmts entry added when a module only uses buffered printf
+/// with non-constant format strings.
+inline constexpr StringLiteral AMDGPUPrintfNonConstFormatMetadata =
+ "0:0:ffffffff,\"Non const format string\"";
+
+/// Splits the constant string argument \p Str, including its terminating NUL,
+/// into the little-endian 32-bit words that buffered printf stores for it,
+/// padded to a multiple of 8 bytes.
+LLVM_ABI void packAMDGPUPrintfConstantString(StringRef Str,
+ SmallVectorImpl<uint32_t> &Words);
+
} // end namespace llvm
#endif // LLVM_TRANSFORMS_UTILS_AMDGPUEMITPRINTF_H
diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
index 93fc21714ecdf..07a0fda28acb7 100644
--- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
+++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp
@@ -180,14 +180,14 @@ static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg,
// Scan the format string to locate all specifiers, and mark the ones that
// specify a string, i.e, the "%s" specifier with optional '*' characters.
-static void locateCStrings(SparseBitVector<8> &BV, StringRef Str) {
+void llvm::locateAMDGPUPrintfCStrings(SparseBitVector<8> &BV, StringRef Str) {
static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn";
size_t SpecPos = 0;
// Skip the first argument, the format string.
unsigned ArgIdx = 1;
while ((SpecPos = Str.find_first_of('%', SpecPos)) != StringRef::npos) {
- if (Str[SpecPos + 1] == '%') {
+ if (SpecPos + 1 < Str.size() && Str[SpecPos + 1] == '%') {
SpecPos += 2;
continue;
}
@@ -305,10 +305,25 @@ static Value *callBufferedPrintfStart(
return Builder.CreateCall(PrintfAllocFn, Alloc_args, "printf_alloc_fn");
}
-// Prepare constant string argument to push onto the buffer
-static void processConstantStringArg(StringData *SD, IRBuilder<> &Builder,
- SmallVectorImpl<Value *> &WhatToStore) {
- std::string Str(SD->Str.str() + '\0');
+uint64_t llvm::getAMDGPUPrintfFormatHash(StringRef Fmt) {
+ MD5 Hasher;
+ MD5::MD5Result Hash;
+ Hasher.update(Fmt);
+ Hasher.final(Hash);
+ return Hash.low();
+}
+
+std::string llvm::getAMDGPUPrintfFormatMetadata(StringRef Fmt) {
+ // Try sticking to llvm.printf.fmts format, although we are not going to
+ // use the ID and argument size fields while printing,
+ return "0:0:" +
+ llvm::utohexstr(getAMDGPUPrintfFormatHash(Fmt), /*LowerCase=*/true) +
+ "," + Fmt.str();
+}
+
+void llvm::packAMDGPUPrintfConstantString(StringRef S,
+ SmallVectorImpl<uint32_t> &Words) {
+ std::string Str(S.str() + '\0');
DataExtractor Extractor(Str, /*IsLittleEndian=*/true);
DataExtractor::Cursor Offset(0);
@@ -333,20 +348,21 @@ static void processConstantStringArg(StringData *SD, IRBuilder<> &Builder,
break;
}
cantFail(Offset.takeError(), "failed to read bytes from constant array");
-
- APInt IntVal(8 * ReadSize, ReadBytes);
-
- // TODO: Should not bother aligning up.
- if (ReadNow < ReadSize)
- IntVal = IntVal.zext(8 * ReadSize);
-
- Type *IntTy = Type::getIntNTy(Builder.getContext(), IntVal.getBitWidth());
- WhatToStore.push_back(ConstantInt::get(IntTy, IntVal));
+ Words.push_back(static_cast<uint32_t>(ReadBytes));
}
// Additional padding for 8 byte alignment
int Rem = (Str.size() % 8);
if (Rem > 0 && Rem <= 4)
- WhatToStore.push_back(ConstantInt::get(Builder.getInt32Ty(), 0));
+ Words.push_back(0);
+}
+
+// Prepare constant string argument to push onto the buffer
+static void processConstantStringArg(StringData *SD, IRBuilder<> &Builder,
+ SmallVectorImpl<Value *> &WhatToStore) {
+ SmallVector<uint32_t, 16> Words;
+ packAMDGPUPrintfConstantString(SD->Str, Words);
+ for (uint32_t Word : Words)
+ WhatToStore.push_back(Builder.getInt32(Word));
}
static Value *processNonStringArg(Value *Arg, IRBuilder<> &Builder) {
@@ -432,7 +448,7 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
StringRef FmtStr;
if (getConstantStringInfo(Fmt, FmtStr))
- locateCStrings(SpecIsCString, FmtStr);
+ locateAMDGPUPrintfCStrings(SpecIsCString, FmtStr);
if (IsBuffered) {
SmallVector<StringData, 8> StringContents;
@@ -479,21 +495,13 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
// same onto buffer and metadata.
NamedMDNode *metaD = M->getOrInsertNamedMetadata("llvm.printf.fmts");
if (IsConstFmtStr) {
- MD5 Hasher;
- MD5::MD5Result Hash;
- Hasher.update(FmtStr);
- Hasher.final(Hash);
-
- // Try sticking to llvm.printf.fmts format, although we are not going to
- // use the ID and argument size fields while printing,
- std::string MetadataStr =
- "0:0:" + llvm::utohexstr(Hash.low(), /*LowerCase=*/true) + "," +
- FmtStr.str();
- MDString *fmtStrArray = MDString::get(Ctx, MetadataStr);
+ MDString *fmtStrArray =
+ MDString::get(Ctx, getAMDGPUPrintfFormatMetadata(FmtStr));
MDNode *myMD = MDNode::get(Ctx, fmtStrArray);
metaD->addOperand(myMD);
- Builder.CreateStore(Builder.getInt64(Hash.low()), Ptr);
+ Builder.CreateStore(Builder.getInt64(getAMDGPUPrintfFormatHash(FmtStr)),
+ Ptr);
Ptr = Builder.CreateConstInBoundsGEP1_32(Int8Ty, Ptr, 8);
} else {
// Include a dummy metadata instance in case of only non constant
@@ -501,7 +509,7 @@ Value *llvm::emitAMDGPUPrintfCall(IRBuilder<> &Builder, ArrayRef<Value *> Args,
// be done for completeness
if (metaD->getNumOperands() == 0) {
MDString *fmtStrArray =
- MDString::get(Ctx, "0:0:ffffffff,\"Non const format string\"");
+ MDString::get(Ctx, AMDGPUPrintfNonConstFormatMetadata);
MDNode *myMD = MDNode::get(Ctx, fmtStrArray);
metaD->addOperand(myMD);
}
>From 905b41487717faa68af2d6d8dafdd6501e239396 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 29 Sep 2026 06:26:19 -0500
Subject: [PATCH 10/10] Fix test
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
---
clang/test/CIR/Tools/offload-printf.cir | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/test/CIR/Tools/offload-printf.cir b/clang/test/CIR/Tools/offload-printf.cir
index bc1e96534c8b1..63cb6363637fd 100644
--- a/clang/test/CIR/Tools/offload-printf.cir
+++ b/clang/test/CIR/Tools/offload-printf.cir
@@ -43,7 +43,7 @@ module attributes {cir.amdgpu_printf_kind = "hostcall"} {
!s8i = !cir.int<s, 8>
module attributes {cir.amdgpu_printf_kind = "buffered"} {
- cir.global "private" constant cir_private @".str" = #cir.const_array<"%d %s\0A" : !cir.array<!s8i x 6>, trailing_zeros> : !cir.array<!s8i x 7> {alignment = 1 : i64}
+ cir.global "private" constant cir_private @".str" = #cir.const_array<"%d %s\0A" : !cir.array<!s8i x 6>, trailing_zeros> : !cir.array<!s8i x 7>
cir.func @buffered(%v: !s32i, %str: !cir.ptr<!s8i>) -> !s32i {
%0 = cir.get_global @".str" : !cir.ptr<!cir.array<!s8i x 7>>
%1 = cir.cast array_to_ptrdecay %0 : !cir.ptr<!cir.array<!s8i x 7>> -> !cir.ptr<!s8i>
More information about the llvm-commits
mailing list