MLIR
22.0.0git
|
A class to build PTX assembly automatically. More...
#include "mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.h"
Public Member Functions | |
PtxBuilder (Operation *op, PatternRewriter &rewriter, bool needsManualRegisterMapping=false) | |
Single constructor that only initializes members. More... | |
LogicalResult | insertValue (Value v, PTXRegisterMod itype=PTXRegisterMod::Read) |
Add an operand with the read/write input type. More... | |
LLVM::InlineAsmOp | build () |
Builds the inline assembly Op and returns it. More... | |
void | buildAndReplaceOp () |
Shortcut to build the inline assembly Op and replace or erase the original op with. More... | |
A class to build PTX assembly automatically.
It is used by BasicPtxBuilderInterface.
Definition at line 62 of file BasicPtxBuilderInterface.h.
|
inline |
Single constructor that only initializes members.
Definition at line 80 of file BasicPtxBuilderInterface.h.
LLVM::InlineAsmOp PtxBuilder::build | ( | ) |
Builds the inline assembly Op and returns it.
The insertValue
needs to be called to pass operands before building the PTX.
Definition at line 426 of file BasicPtxBuilderInterface.cpp.
References canonicalizeRegisterConstraints(), mlir::get(), None, packResultTypes(), and rewriteAsmPlaceholders().
Referenced by buildAndReplaceOp().
void PtxBuilder::buildAndReplaceOp | ( | ) |
Shortcut to build the inline assembly Op and replace or erase the original op with.
Definition at line 467 of file BasicPtxBuilderInterface.cpp.
References build(), mlir::RewriterBase::eraseOp(), extractStructElements(), mlir::detail::IROperandBase::getOwner(), needsPackUnpack(), mlir::NVVM::ReadWrite, and mlir::RewriterBase::replaceOp().
LogicalResult PtxBuilder::insertValue | ( | Value | v, |
PTXRegisterMod | itype = PTXRegisterMod::Read |
||
) |
Add an operand with the read/write input type.
Definition at line 138 of file BasicPtxBuilderInterface.cpp.
References mlir::detail::enumerate(), mlir::remark::failed(), getRegisterType(), mlir::Value::getType(), mlir::RewriterBase::notifyMatchFailure(), mlir::NVVM::Read, mlir::NVVM::ReadWrite, and mlir::NVVM::Write.