https://github.com/steffenlarsen updated https://github.com/llvm/llvm-project/pull/231563
>From 219af38cf63352bd50b8b81ede9f92b86ae6367b Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Tue, 13 Oct 2026 10:26:33 -0500 Subject: [PATCH 1/2] [CIR][AMDGPU] Lower cir.offload.printf to AMDGPU printf CUDA printf lowering to use it, while leaving the AMDGPU lowering unsupported. Add lowering for cir.offload.printf to the AMDGPU printf logic employed by OGCG. Most of the functions added in clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp are CIR equivalents to the functions in llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp. Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 2 +- .../CIR/Lowering/DirectToLLVM/CMakeLists.txt | 1 + .../DirectToLLVM/LowerToLLVMOffloadPrintf.cpp | 684 +++++++++++++++++- clang/test/CIR/CodeGenHIP/device-printf.hip | 258 +++++++ .../offload-printf-amdgpu-unsupported.cir | 17 + .../CIR/Lowering/offload-printf-amdgpu.cir | 86 +++ .../llvm/Transforms/Utils/AMDGPUEmitPrintf.h | 28 + .../lib/Transforms/Utils/AMDGPUEmitPrintf.cpp | 66 +- 8 files changed, 1110 insertions(+), 32 deletions(-) create mode 100644 clang/test/CIR/CodeGenHIP/device-printf.hip create mode 100644 clang/test/CIR/Lowering/offload-printf-amdgpu-unsupported.cir create mode 100644 clang/test/CIR/Lowering/offload-printf-amdgpu.cir diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 5bbaa1828df56a..670277e3db5bc8 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -3191,7 +3191,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(emitDevicePrintfCallExpr(e)); } break; case Builtin::BI__builtin_canonicalize: diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt b/clang/lib/CIR/Lowering/DirectToLLVM/CMakeLists.txt index 009d437da0f166..7dfca636ea8825 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 diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp index e01a0f027539a3..ab64b5a83049b1 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp @@ -7,8 +7,7 @@ //===----------------------------------------------------------------------===// // // 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. +// device target. // //===----------------------------------------------------------------------===// @@ -17,7 +16,10 @@ #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 { @@ -123,8 +125,676 @@ lowerNVPTXPrintf(cir::OffloadPrintfOp op, mlir::Value format, return mlir::success(); } +std::optional<int64_t> getConstantInt(mlir::Value value) { + auto constOp = value.getDefiningOp<cir::ConstantOp>(); + auto intAttr = + constOp ? mlir::dyn_cast<cir::IntAttr>(constOp.getValue()) : nullptr; + if (!intAttr) + return std::nullopt; + const llvm::APInt &v = intAttr.getValue(); + if (mlir::cast<cir::IntType>(intAttr.getType()).isSigned() + ? v.getSignificantBits() > 64 + : v.getActiveBits() > 63) + return std::nullopt; + return v.getSExtValue(); +} + +bool isByteType(mlir::Type ty) { + auto intTy = mlir::dyn_cast<cir::IntType>(ty); + return intTy && intTy.getWidth() == 8; +} + +bool getConstantStringInfo(mlir::Value ptr, mlir::ModuleOp mod, + mlir::SymbolTableCollection &symbolTables, + std::string &str) { + int64_t offset = 0; + while (mlir::Operation *def = ptr.getDefiningOp()) { + if (auto cast = mlir::dyn_cast<cir::CastOp>(def)) { + if (cast.getKind() != cir::CastKind::array_to_ptrdecay && + cast.getKind() != cir::CastKind::bitcast && + cast.getKind() != cir::CastKind::address_space) + break; + ptr = cast.getSrc(); + } else if (auto elt = mlir::dyn_cast<cir::GetElementOp>(def)) { + std::optional<int64_t> index = getConstantInt(elt.getIndex()); + if (!index || !isByteType(elt.getElementType()) || + llvm::AddOverflow(offset, *index, offset)) + return false; + ptr = elt.getBase(); + } else if (auto stride = mlir::dyn_cast<cir::PtrStrideOp>(def)) { + std::optional<int64_t> index = getConstantInt(stride.getStride()); + if (!index || !isByteType(stride.getElementType()) || + llvm::AddOverflow(offset, *index, offset)) + return false; + ptr = stride.getBase(); + } else { + break; + } + } + + auto getGlobal = ptr.getDefiningOp<cir::GetGlobalOp>(); + if (!getGlobal) + return false; + + // The global may or may not have been converted already. + std::string data; + mlir::Operation *symbol = + symbolTables.lookupSymbolIn(mod, getGlobal.getNameAttr()); + if (auto global = mlir::dyn_cast_if_present<cir::GlobalOp>(symbol)) { + std::optional<mlir::Attribute> init = global.getInitialValue(); + auto arrayTy = mlir::dyn_cast<cir::ArrayType>(global.getSymType()); + if (!global.getConstant() || !init || !arrayTy || + !isByteType(arrayTy.getElementType())) + return false; + if (auto constArr = mlir::dyn_cast<cir::ConstArrayAttr>(*init)) { + auto strAttr = mlir::dyn_cast<mlir::StringAttr>(constArr.getElts()); + if (!strAttr) + return false; + data = strAttr.getValue().str(); + } else if (!mlir::isa<cir::ZeroAttr>(*init)) { + return false; + } + // Add back the trailing NULs the string attribute leaves out. + data.resize(arrayTy.getSize(), '\0'); + } 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().str(); + } else if (auto dense = + mlir::dyn_cast_if_present<mlir::DenseIntElementsAttr>(value); + dense && dense.getElementType().isInteger(8)) { + for (const llvm::APInt &c : dense.getValues<llvm::APInt>()) + data.push_back(static_cast<char>(c.getZExtValue())); + } else { + return false; + } + } else { + return false; + } + + if (offset < 0 || static_cast<uint64_t>(offset) > data.size()) + return false; + llvm::StringRef rest = llvm::StringRef(data).drop_front(offset); + str = rest.substr(0, rest.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; +}; + +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 gepInBounds(mlir::Value ptr, mlir::LLVM::GEPArg offset); + void addPrintfFormatMetadata(llvm::StringRef entry, bool onlyIfEmpty); + + mlir::Value fitArgInto64Bits(mlir::Value arg); + mlir::Value callPrintfBegin(mlir::Value version); + 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); + bool isCStringArg(mlir::ValueRange args, size_t i) const; + mlir::Value callBufferedPrintfStart( + mlir::ValueRange args, mlir::Value fmt, 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); + mlir::Value emitAMDGPUPrintfCall(mlir::ValueRange args, bool isBuffered); + + cir::OffloadPrintfOp op; + mlir::ConversionPatternRewriter &rewriter; + const mlir::DataLayout &dataLayout; + mlir::SymbolTableCollection &symbolTables; + mlir::ModuleOp mod; + + /// The rest of the op's block, which the expansion branches to. + 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::gepInBounds(mlir::Value ptr, + mlir::LLVM::GEPArg offset) { + return mlir::LLVM::GEPOp::create(rewriter, op.getLoc(), ptr.getType(), + rewriter.getI8Type(), ptr, + llvm::ArrayRef<mlir::LLVM::GEPArg>{offset}, + mlir::LLVM::GEPNoWrapFlags::inbounds); +} + +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::fitArgInto64Bits(mlir::Value arg) { + mlir::Type int64Ty = rewriter.getI64Type(); + 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(), int64Ty, arg); + case 64: + return arg; + } + } + + if (ty.isF64()) { + return mlir::LLVM::BitcastOp::create(rewriter, op.getLoc(), int64Ty, arg); + } + + if (mlir::isa<mlir::LLVM::LLVMPointerType>(ty)) { + return mlir::LLVM::PtrToIntOp::create(rewriter, op.getLoc(), int64Ty, arg); + } + + llvm_unreachable("unexpected type"); +} + +mlir::Value AMDGPUPrintfLowering::callPrintfBegin(mlir::Value version) { + return createCall("__ockl_printf_begin", rewriter.getI64Type(), version); +} + +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::Value charZero = mlir::LLVM::ConstantOp::create( + rewriter, op.getLoc(), rewriter.getI8Type(), llvm::APInt(8, 0)); + mlir::Value one = getI64(1); + mlir::Value zero = getI64(0); + mlir::Type int64Ty = rewriter.getI64Type(); + + // The length is either zero for a null pointer, or the computed value for an + // actual string. We need a join block for a phi that 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 *join = rewriter.createBlock(region, cont->getIterator(), + {int64Ty}, {op.getLoc()}); + mlir::Block *whileBlock = rewriter.createBlock( + region, join->getIterator(), {str.getType()}, {op.getLoc()}); + mlir::Block *whileDone = rewriter.createBlock(region, join->getIterator()); + + // Emit an early return for when the pointer is null. + rewriter.setInsertionPointToEnd(prev); + mlir::Value cmpNull = mlir::LLVM::ICmpOp::create( + rewriter, op.getLoc(), mlir::LLVM::ICmpPredicate::eq, str, + mlir::LLVM::ZeroOp::create(rewriter, op.getLoc(), str.getType())); + 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(), str.getType(), rewriter.getI8Type(), ptrPhi, + llvm::ArrayRef<mlir::LLVM::GEPArg>{one}); + + // 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, charZero); + 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(str.getType())); + 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 != int64Ty) + len = mlir::LLVM::ZExtOp::create(rewriter, op.getLoc(), int64Ty, len); + len = mlir::LLVM::AddOp::create(rewriter, op.getLoc(), len, one); + + // 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"); + // A call must match its callee's type, which takes the pointer type of the + // first string it was declared for. + 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); +} + +// As in processArg, a non-pointer argument for a string is sent as a scalar. +bool AMDGPUPrintfLowering::isCStringArg(mlir::ValueRange args, size_t i) const { + return specIsCString.test(i) && + mlir::isa<mlir::LLVM::LLVMPointerType>(args[i].getType()); +} + +// 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, mlir::Value fmt, bool isConstFmtStr, + llvm::SmallVectorImpl<StringData> &stringContents, mlir::Value &argSize) { + mlir::Value nonConstStrLen; + mlir::Value lenWithNull; + mlir::Value lenWithNullAligned; + mlir::Value tempAdd; + + // First 4 bytes to be reserved for control dword + size_t bufSize = 4; + if (isConstFmtStr) + // First 8 bytes of MD5 hash + bufSize += 8; + else { + lenWithNull = getStrlenWithNull(fmt); + + // Align the computed length to next 8 byte boundary + tempAdd = mlir::LLVM::AddOp::create(rewriter, op.getLoc(), lenWithNull, + getI64(7U)); + nonConstStrLen = + mlir::LLVM::AndOp::create(rewriter, op.getLoc(), tempAdd, getI64(~7U)); + + 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)) { + auto alignedLen = llvm::alignTo(argStr.size() + 1, 8); + stringContents.push_back({argStr, /*realSize=*/{}, /*alignedSize=*/{}, + /*isConst=*/true}); + bufSize += alignedLen; + } else { + lenWithNull = getStrlenWithNull(args[i]); + + // Align the computed length to next 8 byte boundary + tempAdd = mlir::LLVM::AddOp::create(rewriter, op.getLoc(), lenWithNull, + getI64(7U)); + lenWithNullAligned = mlir::LLVM::AndOp::create(rewriter, op.getLoc(), + tempAdd, getI64(~7U)); + + if (nonConstStrLen) { + mlir::Value val = mlir::LLVM::AddOp::create( + rewriter, op.getLoc(), lenWithNullAligned, nonConstStrLen); + nonConstStrLen = val; + } 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 + mlir::Value sizeToReserve = getI64(bufSize); + if (nonConstStrLen) + sizeToReserve = mlir::LLVM::AddOp::create(rewriter, op.getLoc(), + nonConstStrLen, sizeToReserve); + + argSize = mlir::LLVM::TruncOp::create(rewriter, op.getLoc(), + rewriter.getI32Type(), sizeToReserve); + + // 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); + 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) { + 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. + auto memcpy = mlir::LLVM::MemcpyOp::create( + rewriter, op.getLoc(), ptrToStore, args[i], strIt->realSize, + /*isVolatile=*/false); + mlir::NamedAttribute align1(mlir::LLVM::LLVMDialect::getAlignAttrName(), + rewriter.getI64IntegerAttr(1)); + memcpy.setArgAttrsAttr( + rewriter.getArrayAttr({rewriter.getDictionaryAttr(align1), + rewriter.getDictionaryAttr(align1), + rewriter.getDictionaryAttr({})})); + + 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()); + } + } +} + +mlir::Value AMDGPUPrintfLowering::emitAMDGPUPrintfCall(mlir::ValueRange args, + bool isBuffered) { + auto numOps = args.size(); + assert(numOps >= 1); + + mlir::Value fmt = args[0]; + + if (getConstantStringInfo(op.getFormat(), mod, symbolTables, fmtStr)) + llvm::locateAMDGPUPrintfCStrings(specIsCString, fmtStr); + + if (isBuffered) { + llvm::SmallVector<StringData, 8> stringContents; + bool isConstFmtStr = !fmtStr.empty(); + + mlir::Value argSize; + mlir::Value ptr = callBufferedPrintfStart(args, fmt, 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 rest of the op's block is the end block. + mlir::Block *end = cont; + mlir::Block *argPush = + rewriter.createBlock(cont->getParent(), cont->getIterator()); + + rewriter.setInsertionPointToEnd(cmp.getParentBlock()); + mlir::LLVM::CondBrOp::create(rewriter, op.getLoc(), cmp, argPush, end); + 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 constantTwo = getI32(2); + mlir::Value controlDWord = + mlir::LLVM::ShlOp::create(rewriter, op.getLoc(), argSize, constantTwo); + if (isConstFmtStr) + controlDWord = mlir::LLVM::OrOp::create(rewriter, op.getLoc(), + controlDWord, constantTwo); + + mlir::LLVM::StoreOp::create(rewriter, op.getLoc(), controlDWord, ptr); + + ptr = gepInBounds(ptr, 4); + + // 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 = gepInBounds(ptr, 8); + } 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(), end); + rewriter.setInsertionPointToStart(end); + 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::Value desc = callPrintfBegin(getI64(0)); + desc = appendString(desc, fmt, 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 (unsigned int 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); +} + +mlir::LogicalResult AMDGPUPrintfLowering::lower(mlir::ValueRange args, + bool isBuffered) { + // Unlike fitArgInto64Bits's unreachable, report a type hostcall cannot send. + 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"; + } + } + + mlir::Block *block = op->getBlock(); + cont = rewriter.splitBlock(block, mlir::Block::iterator(op)); + rewriter.setInsertionPointToEnd(block); + + mlir::Value result = emitAMDGPUPrintfCall(args, isBuffered); + if (rewriter.getInsertionBlock() != cont) + 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 { @@ -139,6 +809,16 @@ mlir::LogicalResult CIRToLLVMOffloadPrintfOpLowering::matchAndRewrite( return lowerNVPTXPrintf(op, adaptor.getFormat(), adaptor.getArgs(), rewriter, dataLayout, symbolTables); + 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 << "'"; } diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip new file mode 100644 index 00000000000000..673fdb63217c14 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -0,0 +1,258 @@ +// 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 + +// Device printf is emitted as cir.offload.printf, which AMDGPU expands into the +// OpenCL printf runtime: hostcall by default, or the buffered kind with +// -mprintf-kind=buffered. Both must match classic codegen. + +#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{{(, target_address_space\(4\))?}}>, !s32i +// CIR: cir.offload.printf(%[[FMT]], %[[V]]) : (!cir.ptr<!s8i{{(, target_address_space\(4\))?}}>, !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. 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.offload.printf({{.*}}) +// 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 +// LLVM-DAG: {{store|ret}} i32 [[RET]] +// 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 +// invoke with an exception-handling path. +struct D { __device__ ~D(); }; +// CIR-LABEL: cir.func {{.*}} @_Z2ehi( +// CIR: cir.cleanup.scope { +// 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}} +// BUFFERED-LABEL: @_Z2ehi( +// BUFFERED-NOT: personality +// BUFFERED: call{{.*}} ptr addrspace(1) @__printf_alloc(i32 +// BUFFERED-NOT: {{invoke|landingpad|resume}} +__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{{(, target_address_space\(4\))?}}>, !cir.ptr<!s8i{{(, target_address_space\(4\))?}}>, !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{{(, target_address_space\(4\))?}}>) -> !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{{(, target_address_space\(4\))?}}>, !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); +} + +// A constant offset into a string literal still yields a constant string, so +// the format's %s is found and the argument is passed as a string, not as a +// pointer value. +// CIR-LABEL: cir.func {{.*}} @_Z3offv( +// CIR: %[[FMT:.*]] = cir.get_element +// CIR: %[[STR:.*]] = cir.ptr_stride +// CIR: cir.offload.printf(%[[FMT]], %[[STR]]) +// LLVM-LABEL: @_Z3offv( +// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0) +// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}} getelementptr {{.*}}@.str.5{{.*}}, i64 2), i64 {{%.*}}, i32 0) +// LLVM-NOT: @__ockl_printf_append_args +// LLVM: call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D1]], ptr{{.*}} getelementptr {{.*}}@.str.6{{.*}}, i64 5), i64 {{%.*}}, i32 1) +// BUFFERED-LABEL: @_Z3offv( +// BUFFERED: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 20) +// BUFFERED: store i32 82, ptr addrspace(1) [[BUF]], align 4 +// BUFFERED: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4 +// BUFFERED: store i64 748964781523690102, ptr addrspace(1) [[P0]], align 8 +// BUFFERED: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i32 8 +// "me\0" as a little-endian i32 chunk, then padding. +// BUFFERED: store i32 25965, ptr addrspace(1) [[P1]], align 4 +// BUFFERED: [[P2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P1]], i32 4 +// BUFFERED: store i32 0, ptr addrspace(1) [[P2]], align 4 +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wstring-plus-int" +__device__ void off() { printf(&"xx%s|\n"[2], "skip-me" + 5); } +#pragma clang diagnostic pop + +// 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]+]], [[F4:![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"} +// BUFFERED-DAG: [[F4]] = !{!"0:0:a64db8027ff6276,%s|\0A"} diff --git a/clang/test/CIR/Lowering/offload-printf-amdgpu-unsupported.cir b/clang/test/CIR/Lowering/offload-printf-amdgpu-unsupported.cir new file mode 100644 index 00000000000000..131ffc8e0e77ee --- /dev/null +++ b/clang/test/CIR/Lowering/offload-printf-amdgpu-unsupported.cir @@ -0,0 +1,17 @@ +// RUN: cir-opt %s --cir-to-llvm -verify-diagnostics + +// Hostcall printf passes each argument as a 64-bit word, so it has no encoding +// for a vector. + +!s32i = !cir.int<s, 32> +!s8i = !cir.int<s, 8> + +module attributes {cir.triple = "amdgcn-amd-amdhsa", + cir.amdgpu_printf_kind = "hostcall"} { + cir.func @p(%fmt: !cir.ptr<!s8i>, %v: !cir.vector<4 x !s32i>) -> !s32i { + // expected-error @below {{unsupported argument type '!cir.vector<4 x !cir.int<s, 32>>' for AMDGPU printf}} + // expected-error @below {{failed to legalize operation 'cir.offload.printf'}} + %0 = cir.offload.printf(%fmt, %v) : (!cir.ptr<!s8i>, !cir.vector<4 x !s32i>) -> !s32i + cir.return %0 : !s32i + } +} diff --git a/clang/test/CIR/Lowering/offload-printf-amdgpu.cir b/clang/test/CIR/Lowering/offload-printf-amdgpu.cir new file mode 100644 index 00000000000000..955cfeccdb9467 --- /dev/null +++ b/clang/test/CIR/Lowering/offload-printf-amdgpu.cir @@ -0,0 +1,86 @@ +// RUN: cir-translate %s -cir-to-llvmir --target amdgcn-amd-amdhsa --disable-cc-lowering --split-input-file -o - \ +// RUN: | FileCheck %s --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> + 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) align 1 [[P3]], ptr align 1 %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 e7adf91b3f5695..579dabd08e757e 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 90c48fbd9127cc..2249a7a7c864de 100644 --- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp +++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp @@ -180,7 +180,7 @@ 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. @@ -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 e7cbe76e8903ce5d866fc6b7d6321742b7f97a04 Mon Sep 17 00:00:00 2001 From: Steffen Holst Larsen <[email protected]> Date: Wed, 14 Oct 2026 08:41:49 -0500 Subject: [PATCH 2/2] Add const char array case Signed-off-by: Steffen Holst Larsen <[email protected]> --- clang/test/CIR/CodeGenHIP/device-printf.hip | 63 +++++++++++++++++---- 1 file changed, 52 insertions(+), 11 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/device-printf.hip b/clang/test/CIR/CodeGenHIP/device-printf.hip index 673fdb63217c14..8e70bb195a8f77 100644 --- a/clang/test/CIR/CodeGenHIP/device-printf.hip +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -2,30 +2,30 @@ // 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 -fcxx-exceptions -fexceptions -fclangir -fcuda-is-device -emit-llvm %s -o - | FileCheck --check-prefixes=LLVM,LLVM-CIR %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: | FileCheck --check-prefixes=LLVM,LLVM-OGCG %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: | FileCheck --check-prefixes=BUFFERED,BUFFERED-CIR %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: | FileCheck --check-prefixes=BUFFERED,BUFFERED-OGCG %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: | FileCheck --check-prefixes=LLVM,LLVM-CIR %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: | FileCheck --check-prefixes=LLVM,LLVM-OGCG %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: | FileCheck --check-prefixes=BUFFERED,BUFFERED-CIR %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 +// RUN: | FileCheck --check-prefixes=BUFFERED,BUFFERED-OGCG %s // Device printf is emitted as cir.offload.printf, which AMDGPU expands into the // OpenCL printf runtime: hostcall by default, or the buffered kind with -// -mprintf-kind=buffered. Both must match classic codegen. +// -mprintf-kind=buffered. #define __device__ __attribute__((device)) extern "C" __device__ int printf(const char *, ...); @@ -247,12 +247,53 @@ __device__ void many(int a, long b, double c, void *d, float e, char f, __device__ void off() { printf(&"xx%s|\n"[2], "skip-me" + 5); } #pragma clang diagnostic pop +// A constant format in a deferred global. Classic codegen expands printf before +// the global's initializer is emitted, so it treats the format as non-constant +// and passes the %s argument as a pointer value. +// CIR-LABEL: cir.func {{.*}} @_Z2bfv( +// CIR: cir.get_global @_ZL9brace_fmt +// CIR: %[[FMT:.*]] = cir.cast array_to_ptrdecay +// CIR: %[[STR:.*]] = cir.cast array_to_ptrdecay +// CIR: cir.offload.printf(%[[FMT]], %[[STR]]) +// LLVM-LABEL: @_Z2bfv( +// LLVM: [[D0:%.*]] = call{{.*}} i64 @__ockl_printf_begin(i64 0) +// LLVM: [[D1:%.*]] = call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D0]], ptr{{.*}}@_ZL9brace_fmt{{.*}}, i64 {{%.*}}, i32 0) +// LLVM-CIR-NOT: @__ockl_printf_append_args +// LLVM-CIR: call{{.*}} i64 @__ockl_printf_append_string_n(i64 [[D1]], ptr{{.*}}@.str.7{{.*}}, i64 {{%.*}}, i32 1) +// LLVM-OGCG: call{{.*}} i64 @__ockl_printf_append_args(i64 [[D1]], i32 1, i64 ptrtoint ({{.*}}@.str.7{{.*}}), i64 0, i64 0, i64 0, i64 0, i64 0, i64 0, i32 1) +// BUFFERED-LABEL: @_Z2bfv( +// BUFFERED-CIR: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 20) +// BUFFERED-CIR: store i32 82, ptr addrspace(1) [[BUF]], align 4 +// BUFFERED-CIR: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4 +// BUFFERED-CIR: store i64 -88937716022153584, ptr addrspace(1) [[P0]], align 8 +// BUFFERED-CIR: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i32 8 +// BUFFERED-CIR: store i32 120, ptr addrspace(1) [[P1]], align 4 +// BUFFERED-CIR: [[P2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P1]], i32 4 +// BUFFERED-CIR: store i32 0, ptr addrspace(1) [[P2]], align 4 +// BUFFERED-OGCG: [[LEN:%.*]] = phi i64 +// BUFFERED-OGCG: [[ADD:%.*]] = add i64 [[LEN]], 7 +// BUFFERED-OGCG: [[ALIGNED:%.*]] = and i64 [[ADD]], 4294967288 +// BUFFERED-OGCG: [[TOTAL:%.*]] = add i64 [[ALIGNED]], 12 +// BUFFERED-OGCG: [[SIZE:%.*]] = trunc i64 [[TOTAL]] to i32 +// BUFFERED-OGCG: [[BUF:%.*]] = call{{.*}} ptr addrspace(1) @__printf_alloc(i32 [[SIZE]]) +// BUFFERED-OGCG: [[CTRL:%.*]] = shl i32 [[SIZE]], 2 +// BUFFERED-OGCG-NEXT: store i32 [[CTRL]], ptr addrspace(1) [[BUF]], align 4 +// BUFFERED-OGCG-NEXT: [[P0:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4 +// BUFFERED-OGCG-NEXT: call{{.*}} void @llvm.memcpy.p1.{{p0|p4}}.i64(ptr addrspace(1){{.*}} [[P0]], ptr{{.*}}@_ZL9brace_fmt{{.*}}, i64 [[LEN]], i1 false) +// BUFFERED-OGCG-NEXT: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P0]], i64 [[ALIGNED]] +// BUFFERED-OGCG-NEXT: store ptr{{.*}}@.str.7{{.*}}, ptr addrspace(1) [[P1]], align 8 +const char brace_fmt[] = {'%', 's', '\n', 0}; +__device__ void bf() { printf(brace_fmt, "x"); } + // Constant formats are recorded with the low 64 bits of their MD5 hash, in -// emission order. The runtime format in nf adds no entry. +// emission order. The runtime format in nf adds no entry, nor does brace_fmt +// in classic codegen. // LLVM-NOT: !llvm.printf.fmts -// BUFFERED: !llvm.printf.fmts = !{[[F0:![0-9]+]], [[F0]], [[F0]], [[F1:![0-9]+]], [[F2:![0-9]+]], [[F3:![0-9]+]], [[F4:![0-9]+]]} +// BUFFERED-CIR: !llvm.printf.fmts = !{[[F0:![0-9]+]], [[F0]], [[F0]], [[F1:![0-9]+]], [[F2:![0-9]+]], [[F3:![0-9]+]], [[F4:![0-9]+]], [[F5:![0-9]+]]} +// BUFFERED-OGCG: !llvm.printf.fmts = !{[[F0:![0-9]+]], [[F0]], [[F0]], [[F1:![0-9]+]], [[F2:![0-9]+]], [[F3:![0-9]+]], [[F4:![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"} // BUFFERED-DAG: [[F4]] = !{!"0:0:a64db8027ff6276,%s|\0A"} +// BUFFERED-CIR-DAG: [[F5]] = !{!"0:0:fec4079e555d6e90,%s\0A"} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
