Author: Steffen Larsen Date: 2026-10-30T14:14:38+01:00 New Revision: 050a017f870be148f745381d625fb7c9f3cde174
URL: https://github.com/llvm/llvm-project/commit/050a017f870be148f745381d625fb7c9f3cde174 DIFF: https://github.com/llvm/llvm-project/commit/050a017f870be148f745381d625fb7c9f3cde174.diff LOG: [CIR][AMDGPU] Lower cir.offload.printf to AMDGPU printf (#231563) 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. Assisted-by: Claude Opus 5.5 --------- Signed-off-by: Steffen Holst Larsen <[email protected]> Added: clang/test/CIR/CodeGenHIP/device-printf.hip clang/test/CIR/Lowering/offload-printf-amdgpu-unsupported.cir clang/test/CIR/Lowering/offload-printf-amdgpu.cir llvm/include/llvm/Support/AMDGPUPrintf.h llvm/lib/Support/AMDGPUPrintf.cpp Modified: clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp llvm/lib/Support/CMakeLists.txt llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 21a86114420e19..55941a8480cf1d 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -3226,14 +3226,8 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, if (getTarget().getTriple().isNVPTX() || getTarget().getTriple().isAMDGCN() || (getTarget().getTriple().isSPIRV() && - getTarget().getTriple().getVendor() == llvm::Triple::AMD)) { - if (getTarget().getTriple().isNVPTX()) - return RValue::get(emitDevicePrintfCallExpr(e)); - if ((getTarget().getTriple().isAMDGCN() || - getTarget().getTriple().isSPIRV()) && - getLangOpts().HIP) - return errorBuiltinNYI(*this, e, builtinID); - } + getTarget().getTriple().getVendor() == llvm::Triple::AMD)) + return RValue::get(emitDevicePrintfCallExpr(e)); break; case Builtin::BI__builtin_canonicalize: case Builtin::BI__builtin_canonicalizef: diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp index e01a0f027539a3..769b52a451059a 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVMOffloadPrintf.cpp @@ -7,16 +7,19 @@ //===----------------------------------------------------------------------===// // // 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. // //===----------------------------------------------------------------------===// #include "LowerToLLVM.h" +#include "TargetLoweringInfo.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" #include "mlir/IR/BuiltinOps.h" #include "clang/CIR/Dialect/IR/CIRDialect.h" +#include "llvm/ADT/SparseBitVector.h" +#include "llvm/Support/AMDGPUPrintf.h" +#include "llvm/Support/MathExtras.h" #include "llvm/TargetParser/Triple.h" namespace cir { @@ -123,8 +126,681 @@ 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, + unsigned genericAddrSpace) + : op(op), rewriter(rewriter), dataLayout(dataLayout), + symbolTables(symbolTables), mod(op->getParentOfType<mlir::ModuleOp>()), + genericAddrSpace(genericAddrSpace) {} + + 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; + unsigned genericAddrSpace; + + /// 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); + // __ockl_printf_append_string_n takes a generic pointer, which can address a + // string in any address space. + auto ptrTy = + mlir::LLVM::LLVMPointerType::get(rewriter.getContext(), genericAddrSpace); + if (str.getType() != ptrTy) + str = + mlir::LLVM::AddrSpaceCastOp::create(rewriter, op.getLoc(), ptrTy, str); + return createCall("__ockl_printf_append_string_n", 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.getTypeAllocSize(args[i].getType()).getFixedValue(); + // We end up expanding non string arguments to 8 bytes + // (args smaller than 8 bytes) + bufSize += llvm::alignTo(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.getTypeAllocSize(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++) { + if (i != 0 && !isCStringArg(args, i)) { + // The runtime reads each non-string argument from a slot of at least 8 + // bytes, also when the stored value is smaller. + mlir::Value toStore = processNonStringArg(args[i]); + mlir::LLVM::StoreOp::create(rewriter, op.getLoc(), toStore, ptrToStore); + uint64_t allocSize = + dataLayout.getTypeAllocSize(toStore.getType()).getFixedValue(); + ptrToStore = gepInBounds(ptrToStore, llvm::alignTo(allocSize, 8)); + continue; + } + + if (!strIt->isConst) { + // 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; + } + + llvm::SmallVector<mlir::Value, 32> whatToStore; + processConstantStringArg(*strIt, whatToStore); + strIt++; + + for (mlir::Value toStore : whatToStore) { + mlir::LLVM::StoreOp::create(rewriter, op.getLoc(), toStore, ptrToStore); + ptrToStore = gepInBounds( + ptrToStore, + dataLayout.getTypeAllocSize(toStore.getType()).getFixedValue()); + } + } +} + +mlir::Value AMDGPUPrintfLowering::emitAMDGPUPrintfCall(mlir::ValueRange args, + bool isBuffered) { + size_t 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 +815,22 @@ 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"; + std::unique_ptr<cir::TargetLoweringInfo> targetInfo = + triple.isSPIRV() ? cir::createSPIRVTargetLoweringInfo() + : cir::createAMDGPUTargetLoweringInfo(); + unsigned genericAddrSpace = targetInfo->getTargetAddrSpaceFromCIRAddrSpace( + cir::LangAddressSpace::OffloadGeneric); + return AMDGPUPrintfLowering(op, rewriter, dataLayout, symbolTables, + genericAddrSpace) + .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..8e70bb195a8f77 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/device-printf.hip @@ -0,0 +1,299 @@ +// 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-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-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-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-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-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-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-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-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. + +#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 diff ers 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 + +// 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, nor does brace_fmt +// in classic codegen. +// LLVM-NOT: !llvm.printf.fmts +// 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"} 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..82fa0b9c23b0bd --- /dev/null +++ b/clang/test/CIR/Lowering/offload-printf-amdgpu.cir @@ -0,0 +1,134 @@ +// 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 = "hostcall"} { + cir.func @hostcall_constant_as(%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 + } +} + +// The length of a string outside the generic address space is computed in its +// own address space, but the string is passed to the device library as a +// generic pointer, so the declaration matches the device library. +// HOSTCALL: declare i64 @__ockl_printf_append_string_n(i64, ptr, i64, i32) +// HOSTCALL-LABEL: define{{.*}} i32 @hostcall_constant_as(ptr addrspace(4) %0) +// HOSTCALL: icmp eq ptr addrspace(4) %0, null +// HOSTCALL: load i8, ptr addrspace(4) %{{.*}}, align 1 +// HOSTCALL: [[STR:%.*]] = addrspacecast ptr addrspace(4) %0 to ptr +// HOSTCALL: call i64 @__ockl_printf_append_string_n(i64 %{{.*}}, ptr [[STR]], i64 %{{.*}}, i32 1) + +// ----- + +!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"} + +// ----- + +!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<"%p %d\0A" : !cir.array<!s8i x 6>, trailing_zeros> : !cir.array<!s8i x 7> + cir.func @buffered_small_arg(%p: !cir.ptr<!s32i, target_address_space(3)>, %v: !s32i) -> !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, %p, %v) : (!cir.ptr<!s8i>, !cir.ptr<!s32i, target_address_space(3)>, !s32i) -> !s32i + cir.return %2 : !s32i + } +} + +// An argument stored in fewer than 8 bytes, here a 32-bit pointer, still takes +// an 8-byte slot, so the next argument starts 8 bytes after it. +// BUFFERED-LABEL: define{{.*}} i32 @buffered_small_arg(ptr addrspace(3) %0, i32 %1) +// BUFFERED: [[BUF:%.*]] = call ptr addrspace(1) @__printf_alloc(i32 28) +// BUFFERED: [[P1:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[BUF]], i32 4 +// BUFFERED: [[P2:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P1]], i32 8 +// BUFFERED: store ptr addrspace(3) %0, ptr addrspace(1) [[P2]], align 4 +// BUFFERED: [[P3:%.*]] = getelementptr inbounds i8, ptr addrspace(1) [[P2]], i32 8 +// BUFFERED: [[ARG:%.*]] = zext i32 %1 to i64 +// BUFFERED: store i64 [[ARG]], ptr addrspace(1) [[P3]], align 8 diff --git a/llvm/include/llvm/Support/AMDGPUPrintf.h b/llvm/include/llvm/Support/AMDGPUPrintf.h new file mode 100644 index 00000000000000..ba382e470482b7 --- /dev/null +++ b/llvm/include/llvm/Support/AMDGPUPrintf.h @@ -0,0 +1,53 @@ +//===- AMDGPUPrintf.h -------------------------------------------*- C++ -*-===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +/// \file +/// Helpers for the AMDGPU printf runtime interface, shared by the lowerings of +/// printf calls on the AMDGPU target. +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_SUPPORT_AMDGPUPRINTF_H +#define LLVM_SUPPORT_AMDGPUPRINTF_H + +#include "llvm/ADT/SmallVector.h" +#include "llvm/ADT/SparseBitVector.h" +#include "llvm/ADT/StringRef.h" +#include "llvm/Support/Compiler.h" +#include <cstdint> +#include <string> + +namespace llvm { + +/// 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_SUPPORT_AMDGPUPRINTF_H diff --git a/llvm/lib/Support/AMDGPUPrintf.cpp b/llvm/lib/Support/AMDGPUPrintf.cpp new file mode 100644 index 00000000000000..a596b3374796be --- /dev/null +++ b/llvm/lib/Support/AMDGPUPrintf.cpp @@ -0,0 +1,96 @@ +//===- AMDGPUPrintf.cpp ---------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// +// Helpers for the AMDGPU printf runtime interface, shared by the lowerings of +// printf calls on the AMDGPU target. +// +//===----------------------------------------------------------------------===// + +#include "llvm/Support/AMDGPUPrintf.h" +#include "llvm/ADT/StringExtras.h" +#include "llvm/Support/DataExtractor.h" +#include "llvm/Support/ErrorHandling.h" +#include "llvm/Support/MD5.h" + +using namespace llvm; + +// Scan the format string to locate all specifiers, and mark the ones that +// specify a string, i.e, the "%s" specifier with optional '*' characters. +void llvm::locateAMDGPUPrintfCStrings(SparseBitVector<8> &BV, StringRef Str) { + static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn"; + // Skip the first argument, the format string. + unsigned ArgIdx = 1; + + size_t SpecPos; + while ((SpecPos = Str.find('%')) != StringRef::npos) { + Str = Str.drop_front(SpecPos + 1); + // "%%" prints a '%' and takes no argument. + if (Str.consume_front("%")) + continue; + size_t SpecEnd = Str.find_first_of(ConvSpecifiers); + if (SpecEnd == StringRef::npos) + return; + // Each '*' in the specifier takes an argument of its own. + ArgIdx += Str.take_front(SpecEnd).count('*'); + if (Str[SpecEnd] == 's') + BV.set(ArgIdx); + Str = Str.drop_front(SpecEnd + 1); + ++ArgIdx; + } +} + +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); + while (Offset && Offset.tell() < Str.size()) { + const uint64_t ReadSize = 4; + uint64_t ReadNow = std::min(ReadSize, Str.size() - Offset.tell()); + uint64_t ReadBytes = 0; + switch (ReadNow) { + default: + llvm_unreachable("min(4, X) > 4?"); + case 1: + ReadBytes = Extractor.getU8(Offset); + break; + case 2: + ReadBytes = Extractor.getU16(Offset); + break; + case 3: + ReadBytes = Extractor.getU24(Offset); + break; + case 4: + ReadBytes = Extractor.getU32(Offset); + break; + } + cantFail(Offset.takeError(), "failed to read bytes from constant array"); + Words.push_back(static_cast<uint32_t>(ReadBytes)); + } + // Additional padding for 8 byte alignment + int Rem = (Str.size() % 8); + if (Rem > 0 && Rem <= 4) + Words.push_back(0); +} diff --git a/llvm/lib/Support/CMakeLists.txt b/llvm/lib/Support/CMakeLists.txt index 3d3e27d89ef352..a235991aa1689e 100644 --- a/llvm/lib/Support/CMakeLists.txt +++ b/llvm/lib/Support/CMakeLists.txt @@ -156,6 +156,7 @@ add_subdirectory(BLAKE3) add_llvm_component_library(LLVMSupport ABIBreak.cpp AMDGPUMetadata.cpp + AMDGPUPrintf.cpp APFixedPoint.cpp APFloat.cpp APInt.cpp diff --git a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp index 11fb00f9b6a373..925e12379d9bb5 100644 --- a/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp +++ b/llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp @@ -16,11 +16,9 @@ #include "llvm/Transforms/Utils/AMDGPUEmitPrintf.h" #include "llvm/ADT/SparseBitVector.h" -#include "llvm/ADT/StringExtras.h" #include "llvm/Analysis/ValueTracking.h" #include "llvm/IR/Module.h" -#include "llvm/Support/DataExtractor.h" -#include "llvm/Support/MD5.h" +#include "llvm/Support/AMDGPUPrintf.h" #include "llvm/Support/MathExtras.h" using namespace llvm; @@ -178,31 +176,6 @@ static Value *processArg(IRBuilder<> &Builder, Value *Desc, Value *Arg, return appendArg(Builder, Desc, Arg, IsLast); } -// 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) { - static const char ConvSpecifiers[] = "diouxXfFeEgGaAcspn"; - // Skip the first argument, the format string. - unsigned ArgIdx = 1; - - size_t SpecPos; - while ((SpecPos = Str.find('%')) != StringRef::npos) { - Str = Str.drop_front(SpecPos + 1); - // "%%" prints a '%' and takes no argument. - if (Str.consume_front("%")) - continue; - size_t SpecEnd = Str.find_first_of(ConvSpecifiers); - if (SpecEnd == StringRef::npos) - return; - // Each '*' in the specifier takes an argument of its own. - ArgIdx += Str.take_front(SpecEnd).count('*'); - if (Str[SpecEnd] == 's') - BV.set(ArgIdx); - Str = Str.drop_front(SpecEnd + 1); - ++ArgIdx; - } -} - // helper struct to package the string related data struct StringData { StringRef Str; @@ -307,45 +280,10 @@ static Value *callBufferedPrintfStart( // 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'); - - DataExtractor Extractor(Str, /*IsLittleEndian=*/true); - DataExtractor::Cursor Offset(0); - while (Offset && Offset.tell() < Str.size()) { - const uint64_t ReadSize = 4; - uint64_t ReadNow = std::min(ReadSize, Str.size() - Offset.tell()); - uint64_t ReadBytes = 0; - switch (ReadNow) { - default: - llvm_unreachable("min(4, X) > 4?"); - case 1: - ReadBytes = Extractor.getU8(Offset); - break; - case 2: - ReadBytes = Extractor.getU16(Offset); - break; - case 3: - ReadBytes = Extractor.getU24(Offset); - break; - case 4: - ReadBytes = Extractor.getU32(Offset); - 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)); - } - // Additional padding for 8 byte alignment - int Rem = (Str.size() % 8); - if (Rem > 0 && Rem <= 4) - WhatToStore.push_back(ConstantInt::get(Builder.getInt32Ty(), 0)); + 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) { @@ -431,7 +369,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; @@ -478,21 +416,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 @@ -500,7 +430,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); } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
