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	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...

unsigned	getPermutedPosition (unsigned input) const
	Extracts the permuted position where given input index 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)

Detailed Description

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 42 of file AffineMap.h.

Member Typedef Documentation

◆ ImplType

using mlir::AffineMap::ImplType = detail::AffineMapStorage

Definition at line 44 of file AffineMap.h.

Constructor & Destructor Documentation

◆ AffineMap() [1/2]

constexpr mlir::AffineMap::AffineMap ( )

default

◆ AffineMap() [2/2]

mlir::AffineMap::AffineMap ( ImplType * map )

inlineexplicit

Definition at line 47 of file AffineMap.h.

References inferFromExprList().

Member Function Documentation

◆ compose() [1/2]

AffineMap AffineMap::compose ( AffineMap map ) const

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 451 of file AffineMap.cpp.

References mlir::AffineExpr::compose(), get(), mlir::getAffineDimExpr(), mlir::getAffineSymbolExpr(), getContext(), getNumDims(), getNumResults(), getNumSymbols(), and replaceDimsAndSymbols().

Referenced by foldTransferInBoundsAttribute(), mlir::linalg::fuseElementwiseOps(), getIndexingMapOfProducerOperandsInCoordinatesOfFusedOp(), getPermutationMapAttr(), getResultsPositionInLoopsToShapeMap(), insertParallelDim(), mlir::linalg::interchangeGenericOp(), TransferReadPermutationLowering::matchAndRewrite(), scaleReductionDim(), vectorizeAsLinalgGeneric(), and mlir::linalg::detail::verifyStructuredOpInterface().

◆ compose() [2/2]

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 475 of file AffineMap.cpp.

References get(), and mlir::getAffineConstantExpr().

◆ constantFold()

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 343 of file AffineMap.cpp.

References mlir::failure(), and mlir::success().

Referenced by foldLoopBounds().

◆ dropResult()

AffineMap mlir::AffineMap::dropResult ( int64_t pos )

inline

Returns a new AffineMap with the same number of dims and symbols and one less result at pos, dropped.

Definition at line 246 of file AffineMap.h.

Referenced by dropResults().

◆ dropResults()

AffineMap mlir::AffineMap::dropResults ( ArrayRef< int64_t > positions )

inline

Definition at line 254 of file AffineMap.h.

References dropResult().

Referenced by getGenericEffectsImpl().

◆ dump()

void AffineMap::dump ( ) const

Definition at line 3324 of file AsmPrinter.cpp.

References print().

Referenced by mlir::MemRefRegion::compute(), and mlir::FlatAffineValueConstraints::getSliceBounds().

◆ get() [1/4]

AffineMap AffineMap::get ( MLIRContext * context )

static

Returns a zero result affine map with no dimensions or symbols: () -> ().

Definition at line 1022 of file MLIRContext.cpp.

References mlir::MLIRContext::getImpl().

Referenced by addConstToResults(), addLoopRangeConstraints(), adjustMap(), augmentMapAndBounds(), mlir::tensor::bubbleUpPadSlice(), calculateImplicitMap(), mlir::canonicalizeStridedLayout(), mlir::coalesceLoops(), compose(), composeMultiResultAffineMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), compressUnusedImpl(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), mlir::createAffineComputationSlice(), createDepthwiseConvCollapseMap(), createMask(), createNewDynamicSizes(), createPrivateMemRef(), createReassociationMapsForCollapse(), doubleBuffer(), elementwiseMatchAndRewriteHelper(), mlir::linalg::extractOrIdentityMap(), findAncestorOpInRegion(), foldMemRefCast(), fuseWithReshapeByExpansion(), generateCopy(), generatePointWiseCopy(), generateTileLoopNest(), genVectorMask(), mlir::MutableAffineMap::getAffineMap(), getCleanupLoopLowerBound(), getCollapsedOpIndexingMap(), getCollapsedOutputDimFromInputShape(), mlir::Builder::getConstantAffineMap(), mlir::Builder::getDimIdentityMap(), mlir::Builder::getEmptyAffineMap(), getExpandedOutputDimFromInputShape(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::FlatAffineValueConstraints::getIneqAsAffineValueMap(), mlir::nvgpu::getLaneIdAndValueIdToOperandCoord(), mlir::nvgpu::getLaneIdToLdMatrixMatrixCoord(), getLocalInvocationDimSize(), mlir::FlatAffineValueConstraints::getLowerAndUpperBound(), getMinorIdentityMap(), mlir::Builder::getMultiDimIdentityMap(), getPermutationMap(), getPermutationMapAttr(), mlir::nvgpu::getRegisterIndexToTileOffsetMap(), mlir::Builder::getShiftedAffineMap(), mlir::Builder::getSingleDimShiftAffineMap(), mlir::FlatAffineValueConstraints::getSliceBounds(), getSliceMap(), getSubMap(), mlir::Builder::getSymbolIdentityMap(), mlir::getSymbolLessAffineMaps(), mlir::vector::getTransferMinorIdentityMap(), mlir::getTripCountMapAndOperands(), inferFromExprList(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), mlir::isColumnMajorMatmul(), mlir::isRowMajorBatchMatmul(), mlir::isRowMajorMatmul(), mlir::loopUnrollByFactor(), mlir::loopUnrollJamByFactor(), makePermutationMap(), mlir::makeStridedLinearLayoutMap(), TransferOpReduceRank::matchAndRewrite(), mlir::linalg::PadOpTransformationPattern::matchAndRewrite(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), mlirAffineMapEmptyGet(), mlirAffineMapGet(), mlirAffineMapZeroResultGet(), mlir::memref::multiBuffer(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), peelForLoop(), processParallelLoop(), ReifyExpandOrCollapseShapeOp< OpTy >::reifyResultShapes(), mlir::removeDuplicateExprs(), replace(), mlir::replaceAllMemRefUsesWith(), replaceUnitExtents(), resolveSourceIndicesCollapseShape(), resolveSourceIndicesExpandShape(), resolveSourceIndicesSubView(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), sliceTransferIndices(), splatZero(), mlir::linalg::splitReduction(), mlir::linalg::splitReductionByScaling(), mlir::scf::tileParallelLoop(), and transferReadSupportsMMAMatrixType().

◆ get() [2/4]

AffineMap AffineMap::get	(	unsigned	dimCount,
		unsigned	symbolCount,
		MLIRContext *	context
	)

static

Returns a zero result affine map with dimCount dimensions and symbolCount symbols, e.g.

: (...) -> ().

Definition at line 1026 of file MLIRContext.cpp.

References mlir::MLIRContext::getImpl().

◆ get() [3/4]

AffineMap AffineMap::get	(	unsigned	dimCount,
		unsigned	symbolCount,
		AffineExpr	result
	)

static

Returns an affine map with dimCount dimensions and symbolCount mapping to a single output dimension.

Definition at line 1031 of file MLIRContext.cpp.

References mlir::AffineExpr::getContext(), mlir::MLIRContext::getImpl(), and willBeValidAffineMap().

◆ get() [4/4]

AffineMap AffineMap::get	(	unsigned	dimCount,
		unsigned	symbolCount,
		ArrayRef< AffineExpr >	results,
		MLIRContext *	context
	)

static

Returns an affine map with dimCount dimensions and symbolCount mapping to the given results.

Definition at line 1037 of file MLIRContext.cpp.

References mlir::MLIRContext::getImpl(), and willBeValidAffineMap().

◆ getAsOpaquePointer()

const void* mlir::AffineMap::getAsOpaquePointer ( ) const

inline

Methods supporting C API.

Definition at line 339 of file AffineMap.h.

◆ getConstantMap()

AffineMap AffineMap::getConstantMap	(	int64_t	val,
		MLIRContext *	context
	)

static

Returns a single constant result affine map.

Definition at line 96 of file AffineMap.cpp.

References mlir::getAffineConstantExpr().

Referenced by mlir::AffineDmaWaitOp::fold(), mlir::getTripCountMapAndOperands(), mlir::linalg::getUpperBoundForIndex(), hasTrivialZeroTripCount(), and mlirAffineMapConstantGet().

◆ getConstantResults()

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 298 of file AffineMap.cpp.

References mlir::AffineExpr::cast(), and mlir::AffineConstantExpr::getValue().

◆ getContext()

MLIRContext * AffineMap::getContext ( ) const

Definition at line 265 of file AffineMap.cpp.

Referenced by addConstToResults(), mlir::alignAffineMapWithValues(), compose(), composeAffineMapAndOperands(), composeMultiResultAffineMap(), mlir::compressDims(), mlir::compressSymbols(), getCollapsedOpIndexingMap(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), mlir::removeDuplicateExprs(), replaceDimOrSym(), mlir::MutableAffineMap::reset(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), and mlir::vector::detail::transferMaskType().

◆ getDimPosition()

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 327 of file AffineMap.cpp.

Referenced by adjustMap(), getDimPosition(), getReductionIndex(), getResultIndex(), getTiledProducerLoops(), mlir::inverseAndBroadcastProjectedPermutation(), isFusableWithReshapeByDimExpansion(), permute(), and mlir::linalg::splitReduction().

◆ getFromOpaquePointer()

static AffineMap mlir::AffineMap::getFromOpaquePointer ( const void * pointer )

inlinestatic

Definition at line 342 of file AffineMap.h.

◆ getLargestKnownDivisorOfMapExprs()

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 245 of file AffineMap.cpp.

References mlir::presburger::gcd(), and max().

◆ getMajorSubMap()

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 539 of file AffineMap.cpp.

◆ getMinorIdentityMap()

AffineMap AffineMap::getMinorIdentityMap	(	unsigned	dims,
		unsigned	results,
		MLIRContext *	context
	)

static

Returns an identity affine map (d0, ..., dn) -> (dp, ..., dn) on the most minor dimensions.

Definition at line 103 of file AffineMap.cpp.

References get(), and getMultiDimIdentityMap().

Referenced by mlir::vector::getTransferMinorIdentityMap(), TransferWritePermutationLowering::matchAndRewrite(), and mlirAffineMapMinorIdentityGet().

◆ getMinorSubMap()

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 547 of file AffineMap.cpp.

◆ getMultiDimIdentityMap()

AffineMap AffineMap::getMultiDimIdentityMap	(	unsigned	numDims,
		MLIRContext *	context
	)

static

Returns an AffineMap with 'numDims' identity result dim exprs.

Definition at line 256 of file AffineMap.cpp.

References mlir::getAffineDimExpr().

Referenced by mlir::tensor::bubbleUpPadSlice(), buildMax(), buildMin(), mlir::linalg::extractOrIdentityMap(), getGenericEffectsImpl(), mlir::detail::getMemorySpaceAsInt(), getMinorIdentityMap(), getUnorderedCOOFromType(), mlir::linalg::getUpperBoundForIndex(), insertParallelDim(), mlir::linalg::makeMemRefCopyOp(), mlir::linalg::makeTransposeOp(), mlirAffineMapMultiDimIdentityGet(), scaleReductionDim(), tileLinalgOpImpl(), and tilePadOp().

◆ getNumDims()

unsigned AffineMap::getNumDims ( ) const

Definition at line 306 of file AffineMap.cpp.

Referenced by acceptBitWidth(), mlir::FlatAffineValueConstraints::addBound(), addConstToResults(), adjustMap(), mlir::alignAffineMapWithValues(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), compose(), composeAffineMapAndOperands(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), mlir::MemRefRegion::compute(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), mlir::expandAffineMap(), getCleanupLoopLowerBound(), getConstDifference(), mlir::getFlattenedAffineExprs(), getIndicesForAccess(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumDims(), mlir::getRelationFromMap(), getResultsPositionInLoopsToShapeMap(), mlir::Builder::getShiftedAffineMap(), getStridesAndOffset(), getXferIndices(), hasTrivialZeroTripCount(), mlir::inversePermutation(), isFusableWithReshapeByDimExpansion(), TransferWritePermutationLowering::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), mlir::AsmPrinter::Impl::printAffineMap(), printBound(), printMinMaxBound(), remainsLegalAfterInline(), mlir::removeDuplicateExprs(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), splatZero(), mlir::linalg::splitReduction(), transferReadSupportsMMAMatrixType(), mlir::detail::verifyAffineMapAsLayout(), verifyDimAndSymbolIdentifiers(), and mlir::linalg::detail::verifyStructuredOpInterface().

◆ getNumInputs()

unsigned AffineMap::getNumInputs ( ) const

Definition at line 315 of file AffineMap.cpp.

Referenced by mlir::FlatAffineValueConstraints::addBound(), mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), mlir::alignAffineMapWithValues(), mlir::applyPermutationMap(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeSetAndOperands(), mlir::FlatAffineValueConstraints::computeAlignedMap(), createNewDynamicSizes(), mlir::AffineDmaWaitOp::fold(), 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(), hasTrivialZeroTripCount(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), mlir::AffineDmaStartOp::isStrided(), makeCanonicalAffineApplies(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), printBound(), mlir::replaceAllMemRefUsesWith(), simplifyMapWithOperands(), verifyMemoryOpIndexing(), verifyPermutationMap(), verifyTransferOp(), and verifyVectorMemoryOp().

◆ getNumResults()

unsigned AffineMap::getNumResults ( ) const

Definition at line 314 of file AffineMap.cpp.

Referenced by acceptBitWidth(), mlir::FlatAffineValueConstraints::addBound(), mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), adjustMap(), mlir::applyPermutationMap(), mlir::linalg::areElementwiseOpsFusable(), mlir::buildSliceTripCountMap(), calculateExpandedAccessIndices(), canonicalizeMapExprAndTermOrder(), canonicalizeMinMaxOp(), mlir::canonicalizeStridedLayout(), compose(), composeAffineMapAndOperands(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeMultiResultAffineMap(), mlir::compressDims(), mlir::compressSymbols(), constructTiledIndexSetHyperRect(), createMask(), mlir::ComputationSliceState::dump(), foldLoopBounds(), getCleanupLoopLowerBound(), getCollapsedOpOperand(), getConstDifference(), getDimPosition(), mlir::getFlattenedAffineExprs(), getIndicesForAccess(), mlir::getLargestDivisorOfTripCount(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumResults(), getOperandReassociation(), getReductionIndex(), mlir::getRelationFromMap(), getResultIndex(), getShapeDefiningLoopRange(), mlir::Builder::getShiftedAffineMap(), mlir::FlatAffineValueConstraints::getSliceBounds(), getStridesAndOffset(), getTiledProducerLoops(), hasTrivialZeroTripCount(), mlir::inverseAndBroadcastProjectedPermutation(), mlir::inversePermutation(), isDimSequencePreserved(), isFusableWithReshapeByDimExpansion(), isTiled(), mlir::makeComposedFoldedAffineApply(), mlir::makeComposedFoldedMultiResultAffineApply(), mlir::linalg::makeTiledLoopRanges(), TransferReadPermutationLowering::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), mlir::normalizeMemRefType(), permute(), printBound(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), simplifyMapWithOperands(), mlir::linalg::splitReduction(), verifyDimAndSymbolIdentifiers(), verifyMemoryOpIndexing(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().

◆ getNumSymbols()

unsigned AffineMap::getNumSymbols ( ) const

Definition at line 310 of file AffineMap.cpp.

Referenced by mlir::FlatAffineValueConstraints::addBound(), addConstToResults(), mlir::alignAffineMapWithValues(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::canonicalizeStridedLayout(), compose(), mlir::FlatAffineValueConstraints::composeMatchingMap(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), mlir::compressUnusedSymbols(), mlir::MemRefRegion::compute(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), createNewDynamicSizes(), getCleanupLoopLowerBound(), getConstDifference(), mlir::getFlattenedAffineExprs(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::getMultiAffineFunctionFromMap(), mlir::AffineValueMap::getNumSymbols(), mlir::Builder::getShiftedAffineMap(), getStridesAndOffset(), hasTrivialZeroTripCount(), mlir::inversePermutation(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), mlir::AsmPrinter::Impl::printAffineMap(), printBound(), remainsLegalAfterInline(), mlir::removeDuplicateExprs(), mlir::replaceAllMemRefUsesWith(), replaceDimOrSym(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), splatZero(), verifyDimAndSymbolIdentifiers(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().

◆ getPermutationMap()

AffineMap AffineMap::getPermutationMap	(	ArrayRef< unsigned >	permutation,
		MLIRContext *	context
	)

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 206 of file AffineMap.cpp.

References get(), and mlir::getAffineDimExpr().

Referenced by convertTransferReadOp(), foldTransferInBoundsAttribute(), mlir::linalg::interchangeGenericOp(), interchangeGenericOpPrecondition(), mlir::linalg::makeTransposeOp(), TransferReadPermutationLowering::matchAndRewrite(), TransferWritePermutationLowering::matchAndRewrite(), TransferOpReduceRank::matchAndRewrite(), mlirAffineMapPermutationGet(), permute(), splatZero(), tileLinalgOpImpl(), and transferReadSupportsMMAMatrixType().

◆ getPermutedPosition()

unsigned AffineMap::getPermutedPosition ( unsigned input ) const

Extracts the permuted position where given input index resides.

Fails when called on a non-permutation.

Definition at line 331 of file AffineMap.cpp.

References getDimPosition(), and mlir::linalg::isPermutation().

◆ getResult()

AffineExpr AffineMap::getResult ( unsigned idx ) const

Definition at line 323 of file AffineMap.cpp.

Referenced by mlir::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::FlatAffineValueConstraints::addSliceBounds(), mlir::canonicalizeStridedLayout(), mlir::ComputationSliceState::dump(), foldTransferInBoundsAttribute(), mlir::linalg::generateLibraryCallName(), getCleanupLoopLowerBound(), getConstDifference(), mlir::nvgpu::getLaneIdAndValueIdToOperandCoord(), getOperandReassociation(), mlir::AffineValueMap::getResult(), getResultIndex(), getStridesAndOffset(), getXferIndices(), mlir::inverseAndBroadcastProjectedPermutation(), isDimSequencePreserved(), mlir::AffineValueMap::isFunctionOf(), isTiled(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), printBound(), printMinMaxBound(), mlir::vector::detail::transferMaskType(), mlir::linalg::detail::verifyStructuredOpInterface(), and verifyTransferOp().

◆ getResults()

ArrayRef< AffineExpr > AffineMap::getResults ( ) const

Definition at line 319 of file AffineMap.cpp.

Referenced by addConstToResults(), mlir::applyPermutationMap(), augmentMapAndBounds(), canonicalizeMapExprAndTermOrder(), mlir::coalesceLoops(), mlir::AffineExpr::compose(), composeSetAndOperands(), mlir::compressDims(), mlir::compressSymbols(), computeMemoryOpIndices(), constructTiledIndexSetHyperRect(), createNewDynamicSizes(), mlir::expandAffineMap(), fuse(), getCollapsedOpIndexingMap(), getCollapsedOutputDimFromInputShape(), mlir::getConstantTripCount(), mlir::linalg::getConstantUpperBoundForIndex(), getDomainReassociation(), getExpandedType(), mlir::getFlattenedAffineExprs(), getIndexingMapInExpandedOp(), getIndicesForAccess(), mlir::getLargestDivisorOfTripCount(), getPreservedDims(), getReassociationForExpansion(), getResultsPositionInLoopsToShapeMap(), getShapeDefiningLoopRange(), mlir::Builder::getShiftedAffineMap(), getTiledSliceDims(), getTileSizePos(), inlinePayload(), mlir::inversePermutation(), isDimSequencePreserved(), isFusableWithReshapeByDimExpansion(), isNormalizedMemRefDynamicDim(), TransferOpReduceRank::matchAndRewrite(), DeduplicateAffineMinMaxExpressions< T >::matchAndRewrite(), MergeAffineMinMaxOp< T >::matchAndRewrite(), CanonicalizeSingleResultAffineMinMaxOp< T >::matchAndRewrite(), mlir::normalizeAffineFor(), mlir::normalizeAffineParallel(), padOperandToSmallestStaticBoundingBox(), permuteValues(), mlir::AsmPrinter::Impl::printAffineMap(), mlir::removeDuplicateExprs(), mlir::replaceAllMemRefUsesWith(), replaceUnitDimIndexOps(), replaceUnitExtents(), mlir::MutableAffineMap::reset(), setInterTileBoundsParametric(), setIntraTileBoundsParametric(), mlir::simplifyAffineMap(), simplifyMapWithOperands(), sliceTransferIndices(), splatZero(), and verifyPermutationMap().

◆ getSingleConstantResult()

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 293 of file AffineMap.cpp.

◆ getSliceMap()

AffineMap AffineMap::getSliceMap	(	unsigned	start,
		unsigned	length
	)		const

Returns the map consisting of length expressions starting from start.

Definition at line 534 of file AffineMap.cpp.

References get().

Referenced by getResultsPositionInLoopsToShapeMap(), and printMinMaxBound().

◆ getSubMap()

AffineMap AffineMap::getSubMap ( ArrayRef< unsigned > resultPos ) const

Returns the map consisting of the resultPos subset.

Definition at line 526 of file AffineMap.cpp.

References get().

Referenced by calculateExpandedAccessIndices(), canonicalizeMinMaxOp(), composeMultiResultAffineMap(), mlir::linalg::computeSliceParameters(), generateFusedElementwiseOpRegion(), getTiledProducerLoops(), mlir::linalg::interchangeGenericOp(), and mlir::makeComposedFoldedMultiResultAffineApply().

◆ inferFromExprList() [1/2]

SmallVector< AffineMap, 4 > AffineMap::inferFromExprList ( ArrayRef< ArrayRef< AffineExpr >> exprsList )

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 236 of file AffineMap.cpp.

References inferFromExprList().

Referenced by mlir::linalg::computeSliceParameters(), contractSupportsMMAMatrixType(), createReassociationMapsForCollapse(), createSubViewIntersection(), mlir::StructuredGenerator< StructuredOpInterface >::layout(), mlir::makeCanonicalStridedLayoutExpr(), mlir::makeComposedAffineApply(), mlir::makeComposedFoldedAffineApply(), reduceMatchAndRewriteHelper(), replace(), and mlir::linalg::splitOp().

◆ inferFromExprList() [2/2]

SmallVector< AffineMap, 4 > AffineMap::inferFromExprList ( ArrayRef< SmallVector< AffineExpr, 4 >> exprsList )

static

Definition at line 241 of file AffineMap.cpp.

References inferFromExprList().

◆ insertResult()

AffineMap mlir::AffineMap::insertResult	(	AffineExpr	expr,
		unsigned	pos
	)

inline

Returns a new AffineMap with the same number of dims and symbols and an extra result inserted at pos.

Definition at line 263 of file AffineMap.h.

References mlir::hash_value(), and mlir::linalg::isPermutation().

Referenced by insertParallelDim().

◆ isConstant()

bool AffineMap::isConstant ( ) const

Returns true if this affine map has only constant results.

Definition at line 287 of file AffineMap.cpp.

References mlir::AffineExpr::isa().

◆ isEmpty()

bool AffineMap::isEmpty ( ) const

Returns true if this affine map is an empty map, i.e., () -> ().

Definition at line 279 of file AffineMap.cpp.

Referenced by mlir::linalg::interchangeGenericOp(), mlir::inversePermutation(), and makeCanonicalAffineApplies().

◆ isFunctionOfDim()

bool mlir::AffineMap::isFunctionOfDim ( unsigned position ) const

inline

Return true if any affine expression involves AffineDimExpr position.

Definition at line 174 of file AffineMap.h.

References mlir::AffineExpr::isFunctionOfDim().

Referenced by composeAffineMapAndOperands().

◆ isFunctionOfSymbol()

bool mlir::AffineMap::isFunctionOfSymbol ( unsigned position ) const

inline

Return true if any affine expression involves AffineSymbolExpr position.

Definition at line 181 of file AffineMap.h.

References mlir::AffineExpr::isFunctionOfSymbol().

Referenced by composeAffineMapAndOperands().

◆ isIdentity()

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 267 of file AffineMap.cpp.

Referenced by mlir::canonicalizeStridedLayout(), getStridesAndOffset(), and TransferReadPermutationLowering::matchAndRewrite().

◆ isMinorIdentity()

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 110 of file AffineMap.cpp.

Referenced by TransferWritePermutationLowering::matchAndRewrite(), and transferReadSupportsMMAMatrixType().

◆ isMinorIdentityWithBroadcasting()

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 118 of file AffineMap.cpp.

References mlir::AffineExpr::dyn_cast(), and mlir::detail::enumerate().

Referenced by TransferOpReduceRank::matchAndRewrite().

◆ isPermutation()

bool AffineMap::isPermutation ( ) const

Returns true if the AffineMap represents a symbol-less permutation map.

Definition at line 520 of file AffineMap.cpp.

Referenced by acceptBitWidth(), mlir::linalg::areElementwiseOpsFusable(), permuteValues(), and reductionPreconditions().

◆ isPermutationOfMinorIdentityWithBroadcasting()

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:

(d0, d1, d2) -> (0, d1) maps to minor identity (d1, 0 = d2) with perm = [1, 0] and broadcast d2
(d0, d1, d2) -> (d0, 0) cannot be mapped to a minor identity by permutation + broadcast
(d0, d1, d2, d3) -> (0, d1, d3) maps to minor identity (d1, 0 = d2, d3) with perm = [1, 0, 2] and broadcast d2
(d0, d1) -> (d1, 0, 0, d0) maps to minor identity (d0, d1) with extra leading broadcat dimensions. The map returned would be (0, 0, d0, d1) with perm = [3, 0, 1, 2]

Definition at line 158 of file AffineMap.cpp.

References mlir::AffineExpr::dyn_cast(), mlir::detail::enumerate(), and max().

Referenced by TransferReadPermutationLowering::matchAndRewrite(), and TransferWritePermutationLowering::matchAndRewrite().

◆ isProjectedPermutation()

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 490 of file AffineMap.cpp.

References mlir::AffineExpr::dyn_cast().

Referenced by mlir::applyPermutationMap(), getCollapsedOpIndexingMap(), getDomainReassociation(), getPreservedDims(), getTiledProducerLoops(), mlir::inverseAndBroadcastProjectedPermutation(), isContractionInterfaceImpl(), isDimSequencePreserved(), and reindexIndexingMap().

◆ isSingleConstant()

bool AffineMap::isSingleConstant ( ) const

Returns true if this affine map is a single result constant function.

Definition at line 283 of file AffineMap.cpp.

Referenced by hasTrivialZeroTripCount().

◆ operator bool()

mlir::AffineMap::operator bool ( ) const

inlineexplicit

Definition at line 98 of file AffineMap.h.

◆ operator!=()

bool mlir::AffineMap::operator!= ( AffineMap other ) const

inline

Definition at line 100 of file AffineMap.h.

References getDimPosition(), and print().

◆ operator==()

bool mlir::AffineMap::operator== ( AffineMap other ) const

inline

Definition at line 99 of file AffineMap.h.

◆ partialConstantFold()

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 362 of file AffineMap.cpp.

References mlir::getAffineConstantExpr().

◆ print()

void AffineMap::print ( raw_ostream & os ) const

Definition at line 3348 of file AsmPrinter.cpp.

References mlir::AsmState::getImpl().

Referenced by mlir::operator<<().

◆ replace() [1/3]

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 419 of file AffineMap.cpp.

References get().

Referenced by replaceDimOrSym(), and scaleReductionDim().

◆ replace() [2/3]

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 443 of file AffineMap.cpp.

References inferFromExprList().

◆ replace() [3/3]

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 432 of file AffineMap.cpp.

References get().

◆ replaceDimsAndSymbols()

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 404 of file AffineMap.cpp.

References mlir::AffineExpr::replaceDimsAndSymbols().

Referenced by mlir::alignAffineMapWithValues(), compose(), and composeAffineMapAndOperands().

◆ shiftDims()

AffineMap mlir::AffineMap::shiftDims	(	unsigned	shift,
		unsigned	offset = `0`
	)		const

inline

Replace dims[offset ...

numDims) by dims[offset + shift ... shift + numDims).

Definition at line 220 of file AffineMap.h.

References mlir::AffineExpr::shiftDims().

Referenced by canonicalizeMinMaxOp(), insertParallelDim(), replaceDimOrSym(), and scaleReductionDim().

◆ shiftSymbols()

AffineMap mlir::AffineMap::shiftSymbols	(	unsigned	shift,
		unsigned	offset = `0`
	)		const

inline

Replace symbols[offset ...

numSymbols) by symbols[offset + shift ... shift + numSymbols).

Definition at line 233 of file AffineMap.h.

References mlir::AffineExpr::shiftSymbols().

◆ walkExprs()

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 394 of file AffineMap.cpp.

References mlir::AffineExpr::walk().

Referenced by mlir::compressUnusedSymbols().

Friends And Related Function Documentation

◆ hash_value

::llvm::hash_code hash_value ( AffineMap arg )

friend

Definition at line 354 of file AffineMap.h.

The documentation for this class was generated from the following files:

include/mlir/IR/AffineMap.h
lib/IR/AffineMap.cpp
lib/IR/AsmPrinter.cpp
lib/IR/MLIRContext.cpp

Public Types

Public Member Functions

Static Public Member Functions

Friends

Detailed Description

Member Typedef Documentation

◆ ImplType

Constructor & Destructor Documentation

◆ AffineMap() [1/2]

◆ AffineMap() [2/2]

Member Function Documentation

◆ compose() [1/2]

◆ compose() [2/2]

◆ constantFold()

◆ dropResult()

◆ dropResults()

◆ dump()

◆ get() [1/4]

◆ get() [2/4]

◆ get() [3/4]

◆ get() [4/4]

◆ getAsOpaquePointer()

◆ getConstantMap()

◆ getConstantResults()

◆ getContext()

◆ getDimPosition()

◆ getFromOpaquePointer()

◆ getLargestKnownDivisorOfMapExprs()

◆ getMajorSubMap()

◆ getMinorIdentityMap()

◆ getMinorSubMap()

◆ getMultiDimIdentityMap()

◆ getNumDims()

◆ getNumInputs()

◆ getNumResults()

◆ getNumSymbols()

◆ getPermutationMap()

◆ getPermutedPosition()

◆ getResult()

◆ getResults()

◆ getSingleConstantResult()

◆ getSliceMap()

◆ getSubMap()

◆ inferFromExprList() [1/2]

◆ inferFromExprList() [2/2]

◆ insertResult()

◆ isConstant()

◆ isEmpty()

◆ isFunctionOfDim()

◆ isFunctionOfSymbol()

◆ isIdentity()

◆ isMinorIdentity()

◆ isMinorIdentityWithBroadcasting()

◆ isPermutation()

◆ isPermutationOfMinorIdentityWithBroadcasting()

◆ isProjectedPermutation()

◆ isSingleConstant()

◆ operator bool()

◆ operator!=()

◆ operator==()

◆ partialConstantFold()

◆ print()

◆ replace() [1/3]

◆ replace() [2/3]

◆ replace() [3/3]

◆ replaceDimsAndSymbols()

◆ shiftDims()

◆ shiftSymbols()

◆ walkExprs()

Friends And Related Function Documentation

◆ hash_value