MLIR
17.0.0git
|
A multi-dimensional affine map Affine map's are immutable like Type's, and they are uniqued. More...
#include "mlir/IR/AffineMap.h"
Public Types | |
using | ImplType = detail::AffineMapStorage |
Public Member Functions | |
constexpr | AffineMap ()=default |
AffineMap (ImplType *map) | |
MLIRContext * | getContext () const |
operator bool () const | |
bool | operator== (AffineMap other) const |
bool | operator!= (AffineMap other) const |
bool | isIdentity () const |
Returns true if this affine map is an identity affine map. More... | |
bool | isSymbolIdentity () const |
Returns true if this affine map is an identity affine map on the symbol identifiers. More... | |
bool | isMinorIdentity () const |
Returns true if this affine map is a minor identity, i.e. More... | |
bool | isMinorIdentityWithBroadcasting (SmallVectorImpl< unsigned > *broadcastedDims=nullptr) const |
Returns true if this affine map is a minor identity up to broadcasted dimensions which are indicated by value 0 in the result. More... | |
bool | isPermutationOfMinorIdentityWithBroadcasting (SmallVectorImpl< unsigned > &permutedDims) const |
Return true if this affine map can be converted to a minor identity with broadcast by doing a permute. More... | |
bool | isEmpty () const |
Returns true if this affine map is an empty map, i.e., () -> (). More... | |
bool | isSingleConstant () const |
Returns true if this affine map is a single result constant function. More... | |
bool | isConstant () const |
Returns true if this affine map has only constant results. More... | |
int64_t | getSingleConstantResult () const |
Returns the constant result of this map. More... | |
SmallVector< int64_t > | getConstantResults () const |
Returns the constant results of this map. More... | |
void | print (raw_ostream &os) const |
void | dump () const |
unsigned | getNumDims () const |
unsigned | getNumSymbols () const |
unsigned | getNumResults () const |
unsigned | getNumInputs () const |
ArrayRef< AffineExpr > | getResults () const |
AffineExpr | getResult (unsigned idx) const |
unsigned | getDimPosition (unsigned idx) const |
Extracts the position of the dimensional expression at the given result, when the caller knows it is safe to do so. More... | |
std::optional< unsigned > | getResultPosition (AffineExpr input) const |
Extracts the first result position where input dimension resides. More... | |
bool | isFunctionOfDim (unsigned position) const |
Return true if any affine expression involves AffineDimExpr position . More... | |
bool | isFunctionOfSymbol (unsigned position) const |
Return true if any affine expression involves AffineSymbolExpr position . More... | |
void | walkExprs (llvm::function_ref< void(AffineExpr)> callback) const |
Walk all of the AffineExpr's in this mapping. More... | |
AffineMap | replaceDimsAndSymbols (ArrayRef< AffineExpr > dimReplacements, ArrayRef< AffineExpr > symReplacements, unsigned numResultDims, unsigned numResultSyms) const |
This method substitutes any uses of dimensions and symbols (e.g. More... | |
AffineMap | replace (AffineExpr expr, AffineExpr replacement, unsigned numResultDims, unsigned numResultSyms) const |
Sparse replace method. More... | |
AffineMap | replace (const DenseMap< AffineExpr, AffineExpr > &map) const |
Sparse replace method. More... | |
AffineMap | replace (const DenseMap< AffineExpr, AffineExpr > &map, unsigned numResultDims, unsigned numResultSyms) const |
Sparse replace method. More... | |
AffineMap | shiftDims (unsigned shift, unsigned offset=0) const |
Replace dims[offset ... More... | |
AffineMap | shiftSymbols (unsigned shift, unsigned offset=0) const |
Replace symbols[offset ... More... | |
AffineMap | dropResult (int64_t pos) |
Returns a new AffineMap with the same number of dims and symbols and one less result at pos , dropped. More... | |
AffineMap | dropResults (ArrayRef< int64_t > positions) |
AffineMap | insertResult (AffineExpr expr, unsigned pos) |
Returns a new AffineMap with the same number of dims and symbols and an extra result inserted at pos . More... | |
LogicalResult | constantFold (ArrayRef< Attribute > operandConstants, SmallVectorImpl< Attribute > &results) const |
Folds the results of the application of an affine map on the provided operands to a constant if possible. More... | |
AffineMap | partialConstantFold (ArrayRef< Attribute > operandConstants, SmallVectorImpl< int64_t > *results=nullptr) const |
Propagates the constant operands into this affine map. More... | |
AffineMap | compose (AffineMap map) const |
Returns the AffineMap resulting from composing this with map . More... | |
SmallVector< int64_t, 4 > | compose (ArrayRef< int64_t > values) const |
Applies composition by the dims of this to the integer values and returns the resulting values. More... | |
bool | isProjectedPermutation (bool allowZeroInResults=false) const |
Returns true if the AffineMap represents a subset (i.e. More... | |
bool | isPermutation () const |
Returns true if the AffineMap represents a symbol-less permutation map. More... | |
AffineMap | getSubMap (ArrayRef< unsigned > resultPos) const |
Returns the map consisting of the resultPos subset. More... | |
AffineMap | getSliceMap (unsigned start, unsigned length) const |
Returns the map consisting of length expressions starting from start . More... | |
AffineMap | getMajorSubMap (unsigned numResults) const |
Returns the map consisting of the most major numResults results. More... | |
AffineMap | getMinorSubMap (unsigned numResults) const |
Returns the map consisting of the most minor numResults results. More... | |
uint64_t | getLargestKnownDivisorOfMapExprs () |
Get the largest known divisor of all map expressions. More... | |
const void * | getAsOpaquePointer () const |
Methods supporting C API. More... | |
Static Public Member Functions | |
static AffineMap | get (MLIRContext *context) |
Returns a zero result affine map with no dimensions or symbols: () -> (). More... | |
static AffineMap | get (unsigned dimCount, unsigned symbolCount, MLIRContext *context) |
Returns a zero result affine map with dimCount dimensions and symbolCount symbols, e.g. More... | |
static AffineMap | get (unsigned dimCount, unsigned symbolCount, AffineExpr result) |
Returns an affine map with dimCount dimensions and symbolCount mapping to a single output dimension. More... | |
static AffineMap | get (unsigned dimCount, unsigned symbolCount, ArrayRef< AffineExpr > results, MLIRContext *context) |
Returns an affine map with dimCount dimensions and symbolCount mapping to the given results. More... | |
static AffineMap | getConstantMap (int64_t val, MLIRContext *context) |
Returns a single constant result affine map. More... | |
static AffineMap | getMultiDimIdentityMap (unsigned numDims, MLIRContext *context) |
Returns an AffineMap with 'numDims' identity result dim exprs. More... | |
static AffineMap | getMinorIdentityMap (unsigned dims, unsigned results, MLIRContext *context) |
Returns an identity affine map (d0, ..., dn) -> (dp, ..., dn) on the most minor dimensions. More... | |
static AffineMap | getPermutationMap (ArrayRef< unsigned > permutation, MLIRContext *context) |
Returns an AffineMap representing a permutation. More... | |
static SmallVector< AffineMap, 4 > | inferFromExprList (ArrayRef< ArrayRef< AffineExpr >> exprsList) |
Returns a vector of AffineMaps; each with as many results as exprs.size() , as many dims as the largest dim in exprs and as many symbols as the largest symbol in exprs . More... | |
static SmallVector< AffineMap, 4 > | inferFromExprList (ArrayRef< SmallVector< AffineExpr, 4 >> exprsList) |
static AffineMap | getFromOpaquePointer (const void *pointer) |
Friends | |
::llvm::hash_code | hash_value (AffineMap arg) |
A multi-dimensional affine map Affine map's are immutable like Type's, and they are uniqued.
Eg: (d0, d1) -> (d0/128, d0 mod 128, d1) The names used (d0, d1) don't matter - it's the mathematical function that is unique to this affine map.
Definition at line 43 of file AffineMap.h.
Definition at line 45 of file AffineMap.h.
|
constexprdefault |
Referenced by getFromOpaquePointer(), getMajorSubMap(), and getMinorSubMap().
|
inlineexplicit |
Definition at line 48 of file AffineMap.h.
Returns the AffineMap resulting from composing this
with map
.
The resulting AffineMap has as many AffineDimExpr as map
and as many AffineSymbolExpr as the concatenation of this
and map
(in which case the symbols of this
map come first).
Prerequisites: The maps are composable, i.e. that the number of AffineDimExpr of this
matches the number of results of map
.
Example: map1: (d0, d1)[s0, s1] -> (d0 + 1 + s1, d1 - 1 - s0)
map2: (d0)[s0] -> (d0 + s0, d0 - s0)
map1.compose(map2): (d0)[s0, s1, s2] -> (d0 + s1 + s2 + 1, d0 - s0 - s2 - 1)
Definition at line 468 of file AffineMap.cpp.
References mlir::AffineExpr::compose(), get(), mlir::getAffineDimExpr(), mlir::getAffineSymbolExpr(), getContext(), getNumDims(), getNumResults(), getNumSymbols(), getResults(), and replaceDimsAndSymbols().
Referenced by buildVectorWrite(), compose(), mlir::linalg::fuseElementwiseOps(), getIndexingMapOfProducerOperandsInCoordinatesOfFusedOp(), getPermutationMapAttr(), inferTransferOpMaskType(), insertParallelDim(), mlir::linalg::interchangeGenericOp(), TransferReadPermutationLowering::matchAndRewrite(), scaleReductionDim(), transposeOneLinalgOperandAndReplace(), vectorizeAsLinalgGeneric(), verifyOutputShape(), and mlir::linalg::detail::verifyStructuredOpInterface().
SmallVector< int64_t, 4 > AffineMap::compose | ( | ArrayRef< int64_t > | values | ) | const |
Applies composition by the dims of this
to the integer values
and returns the resulting values.
this
must be symbol-less.
Definition at line 492 of file AffineMap.cpp.
References compose(), get(), mlir::getAffineConstantExpr(), getContext(), getNumSymbols(), and mlir::AffineConstantExpr::getValue().
LogicalResult AffineMap::constantFold | ( | ArrayRef< Attribute > | operandConstants, |
SmallVectorImpl< Attribute > & | results | ||
) | const |
Folds the results of the application of an affine map on the provided operands to a constant if possible.
Returns false if the folding happens, true otherwise.
Definition at line 360 of file AffineMap.cpp.
References mlir::failure(), getContext(), partialConstantFold(), and mlir::success().
Referenced by foldLoopBounds().
|
inline |
Returns a new AffineMap with the same number of dims and symbols and one less result at pos
, dropped.
Definition at line 252 of file AffineMap.h.
References dropResults().
Referenced by mlir::linalg::splitReductionByScaling().
Definition at line 256 of file AffineMap.h.
References get(), getContext(), getNumDims(), getNumSymbols(), and getResults().
Referenced by dropResult(), and vectorizeAsLinalgGeneric().
void AffineMap::dump | ( | ) | const |
Definition at line 3562 of file AsmPrinter.cpp.
References print().
Referenced by mlir::MemRefRegion::compute(), and mlir::FlatAffineValueConstraints::getSliceBounds().
|
static |
Returns a zero result affine map with no dimensions or symbols: () -> ().
Definition at line 1047 of file MLIRContext.cpp.
References mlir::OperationName::getImpl().
Referenced by addConstToResults(), addLoopRangeConstraints(), adjustMap(), augmentMapAndBounds(), mlir::tensor::bubbleUpPadSlice(), calculateImplicitMap(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), mlir::coalesceLoops(), compose(), composeMultiResultAffineMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), compressUnusedImpl(), computeMemoryOpIndices(), mlir::concatAffineMaps(), constructTiledIndexSetHyperRect(), mlir::createAffineComputationSlice(), createMask(), createPrivateMemRef(), mlir::AffineValueMap::difference(), doubleBuffer(), dropResults(), elementwiseMatchAndRewriteHelper(), mlir::linalg::extractOrIdentityMap(), fuseWithReshapeByExpansion(), generateCopy(), generatePointWiseCopy(), mlir::MutableAffineMap::getAffineMap(), getBoundedTileSize(), getCleanupLoopLowerBound(), getCollapsedOpIndexingMap(), getCollapsedOutputDimFromInputShape(), mlir::Builder::getConstantAffineMap(), getConstantMap(), mlir::Builder::getDimIdentityMap(), mlir::Builder::getEmptyAffineMap(), getExpandedOutputDimFromInputShape(), mlir::FlatAffineValueConstraints::getHyperrectangular(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::FlatAffineValueConstraints::getIneqAsAffineValueMap(), mlir::nvgpu::getLaneIdAndValueIdToOperandCoord(), mlir::nvgpu::getLaneIdToLdMatrixMatrixCoord(), mlir::FlatAffineValueConstraints::getLowerAndUpperBound(), getMinorIdentityMap(), getMultiDimIdentityMap(), mlir::Builder::getMultiDimIdentityMap(), getPermutationMap(), getPermutationMapAttr(), getRegisterIndexToTileOffsetMap(), mlir::Builder::getShiftedAffineMap(), mlir::Builder::getSingleDimShiftAffineMap(), mlir::FlatAffineValueConstraints::getSliceBounds(), getSliceMap(), getSubMap(), mlir::Builder::getSymbolIdentityMap(), mlir::getSymbolLessAffineMaps(), mlir::vector::getTransferMinorIdentityMap(), mlir::getTripCountMapAndOperands(), inferFromExprList(), insertResult(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), mlir::isColumnMajorMatmul(), mlir::isRowMajorBatchMatmul(), mlir::isRowMajorMatmul(), isTransposeMatrixLoadMap(), mlir::loopUnrollJamByFactor(), makeCanonicalAffineApplies(), makePermutationMap(), mlir::makeStridedLinearLayoutMap(), mlir::mapLoopToProcessorIds(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), mlir::linalg::PadOpTransformationPattern::matchAndRewrite(), mlir::vector::ContractionOpToMatmulOpLowering::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), TransferWriteNonPermutationLowering::matchAndRewrite(), mlirAffineMapEmptyGet(), mlirAffineMapGet(), mlirAffineMapZeroResultGet(), mlir::memref::multiBuffer(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), parseAffineMapWithMinMax(), partialConstantFold(), peelForLoop(), processParallelLoop(), mlir::removeDuplicateExprs(), replace(), mlir::replaceAllMemRefUsesWith(), replaceDimsAndSymbols(), replaceUnitExtents(), resolveSourceIndicesCollapseShape(), resolveSourceIndicesExpandShape(), resolveSourceIndicesSubView(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), shiftDims(), shiftSymbols(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), sliceTransferIndices(), mlir::linalg::splitReduction(), mlir::linalg::splitReductionByScaling(), mlir::scf::tileParallelLoop(), transferReadSupportsMMAMatrixType(), and verifyOutputShape().
|
static |
Returns an affine map with dimCount
dimensions and symbolCount
mapping to a single output dimension.
Definition at line 1056 of file MLIRContext.cpp.
References mlir::AffineExpr::getContext(), mlir::OperationName::getImpl(), and willBeValidAffineMap().
|
static |
Returns an affine map with dimCount
dimensions and symbolCount
mapping to the given results.
Definition at line 1062 of file MLIRContext.cpp.
References mlir::OperationName::getImpl(), and willBeValidAffineMap().
|
static |
Returns a zero result affine map with dimCount
dimensions and symbolCount
symbols, e.g.
: (...) -> ()
.
Definition at line 1051 of file MLIRContext.cpp.
References mlir::OperationName::getImpl().
|
inline |
Methods supporting C API.
Definition at line 344 of file AffineMap.h.
|
static |
Returns a single constant result affine map.
Definition at line 97 of file AffineMap.cpp.
References get(), and mlir::getAffineConstantExpr().
Referenced by mlir::getTripCountMapAndOperands(), mlir::linalg::getUpperBoundForIndex(), and mlirAffineMapConstantGet().
SmallVector< int64_t > AffineMap::getConstantResults | ( | ) | const |
Returns the constant results of this map.
This method asserts that the map has all constant results.
Definition at line 311 of file AffineMap.cpp.
References mlir::AffineExpr::cast(), getResults(), mlir::AffineConstantExpr::getValue(), and isConstant().
MLIRContext * AffineMap::getContext | ( | ) | const |
Definition at line 266 of file AffineMap.cpp.
References mlir::detail::AffineMapStorage::context.
Referenced by addConstToResults(), mlir::alignAffineMapWithValues(), canonicalizeMapExprAndTermOrder(), compose(), composeAffineMapAndOperands(), composeMultiResultAffineMap(), mlir::compressDims(), mlir::compressSymbols(), constantFold(), dropResults(), getCollapsedOpIndexingMap(), getSliceMap(), getSubMap(), inferTransferOpMaskType(), insertResult(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), isMinorIdentity(), packLinalgMetadataOnce(), partialConstantFold(), mlir::removeDuplicateExprs(), replace(), replaceDimOrSym(), replaceDimsAndSymbols(), mlir::MutableAffineMap::reset(), shiftDims(), shiftSymbols(), mlir::simplifyAffineMap(), and simplifyMapWithOperands().
unsigned AffineMap::getDimPosition | ( | unsigned | idx | ) | const |
Extracts the position of the dimensional expression at the given result, when the caller knows it is safe to do so.
Definition at line 340 of file AffineMap.cpp.
References mlir::AffineExpr::cast(), and getResult().
Referenced by adjustMap(), getDimPosition(), getReductionIndex(), getResultIndex(), mlir::inverseAndBroadcastProjectedPermutation(), and mlir::linalg::splitReduction().
|
inlinestatic |
Definition at line 347 of file AffineMap.h.
References AffineMap().
uint64_t AffineMap::getLargestKnownDivisorOfMapExprs | ( | ) |
Get the largest known divisor of all map expressions.
For eg: for (d0, d1) -> (8*d0 + 4, 4*d1 + 2), the result is 2. In the case of maps with no expressions or all zero constant expressions, the largest known divisor is trivially the max uint64_t value.
Definition at line 246 of file AffineMap.cpp.
References mlir::presburger::gcd(), getResults(), and max().
AffineMap AffineMap::getMajorSubMap | ( | unsigned | numResults | ) | const |
Returns the map consisting of the most major numResults
results.
Returns the null AffineMap if numResults
== 0. Returns *this
if numResults
>= this->getNumResults()
.
Definition at line 556 of file AffineMap.cpp.
References AffineMap(), getNumResults(), and getSliceMap().
|
static |
Returns an identity affine map (d0, ..., dn) -> (dp, ..., dn) on the most minor dimensions.
Definition at line 104 of file AffineMap.cpp.
References get(), getMultiDimIdentityMap(), and getResults().
Referenced by mlir::vector::getTransferMinorIdentityMap(), isMinorIdentity(), TransferWritePermutationLowering::matchAndRewrite(), and mlirAffineMapMinorIdentityGet().
AffineMap AffineMap::getMinorSubMap | ( | unsigned | numResults | ) | const |
Returns the map consisting of the most minor numResults
results.
Returns the null AffineMap if numResults
== 0. Returns *this
if numResults
>= this->getNumResults()
.
Definition at line 564 of file AffineMap.cpp.
References AffineMap(), getNumResults(), and getSliceMap().
|
static |
Returns an AffineMap with 'numDims' identity result dim exprs.
Definition at line 257 of file AffineMap.cpp.
References get(), and mlir::getAffineDimExpr().
Referenced by mlir::tensor::bubbleUpPadSlice(), buildMax(), buildMin(), mlir::linalg::extractOrIdentityMap(), getMinorIdentityMap(), getUnorderedCOOFromType(), mlir::linalg::getUpperBoundForIndex(), hasSameDimOrdering(), insertParallelDim(), mlir::linalg::makeMemRefCopyOp(), mlir::linalg::makeTransposeOp(), mlir::AffineBuilder::max(), mlir::AffineBuilder::min(), mlirAffineMapMultiDimIdentityGet(), scaleReductionDim(), tileLinalgOpImpl(), and tilePadOp().
unsigned AffineMap::getNumDims | ( | ) | const |
Definition at line 319 of file AffineMap.cpp.
References mlir::detail::AffineMapStorage::numDims.
Referenced by mlir::FlatAffineValueConstraints::addBound(), addConstToResults(), adjustMap(), mlir::alignAffineMapWithValues(), mlir::applyMapToValues(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), compose(), composeAffineMapAndOperands(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeMultiResultAffineMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), compressUnusedImpl(), mlir::MemRefRegion::compute(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), dropResults(), mlir::expandAffineMap(), getCleanupLoopLowerBound(), getConstDifference(), mlir::getFlattenedAffineExprs(), getIndicesForAccess(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumDims(), mlir::getRelationFromMap(), mlir::Builder::getShiftedAffineMap(), getSliceMap(), getStridesAndOffset(), getSubMap(), getXferIndices(), insertParallelDim(), insertResult(), mlir::inversePermutation(), isEmpty(), isIdentity(), isMinorIdentity(), isMinorIdentityWithBroadcasting(), isPermutation(), isTransposeMatrixLoadMap(), makeCanonicalAffineApplies(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), TransferWriteNonPermutationLowering::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), packLinalgMetadataOnce(), parseBound(), partialConstantFold(), permute(), mlir::AsmPrinter::Impl::printAffineMap(), printBound(), printMinMaxBound(), remainsLegalAfterInline(), mlir::removeDuplicateExprs(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), scaleReductionDim(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), shiftDims(), shiftSymbols(), mlir::simplifyAffineMap(), mlir::simplifyConstrainedMinMaxOp(), simplifyMapWithOperands(), mlir::linalg::splitReduction(), transferReadSupportsMMAMatrixType(), mlir::detail::verifyAffineMapAsLayout(), and mlir::linalg::detail::verifyStructuredOpInterface().
unsigned AffineMap::getNumInputs | ( | ) | const |
Definition at line 328 of file AffineMap.cpp.
References mlir::detail::AffineMapStorage::numDims, and mlir::detail::AffineMapStorage::numSymbols.
Referenced by mlir::FlatAffineValueConstraints::addBound(), mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), mlir::alignAffineMapWithValues(), mlir::applyPermutationMap(), mlir::FlatAffineValueConstraints::composeMatchingMap(), mlir::FlatAffineValueConstraints::computeAlignedMap(), createNewDynamicSizes(), generatePointWiseCopy(), generateShiftedLoop(), mlir::AffineDmaStartOp::getDstIndices(), mlir::AffineDmaStartOp::getDstMemRefOperandIndex(), mlir::MemRefRegion::getLowerAndUpperBound(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineDmaStartOp::getSrcIndices(), mlir::AffineDmaStartOp::getTagIndices(), mlir::AffineDmaWaitOp::getTagIndices(), mlir::AffineDmaStartOp::getTagMemRefOperandIndex(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), isPermutationOfMinorIdentityWithBroadcasting(), isProjectedPermutation(), mlir::AffineDmaStartOp::isStrided(), makeCanonicalAffineApplies(), partialConstantFold(), mlir::replaceAllMemRefUsesWith(), simplifyMapWithOperands(), mlir::AffineDmaStartOp::verifyInvariantsImpl(), verifyMemoryOpIndexing(), verifyOutputShape(), verifyPermutationMap(), and verifyTransferOp().
unsigned AffineMap::getNumResults | ( | ) | const |
Definition at line 327 of file AffineMap.cpp.
References getResults().
Referenced by mlir::FlatAffineValueConstraints::addBound(), mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), adjustMap(), mlir::applyMapToValues(), mlir::applyPermutationMap(), mlir::linalg::areElementwiseOpsFusable(), mlir::buildSliceTripCountMap(), calculateExpandedAccessIndices(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), compose(), composeAffineMapAndOperands(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeMultiResultAffineMap(), mlir::compressDims(), mlir::compressSymbols(), mlir::concatAffineMaps(), constructTiledIndexSetHyperRect(), createMask(), foldLoopBounds(), getCleanupLoopLowerBound(), getCollapsedOpOperand(), getConstDifference(), getDimPosition(), getFirstResultIndexFunctionOf(), mlir::getFlattenedAffineExprs(), getIndicesForAccess(), mlir::getLargestDivisorOfTripCount(), getMajorSubMap(), getMinorSubMap(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumResults(), getOperandReassociation(), getReductionIndex(), mlir::getRelationFromMap(), getResultIndex(), getResultPosition(), getShapeDefiningLoopRange(), mlir::Builder::getShiftedAffineMap(), mlir::FlatAffineValueConstraints::getSliceBounds(), getStridesAndOffset(), getTiledProducerLoops(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), isDimSequencePreserved(), isEmpty(), isIdentity(), isMinorIdentity(), isMinorIdentityWithBroadcasting(), isPermutation(), isPermutationOfMinorIdentityWithBroadcasting(), isProjectedPermutation(), isSingleConstant(), isSymbolIdentity(), isTiled(), makeCanonicalAffineApplies(), mlir::makeComposedFoldedAffineApply(), mlir::makeComposedFoldedMultiResultAffineApply(), mlir::linalg::makeTiledLoopRanges(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), TransferReadPermutationLowering::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), mlir::normalizeMemRefType(), packLinalgMetadataOnce(), parseBound(), partialConstantFold(), permute(), printBound(), replace(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), replaceDimsAndSymbols(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::simplifyConstrainedMinMaxOp(), simplifyMapWithOperands(), mlir::linalg::splitReduction(), verifyMemoryOpIndexing(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().
unsigned AffineMap::getNumSymbols | ( | ) | const |
Definition at line 323 of file AffineMap.cpp.
References mlir::detail::AffineMapStorage::numSymbols.
Referenced by mlir::FlatAffineValueConstraints::addBound(), addConstToResults(), mlir::alignAffineMapWithValues(), mlir::applyMapToValues(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), compose(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), compressUnusedImpl(), mlir::compressUnusedSymbols(), mlir::MemRefRegion::compute(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), createNewDynamicSizes(), dropResults(), getCleanupLoopLowerBound(), getConstDifference(), mlir::getFlattenedAffineExprs(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumSymbols(), mlir::Builder::getShiftedAffineMap(), getSliceMap(), getStridesAndOffset(), getSubMap(), insertResult(), mlir::inversePermutation(), isEmpty(), isProjectedPermutation(), isSymbolIdentity(), makeCanonicalAffineApplies(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), parseBound(), partialConstantFold(), mlir::AsmPrinter::Impl::printAffineMap(), printBound(), remainsLegalAfterInline(), mlir::removeDuplicateExprs(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), shiftDims(), shiftSymbols(), mlir::simplifyAffineMap(), mlir::simplifyConstrainedMinMaxOp(), simplifyMapWithOperands(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().
|
static |
Returns an AffineMap representing a permutation.
The permutation is expressed as a non-empty vector of integers. E.g. the permutation (i,j,k) -> (j,k,i)
will be expressed with permutation = [1,2,0]
. All values in permutation
must be integers, in the range 0..permutation.size()-1
without duplications (i.e. [1,1,2]
is an invalid permutation).
Definition at line 207 of file AffineMap.cpp.
References get(), and mlir::getAffineDimExpr().
Referenced by convertTransferReadOp(), foldTransferInBoundsAttribute(), mlir::linalg::interchangeGenericOp(), interchangeGenericOpPrecondition(), mlir::linalg::makeTransposeOp(), TransferReadPermutationLowering::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), TransferWriteNonPermutationLowering::matchAndRewrite(), mlirAffineMapPermutationGet(), permute(), tileLinalgOpImpl(), transferReadSupportsMMAMatrixType(), and transposeOneLinalgOperandAndReplace().
AffineExpr AffineMap::getResult | ( | unsigned | idx | ) | const |
Definition at line 336 of file AffineMap.cpp.
References getResults().
Referenced by mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), composeMultiResultAffineMap(), convertTransferReadOp(), foldTransferInBoundsAttribute(), getCleanupLoopLowerBound(), getConstDifference(), getDimPosition(), getFirstResultIndexFunctionOf(), mlir::nvgpu::getLaneIdAndValueIdToOperandCoord(), getOperandReassociation(), mlir::AffineValueMap::getResult(), getResultIndex(), getResultPosition(), getSingleConstantResult(), getStridesAndOffset(), getSubMap(), getXferIndices(), mlir::inverseAndBroadcastProjectedPermutation(), isDimSequencePreserved(), mlir::AffineValueMap::isFunctionOf(), isSingleConstant(), isTiled(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), packLinalgMetadataOnce(), permute(), printBound(), printMinMaxBound(), replaceDimOrSym(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().
std::optional< unsigned > AffineMap::getResultPosition | ( | AffineExpr | input | ) | const |
Extracts the first result position where input
dimension resides.
Returns std::nullopt
if input
is not a dimension expression or cannot be found in results.
Definition at line 344 of file AffineMap.cpp.
References getNumResults(), getResult(), and mlir::AffineExpr::isa().
ArrayRef< AffineExpr > AffineMap::getResults | ( | ) | const |
Definition at line 332 of file AffineMap.cpp.
References mlir::detail::AffineMapStorage::results().
Referenced by addConstToResults(), mlir::applyMapToValues(), mlir::applyPermutationMap(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::coalesceLoops(), mlir::AffineExpr::compose(), compose(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), compressUnusedImpl(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), createNewDynamicSizes(), dropResults(), mlir::expandAffineMap(), findPermutationsIndexingOperand(), fuse(), getCollapsedOpIndexingMap(), getCollapsedOutputDimFromInputShape(), getConstantResults(), mlir::getConstantTripCount(), mlir::linalg::getConstantUpperBoundForIndex(), getDomainReassociation(), getExpandedType(), mlir::getFlattenedAffineExprs(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::getLargestDivisorOfTripCount(), getLargestKnownDivisorOfMapExprs(), getMinorIdentityMap(), getNumCompoundAffineOnSparseDims(), getNumResults(), getPreservedDims(), getReassociationForExpansion(), getResult(), getShapeDefiningLoopRange(), mlir::Builder::getShiftedAffineMap(), getSliceMap(), getTiledSliceDims(), getTileSizePos(), hasAtMostOneResultFunctionOfDim(), inferTransposeResultType(), insertResult(), mlir::inversePermutation(), isConstant(), isDimSequencePreserved(), isFunctionOfDim(), isFunctionOfSymbol(), isIdentity(), isMinorIdentityWithBroadcasting(), isNormalizedMemRefDynamicDim(), isPermutationOfMinorIdentityWithBroadcasting(), isProjectedPermutation(), isSymbolIdentity(), makeCanonicalAffineApplies(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), TransferWriteNonPermutationLowering::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), mlir::MemRefAccess::operator==(), padOperandToSmallestStaticBoundingBox(), partialConstantFold(), permuteValues(), mlir::AsmPrinter::Impl::printAffineMap(), mlir::removeDuplicateExprs(), replace(), mlir::replaceAllMemRefUsesWith(), replaceDimsAndSymbols(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), shiftDims(), shiftSymbols(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), sliceTransferIndices(), vectorizeAsLinalgGeneric(), verifyOutputShape(), verifyPermutationMap(), and walkExprs().
int64_t AffineMap::getSingleConstantResult | ( | ) | const |
Returns the constant result of this map.
This methods asserts that the map has a single constant result.
Definition at line 306 of file AffineMap.cpp.
References mlir::AffineExpr::cast(), getResult(), and isSingleConstant().
Referenced by mlir::FlatAffineValueConstraints::addAffineParallelOpDomain().
AffineMap AffineMap::getSliceMap | ( | unsigned | start, |
unsigned | length | ||
) | const |
Returns the map consisting of length
expressions starting from start
.
Definition at line 551 of file AffineMap.cpp.
References get(), getContext(), getNumDims(), getNumSymbols(), and getResults().
Referenced by getMajorSubMap(), getMinorSubMap(), and printMinMaxBound().
Returns the map consisting of the resultPos
subset.
Definition at line 543 of file AffineMap.cpp.
References get(), getContext(), getNumDims(), getNumSymbols(), and getResult().
Referenced by calculateExpandedAccessIndices(), composeMultiResultAffineMap(), mlir::linalg::computeSliceParameters(), generateFusedElementwiseOpRegion(), getTiledProducerLoops(), mlir::linalg::interchangeGenericOp(), and mlir::simplifyConstrainedMinMaxOp().
|
static |
Returns a vector of AffineMaps; each with as many results as exprs.size()
, as many dims as the largest dim in exprs
and as many symbols as the largest symbol in exprs
.
Definition at line 237 of file AffineMap.cpp.
References inferFromExprList().
Referenced by mlir::linalg::computeSliceParameters(), contractSupportsMMAMatrixType(), createSubViewIntersection(), mlir::StructuredGenerator< StructuredOpInterface, IteratorTypeT >::layout(), mlir::makeCanonicalStridedLayoutExpr(), mlir::makeComposedAffineApply(), mlir::makeComposedFoldedAffineApply(), mlir::nvgpu::PrepareContractToGPUMMASync::matchAndRewrite(), mlir::vector::ContractionOpToDotLowering::matchAndRewrite(), reduceMatchAndRewriteHelper(), replace(), and mlir::linalg::splitOp().
|
static |
Definition at line 242 of file AffineMap.cpp.
References inferFromExprList().
|
inline |
Returns a new AffineMap with the same number of dims and symbols and an extra result inserted at pos
.
Definition at line 268 of file AffineMap.h.
References get(), getContext(), getNumDims(), getNumSymbols(), and getResults().
Referenced by insertParallelDim(), and packLinalgMetadataOnce().
bool AffineMap::isConstant | ( | ) | const |
Returns true if this affine map has only constant results.
Definition at line 300 of file AffineMap.cpp.
References getResults(), and mlir::AffineExpr::isa().
Referenced by mlir::FlatAffineValueConstraints::addAffineParallelOpDomain(), and getConstantResults().
bool AffineMap::isEmpty | ( | ) | const |
Returns true if this affine map is an empty map, i.e., () -> ().
Definition at line 292 of file AffineMap.cpp.
References getNumDims(), getNumResults(), and getNumSymbols().
Referenced by mlir::linalg::interchangeGenericOp(), mlir::inversePermutation(), and makeCanonicalAffineApplies().
|
inline |
Return true if any affine expression involves AffineDimExpr position
.
Definition at line 180 of file AffineMap.h.
References getResults(), and mlir::AffineExpr::isFunctionOfDim().
Referenced by composeAffineMapAndOperands().
|
inline |
Return true if any affine expression involves AffineSymbolExpr position
.
Definition at line 187 of file AffineMap.h.
References getResults(), and mlir::AffineExpr::isFunctionOfSymbol().
Referenced by composeAffineMapAndOperands().
bool AffineMap::isIdentity | ( | ) | const |
Returns true if this affine map is an identity affine map.
An identity affine map corresponds to an identity affine function on the dimensional identifiers.
Definition at line 268 of file AffineMap.cpp.
References getNumDims(), getNumResults(), and getResults().
Referenced by mlir::canonicalizeStridedLayout(), getStridesAndOffset(), and TransferReadPermutationLowering::matchAndRewrite().
bool AffineMap::isMinorIdentity | ( | ) | const |
Returns true if this affine map is a minor identity, i.e.
an identity affine map (d0, ..., dn) -> (dp, ..., dn) on the most minor dimensions.
Definition at line 111 of file AffineMap.cpp.
References getContext(), getMinorIdentityMap(), getNumDims(), and getNumResults().
Referenced by TransferWritePermutationLowering::matchAndRewrite(), and transferReadSupportsMMAMatrixType().
bool AffineMap::isMinorIdentityWithBroadcasting | ( | SmallVectorImpl< unsigned > * | broadcastedDims = nullptr | ) | const |
Returns true if this affine map is a minor identity up to broadcasted dimensions which are indicated by value 0 in the result.
If broadcastedDims
is not null, it will be populated with the indices of the broadcasted dimensions in the result array. Example: affine_map<(d0, d1, d2, d3, d4) -> (0, d2, 0, d4)> (broadcastedDims
will contain [0, 2])
Definition at line 119 of file AffineMap.cpp.
References mlir::AffineExpr::dyn_cast(), mlir::detail::enumerate(), getNumDims(), getNumResults(), and getResults().
Referenced by TransferOpReduceRank::matchAndRewrite().
bool AffineMap::isPermutation | ( | ) | const |
Returns true if the AffineMap represents a symbol-less permutation map.
Definition at line 537 of file AffineMap.cpp.
References getNumDims(), getNumResults(), and isProjectedPermutation().
Referenced by mlir::linalg::areElementwiseOpsFusable(), permuteValues(), and reductionPreconditions().
bool AffineMap::isPermutationOfMinorIdentityWithBroadcasting | ( | SmallVectorImpl< unsigned > & | permutedDims | ) | const |
Return true if this affine map can be converted to a minor identity with broadcast by doing a permute.
Return a permutation (there may be several) to apply to get to a minor identity with broadcasts. Ex:
Definition at line 159 of file AffineMap.cpp.
References mlir::AffineExpr::dyn_cast(), mlir::detail::enumerate(), getNumInputs(), getNumResults(), getResults(), and max().
Referenced by TransferReadPermutationLowering::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), and TransferWriteNonPermutationLowering::matchAndRewrite().
bool AffineMap::isProjectedPermutation | ( | bool | allowZeroInResults = false | ) | const |
Returns true if the AffineMap represents a subset (i.e.
a projection) of a symbol-less permutation map. allowZeroInResults
allows projected permutation maps with constant zero result expressions. TODO: Remove allowZeroInResults
when constant zero result expressions are broadly supported.
Definition at line 507 of file AffineMap.cpp.
References mlir::AffineExpr::dyn_cast(), getNumInputs(), getNumResults(), getNumSymbols(), and getResults().
Referenced by mlir::applyPermutationMap(), getCollapsedOpIndexingMap(), getDomainReassociation(), getPreservedDims(), getTiledProducerLoops(), mlir::inverseAndBroadcastProjectedPermutation(), isContractionInterfaceImpl(), isDimSequencePreserved(), isPermutation(), and reindexIndexingMap().
bool AffineMap::isSingleConstant | ( | ) | const |
Returns true if this affine map is a single result constant function.
Definition at line 296 of file AffineMap.cpp.
References getNumResults(), getResult(), and mlir::AffineExpr::isa().
Referenced by getSingleConstantResult().
bool AffineMap::isSymbolIdentity | ( | ) | const |
Returns true if this affine map is an identity affine map on the symbol identifiers.
Definition at line 280 of file AffineMap.cpp.
References getNumResults(), getNumSymbols(), and getResults().
|
inlineexplicit |
Definition at line 99 of file AffineMap.h.
|
inline |
Definition at line 101 of file AffineMap.h.
|
inline |
Definition at line 100 of file AffineMap.h.
AffineMap AffineMap::partialConstantFold | ( | ArrayRef< Attribute > | operandConstants, |
SmallVectorImpl< int64_t > * | results = nullptr |
||
) | const |
Propagates the constant operands into this affine map.
Operands are allowed to be null, at which point they are treated as non-constant. This does not change the number of symbols and dimensions. Returns a new map, which may be equal to the old map if no folding happened. If results
is provided and if all expressions in the map were folded to constants, results
will contain the values of these constants.
Definition at line 379 of file AffineMap.cpp.
References get(), mlir::getAffineConstantExpr(), getContext(), getNumDims(), getNumInputs(), getNumResults(), getNumSymbols(), and getResults().
Referenced by constantFold().
void AffineMap::print | ( | raw_ostream & | os | ) | const |
Definition at line 3586 of file AsmPrinter.cpp.
References mlir::AsmState::getImpl(), and mlir::AsmPrinter::Impl::printAffineMap().
Referenced by mlir::operator<<().
AffineMap AffineMap::replace | ( | AffineExpr | expr, |
AffineExpr | replacement, | ||
unsigned | numResultDims, | ||
unsigned | numResultSyms | ||
) | const |
Sparse replace method.
Apply AffineExpr::replace(expr
, replacement
) to each of the results and return a new AffineMap with the new results and with the specified number of dims and symbols.
Definition at line 436 of file AffineMap.cpp.
References get(), getContext(), getNumResults(), and getResults().
Referenced by replaceDimOrSym(), scaleReductionDim(), and mlir::simplifyConstrainedMinMaxOp().
AffineMap AffineMap::replace | ( | const DenseMap< AffineExpr, AffineExpr > & | map | ) | const |
Sparse replace method.
Apply AffineExpr::replace(map
) to each of the results and return a new AffineMap with the new results and with inferred number of dims and symbols.
Definition at line 460 of file AffineMap.cpp.
References getNumResults(), getResults(), and inferFromExprList().
AffineMap AffineMap::replace | ( | const DenseMap< AffineExpr, AffineExpr > & | map, |
unsigned | numResultDims, | ||
unsigned | numResultSyms | ||
) | const |
Sparse replace method.
Apply AffineExpr::replace(map
) to each of the results and return a new AffineMap with the new results and with the specified number of dims and symbols.
Definition at line 449 of file AffineMap.cpp.
References get(), getContext(), getNumResults(), and getResults().
AffineMap AffineMap::replaceDimsAndSymbols | ( | ArrayRef< AffineExpr > | dimReplacements, |
ArrayRef< AffineExpr > | symReplacements, | ||
unsigned | numResultDims, | ||
unsigned | numResultSyms | ||
) | const |
This method substitutes any uses of dimensions and symbols (e.g.
dim#0 with dimReplacements[0]) in subexpressions and returns the modified expression mapping. Because this can be used to eliminate dims and symbols, the client needs to specify the number of dims and symbols in the result. The returned map always has the same number of results.
Definition at line 421 of file AffineMap.cpp.
References get(), getContext(), getNumResults(), getResults(), and mlir::AffineExpr::replaceDimsAndSymbols().
Referenced by mlir::alignAffineMapWithValues(), compose(), and composeAffineMapAndOperands().
|
inline |
Replace dims[offset ...
numDims) by dims[offset + shift ... shift + numDims).
Definition at line 226 of file AffineMap.h.
References get(), getContext(), getNumDims(), getNumSymbols(), getResults(), and mlir::AffineExpr::shiftDims().
Referenced by composeMultiResultAffineMap(), mlir::AffineValueMap::difference(), insertParallelDim(), packLinalgMetadataOnce(), replaceDimOrSym(), scaleReductionDim(), and mlir::simplifyConstrainedMinMaxOp().
|
inline |
Replace symbols[offset ...
numSymbols) by symbols[offset + shift ... shift + numSymbols).
Definition at line 239 of file AffineMap.h.
References get(), getContext(), getNumDims(), getNumSymbols(), getResults(), and mlir::AffineExpr::shiftSymbols().
Referenced by composeMultiResultAffineMap(), mlir::AffineValueMap::difference(), and replaceDimOrSym().
void AffineMap::walkExprs | ( | llvm::function_ref< void(AffineExpr)> | callback | ) | const |
Walk all of the AffineExpr's in this mapping.
Each node in an expression tree is visited in postorder.
Definition at line 411 of file AffineMap.cpp.
References getResults(), and mlir::AffineExpr::walk().
Referenced by mlir::compressUnusedSymbols().
|
friend |
Definition at line 359 of file AffineMap.h.