https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/230517
>From 88e019cefbf338fe58e53171fe10513813c5801c Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Mon, 5 Oct 2026 10:05:09 -0500 Subject: [PATCH 1/2] [CIR][CUDA] Add offload.printf op and CUDA target for it Previously the a printf call in CIR on device targets were left for each target to lower. This meant that the printf calls would be processed differently for each target, sometimes even before lowering to LLVM. This commit adds a new offload.printf op that is lowered to the correct target specific printf call during lowering to LLVM. This ensures that the printf calls are handled consistently across all targets, until reaching the target-specific lowering. As a first target, this commit migrates the NVPTX target printf support to use the new offload.printf op. Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/include/clang/CIR/Dialect/IR/CIROps.td | 47 ++++++ clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 44 +++++- clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp | 106 ------------- clang/lib/CIR/CodeGen/CIRGenFunction.h | 4 +- .../CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 + .../CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp | 5 +- .../CIR/Lowering/DirectToLLVM/LowerToLLVM.h | 9 ++ .../DirectToLLVM/LowerToLLVMOffloadPrintf.cpp | 146 ++++++++++++++++++ clang/test/CIR/CodeGenCUDA/device-printf.cu | 101 +++++++++++- clang/test/CIR/IR/offload-printf.cir | 17 ++ .../test/CIR/Lowering/offload-printf-nyi.cir | 13 ++ clang/test/CIR/Lowering/offload-printf.cir | 51 ++++++ 12 files changed, 425 insertions(+), 119 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 create mode 100644 clang/test/CIR/Lowering/offload-printf.cir diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index 8082f31f24187..99748eec557b6 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -9784,6 +9784,53 @@ def CIR_MemChrOp : CIR_Op<"libc.memchr", let hasVerifier = 1; } +//===----------------------------------------------------------------------===// +// OffloadPrintfOp +//===----------------------------------------------------------------------===// + +def CIR_OffloadPrintfOp : CIR_Op<"offload.printf"> { + let summary = "Device-side printf"; + let description = [{ + Represents a call to `printf` in device code. The target is responsible for + the appropriate lowering. + + `format` is the format string and `args` are the remaining arguments, + after the language's default argument promotions. `result` is the value + `printf` returns on the target, which follows the target's conventions + rather than C's. + + Like the GPU dialect's `gpu.printf`, the op represents a device printf + independently of the target, and each target expands it when lowering. + Unlike `gpu.printf`, it must represent any call to C `printf`, so: + + - `format` is a value rather than an attribute, as the format string need + not be a constant. + - `args` may also be pointers, such as strings printed with `%s`, + pointers to data members, and vectors. + - The op returns the result of `printf`. + - The op has no memory effect annotations, as printing may read memory + through `format` and `args`. + + 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, + CIR_AnyDataMemberType, CIR_AnyVectorType]>>:$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/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index c47165ec153a7..e396389945ea5 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -3167,7 +3167,7 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, (getTarget().getTriple().isSPIRV() && getTarget().getTriple().getVendor() == llvm::Triple::AMD)) { if (getTarget().getTriple().isNVPTX()) - return RValue::get(emitNVPTXDevicePrintfCallExpr(e)); + return RValue::get(emitDevicePrintfCallExpr(e)); if ((getTarget().getTriple().isAMDGCN() || getTarget().getTriple().isSPIRV()) && getLangOpts().HIP) @@ -3573,3 +3573,45 @@ mlir::Value CIRGenFunction::evaluateOrEmitBuiltinObjectSize( *objectSize); return emitBuiltinObjectSize(e, type, resType, emittedE, isDynamic); } + +// Emit a device printf as a cir.offload.printf. Unlike classic codegen, which +// emits each target's printf runtime interface here, CIR leaves that to the +// target-specific lowering of the op. +mlir::Value CIRGenFunction::emitDevicePrintfCallExpr(const CallExpr *expr) { + assert(cgm.getTriple().isNVPTX() || 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); // printf always has at least one arg. + CallArgList args; + emitCallArgs(args, + expr->getDirectCallee()->getType()->getAs<FunctionProtoType>(), + expr->arguments(), expr->getDirectCallee()); + + mlir::Location loc = getLoc(expr->getBeginLoc()); + + // We don't know how to emit non-scalar varargs, nor scalars that + // cir.offload.printf does not accept. + bool hasNonScalar = + llvm::any_of(llvm::drop_begin(args), [&](const CallArg &a) { + if (a.hasLValue() || !a.getKnownRValue().isScalar()) + return true; + mlir::Type ty = a.getKnownRValue().getValue().getType(); + return !mlir::isa<cir::IntType, cir::PointerType, cir::DataMemberType, + cir::VectorType>(ty) && + !cir::isAnyFloatingPointType(ty); + }); + 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()); + + return cir::OffloadPrintfOp::create(builder, loc, builder.getSInt32Ty(), + callArgs.front(), + llvm::drop_begin(callArgs)); +} diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp index 5b3618c137fb7..e800f9e9ae1d7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinNVPTX.cpp @@ -14,7 +14,6 @@ #include "mlir/IR/Value.h" #include "clang/Basic/TargetBuiltins.h" -#include "clang/CIR/Dialect/IR/CIRDataLayout.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" #include "llvm/Support/NVPTXAddrSpace.h" @@ -989,108 +988,3 @@ CIRGenFunction::emitNVPTXBuiltinExpr(unsigned builtinId, const CallExpr *expr) { return std::nullopt; } } - -// vprintf takes two args: A format string, and a pointer to a buffer containing -// the varargs. -// -// For example, the call -// -// printf("format string", arg1, arg2, arg3); -// -// is converted into something resembling -// -// struct Tmp { -// Arg1 a1; -// Arg2 a2; -// Arg3 a3; -// }; -// char* buf = alloca(sizeof(Tmp)); -// *(Tmp*)buf = {a1, a2, a3}; -// vprintf("format string", buf); -// -// `buf` is aligned to the max of {alignof(Arg1), ...}. Furthermore, each of -// the args is itself aligned to its preferred alignment. -// -// Note that by the time this function runs, the arguments have already -// undergone the standard C vararg promotion (short -> int, float -> double -// etc). In this function we pack the arguments into the buffer described above. -static mlir::Value packArgsIntoNVPTXFormatBuffer(CIRGenFunction &cgf, - const CallArgList &args, - mlir::Location loc) { - const cir::CIRDataLayout dataLayout = cgf.cgm.getDataLayout(); - CIRGenBuilderTy &builder = cgf.getBuilder(); - - if (args.size() <= 1) - // If there are no arguments other than the format string, - // pass a nullptr to vprintf. - return builder.getNullPtr(builder.getVoidPtrTy(), loc); - - llvm::SmallVector<mlir::Type> argTypes; - for (const auto &arg : llvm::drop_begin(args)) - argTypes.push_back(arg.getKnownRValue().getValue().getType()); - - // We can directly store the arguments into a struct, and the alignment - // would automatically be correct. That's because vprintf does not - // accept aggregates. - mlir::Type allocaTy = builder.getAnonRecordTy( - argTypes, /*packed=*/false, cir::RecordType::getAllDataKinds(argTypes)); - auto allocaAlign = clang::CharUnits::fromQuantity( - dataLayout.getABITypeAlign(allocaTy).value()); - Address allocaAddr = - cgf.createTempAlloca(allocaTy, allocaAlign, loc, "printf_args"); - mlir::Value alloca = allocaAddr.getPointer(); - - for (auto [i, arg] : llvm::enumerate(llvm::drop_begin(args))) { - mlir::Value member = builder.createGetMember( - loc, cir::PointerType::get(argTypes[i], allocaAddr.getAddressSpace()), - alloca, /*name=*/"", - /*index=*/i); - auto abiAlign = clang::CharUnits::fromQuantity( - dataLayout.getABITypeAlign(argTypes[i]).value()); - cir::StoreOp::create(builder, loc, arg.getKnownRValue().getValue(), member, - /*is_volatile=*/false, - /*isNontemporal=*/false, - builder.getAlignmentAttr(abiAlign), - /*sync_scope=*/cir::SyncScopeKindAttr{}, - /*mem_order=*/cir::MemOrderAttr{}); - } - - return builder.createBitcast(alloca, builder.getVoidPtrTy()); -} - -mlir::Value -CIRGenFunction::emitNVPTXDevicePrintfCallExpr(const CallExpr *expr) { - assert(cgm.getTriple().isNVPTX()); - assert(expr->getBuiltinCallee() == Builtin::BIprintf || - expr->getBuiltinCallee() == Builtin::BI__builtin_printf); - assert(expr->getNumArgs() >= 1); // printf always has at least one arg. - CallArgList args; - emitCallArgs(args, - expr->getDirectCallee()->getType()->getAs<FunctionProtoType>(), - 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(llvm::drop_begin(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); - } - - mlir::Value packedData = packArgsIntoNVPTXFormatBuffer(*this, args, loc); - - // int vprintf(char *format, void *packedData); - auto vprintf = cgm.createRuntimeFunction( - cir::FuncType::get( - {cir::PointerType::get(builder.getSInt8Ty()), builder.getVoidPtrTy()}, - builder.getSInt32Ty()), - "vprintf"); - auto formatString = args[0].getKnownRValue().getValue(); - return builder - .createCallOp(loc, vprintf, mlir::ValueRange{formatString, packedData}) - .getResult(); -} diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h index c6427d60a3792..70a2c779c85d7 100644 --- a/clang/lib/CIR/CodeGen/CIRGenFunction.h +++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h @@ -2388,8 +2388,8 @@ class CIRGenFunction : public CIRGenTypeCache { std::optional<mlir::Value> emitNVPTXBuiltinExpr(unsigned builtinID, const CallExpr *expr); - /// Emit a device-side printf call for NVPTX targets. - mlir::Value emitNVPTXDevicePrintfCallExpr(const CallExpr *expr); + /// Emit a device-side printf call for GPU targets. + mlir::Value emitDevicePrintfCallExpr(const CallExpr *expr); LValue emitOpaqueValueLValue(const OpaqueValueExpr *e); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt index d3d0805afb4d8..009d437da0f16 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt +++ b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt @@ -6,6 +6,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 1871da4ec0607..94d1c4d158614 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -4677,9 +4677,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)); 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..1bac9b6c4c878 --- /dev/null +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp @@ -0,0 +1,146 @@ +//===- LowerToLLVMOffloadPrintf.cpp - Lower cir.offload.printf ------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// This file lowers cir.offload.printf to the printf runtime interface of each +// device target. The expansions follow the lowerings of the GPU dialect's +// gpu.printf in mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp. +// +//===----------------------------------------------------------------------===// + +#include "LowerToLLVM.h" + +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/BuiltinOps.h" +#include "clang/CIR/Dialect/IR/CIRDialect.h" +#include "llvm/TargetParser/Triple.h" + +namespace cir { +namespace direct { + +namespace { + +// vprintf takes two args: A format string, and a pointer to a buffer containing +// the varargs. +// +// For example, the call +// +// printf("format string", arg1, arg2, arg3); +// +// is converted into something resembling +// +// struct Tmp { +// Arg1 a1; +// Arg2 a2; +// Arg3 a3; +// }; +// char* buf = alloca(sizeof(Tmp)); +// *(Tmp*)buf = {a1, a2, a3}; +// vprintf("format string", buf); +// +// `buf` is aligned to the max of {alignof(Arg1), ...}. Furthermore, each of +// the args is itself aligned to its ABI alignment. +// +// Note that by the time this function runs, the arguments have already +// undergone the standard C vararg promotion (short -> int, float -> double +// etc). In this function we pack the arguments into the buffer described above. +mlir::Value +packArgsIntoNVPTXFormatBuffer(cir::OffloadPrintfOp op, mlir::ValueRange args, + mlir::ConversionPatternRewriter &rewriter, + const mlir::DataLayout &dataLayout) { + mlir::Location loc = op.getLoc(); + auto ptrTy = mlir::LLVM::LLVMPointerType::get(rewriter.getContext()); + + if (args.empty()) + // If there are no arguments other than the format string, + // pass a nullptr to vprintf. + return mlir::LLVM::ZeroOp::create(rewriter, loc, ptrTy); + + // We can directly store the arguments into a struct, and the alignment + // would automatically be correct. That's because vprintf does not + // accept aggregates. + auto allocaTy = mlir::LLVM::LLVMStructType::getLiteral( + rewriter.getContext(), llvm::to_vector(args.getTypes())); + + // Allocate the buffer in the entry block, so that a printf in a loop does + // not grow the stack on every iteration. + mlir::Value alloca; + { + mlir::OpBuilder::InsertionGuard guard(rewriter); + auto fn = op->getParentOfType<mlir::LLVM::LLVMFuncOp>(); + rewriter.setInsertionPointToStart(&fn.getBody().front()); + mlir::Value one = + mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(), 1); + alloca = + mlir::LLVM::AllocaOp::create(rewriter, loc, ptrTy, allocaTy, one, + dataLayout.getTypeABIAlignment(allocaTy)); + } + + // Member accesses are inbounds and nuw, as for cir.get_member. + mlir::LLVM::GEPNoWrapFlags flags = + mlir::LLVM::GEPNoWrapFlags::inbounds | mlir::LLVM::GEPNoWrapFlags::nuw; + for (auto [i, arg] : llvm::enumerate(args)) { + mlir::Value member = mlir::LLVM::GEPOp::create( + rewriter, loc, ptrTy, allocaTy, alloca, + llvm::ArrayRef<mlir::LLVM::GEPArg>{0, static_cast<int32_t>(i)}, flags); + mlir::LLVM::StoreOp::create(rewriter, loc, arg, member, + dataLayout.getTypeABIAlignment(arg.getType())); + } + + return alloca; +} + +// Lowers a printf to a call to vprintf, as GPUPrintfOpToVPrintfLowering does +// for gpu.printf. +mlir::LogicalResult +lowerNVPTXPrintf(cir::OffloadPrintfOp op, mlir::Value format, + mlir::ValueRange args, + mlir::ConversionPatternRewriter &rewriter, + const mlir::DataLayout &dataLayout, + mlir::SymbolTableCollection &symbolTables) { + mlir::Value packedData = + packArgsIntoNVPTXFormatBuffer(op, args, rewriter, dataLayout); + + // int vprintf(char *format, void *packedData); + auto ptrTy = mlir::LLVM::LLVMPointerType::get(rewriter.getContext()); + mlir::Type i32Ty = rewriter.getI32Type(); + const llvm::StringRef fnName = "vprintf"; + createLLVMFuncOpIfNotExist( + rewriter, symbolTables, op, fnName, + mlir::LLVM::LLVMFunctionType::get(i32Ty, {ptrTy, ptrTy})); + // vprintf takes a generic pointer, which can address a format string in any + // address space. + if (format.getType() != ptrTy) + format = mlir::LLVM::AddrSpaceCastOp::create(rewriter, op.getLoc(), ptrTy, + format); + rewriter.replaceOpWithNewOp<mlir::LLVM::CallOp>( + op, mlir::TypeRange{i32Ty}, fnName, mlir::ValueRange{format, packedData}); + return mlir::success(); +} + +} // namespace + +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.isNVPTX()) + return lowerNVPTXPrintf(op, adaptor.getFormat(), adaptor.getArgs(), + rewriter, dataLayout, symbolTables); + + return op.emitError() << "offload printf lowering is NYI for target '" + << tripleStr << "'"; +} + +} // namespace direct +} // namespace cir diff --git a/clang/test/CIR/CodeGenCUDA/device-printf.cu b/clang/test/CIR/CodeGenCUDA/device-printf.cu index ed49dc133943b..eeecdcb2ef626 100644 --- a/clang/test/CIR/CodeGenCUDA/device-printf.cu +++ b/clang/test/CIR/CodeGenCUDA/device-printf.cu @@ -12,17 +12,23 @@ // RUN: -fcuda-is-device -emit-llvm %s -o %t-og.ll // RUN: FileCheck --check-prefix=LLVM --input-file=%t-og.ll %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions -fclangir \ +// RUN: -emit-llvm %s -o %t-eh.ll +// RUN: FileCheck --check-prefix=EH --input-file=%t-eh.ll %s + +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_80 -x cuda \ +// RUN: -fcuda-is-device -fcxx-exceptions -fexceptions \ +// RUN: -emit-llvm %s -o %t-eh-og.ll +// RUN: FileCheck --check-prefix=EH --input-file=%t-eh-og.ll %s + __device__ void print_int() { printf("%d", 42); } // CIR: cir.func no_inline dso_local @_Z9print_intv() -// CIR: %[[#ALLOCA:]] = cir.alloca {{.*}} : !cir.ptr<!rec_anon_struct> // CIR: %[[#VAL:]] = cir.const #cir.int<42> : !s32i -// CIR: %[[#FIELD:]] = cir.get_member %[[#ALLOCA]][0] -// CIR: cir.store align(4) %[[#VAL]], %[[#FIELD]] -// CIR: %[[#BUF:]] = cir.cast bitcast %[[#ALLOCA]] : !cir.ptr<!rec_anon_struct> -> !cir.ptr<!void> -// CIR: cir.call @vprintf(%{{.+}}, %[[#BUF]]) +// CIR: cir.offload.printf(%{{.+}}, %[[#VAL]]) : (!cir.ptr<!s8i>, !s32i) -> !s32i // CIR: cir.return // LLVM: define dso_local void @_Z9print_intv() @@ -37,10 +43,91 @@ __device__ void print_no_args() { } // CIR: cir.func no_inline dso_local @_Z13print_no_argsv() -// CIR: %[[#NULL:]] = cir.const #cir.ptr<null> : !cir.ptr<!void> -// CIR: cir.call @vprintf(%{{.+}}, %[[#NULL]]) +// CIR: cir.offload.printf(%{{.+}}) : (!cir.ptr<!s8i>) -> !s32i // CIR: cir.return // LLVM: define dso_local void @_Z13print_no_argsv() // LLVM: call i32 @vprintf(ptr @{{.*}}, ptr null) // LLVM: ret void + +// The char is promoted to int. Each argument is stored with its own alignment. +__device__ void print_mixed(char c, double d, const char *s) { + printf("%c %f %s", c, d, s); +} + +// CIR: cir.func no_inline dso_local @_Z11print_mixedcdPKc( +// CIR: cir.offload.printf(%{{.+}}, %{{.+}}, %{{.+}}, %{{.+}}) : (!cir.ptr<!s8i>, !s32i, !cir.double, !cir.ptr<!s8i>) -> !s32i + +// LLVM: define dso_local void @_Z11print_mixedcdPKc( +// LLVM: %[[PACKED:.*]] = alloca {{%printf_args.*|. i32, double, ptr .}}, align 8 +// LLVM: %[[GEP0:.*]] = getelementptr inbounds nuw {{.*}}, ptr %[[PACKED]], i32 0, i32 0 +// LLVM: store i32 %{{.*}}, ptr %[[GEP0]], align 4 +// LLVM: %[[GEP1:.*]] = getelementptr inbounds nuw {{.*}}, ptr %[[PACKED]], i32 0, i32 1 +// LLVM: store double %{{.*}}, ptr %[[GEP1]], align 8 +// LLVM: %[[GEP2:.*]] = getelementptr inbounds nuw {{.*}}, ptr %[[PACKED]], i32 0, i32 2 +// LLVM: store ptr %{{.*}}, ptr %[[GEP2]], align 8 +// LLVM: call i32 @vprintf(ptr @{{.*}}, ptr %[[PACKED]]) + +// A pointer to a data member is passed as its offset, which the C++ ABI +// lowering has substituted by the time CIR is emitted. +struct Pair { + int first, second; +}; + +__device__ void print_member_pointer() { + printf("%ld", &Pair::second); +} + +// CIR: cir.func no_inline dso_local @_Z20print_member_pointerv() +// CIR: %[[#OFFSET:]] = cir.const #cir.int<4> : !s64i +// CIR: cir.offload.printf(%{{.+}}, %[[#OFFSET]]) : (!cir.ptr<!s8i>, !s64i) -> !s32i + +// LLVM: define dso_local void @_Z20print_member_pointerv() +// LLVM: %[[PACKED:.*]] = alloca {{%printf_args.*|. i64 .}}, align 8 +// LLVM: %[[GEP:.*]] = getelementptr inbounds nuw {{.*}}, ptr %[[PACKED]], i32 0, i32 0 +// LLVM: store i64 4, ptr %[[GEP]], align 8 +// LLVM: call i32 @vprintf(ptr @{{.*}}, ptr %[[PACKED]]) + +typedef int int4v __attribute__((ext_vector_type(4))); + +__device__ void print_vector(int4v v) { + printf("%v4d", v); +} + +// CIR: cir.func no_inline dso_local @_Z12print_vectorDv4_i( +// CIR: cir.offload.printf(%{{.+}}, %{{.+}}) : (!cir.ptr<!s8i>, !cir.vector<4 x !s32i>) -> !s32i + +// LLVM: define dso_local void @_Z12print_vectorDv4_i( +// LLVM: %[[PACKED:.*]] = alloca {{%printf_args.*|. <4 x i32> .}}, align 16 +// LLVM: %[[GEP:.*]] = getelementptr inbounds nuw {{.*}}, ptr %[[PACKED]], i32 0, i32 0 +// LLVM: store <4 x i32> %{{.*}}, ptr %[[GEP]], align 16 +// LLVM: call i32 @vprintf(ptr @{{.*}}, ptr %[[PACKED]]) + +// The buffer is allocated in the entry block, not on every iteration. +__device__ void print_in_loop(int n) { + for (int i = 0; i < n; ++i) + printf("%d", i); +} + +// LLVM: define dso_local void @_Z13print_in_loopi( +// LLVM-NOT: br +// LLVM: %[[PACKED:.*]] = alloca {{%printf_args.*|. i32 .}}, align +// LLVM: br +// LLVM: call i32 @vprintf(ptr @{{.*}}, ptr %[[PACKED]]) + +// vprintf is not called with invoke, even when a cleanup is active. +struct WithDtor { + __device__ ~WithDtor(); +}; + +__device__ void print_with_cleanup() { + WithDtor w; + printf("%d", 1); +} + +// EH: define dso_local void @_Z18print_with_cleanupv() #{{[0-9]+}} { +// EH-NOT: invoke +// EH: call i32 @vprintf( +// EH-NOT: landingpad +// EH: call void @_ZN8WithDtorD1Ev( +// EH: ret void diff --git a/clang/test/CIR/IR/offload-printf.cir b/clang/test/CIR/IR/offload-printf.cir new file mode 100644 index 0000000000000..98d216cd8bc67 --- /dev/null +++ b/clang/test/CIR/IR/offload-printf.cir @@ -0,0 +1,17 @@ +// RUN: cir-opt %s --verify-roundtrip | FileCheck %s + +!s32i = !cir.int<s, 32> +!s64i = !cir.int<s, 64> +!s8i = !cir.int<s, 8> +!rec_S = !cir.struct<"S" {data !s32i}> +module { + cir.func @f(%fmt : !cir.ptr<!s8i>, %i : !s32i, %l : !s64i, %d : !cir.double, + %p : !cir.ptr<!s8i>, %m : !cir.data_member<!s32i in !rec_S>, + %v : !cir.vector<4 x !s32i>) -> !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>, !cir.data_member<!s32i in !rec_S>, !cir.vector<4 x !s32i>) -> !s32i + %1 = cir.offload.printf(%fmt, %i, %l, %d, %p, %m, %v) : (!cir.ptr<!s8i>, !s32i, !s64i, !cir.double, !cir.ptr<!s8i>, !cir.data_member<!s32i in !rec_S>, !cir.vector<4 x !s32i>) -> !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..7e27aec9b8fe5 --- /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 = "x86_64-unknown-linux-gnu"} { + cir.func @p(%fmt: !cir.ptr<!s8i>, %v: !s32i) -> !s32i { + // expected-error @below {{offload printf lowering is NYI for target 'x86_64-unknown-linux-gnu'}} + // 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/Lowering/offload-printf.cir b/clang/test/CIR/Lowering/offload-printf.cir new file mode 100644 index 0000000000000..f0948bacd271b --- /dev/null +++ b/clang/test/CIR/Lowering/offload-printf.cir @@ -0,0 +1,51 @@ +// RUN: cir-translate %s -cir-to-llvmir --target nvptx64-nvidia-cuda --disable-cc-lowering -o - \ +// RUN: | FileCheck %s + +!s32i = !cir.int<s, 32> +!s8i = !cir.int<s, 8> + +// All printfs share one vprintf declaration. +// CHECK: declare i32 @vprintf(ptr, ptr) +// CHECK-NOT: declare + +module { + // Without arguments, vprintf gets a null buffer. + cir.func @no_args(%fmt: !cir.ptr<!s8i>) -> !s32i { + %0 = cir.offload.printf(%fmt) : (!cir.ptr<!s8i>) -> !s32i + cir.return %0 : !s32i + } + + // CHECK-LABEL: define i32 @no_args(ptr %0) + // CHECK: %[[R:.*]] = call i32 @vprintf(ptr %0, ptr null) + // CHECK: ret i32 %[[R]] + + // The buffer is allocated in the entry block, even when the printf is not. + cir.func @args(%fmt: !cir.ptr<!s8i>, %i: !s32i, %d: !cir.double) -> !s32i { + cir.br ^bb1 + ^bb1: + %0 = cir.offload.printf(%fmt, %i, %d) : (!cir.ptr<!s8i>, !s32i, !cir.double) -> !s32i + cir.return %0 : !s32i + } + + // CHECK-LABEL: define i32 @args(ptr %0, i32 %1, double %2) + // CHECK: %[[BUF:.*]] = alloca { i32, double }, align 8 + // CHECK: br label %[[BB:.*]] + // CHECK: [[BB]]: + // CHECK: %[[F0:.*]] = getelementptr inbounds nuw { i32, double }, ptr %[[BUF]], i32 0, i32 0 + // CHECK: store i32 %1, ptr %[[F0]], align 4 + // CHECK: %[[F1:.*]] = getelementptr inbounds nuw { i32, double }, ptr %[[BUF]], i32 0, i32 1 + // CHECK: store double %2, ptr %[[F1]], align 8 + // CHECK: %[[R:.*]] = call i32 @vprintf(ptr %0, ptr %[[BUF]]) + // CHECK: ret i32 %[[R]] + + // A format string in another address space is cast to the generic one. + cir.func @constant_format(%fmt: !cir.ptr<!s8i, target_address_space(4)>) -> !s32i { + %0 = cir.offload.printf(%fmt) : (!cir.ptr<!s8i, target_address_space(4)>) -> !s32i + cir.return %0 : !s32i + } + + // CHECK-LABEL: define i32 @constant_format(ptr addrspace(4) %0) + // CHECK: %[[FMT:.*]] = addrspacecast ptr addrspace(4) %0 to ptr + // CHECK: %[[R:.*]] = call i32 @vprintf(ptr %[[FMT]], ptr null) + // CHECK: ret i32 %[[R]] +} >From 4e307ef2e99e0603f2e3958c04ea5dd70550ce1c Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 13 Oct 2026 02:08:46 -0500 Subject: [PATCH 2/2] Add func op assert Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp index 1bac9b6c4c878..e01a0f027539a 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp @@ -72,6 +72,7 @@ packArgsIntoNVPTXFormatBuffer(cir::OffloadPrintfOp op, mlir::ValueRange args, { mlir::OpBuilder::InsertionGuard guard(rewriter); auto fn = op->getParentOfType<mlir::LLVM::LLVMFuncOp>(); + assert(fn && "lowering expects printf to be inside a function"); rewriter.setInsertionPointToStart(&fn.getBody().front()); mlir::Value one = mlir::LLVM::ConstantOp::create(rewriter, loc, rewriter.getI32Type(), 1); _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
