Namespaces
namespace	impl
namespace	matcher

Classes
class	AffineBound
	AffineBound represents a lower or upper bound in the for operation. More...
struct	AffineBuilder
	Helper struct to build simple AffineValueExprs with minimal type inference support. More...
struct	AffineCopyOptions
	Explicit copy / DMA generation options for mlir::affineDataCopyGenerate. More...
struct	AffineDataCopyGenerationOptions
class	AffineDmaStartOp
	AffineDmaStartOp starts a non-blocking DMA operation that transfers data from a source memref to a destination memref. More...
class	AffineDmaWaitOp
	AffineDmaWaitOp blocks until the completion of a DMA operation associated with the tag element 'tag[index]'. More...
struct	AffineLoopFusionOptions
struct	AffineLoopNormalizeOptions
struct	AffineLoopTilingOptions
struct	AffineLoopUnrollAndJamOptions
struct	AffineLoopUnrollOptions
struct	AffineParallelizeOptions
struct	AffineValueExpr
class	AffineValueMap
	An AffineValueMap is an affine map plus its ML value operands and results for analysis purposes. More...
struct	AffineVectorizeOptions
struct	ComputationSliceState
	ComputationSliceState aggregates loop IVs, loop bound AffineMaps and their associated operands for a set of loops within a loop nest (typically the set of loops surrounding a store operation). More...
struct	DependenceComponent
struct	DependenceResult
	Checks whether two accesses to the same memref access the same element. More...
class	FlatAffineRelation
	A FlatAffineRelation represents a set of ordered pairs (domain -> range) where "domain" and "range" are tuples of variables. More...
class	FlatAffineValueConstraints
	FlatAffineValueConstraints is an extension of FlatLinearValueConstraints with helper functions for Affine dialect ops. More...
struct	FusionResult
class	FusionStrategy
	Describes the fusion strategy to be used in the Affine loop fusion utilities. More...
struct	LoopNestStateCollector
struct	LoopNestStats
	LoopNestStats aggregates various per-loop statistics (eg. More...
struct	LoopReduction
	A description of a (parallelizable) reduction in an affine loop. More...
struct	MemRefAccess
	Encapsulates a memref load or store access information. More...
struct	MemRefDependenceGraph
struct	MemRefRegion
	A region of a memref's data space; this is typically constructed by analyzing load/store op's on this memref and the index space of loops surrounding such op's. More...
class	NestedMatch
	An NestedPattern captures nested patterns in the IR. More...
class	NestedPattern
class	NestedPatternContext
	RAII structure to transparently manage the bump allocator for NestedPattern and NestedMatch classes. More...
struct	SliceComputationResult
	Enumerates different result statuses of slice computation by computeSliceUnion More...

Typedefs
using	VectorizableLoopFun = std::function<bool(AffineForOp)>
using	FilterFunctionType = std::function<bool(Operation &)>
	A NestedPattern is a nested operation walker that:
using	ReductionLoopMap = DenseMap<Operation *, SmallVector<LoopReduction, 2>>

Enumerations
enum	FusionMode { Greedy , ProducerConsumer , Sibling }
	Fusion mode to attempt. More...

Functions
void	getSupportedReductions (AffineForOp forOp, SmallVectorImpl< LoopReduction > &supportedReductions)
	Populate supportedReductions with descriptors of the supported reductions.
bool	isLoopMemoryParallel (AffineForOp forOp)
	Returns true if ‘forOp’ is a parallel loop.
void	getReachableAffineApplyOps (ArrayRef< Value > operands, SmallVectorImpl< Operation * > &affineApplyOps)
	Returns in affineApplyOps, the sequence of those AffineApplyOp Operations that are reachable via a search starting from operands and ending at those operands that are not the result of an AffineApplyOp.
LogicalResult	getIndexSet (MutableArrayRef< Operation * > ops, FlatAffineValueConstraints *domain)
	Builds a system of constraints with dimensional variables corresponding to the loop IVs of the forOps and AffineIfOp's operands appearing in that order.
DependenceResult	checkMemrefAccessDependence (const MemRefAccess &srcAccess, const MemRefAccess &dstAccess, unsigned loopDepth, FlatAffineValueConstraints dependenceConstraints=nullptr, SmallVector< DependenceComponent, 2 > dependenceComponents=nullptr, bool allowRAR=false)
bool	hasDependence (DependenceResult result)
	Utility function that returns true if the provided DependenceResult corresponds to a dependence result.
bool	noDependence (DependenceResult result)
	Returns true if the provided DependenceResult corresponds to the absence of a dependence.
void	getDependenceComponents (AffineForOp forOp, unsigned maxLoopDepth, std::vector< SmallVector< DependenceComponent, 2 > > *depCompsVec)
	Returns in 'depCompsVec', dependence components for dependences between all load and store ops in loop nest rooted at 'forOp', at loop depths in range [1, maxLoopDepth].
LogicalResult	getRelationFromMap (AffineMap &map, presburger::IntegerRelation &rel)
	Builds a relation from the given AffineMap/AffineValueMap map, containing all pairs of the form operands -> result that satisfy map.
LogicalResult	getRelationFromMap (const AffineValueMap &map, presburger::IntegerRelation &rel)
void	getTripCountMapAndOperands (AffineForOp forOp, AffineMap map, SmallVectorImpl< Value > operands)
	Returns the trip count of the loop as an affine map with its corresponding operands if the latter is expressible as an affine expression, and nullptr otherwise.
std::optional< uint64_t >	getConstantTripCount (AffineForOp forOp)
	Returns the trip count of the loop if it's a constant, std::nullopt otherwise.
uint64_t	getLargestDivisorOfTripCount (AffineForOp forOp)
	Returns the greatest known integral divisor of the trip count.
template<typename LoadOrStoreOp>
bool	isInvariantAccess (LoadOrStoreOp memOp, AffineForOp forOp)
	Checks if an affine read or write operation depends on forOp's IV, i.e., if the memory access is invariant on forOp.
DenseSet< Value, DenseMapInfo< Value > >	getInvariantAccesses (Value iv, ArrayRef< Value > indices)
	Given an induction variable iv of type AffineForOp and indices of type IndexType, returns the set of indices that are independent of iv.
template<typename LoadOrStoreOp>
bool	isContiguousAccess (Value iv, LoadOrStoreOp memoryOp, int *memRefDim)
	Given:
bool	isVectorizableLoopBody (AffineForOp loop, NestedPattern &vectorTransferMatcher)
	Checks whether the loop is structurally vectorizable; i.e.:
bool	isVectorizableLoopBody (AffineForOp loop, int *memRefDim, NestedPattern &vectorTransferMatcher)
	Checks whether the loop is structurally vectorizable and that all the LoadOp and StoreOp matched have access indexing functions that are either:
bool	isOpwiseShiftValid (AffineForOp forOp, ArrayRef< uint64_t > shifts)
	Checks where SSA dominance would be violated if a for op's body operations are shifted by the specified shifts.
bool	isTilingValid (ArrayRef< AffineForOp > loops)
	Checks whether hyper-rectangular loop tiling of the nest represented by loops is valid.
bool	hasCyclicDependence (AffineForOp root)
	Returns true if the affine nest rooted at root has a cyclic dependence among its affine memory accesses.
bool	defaultFilterFunction (Operation &)
void	getAffineForIVs (Operation &op, SmallVectorImpl< AffineForOp > *loops)
	Populates 'loops' with IVs of the affine.for ops surrounding 'op' ordered from the outermost 'affine.for' operation to the innermost one while not traversing outside of the surrounding affine scope.
void	getAffineIVs (Operation &op, SmallVectorImpl< Value > &ivs)
	Populates 'ivs' with IVs of the surrounding affine.for and affine.parallel ops ordered from the outermost one to the innermost while not traversing outside of the surrounding affine scope.
void	getEnclosingAffineOps (Operation &op, SmallVectorImpl< Operation * > *ops)
	Populates 'ops' with affine operations enclosing op ordered from outermost to innermost while stopping at the boundary of the affine scope.
unsigned	getNestingDepth (Operation *op)
	Returns the nesting depth of this operation, i.e., the number of loops surrounding this operation.
bool	isLoopParallelAndContainsReduction (AffineForOp forOp)
	Returns whether a loop is a parallel loop and contains a reduction loop.
void	getSequentialLoops (AffineForOp forOp, llvm::SmallDenseSet< Value, 8 > *sequentialLoops)
	Returns in 'sequentialLoops' all sequential loops in loop nest rooted at 'forOp'.
void	getComputationSliceState (Operation depSourceOp, Operation depSinkOp, const FlatAffineValueConstraints &dependenceConstraints, unsigned loopDepth, bool isBackwardSlice, ComputationSliceState *sliceState)
	Computes the computation slice loop bounds for one loop nest as affine maps of the other loop nest's IVs and symbols, using 'dependenceConstraints' computed between 'depSourceAccess' and 'depSinkAccess'.
uint64_t	getSliceIterationCount (const llvm::SmallDenseMap< Operation *, uint64_t, 8 > &sliceTripCountMap)
	Return the number of iterations for the slicetripCountMap provided.
bool	buildSliceTripCountMap (const ComputationSliceState &slice, llvm::SmallDenseMap< Operation , uint64_t, 8 > tripCountMap)
	Builds a map 'tripCountMap' from AffineForOp to constant trip count for loop nest surrounding represented by slice loop bounds in 'slice'.
SliceComputationResult	computeSliceUnion (ArrayRef< Operation * > opsA, ArrayRef< Operation * > opsB, unsigned loopDepth, unsigned numCommonLoops, bool isBackwardSlice, ComputationSliceState *sliceUnion)
	Computes in 'sliceUnion' the union of all slice bounds computed at 'loopDepth' between all dependent pairs of ops in 'opsA' and 'opsB', and then verifies if it is valid.
AffineForOp	insertBackwardComputationSlice (Operation srcOpInst, Operation dstOpInst, unsigned dstLoopDepth, ComputationSliceState *sliceState)
	Creates a clone of the computation contained in the loop nest surrounding 'srcOpInst', slices the iteration space of src loop based on slice bounds in 'sliceState', and inserts the computation slice at the beginning of the operation block of the loop at 'dstLoopDepth' in the loop nest surrounding 'dstOpInst'.
std::optional< uint64_t >	getIntOrFloatMemRefSizeInBytes (MemRefType memRefType)
	Returns the size of a memref with element type int or float in bytes if it's statically shaped, std::nullopt otherwise.
template<typename LoadOrStoreOpPointer>
LogicalResult	boundCheckLoadOrStoreOp (LoadOrStoreOpPointer loadOrStoreOp, bool emitError=true)
	Checks a load or store op for an out of bound access; returns failure if the access is out of bounds along any of the dimensions, success otherwise.
unsigned	getNumCommonSurroundingLoops (Operation &a, Operation &b)
	Returns the number of surrounding loops common to both A and B.
std::optional< int64_t >	getMemoryFootprintBytes (AffineForOp forOp, int memorySpace=-1)
	Gets the memory footprint of all data touched in the specified memory space in bytes; if the memory space is unspecified, considers all memory spaces.
std::optional< int64_t >	getMemRefIntOrFloatEltSizeInBytes (MemRefType memRefType)
	Returns the memref's element type's size in bytes where the elemental type is an int or float or a vector of such types.
IntegerSet	simplifyIntegerSet (IntegerSet set)
	Simplify the integer set by simplifying the underlying affine expressions by flattening and some simple inference.
unsigned	getInnermostCommonLoopDepth (ArrayRef< Operation * > ops, SmallVectorImpl< AffineForOp > *surroundingLoops=nullptr)
	Returns the innermost common loop depth for the set of operations in 'ops'.
FailureOr< AffineValueMap >	simplifyConstrainedMinMaxOp (Operation *op, FlatAffineValueConstraints constraints)
	Try to simplify the given affine.min or affine.max op to an affine map with a single result and operands, taking into account the specified constraint set.
mlir::Block *	findInnermostCommonBlockInScope (mlir::Operation a, mlir::Operation b)
	Find the innermost common Block of a and b in the affine scope that a and b are part of.
bool	isTopLevelValue (Value value)
	A utility function to check if a value is defined at the top level of an op with trait AffineScope or is a region argument for such an op.
bool	isTopLevelValue (Value value, Region *region)
	A utility function to check if a value is defined at the top level of region or is an argument of region.
Region *	getAffineScope (Operation *op)
	Returns the closest region enclosing op that is held by an operation with trait AffineScope; nullptr if there is no such region.
Region *	getAffineAnalysisScope (Operation *op)
	Returns the closest region enclosing op that is held by a non-affine operation; nullptr if there is no such region.
OpFoldResult	computeProduct (Location loc, OpBuilder &builder, ArrayRef< OpFoldResult > terms)
	Return the product of terms, creating an affine.apply if any of them are non-constant values.
bool	isValidDim (Value value)
	Returns true if the given Value can be used as a dimension id in the region of the closest surrounding op that has the trait AffineScope.
bool	isValidDim (Value value, Region *region)
	Returns true if the given Value can be used as a dimension id in region, i.e., for all its uses in region.
bool	isValidSymbol (Value value)
	Returns true if the given value can be used as a symbol in the region of the closest surrounding op that has the trait AffineScope.
bool	isValidSymbol (Value value, Region *region)
	Returns true if the given Value can be used as a symbol for region, i.e., for all its uses in region.
ParseResult	parseDimAndSymbolList (OpAsmParser &parser, SmallVectorImpl< Value > &operands, unsigned &numDims)
	Parses dimension and symbol list.
void	canonicalizeMapAndOperands (AffineMap map, SmallVectorImpl< Value > operands)
	Modifies both map and operands in-place so as to:
void	canonicalizeSetAndOperands (IntegerSet set, SmallVectorImpl< Value > operands)
	Canonicalizes an integer set the same way canonicalizeMapAndOperands does for affine maps.
AffineApplyOp	makeComposedAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
	Returns a composed AffineApplyOp by composing map and operands with other AffineApplyOps supplying those operands.
AffineApplyOp	makeComposedAffineApply (OpBuilder &b, Location loc, AffineExpr e, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
OpFoldResult	makeComposedFoldedAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
	Constructs an AffineApplyOp that applies map to operands after composing the map with the maps of any other AffineApplyOp supplying the operands, then immediately attempts to fold it.
OpFoldResult	makeComposedFoldedAffineApply (OpBuilder &b, Location loc, AffineExpr expr, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
	Variant of makeComposedFoldedAffineApply that applies to an expression.
SmallVector< OpFoldResult >	makeComposedFoldedMultiResultAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands, bool composeAffineMin=false)
	Variant of makeComposedFoldedAffineApply suitable for multi-result maps.
AffineMinOp	makeComposedAffineMin (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
	Returns an AffineMinOp obtained by composing map and operands with AffineApplyOps supplying those operands.
OpFoldResult	makeComposedFoldedAffineMin (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
	Constructs an AffineMinOp that computes a minimum across the results of applying map to operands, then immediately attempts to fold it.
OpFoldResult	makeComposedFoldedAffineMax (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
	Constructs an AffineMinOp that computes a maximum across the results of applying map to operands, then immediately attempts to fold it.
void	fullyComposeAffineMapAndOperands (AffineMap map, SmallVectorImpl< Value > operands, bool composeAffineMin=false)
	Given an affine map map and its input operands, this method composes into map, maps of AffineApplyOps whose results are the values in operands, iteratively until no more of operands are the result of an AffineApplyOp.
bool	isAffineForInductionVar (Value val)
	Returns true if the provided value is the induction variable of an AffineForOp.
bool	isAffineParallelInductionVar (Value val)
	Returns true if val is the induction variable of an AffineParallelOp.
bool	isAffineInductionVar (Value val)
	Returns true if the provided value is the induction variable of an AffineForOp or AffineParallelOp.
AffineForOp	getForInductionVarOwner (Value val)
	Returns the loop parent of an induction variable.
AffineParallelOp	getAffineParallelInductionVarOwner (Value val)
	Returns true if the provided value is among the induction variables of an AffineParallelOp.
void	extractForInductionVars (ArrayRef< AffineForOp > forInsts, SmallVectorImpl< Value > *ivs)
	Extracts the induction variables from a list of AffineForOps and places them in the output argument ivs.
void	extractInductionVars (ArrayRef< Operation * > affineOps, SmallVectorImpl< Value > &ivs)
	Extracts the induction variables from a list of either AffineForOp or AffineParallelOp and places them in the output argument ivs.
void	buildAffineLoopNest (OpBuilder &builder, Location loc, ArrayRef< int64_t > lbs, ArrayRef< int64_t > ubs, ArrayRef< int64_t > steps, function_ref< void(OpBuilder &, Location, ValueRange)> bodyBuilderFn=nullptr)
	Builds a perfect nest of affine.for loops, i.e., each loop except the innermost one contains only another loop and a terminator.
void	buildAffineLoopNest (OpBuilder &builder, Location loc, ValueRange lbs, ValueRange ubs, ArrayRef< int64_t > steps, function_ref< void(OpBuilder &, Location, ValueRange)> bodyBuilderFn=nullptr)
void	registerValueBoundsOpInterfaceExternalModels (DialectRegistry &registry)
FailureOr< int64_t >	fullyComposeAndComputeConstantDelta (Value value1, Value value2)
	Compute a constant delta of the given two values.
FusionResult	canFuseLoops (AffineForOp srcForOp, AffineForOp dstForOp, unsigned dstLoopDepth, ComputationSliceState *srcSlice, FusionStrategy fusionStrategy=FusionStrategy::Generic)
	Checks the feasibility of fusing the loop nest rooted at 'srcForOp' into the loop nest rooted at 'dstForOp' at 'dstLoopDepth'.
void	fuseLoops (AffineForOp srcForOp, AffineForOp dstForOp, const ComputationSliceState &srcSlice, bool isInnermostSiblingInsertionFusion=false)
	Fuses 'srcForOp' into 'dstForOp' with destination loop block insertion point and source slice loop bounds specified in 'srcSlice'.
bool	getLoopNestStats (AffineForOp forOp, LoopNestStats *stats)
	Collect loop nest statistics (eg.
int64_t	getComputeCost (AffineForOp forOp, LoopNestStats &stats)
	Computes the total cost of the loop nest rooted at 'forOp' using 'stats'.
bool	getFusionComputeCost (AffineForOp srcForOp, LoopNestStats &srcStats, AffineForOp dstForOp, LoopNestStats &dstStats, const ComputationSliceState &slice, int64_t *computeCost)
	Computes and returns in 'computeCost', the total compute cost of fusing the 'slice' of the loop nest rooted at 'srcForOp' into 'dstForOp'.
void	gatherProducerConsumerMemrefs (ArrayRef< Operation * > srcOps, ArrayRef< Operation * > dstOps, DenseSet< Value > &producerConsumerMemrefs)
	Returns in 'producerConsumerMemrefs' the memrefs involved in a producer-consumer dependence between write ops in 'srcOps' and read ops in 'dstOps'.
LogicalResult	loopUnrollFull (AffineForOp forOp)
	Unrolls this for operation completely if the trip count is known to be constant.
LogicalResult	loopUnrollByFactor (AffineForOp forOp, uint64_t unrollFactor, function_ref< void(unsigned, Operation *, OpBuilder)> annotateFn=nullptr, bool cleanUpUnroll=false)
	Unrolls this for operation by the specified unroll factor.
LogicalResult	loopUnrollUpToFactor (AffineForOp forOp, uint64_t unrollFactor)
	Unrolls this loop by the specified unroll factor or its trip count, whichever is lower.
bool	isPerfectlyNested (ArrayRef< AffineForOp > loops)
	Returns true if loops is a perfectly nested loop nest, where loops appear in it from outermost to innermost.
void	getPerfectlyNestedLoops (SmallVectorImpl< AffineForOp > &nestedLoops, AffineForOp root)
	Get perfectly nested sequence of loops starting at root of loop nest (the first op being another AffineFor, and the second op - a terminator).
LogicalResult	loopUnrollJamByFactor (AffineForOp forOp, uint64_t unrollJamFactor)
	Unrolls and jams this loop by the specified factor.
LogicalResult	loopUnrollJamUpToFactor (AffineForOp forOp, uint64_t unrollJamFactor)
	Unrolls and jams this loop by the specified factor or by the trip count (if constant), whichever is lower.
LogicalResult	promoteIfSingleIteration (AffineForOp forOp)
	Promotes the loop body of a AffineForOp to its containing block if the loop was known to have a single iteration.
void	promoteSingleIterationLoops (func::FuncOp f)
	Promotes all single iteration AffineForOp's in the Function, i.e., moves their body into the containing Block.
LogicalResult	affineForOpBodySkew (AffineForOp forOp, ArrayRef< uint64_t > shifts, bool unrollPrologueEpilogue=false)
	Skew the operations in an affine.for's body with the specified operation-wise shifts.
LogicalResult	tilePerfectlyNested (MutableArrayRef< AffineForOp > input, ArrayRef< unsigned > tileSizes, SmallVectorImpl< AffineForOp > *tiledNest=nullptr)
	Tiles the specified band of perfectly nested loops creating tile-space loops and intra-tile loops.
LogicalResult	tilePerfectlyNestedParametric (MutableArrayRef< AffineForOp > input, ArrayRef< Value > tileSizes, SmallVectorImpl< AffineForOp > *tiledNest=nullptr)
	Tiles the specified band of perfectly nested loops creating tile-space loops and intra-tile loops, using SSA values as tiling parameters.
void	interchangeLoops (AffineForOp forOpA, AffineForOp forOpB)
	Performs loop interchange on 'forOpA' and 'forOpB'.
bool	isValidLoopInterchangePermutation (ArrayRef< AffineForOp > loops, ArrayRef< unsigned > loopPermMap)
	Checks if the loop interchange permutation 'loopPermMap', of the perfectly nested sequence of loops in 'loops', would violate dependences (loop 'i' in 'loops' is mapped to location 'j = 'loopPermMap[i]' in the interchange).
unsigned	permuteLoops (ArrayRef< AffineForOp > inputNest, ArrayRef< unsigned > permMap)
	Performs a loop permutation on a perfectly nested loop nest inputNest (where the contained loops appear from outer to inner) as specified by the permutation permMap: loop 'i' in inputNest is mapped to location 'loopPermMap[i]', where positions 0, 1, ... are from the outermost position to inner.
AffineForOp	sinkSequentialLoops (AffineForOp forOp)
SmallVector< SmallVector< AffineForOp, 8 >, 8 >	tile (ArrayRef< AffineForOp > forOps, ArrayRef< uint64_t > sizes, ArrayRef< AffineForOp > targets)
	Performs tiling fo imperfectly nested loops (with interchange) by strip-mining the forOps by sizes and sinking them, in their order of occurrence in forOps, under each of the targets.
SmallVector< AffineForOp, 8 >	tile (ArrayRef< AffineForOp > forOps, ArrayRef< uint64_t > sizes, AffineForOp target)
	Performs tiling (with interchange) by strip-mining the forOps by sizes and sinking them, in their order of occurrence in forOps, under target.
int64_t	numEnclosingInvariantLoops (OpOperand &operand)
	Performs explicit copying for the contiguous sequence of operations in the block iterator range [‘begin’, ‘end’), where ‘end’ can't be past the / terminator of the block (since additional operations are potentially / inserted right before end.
std::unique_ptr<::mlir::Pass >	createAffineParallelize ()
std::unique_ptr<::mlir::Pass >	createAffineParallelize (AffineParallelizeOptions options)
std::unique_ptr<::mlir::Pass >	createAffineVectorize ()
std::unique_ptr<::mlir::Pass >	createAffineVectorize (AffineVectorizeOptions options)
std::unique_ptr<::mlir::Pass >	createSimplifyAffineMinMaxPass ()
std::unique_ptr< OperationPass< func::FuncOp > >	createSimplifyAffineStructuresPass ()
	Creates a simplification pass for affine structures (maps and sets).
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineLoopInvariantCodeMotionPass ()
	Creates a loop invariant code motion pass that hoists loop invariant operations out of affine loops.
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineParallelizePass ()
	Creates a pass to convert all parallel affine.for's into 1-d affine.parallel ops.
std::unique_ptr< OperationPass< func::FuncOp > >	createRaiseMemrefToAffine ()
	Creates a pass that converts some memref operators to affine operators.
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineLoopNormalizePass (bool promoteSingleIter=false)
	Apply normalization transformations to affine loop-like ops.
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineDataCopyGenerationPass (unsigned slowMemorySpace, unsigned fastMemorySpace, unsigned tagMemorySpace=0, int minDmaTransferSize=1024, uint64_t fastMemCapacityBytes=std::numeric_limits< uint64_t >::max())
	Performs packing (or explicit copying) of accessed memref regions into buffers in the specified faster memory space through either pointwise copies or DMA operations.
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineDataCopyGenerationPass ()
	Overload relying on pass options for initialization.
std::unique_ptr< OperationPass< func::FuncOp > >	createAffineScalarReplacementPass ()
	Creates a pass to replace affine memref accesses by scalars using store to load forwarding and redundant load elimination; consequently also eliminate dead allocs.
std::unique_ptr< OperationPass< func::FuncOp > >	createLoopCoalescingPass ()
	Creates a pass that transforms perfectly nested loops with independent bounds into a single loop.
std::unique_ptr< Pass >	createLoopFusionPass (unsigned fastMemorySpace=0, uint64_t localBufSizeThreshold=0, bool maximalFusion=false, enum FusionMode fusionMode=FusionMode::Greedy)
	Creates a loop fusion pass which fuses affine loop nests at the top-level of the operation the pass is created on according to the type of fusion specified in fusionMode.
std::unique_ptr< OperationPass< func::FuncOp > >	createLoopTilingPass (uint64_t cacheSizeBytes)
	Creates a pass to perform tiling on loop nests.
std::unique_ptr< OperationPass< func::FuncOp > >	createLoopTilingPass ()
	Overload relying on pass options for initialization.
std::unique_ptr< InterfacePass< FunctionOpInterface > >	createLoopUnrollPass (int unrollFactor=-1, bool unrollUpToFactor=false, const std::function< unsigned(AffineForOp)> &getUnrollFactor=nullptr)
	Creates a loop unrolling pass with the provided parameters.
std::unique_ptr< InterfacePass< FunctionOpInterface > >	createLoopUnrollAndJamPass (int unrollJamFactor=-1)
	Creates a loop unroll jam pass to unroll jam by the specified factor.
std::unique_ptr< OperationPass< func::FuncOp > >	createPipelineDataTransferPass ()
	Creates a pass to pipeline explicit movement of data across levels of the memory hierarchy.
std::unique_ptr< Pass >	createAffineExpandIndexOpsPass ()
	Creates a pass to expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect).
std::unique_ptr< Pass >	createAffineExpandIndexOpsAsAffinePass ()
	Creates a pass to expand affine index operations into affine.apply operations.
void	registerAffineDataCopyGeneration ()
void	registerAffineDataCopyGenerationPass ()
void	registerAffineExpandIndexOps ()
void	registerAffineExpandIndexOpsPass ()
void	registerAffineExpandIndexOpsAsAffine ()
void	registerAffineExpandIndexOpsAsAffinePass ()
void	registerAffineLoopFusion ()
void	registerAffineLoopFusionPass ()
void	registerAffineLoopInvariantCodeMotion ()
void	registerAffineLoopInvariantCodeMotionPass ()
void	registerAffineLoopNormalize ()
void	registerAffineLoopNormalizePass ()
void	registerAffineLoopTiling ()
void	registerAffineLoopTilingPass ()
void	registerAffineLoopUnroll ()
void	registerAffineLoopUnrollPass ()
void	registerAffineLoopUnrollAndJam ()
void	registerAffineLoopUnrollAndJamPass ()
void	registerAffineParallelize ()
void	registerAffineParallelizePass ()
void	registerAffinePipelineDataTransfer ()
void	registerAffinePipelineDataTransferPass ()
void	registerAffineScalarReplacement ()
void	registerAffineScalarReplacementPass ()
void	registerAffineVectorize ()
void	registerAffineVectorizePass ()
void	registerLoopCoalescing ()
void	registerLoopCoalescingPass ()
void	registerRaiseMemrefDialect ()
void	registerRaiseMemrefDialectPass ()
void	registerSimplifyAffineMinMaxPass ()
void	registerSimplifyAffineMinMaxPassPass ()
void	registerSimplifyAffineStructures ()
void	registerSimplifyAffineStructuresPass ()
void	registerAffinePasses ()
void	registerTransformDialectExtension (DialectRegistry &registry)
LogicalResult	lowerAffineDelinearizeIndexOp (RewriterBase &rewriter, AffineDelinearizeIndexOp op)
	Lowers affine.delinearize_index into a sequence of division and remainder operations.
LogicalResult	lowerAffineLinearizeIndexOp (RewriterBase &rewriter, AffineLinearizeIndexOp op)
	Lowers affine.linearize_index into a sequence of multiplications and additions.
void	populateAffineExpandIndexOpsPatterns (RewritePatternSet &patterns)
	Populate patterns that expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect).
void	populateAffineExpandIndexOpsAsAffinePatterns (RewritePatternSet &patterns)
	Populate patterns that expand affine index operations into their equivalent affine.apply representations.
void	reorderOperandsByHoistability (RewriterBase &rewriter, AffineApplyOp op)
	Helper function to rewrite op's affine map and reorder its operands such that they are in increasing order of hoistability (i.e.
FailureOr< AffineApplyOp >	decompose (RewriterBase &rewriter, AffineApplyOp op)
	Split an "affine.apply" operation into smaller ops.
FailureOr< OpFoldResult >	reifyValueBound (OpBuilder &b, Location loc, presburger::BoundType type, const ValueBoundsConstraintSet::Variable &var, ValueBoundsConstraintSet::StopConditionFn stopCondition, bool closedUB=false)
	Reify a bound for the given variable in terms of SSA values for which stopCondition is met.
FailureOr< OpFoldResult >	reifyIndexValueBound (OpBuilder &b, Location loc, presburger::BoundType type, Value value, ValueBoundsConstraintSet::StopConditionFn stopCondition=nullptr, bool closedUB=false)
	Reify a bound for the given index-typed value in terms of SSA values for which stopCondition is met.
FailureOr< OpFoldResult >	reifyShapedValueDimBound (OpBuilder &b, Location loc, presburger::BoundType type, Value value, int64_t dim, ValueBoundsConstraintSet::StopConditionFn stopCondition=nullptr, bool closedUB=false)
	Reify a bound for the specified dimension of the given shaped value in terms of SSA values for which stopCondition is met.
OpFoldResult	materializeComputedBound (OpBuilder &b, Location loc, AffineMap boundMap, ArrayRef< std::pair< Value, std::optional< int64_t > > > mapOperands)
	Materialize an already computed bound with Affine dialect ops.
bool	simplifyAffineMinOp (RewriterBase &rewriter, AffineMinOp op)
	This transform tries to simplify the affine min operation op, by finding a common lower bound for a set of expressions in the affine map results.
bool	simplifyAffineMaxOp (RewriterBase &rewriter, AffineMaxOp op)
	This transform tries to simplify the affine max operation op, by finding a common upper bound for a set of expressions in the affine map results.
LogicalResult	simplifyAffineMinMaxOps (RewriterBase &rewriter, ArrayRef< Operation * > ops, bool *modified=nullptr)
	This transform applies simplifyAffineMinOp and simplifyAffineMaxOp to all the affine.min or affine.max operations in ops.
LogicalResult	affineParallelize (AffineForOp forOp, ArrayRef< LoopReduction > parallelReductions={}, AffineParallelOp *resOp=nullptr)
	Replaces a parallel affine.for op with a 1-d affine.parallel op.
template<typename EffectType, typename T>
bool	hasNoInterveningEffect (Operation *start, T memOp, llvm::function_ref< bool(Value, Value)> mayAlias)
	Hoists out affine.if/else to as high as possible, i.e., past all invariant affine.fors/parallel's.
LogicalResult	mergeOffsetsSizesAndStrides (OpBuilder &builder, Location loc, ArrayRef< OpFoldResult > producerOffsets, ArrayRef< OpFoldResult > producerSizes, ArrayRef< OpFoldResult > producerStrides, const llvm::SmallBitVector &droppedProducerDims, ArrayRef< OpFoldResult > consumerOffsets, ArrayRef< OpFoldResult > consumerSizes, ArrayRef< OpFoldResult > consumerStrides, SmallVector< OpFoldResult > &combinedOffsets, SmallVector< OpFoldResult > &combinedSizes, SmallVector< OpFoldResult > &combinedStrides)
	Fills the combinedOffsets, combinedSizes and combinedStrides to use when combining a producer slice into a consumer slice.
LogicalResult	mergeOffsetsSizesAndStrides (OpBuilder &builder, Location loc, OffsetSizeAndStrideOpInterface producer, OffsetSizeAndStrideOpInterface consumer, const llvm::SmallBitVector &droppedProducerDims, SmallVector< OpFoldResult > &combinedOffsets, SmallVector< OpFoldResult > &combinedSizes, SmallVector< OpFoldResult > &combinedStrides)
	Fills the combinedOffsets, combinedSizes and combinedStrides to use when combining a producer slice op into a consumer slice op.
void	resolveIndicesIntoOpWithOffsetsAndStrides (RewriterBase &rewriter, Location loc, ArrayRef< OpFoldResult > mixedSourceOffsets, ArrayRef< OpFoldResult > mixedSourceStrides, const llvm::SmallBitVector &rankReducedDims, ArrayRef< OpFoldResult > consumerIndices, SmallVectorImpl< Value > &resolvedIndices)
	Given the 'consumerIndices' of a load/store operation operating on an op with offsets and strides, return the combined indices.
void	resolveIndicesIntoOpWithOffsetsAndStrides (RewriterBase &rewriter, Location loc, ArrayRef< OpFoldResult > mixedSourceOffsets, ArrayRef< OpFoldResult > mixedSourceStrides, const llvm::SmallBitVector &rankReducedDims, ValueRange consumerIndices, SmallVectorImpl< Value > &resolvedIndices)
void	resolveSizesIntoOpWithSizes (ArrayRef< OpFoldResult > sourceSizes, ArrayRef< OpFoldResult > destSizes, const llvm::SmallBitVector &rankReducedSourceDims, SmallVectorImpl< OpFoldResult > &resolvedSizes)
	Given sourceSizes, destSizes and information about which dimensions are dropped by the source: rankReducedSourceDims, compute the resolved sizes that correspond to dest_op(source_op).

Typedef Documentation

◆ FilterFunctionType

using mlir::affine::FilterFunctionType = std::function<bool(Operation &)>

A NestedPattern is a nested operation walker that:

recursively matches a substructure in the tree;
uses a filter function to refine matches with extra semantic constraints (passed via a lambda of type FilterFunctionType);
TODO: optionally applies actions (lambda).

Nested patterns are meant to capture imperfectly nested loops while matching properties over the whole loop nest. For instance, in vectorization we are interested in capturing all the imperfectly nested loops of a certain type and such that all the load and stores have certain access patterns along the loops' induction variables). Such NestedMatches are first captured using the match function and are later processed to analyze properties and apply transformations in a non-greedy way.

The NestedMatches captured in the IR can grow large, especially after aggressive unrolling. As experience has shown, it is generally better to use a plain walk over operations to match flat patterns but the current implementation is competitive nonetheless.

Definition at line 91 of file NestedMatcher.h.

◆ ReductionLoopMap

using mlir::affine::ReductionLoopMap = DenseMap<Operation *, SmallVector<LoopReduction, 2>>

Definition at line 43 of file Utils.h.

◆ VectorizableLoopFun

using mlir::affine::VectorizableLoopFun = std::function<bool(AffineForOp)>

Definition at line 88 of file LoopAnalysis.h.

Enumeration Type Documentation

◆ FusionMode

enum mlir::affine::FusionMode

Fusion mode to attempt.

The default mode Greedy does both producer-consumer and sibling fusion.

Enumerator
Greedy
ProducerConsumer
Sibling

Definition at line 35 of file Passes.h.

Function Documentation

◆ affineForOpBodySkew()

LogicalResult mlir::affine::affineForOpBodySkew	(	AffineForOp	forOp,
		ArrayRef< uint64_t >	shifts,
		bool	unrollPrologueEpilogue = false )

Skew the operations in an affine.for's body with the specified operation-wise shifts.

The shifts are with respect to the original execution order, and are multiplied by the loop 'step' before being applied. If unrollPrologueEpilogue is set, fully unroll the prologue and epilogue loops when possible.

Definition at line 230 of file LoopUtils.cpp.

References mlir::applyOpPatternsGreedily(), b, mlir::ExistingAndNewOps, generateShiftedLoop(), getConstantTripCount(), isOpwiseShiftValid(), loopUnrollFull(), mlir::patterns, and success().

◆ affineParallelize()

LogicalResult mlir::affine::affineParallelize	(	AffineForOp	forOp,
		ArrayRef< LoopReduction >	parallelReductions = {},
		AffineParallelOp *	resOp = nullptr )

Replaces a parallel affine.for op with a 1-d affine.parallel op.

forOp's body is taken by the affine.parallel op and the former is erased. (mlir::isLoopParallel can be used to detect a parallel affine.for op.) The reductions specified in parallelReductions are also parallelized. Parallelization will fail in the presence of loop iteration arguments that are not listed in parallelReductions. resOp if non-null is set to the newly created affine.parallel op.

Definition at line 352 of file Utils.cpp.

References mlir::Value::getDefiningOp(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::Operation::getOperand(), mlir::Block::getOperations(), mlir::Operation::getResult(), mlir::affine::LoopReduction::kind, mlir::Operation::setOperands(), success(), mlir::affine::LoopReduction::value, and ValueRange.

◆ boundCheckLoadOrStoreOp()

template<typename LoadOrStoreOpPointer>

LogicalResult mlir::affine::boundCheckLoadOrStoreOp	(	LoadOrStoreOpPointer	loadOrStoreOp,
		bool	emitError = true )

Checks a load or store op for an out of bound access; returns failure if the access is out of bounds along any of the dimensions, success otherwise.

Emits a diagnostic error (with location information) if emitError is true.

References b, and mlir::emitError().

◆ buildAffineLoopNest() [1/2]

void mlir::affine::buildAffineLoopNest	(	OpBuilder &	builder,
		Location	loc,
		ArrayRef< int64_t >	lbs,
		ArrayRef< int64_t >	ubs,
		ArrayRef< int64_t >	steps,
		function_ref< void(OpBuilder &, Location, ValueRange)>	bodyBuilderFn = nullptr )

Builds a perfect nest of affine.for loops, i.e., each loop except the innermost one contains only another loop and a terminator.

The loops iterate from "lbs" to "ubs" with "steps". The body of the innermost loop is populated by calling "bodyBuilderFn" and providing it with an OpBuilder, a Location and a list of loop induction variables.

Definition at line 3048 of file AffineOps.cpp.

References buildAffineLoopFromConstants(), and buildAffineLoopNestImpl().

Referenced by mlir::linalg::GenerateLoopNest< LoopTy >::doit().

◆ buildAffineLoopNest() [2/2]

void mlir::affine::buildAffineLoopNest	(	OpBuilder &	builder,
		Location	loc,
		ValueRange	lbs,
		ValueRange	ubs,
		ArrayRef< int64_t >	steps,
		function_ref< void(OpBuilder &, Location, ValueRange)>	bodyBuilderFn = nullptr )

Definition at line 3056 of file AffineOps.cpp.

References buildAffineLoopFromValues(), and buildAffineLoopNestImpl().

◆ buildSliceTripCountMap()

bool mlir::affine::buildSliceTripCountMap	(	const ComputationSliceState &	slice,
		llvm::SmallDenseMap< Operation , uint64_t, 8 >	tripCountMap )

Builds a map 'tripCountMap' from AffineForOp to constant trip count for loop nest surrounding represented by slice loop bounds in 'slice'.

Returns true on success, false otherwise (if a non-constant trip count was encountered).

Definition at line 1818 of file Utils.cpp.

References getConstantTripCount(), getConstDifference(), getForInductionVarOwner(), mlir::AffineMap::getNumResults(), mlir::affine::ComputationSliceState::ivs, mlir::affine::ComputationSliceState::lbs, and mlir::affine::ComputationSliceState::ubs.

Referenced by fuseLoops(), getComputationSliceState(), and getFusionComputeCost().

◆ canFuseLoops()

FusionResult mlir::affine::canFuseLoops	(	AffineForOp	srcForOp,
		AffineForOp	dstForOp,
		unsigned	dstLoopDepth,
		ComputationSliceState *	srcSlice,
		FusionStrategy	fusionStrategy = FusionStrategy::Generic )

Checks the feasibility of fusing the loop nest rooted at 'srcForOp' into the loop nest rooted at 'dstForOp' at 'dstLoopDepth'.

Returns FusionResult 'Success' if fusion of the src/dst loop nests is feasible (i.e. they are in the same block and dependences would not be violated). Otherwise returns a FusionResult explaining why fusion is not feasible. NOTE: This function is not feature complete and should only be used in testing.

Definition at line 248 of file LoopFusionUtils.cpp.

◆ canonicalizeMapAndOperands()

void mlir::affine::canonicalizeMapAndOperands	(	AffineMap *	map,
		SmallVectorImpl< Value > *	operands )

Modifies both map and operands in-place so as to:

drop duplicate operands
drop unused dims and symbols from map
promote valid symbols to symbolic operands in case they appeared as dimensional operands
propagate constant operands and drop them

Definition at line 1759 of file AffineOps.cpp.

References canonicalizeMapOrSetAndOperands().

Referenced by mlir::affine::FlatAffineValueConstraints::addBound(), augmentMapAndBounds(), canonicalizeLoopBounds(), composeAffineMapAndOperands(), composeMultiResultAffineMap(), mlir::affine::AffineValueMap::composeSimplifyAndCanonicalize(), createSubApply(), mlir::affine::AffineValueMap::difference(), fuseLoops(), mlir::affine::MemRefAccess::getAccessMap(), getCleanupLoopLowerBound(), makeCanonicalAffineApplies(), materializeComputedBound(), reorderOperandsByHoistability(), replaceDimOrSym(), and simplifyConstrainedMinMaxOp().

◆ canonicalizeSetAndOperands()

void mlir::affine::canonicalizeSetAndOperands	(	IntegerSet *	set,
		SmallVectorImpl< Value > *	operands )

Canonicalizes an integer set the same way canonicalizeMapAndOperands does for affine maps.

Definition at line 1764 of file AffineOps.cpp.

References canonicalizeMapOrSetAndOperands().

Referenced by mlir::affine::FlatAffineValueConstraints::addAffineIfOpDomain(), and createSeparationCondition().

◆ checkMemrefAccessDependence()

DependenceResult mlir::affine::checkMemrefAccessDependence	(	const MemRefAccess &	srcAccess,
		const MemRefAccess &	dstAccess,
		unsigned	loopDepth,
		FlatAffineValueConstraints *	dependenceConstraints = nullptr,
		SmallVector< DependenceComponent, 2 > *	dependenceComponents = nullptr,
		bool	allowRAR = false )

Definition at line 609 of file AffineAnalysis.cpp.

Referenced by computeSliceUnion(), getDependenceComponents(), getMaxLoopDepth(), hasCyclicDependence(), isTilingValid(), mayDependence(), mayHaveEffect(), and mustReachAtInnermost().

◆ computeProduct()

OpFoldResult mlir::affine::computeProduct	(	Location	loc,
		OpBuilder &	builder,
		ArrayRef< OpFoldResult >	terms )

Return the product of terms, creating an affine.apply if any of them are non-constant values.

If any of terms is nullptr, return nullptr.

◆ computeSliceUnion()

SliceComputationResult mlir::affine::computeSliceUnion	(	ArrayRef< Operation * >	opsA,
		ArrayRef< Operation * >	opsB,
		unsigned	loopDepth,
		unsigned	numCommonLoops,
		bool	isBackwardSlice,
		ComputationSliceState *	sliceUnion )

Computes in 'sliceUnion' the union of all slice bounds computed at 'loopDepth' between all dependent pairs of ops in 'opsA' and 'opsB', and then verifies if it is valid.

The parameter 'numCommonLoops' is the number of loops common to the operations in 'opsA' and 'opsB'. If 'isBackwardSlice' is true, computes slice bounds for loop nest surrounding ops in 'opsA', as a function of IVs and symbols of loop nest surrounding ops in 'opsB' at 'loopDepth'. If 'isBackwardSlice' is false, computes slice bounds for loop nest surrounding ops in 'opsB', as a function of IVs and symbols of loop nest surrounding ops in 'opsA' at 'loopDepth'. Returns 'SliceComputationResult::Success' if union was computed correctly, an appropriate 'failure' otherwise.

Returns 'SliceComputationResult::Success' if union was computed correctly, an appropriate failure otherwise.

Definition at line 1633 of file Utils.cpp.

Referenced by canFuseLoops().

◆ createAffineDataCopyGenerationPass() [1/2]

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineDataCopyGenerationPass ( )

Overload relying on pass options for initialization.

Definition at line 97 of file AffineDataCopyGeneration.cpp.

◆ createAffineDataCopyGenerationPass() [2/2]

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineDataCopyGenerationPass	(	unsigned	slowMemorySpace,
		unsigned	fastMemorySpace,
		unsigned	tagMemorySpace = 0,
		int	minDmaTransferSize = 1024,
		uint64_t	fastMemCapacityBytes = std::numeric_limits<uint64_t>::max() )

Performs packing (or explicit copying) of accessed memref regions into buffers in the specified faster memory space through either pointwise copies or DMA operations.

Generates copies for memref's living in 'slowMemorySpace' into newly created buffers in 'fastMemorySpace', and replaces memory operations to the former by the latter.

Definition at line 89 of file AffineDataCopyGeneration.cpp.

◆ createAffineExpandIndexOpsAsAffinePass()

std::unique_ptr< Pass > mlir::affine::createAffineExpandIndexOpsAsAffinePass ( )

Creates a pass to expand affine index operations into affine.apply operations.

Definition at line 95 of file AffineExpandIndexOpsAsAffine.cpp.

◆ createAffineExpandIndexOpsPass()

std::unique_ptr< Pass > mlir::affine::createAffineExpandIndexOpsPass ( )

Creates a pass to expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect).

Definition at line 235 of file AffineExpandIndexOps.cpp.

◆ createAffineLoopInvariantCodeMotionPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineLoopInvariantCodeMotionPass ( )

Creates a loop invariant code motion pass that hoists loop invariant operations out of affine loops.

Definition at line 205 of file AffineLoopInvariantCodeMotion.cpp.

◆ createAffineLoopNormalizePass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineLoopNormalizePass ( bool promoteSingleIter = false )

Apply normalization transformations to affine loop-like ops.

If promoteSingleIter is true, single iteration loops are promoted (i.e., the loop is replaced by its loop body).

Definition at line 53 of file AffineLoopNormalize.cpp.

◆ createAffineParallelize() [1/2]

std::unique_ptr<::mlir::Pass > mlir::affine::createAffineParallelize ( )

Definition at line 765 of file AffineParallelize.cpp.

◆ createAffineParallelize() [2/2]

std::unique_ptr<::mlir::Pass > mlir::affine::createAffineParallelize ( AffineParallelizeOptions options )

Definition at line 769 of file AffineParallelize.cpp.

◆ createAffineParallelizePass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineParallelizePass ( )

Creates a pass to convert all parallel affine.for's into 1-d affine.parallel ops.

References Greedy.

◆ createAffineScalarReplacementPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createAffineScalarReplacementPass ( )

Creates a pass to replace affine memref accesses by scalars using store to load forwarding and redundant load elimination; consequently also eliminate dead allocs.

Definition at line 43 of file AffineScalarReplacement.cpp.

◆ createAffineVectorize() [1/2]

std::unique_ptr<::mlir::Pass > mlir::affine::createAffineVectorize ( )

Definition at line 990 of file SuperVectorize.cpp.

◆ createAffineVectorize() [2/2]

std::unique_ptr<::mlir::Pass > mlir::affine::createAffineVectorize ( AffineVectorizeOptions options )

Definition at line 994 of file SuperVectorize.cpp.

◆ createLoopCoalescingPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createLoopCoalescingPass ( )

Creates a pass that transforms perfectly nested loops with independent bounds into a single loop.

Definition at line 48 of file LoopCoalescing.cpp.

◆ createLoopFusionPass()

std::unique_ptr< Pass > mlir::affine::createLoopFusionPass	(	unsigned	fastMemorySpace = 0,
		uint64_t	localBufSizeThreshold = 0,
		bool	maximalFusion = false,
		enum FusionMode	fusionMode = FusionMode::Greedy )

Creates a loop fusion pass which fuses affine loop nests at the top-level of the operation the pass is created on according to the type of fusion specified in fusionMode.

Buffers of size less than or equal to localBufSizeThreshold are promoted to memory space fastMemorySpace.

Definition at line 1556 of file LoopFusion.cpp.

◆ createLoopTilingPass() [1/2]

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createLoopTilingPass ( )

Overload relying on pass options for initialization.

Definition at line 83 of file LoopTiling.cpp.

◆ createLoopTilingPass() [2/2]

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createLoopTilingPass ( uint64_t cacheSizeBytes )

Creates a pass to perform tiling on loop nests.

Creates a pass to perform loop tiling on all suitable loop nests of a Function.

Definition at line 79 of file LoopTiling.cpp.

◆ createLoopUnrollAndJamPass()

std::unique_ptr< InterfacePass< FunctionOpInterface > > mlir::affine::createLoopUnrollAndJamPass ( int unrollJamFactor = -1 )

Creates a loop unroll jam pass to unroll jam by the specified factor.

A factor of -1 lets the pass use the default factor or the one on the command line if provided.

Definition at line 73 of file LoopUnrollAndJam.cpp.

◆ createLoopUnrollPass()

std::unique_ptr< InterfacePass< FunctionOpInterface > > mlir::affine::createLoopUnrollPass	(	int	unrollFactor = -1,
		bool	unrollUpToFactor = false,
		const std::function< unsigned(AffineForOp)> &	getUnrollFactor = nullptr )

Creates a loop unrolling pass with the provided parameters.

'getUnrollFactor' is a function callback for clients to supply a function that computes an unroll factor - the callback takes precedence over unroll factors supplied through other means. If -1 is passed as the unrollFactor and no callback is provided, anything passed from the command-line (if at all) or the default unroll factor is used (LoopUnroll:kDefaultUnrollFactor).

Definition at line 146 of file LoopUnroll.cpp.

◆ createPipelineDataTransferPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createPipelineDataTransferPass ( )

Creates a pass to pipeline explicit movement of data across levels of the memory hierarchy.

Definition at line 54 of file PipelineDataTransfer.cpp.

◆ createRaiseMemrefToAffine()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createRaiseMemrefToAffine ( )

Creates a pass that converts some memref operators to affine operators.

Definition at line 185 of file RaiseMemrefDialect.cpp.

◆ createSimplifyAffineMinMaxPass()

std::unique_ptr<::mlir::Pass > mlir::affine::createSimplifyAffineMinMaxPass ( )

We declare an explicit private instantiation because Pass classes should only be visible by the current library.

Definition at line 1194 of file SimplifyAffineMinMax.cpp.

◆ createSimplifyAffineStructuresPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createSimplifyAffineStructuresPass ( )

Creates a simplification pass for affine structures (maps and sets).

In addition, this pass also normalizes memrefs to have the trivial (identity) layout map.

Definition at line 85 of file SimplifyAffineStructures.cpp.

◆ decompose()

FailureOr< AffineApplyOp > mlir::affine::decompose	(	RewriterBase &	rewriter,
		AffineApplyOp	op )

Split an "affine.apply" operation into smaller ops.

This reassociates a large AffineApplyOp into an ordered list of smaller AffineApplyOps. This can be used right before lowering affine ops to arith to exhibit more opportunities for CSE and LICM. Return the sink AffineApplyOp on success or failure if op does not decompose into smaller AffineApplyOps. Note that this can be undone by canonicalization which tries to maximally compose chains of AffineApplyOps.

Definition at line 95 of file DecomposeAffineOps.cpp.

References mlir::Add, createSubApply(), mlir::AffineMap::get(), mlir::getAffineBinaryOpExpr(), mlir::getAffineSymbolExpr(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getNumSymbols(), mlir::AffineMap::getResult(), mlir::Mul, mlir::RewriterBase::notifyMatchFailure(), and mlir::RewriterBase::replaceOp().

◆ defaultFilterFunction()

bool mlir::affine::defaultFilterFunction ( Operation & )

inline

Definition at line 92 of file NestedMatcher.h.

◆ extractForInductionVars()

void mlir::affine::extractForInductionVars	(	ArrayRef< AffineForOp >	forInsts,
		SmallVectorImpl< Value > *	ivs )

Extracts the induction variables from a list of AffineForOps and places them in the output argument ivs.

Extracts the induction variables from a list of AffineForOps and returns them.

Definition at line 2962 of file AffineOps.cpp.

Referenced by getFullMemRefAsRegion(), tilePerfectlyNested(), and tilePerfectlyNestedParametric().

◆ extractInductionVars()

void mlir::affine::extractInductionVars	(	ArrayRef< Operation * >	affineOps,
		SmallVectorImpl< Value > &	ivs )

Extracts the induction variables from a list of either AffineForOp or AffineParallelOp and places them in the output argument ivs.

Definition at line 2969 of file AffineOps.cpp.

◆ findInnermostCommonBlockInScope()

Block * mlir::affine::findInnermostCommonBlockInScope	(	mlir::Operation *	a,
		mlir::Operation *	b )

Find the innermost common Block of a and b in the affine scope that a and b are part of.

Return nullptr if they belong to different affine scopes. Also, return nullptr if they do not have a common Block ancestor (for eg., when they are part of the then and else regions of an op that itself starts an affine scope.

Definition at line 2391 of file Utils.cpp.

References b, getAffineAnalysisScope(), mlir::Operation::getBlock(), mlir::Operation::getParentOp(), and mlir::Operation::getParentRegion().

◆ fullyComposeAffineMapAndOperands()

void mlir::affine::fullyComposeAffineMapAndOperands	(	AffineMap *	map,
		SmallVectorImpl< Value > *	operands,
		bool	composeAffineMin = false )

Given an affine map map and its input operands, this method composes into map, maps of AffineApplyOps whose results are the values in operands, iteratively until no more of operands are the result of an AffineApplyOp.

When this function returns, map becomes the composed affine map, and each Value in operands is guaranteed to be either a loop IV or a terminal symbol, i.e., a symbol defined at the top level or a block/function argument.

Definition at line 1400 of file AffineOps.cpp.

References composeAffineMapAndOperands(), and mlir::Value::getDefiningOp().

Referenced by mlir::affine::FlatAffineValueConstraints::addBound(), composeMultiResultAffineMap(), mlir::affine::AffineValueMap::composeSimplifyAndCanonicalize(), mlir::affine::AffineValueMap::difference(), fullyComposeAndComputeConstantDelta(), generateCopy(), mlir::affine::MemRefAccess::getAccessMap(), and getCleanupLoopLowerBound().

◆ fullyComposeAndComputeConstantDelta()

FailureOr< int64_t > mlir::affine::fullyComposeAndComputeConstantDelta	(	Value	value1,
		Value	value2 )

Compute a constant delta of the given two values.

Return "failure" if we cannot determine a constant delta. value1/value2 must be index-typed.

This function is similar to ValueBoundsConstraintSet::computeConstantDistance. To work around limitations in FlatLinearConstraints, this function fully composes value1 and value2 (if they are the result of affine.apply ops) before populating the constraint set. The folding/composing logic can see opportunities for simplifications that the constraint set implementation cannot see.

Definition at line 169 of file ValueBoundsOpInterfaceImpl.cpp.

References b, mlir::ValueBoundsConstraintSet::computeConstantBound(), mlir::presburger::EQ, fullyComposeAffineMapAndOperands(), mlir::AffineMap::get(), mlir::Value::getContext(), mlir::Value::getType(), and mlir::Type::isIndex().

Referenced by mlir::vector::isDisjointTransferIndices().

◆ fuseLoops()

void mlir::affine::fuseLoops	(	AffineForOp	srcForOp,
		AffineForOp	dstForOp,
		const ComputationSliceState &	srcSlice,
		bool	isInnermostSiblingInsertionFusion = false )

Fuses 'srcForOp' into 'dstForOp' with destination loop block insertion point and source slice loop bounds specified in 'srcSlice'.

isInnermostSiblingInsertionFusion enables cleanup of ‘srcForOp that is a single-iteration reduction loop being sibling-fused into a 'dstForOp’.

Definition at line 425 of file LoopFusionUtils.cpp.

References b, buildSliceTripCountMap(), canonicalizeMapAndOperands(), getForInductionVarOwner(), getSliceIterationCount(), mlir::affine::ComputationSliceState::insertPoint, isLoopParallelAndContainsReduction(), mlir::affine::ComputationSliceState::ivs, mlir::affine::ComputationSliceState::lbOperands, mlir::affine::ComputationSliceState::lbs, mlir::IRMapping::lookupOrNull(), promoteIfSingleIteration(), promoteSingleIterReductionLoop(), mlir::affine::ComputationSliceState::ubOperands, and mlir::affine::ComputationSliceState::ubs.

◆ gatherProducerConsumerMemrefs()

void mlir::affine::gatherProducerConsumerMemrefs	(	ArrayRef< Operation * >	srcOps,
		ArrayRef< Operation * >	dstOps,
		DenseSet< Value > &	producerConsumerMemrefs )

Returns in 'producerConsumerMemrefs' the memrefs involved in a producer-consumer dependence between write ops in 'srcOps' and read ops in 'dstOps'.

Definition at line 639 of file LoopFusionUtils.cpp.

Referenced by gatherProducerConsumerMemrefs(), and getMaxLoopDepth().

◆ getAffineAnalysisScope()

Region * mlir::affine::getAffineAnalysisScope ( Operation * op )

Returns the closest region enclosing op that is held by a non-affine operation; nullptr if there is no such region.

This method is meant to be used by affine analysis methods (e.g. dependence analysis) which are only meaningful when performed among/between operations from the same analysis scope.

Definition at line 275 of file AffineOps.cpp.

References mlir::Operation::getParentOp(), and mlir::Operation::getParentRegion().

Referenced by checkMemrefAccessDependence(), findInnermostCommonBlockInScope(), mayHaveEffect(), and mustReachAtInnermost().

◆ getAffineForIVs()

void mlir::affine::getAffineForIVs	(	Operation &	op,
		SmallVectorImpl< AffineForOp > *	loops )

Populates 'loops' with IVs of the affine.for ops surrounding 'op' ordered from the outermost 'affine.for' operation to the innermost one while not traversing outside of the surrounding affine scope.

Definition at line 851 of file Utils.cpp.

References mlir::Operation::getParentOp().

Referenced by getComputationSliceState(), getFullMemRefAsRegion(), getFusionComputeCost(), getInnermostCommonLoopDepth(), getLastDependentOpInRange(), mlir::affine::MemRefDependenceGraph::init(), insertBackwardComputationSlice(), and isOpLoopInvariant().

◆ getAffineIVs()

void mlir::affine::getAffineIVs	(	Operation &	op,
		SmallVectorImpl< Value > &	ivs )

Populates 'ivs' with IVs of the surrounding affine.for and affine.parallel ops ordered from the outermost one to the innermost while not traversing outside of the surrounding affine scope.

Definition at line 2091 of file Utils.cpp.

References mlir::Operation::getParentOp().

Referenced by mlir::affine::MemRefRegion::compute(), and getNumCommonSurroundingLoops().

◆ getAffineParallelInductionVarOwner()

AffineParallelOp mlir::affine::getAffineParallelInductionVarOwner ( Value val )

Returns true if the provided value is among the induction variables of an AffineParallelOp.

Definition at line 2949 of file AffineOps.cpp.

References mlir::Operation::getParentOp().

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), mlir::affine::MemRefRegion::compute(), and isAffineParallelInductionVar().

◆ getAffineScope()

Region * mlir::affine::getAffineScope ( Operation * op )

Returns the closest region enclosing op that is held by an operation with trait AffineScope; nullptr if there is no such region.

Definition at line 265 of file AffineOps.cpp.

Referenced by isValidDim(), isValidSymbol(), verifyDimAndSymbolIdentifiers(), mlir::affine::AffineDmaStartOp::verifyInvariantsImpl(), and mlir::affine::AffineDmaWaitOp::verifyInvariantsImpl().

◆ getComputationSliceState()

void mlir::affine::getComputationSliceState	(	Operation *	depSourceOp,
		Operation *	depSinkOp,
		const FlatAffineValueConstraints &	dependenceConstraints,
		unsigned	loopDepth,
		bool	isBackwardSlice,
		ComputationSliceState *	sliceState )

Computes the computation slice loop bounds for one loop nest as affine maps of the other loop nest's IVs and symbols, using 'dependenceConstraints' computed between 'depSourceAccess' and 'depSinkAccess'.

If 'isBackwardSlice' is true, a backwards slice is computed in which the slice bounds of loop nest surrounding 'depSourceAccess' are computed in terms of loop IVs and symbols of the loop nest surrounding 'depSinkAccess' at 'loopDepth'. If 'isBackwardSlice' is false, a forward slice is computed in which the slice bounds of loop nest surrounding 'depSinkAccess' are computed in terms of loop IVs and symbols of the loop nest surrounding 'depSourceAccess' at 'loopDepth'. The slice loop bounds and associated operands are returned in 'sliceState'.

Definition at line 1870 of file Utils.cpp.

References buildSliceTripCountMap(), getAffineForIVs(), mlir::Operation::getContext(), mlir::presburger::IntegerRelation::getNumDimAndSymbolVars(), getSequentialLoops(), mlir::FlatLinearConstraints::getSliceBounds(), getSliceIterationCount(), mlir::FlatLinearValueConstraints::getValue(), mlir::FlatLinearValueConstraints::getValues(), mlir::affine::ComputationSliceState::insertPoint, isLoopParallelAndContainsReduction(), mlir::affine::ComputationSliceState::isMaximal(), mlir::affine::ComputationSliceState::ivs, kSliceFusionBarrierAttrName, mlir::affine::ComputationSliceState::lbOperands, mlir::affine::ComputationSliceState::lbs, mlir::FlatLinearValueConstraints::projectOut(), mlir::affine::ComputationSliceState::ubOperands, and mlir::affine::ComputationSliceState::ubs.

Referenced by computeSliceUnion().

◆ getComputeCost()

int64_t mlir::affine::getComputeCost	(	AffineForOp	forOp,
		LoopNestStats &	stats )

Computes the total cost of the loop nest rooted at 'forOp' using 'stats'.

Currently, the total cost is computed by counting the total operation instance count (i.e. total number of operations in the loop body * loop trip count) for the entire loop nest.

Definition at line 563 of file LoopFusionUtils.cpp.

References getComputeCostHelper().

Referenced by isFusionProfitable().

◆ getConstantTripCount()

std::optional< uint64_t > mlir::affine::getConstantTripCount ( AffineForOp forOp )

Returns the trip count of the loop if it's a constant, std::nullopt otherwise.

This uses affine expression analysis and is able to determine constant trip count in non-trivial cases.

This method uses affine expression analysis (in turn using getTripCount) and is able to determine constant trip count in non-trivial cases.

Definition at line 216 of file LoopAnalysis.cpp.

References mlir::AffineMap::getResults(), and getTripCountMapAndOperands().

Referenced by adjustToDivisorsOfTripCounts(), affineForOpBodySkew(), buildSliceTripCountMap(), constructTiledIndexSetHyperRect(), getLoopNestStats(), loopUnrollByFactor(), loopUnrollFull(), loopUnrollJamByFactor(), loopUnrollJamUpToFactor(), loopUnrollUpToFactor(), promoteIfSingleIteration(), and promoteSingleIterReductionLoop().

◆ getDependenceComponents()

void mlir::affine::getDependenceComponents	(	AffineForOp	forOp,
		unsigned	maxLoopDepth,
		std::vector< SmallVector< DependenceComponent, 2 > > *	depCompsVec )

Returns in 'depCompsVec', dependence components for dependences between all load and store ops in loop nest rooted at 'forOp', at loop depths in range [1, maxLoopDepth].

Gathers dependence components for dependences between all ops in loop nest rooted at 'forOp' at loop depths in range [1, maxLoopDepth].

Definition at line 697 of file AffineAnalysis.cpp.

References checkMemrefAccessDependence(), hasDependence(), and result.

Referenced by isValidLoopInterchangePermutation(), and sinkSequentialLoops().

◆ getEnclosingAffineOps()

void mlir::affine::getEnclosingAffineOps	(	Operation &	op,
		SmallVectorImpl< Operation * > *	ops )

Populates 'ops' with affine operations enclosing op ordered from outermost to innermost while stopping at the boundary of the affine scope.

affine.for, affine.if, or affine.parallel ops comprise such surrounding affine ops. ops is guaranteed by design to have a successive chain of affine parent ops.

Definition at line 865 of file Utils.cpp.

References mlir::Operation::getParentOp(), and mlir::Operation::hasTrait().

Referenced by findHighestBlockForPlacement(), and getOpIndexSet().

◆ getForInductionVarOwner()

AffineForOp mlir::affine::getForInductionVarOwner ( Value val )

Returns the loop parent of an induction variable.

If the provided value is not an induction variable, then return nullptr.

Definition at line 2938 of file AffineOps.cpp.

Referenced by mlir::affine::FlatAffineValueConstraints::addDomainFromSliceMaps(), mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), addMissingLoopIVBounds(), mlir::affine::FlatAffineValueConstraints::addSliceBounds(), buildSliceTripCountMap(), mlir::affine::MemRefRegion::compute(), mlir::affine::FlatAffineValueConstraints::convertLoopIVSymbolsToDims(), fuseLoops(), mlir::affine::ComputationSliceState::getAsConstraints(), getLargestKnownDivisor(), getLowerBound(), getNumCommonLoops(), mlir::affine::ComputationSliceState::getSourceAsConstraints(), getUpperBound(), isAffineForInductionVar(), mlir::affine::ComputationSliceState::isMaximal(), isNonNegativeBoundedBy(), and isUniformDefinition().

◆ getFusionComputeCost()

bool mlir::affine::getFusionComputeCost	(	AffineForOp	srcForOp,
		LoopNestStats &	srcStats,
		AffineForOp	dstForOp,
		LoopNestStats &	dstStats,
		const ComputationSliceState &	slice,
		int64_t *	computeCost )

Computes and returns in 'computeCost', the total compute cost of fusing the 'slice' of the loop nest rooted at 'srcForOp' into 'dstForOp'.

Currently, the total cost is computed by counting the total operation instance count (i.e. total number of operations in the loop body * loop trip count) for the entire loop nest. Returns true on success, failure otherwise (e.g. non-constant trip counts).

Currently, the total cost is computed by counting the total operation instance count (i.e. total number of operations in the loop body * loop trip count) for the entire loop nest.

Definition at line 574 of file LoopFusionUtils.cpp.

References buildSliceTripCountMap(), getAffineForIVs(), getComputeCostHelper(), getSliceIterationCount(), and mlir::affine::ComputationSliceState::insertPoint.

◆ getIndexSet()

LogicalResult mlir::affine::getIndexSet	(	MutableArrayRef< Operation * >	ops,
		FlatAffineValueConstraints *	domain )

Builds a system of constraints with dimensional variables corresponding to the loop IVs of the forOps and AffineIfOp's operands appearing in that order.

Bounds of the loop are used to add appropriate inequalities. Constraints from the index sets of AffineIfOp are also added. Any symbols founds in the bound operands are added as symbols in the system. Returns failure for the yet unimplemented cases. ops accepts both AffineForOp and AffineIfOp.

Referenced by checkIfHyperRectangular(), createFullTiles(), createSeparationCondition(), and getOpIndexSet().

◆ getInnermostCommonLoopDepth()

unsigned mlir::affine::getInnermostCommonLoopDepth	(	ArrayRef< Operation * >	ops,
		SmallVectorImpl< AffineForOp > *	surroundingLoops = nullptr )

Returns the innermost common loop depth for the set of operations in 'ops'.

Definition at line 1601 of file Utils.cpp.

References getAffineForIVs().

Referenced by computeSliceUnion(), and getMaxLoopDepth().

◆ getIntOrFloatMemRefSizeInBytes()

std::optional< uint64_t > mlir::affine::getIntOrFloatMemRefSizeInBytes ( MemRefType memRefType )

Returns the size of a memref with element type int or float in bytes if it's statically shaped, std::nullopt otherwise.

Returns the size of memref data in bytes if it's statically shaped, std::nullopt otherwise.

If the element of the memref has vector type, takes into account size of the vector as well.

Definition at line 1463 of file Utils.cpp.

References getMemRefIntOrFloatEltSizeInBytes().

Referenced by generateCopy().

◆ getInvariantAccesses()

DenseSet< Value > mlir::affine::getInvariantAccesses	(	Value	iv,
		ArrayRef< Value >	indices )

Given an induction variable iv of type AffineForOp and indices of type IndexType, returns the set of indices that are independent of iv.

Prerequisites (inherited from isAccessInvariant above):

iv and indices of the proper type;
at most one affine.apply is reachable from each index in indices;

Emits a note if it encounters a chain of affine.apply and conservatively those cases.

Definition at line 308 of file LoopAnalysis.cpp.

References indices, and isAccessIndexInvariant().

Referenced by isIVMappedToMultipleIndices(), and makePermutationMap().

◆ getLargestDivisorOfTripCount()

uint64_t mlir::affine::getLargestDivisorOfTripCount ( AffineForOp forOp )

Returns the greatest known integral divisor of the trip count.

Affine expression analysis is used (indirectly through getTripCount), and this method is thus able to determine non-trivial divisors.

Definition at line 243 of file LoopAnalysis.cpp.

References mlir::AffineMap::getNumResults(), mlir::AffineMap::getResults(), and getTripCountMapAndOperands().

Referenced by constructTiledIndexSetHyperRect(), loopUnrollByFactor(), and loopUnrollJamByFactor().

◆ getLoopNestStats()

bool mlir::affine::getLoopNestStats	(	AffineForOp	forOpRoot,
		LoopNestStats *	stats )

Collect loop nest statistics (eg.

loop trip count and operation count) in 'stats' for loop nest rooted at 'forOp'. Returns true on success, returns false otherwise.

Definition at line 474 of file LoopFusionUtils.cpp.

References mlir::WalkResult::advance(), getConstantTripCount(), mlir::WalkResult::interrupt(), mlir::affine::LoopNestStats::loopMap, mlir::affine::LoopNestStats::opCountMap, and mlir::affine::LoopNestStats::tripCountMap.

Referenced by isFusionProfitable().

◆ getMemoryFootprintBytes()

std::optional< int64_t > mlir::affine::getMemoryFootprintBytes	(	AffineForOp	forOp,
		int	memorySpace = -1 )

Gets the memory footprint of all data touched in the specified memory space in bytes; if the memory space is unspecified, considers all memory spaces.

Definition at line 2173 of file Utils.cpp.

Referenced by isFusionProfitable().

◆ getMemRefIntOrFloatEltSizeInBytes()

std::optional< int64_t > mlir::affine::getMemRefIntOrFloatEltSizeInBytes ( MemRefType memRefType )

Returns the memref's element type's size in bytes where the elemental type is an int or float or a vector of such types.

Definition at line 1419 of file Utils.cpp.

Referenced by createPrivateMemRef(), getIntOrFloatMemRefSizeInBytes(), and mlir::affine::MemRefRegion::getRegionSize().

◆ getNestingDepth()

unsigned mlir::affine::getNestingDepth ( Operation * op )

Returns the nesting depth of this operation, i.e., the number of loops surrounding this operation.

Returns the nesting depth of this statement, i.e., the number of loops surrounding this statement.

Definition at line 2063 of file Utils.cpp.

References mlir::Operation::getParentOp().

Referenced by computeSliceUnion(), getMemoryFootprintBytes(), hasCyclicDependence(), and mayDependence().

◆ getNumCommonSurroundingLoops()

unsigned mlir::affine::getNumCommonSurroundingLoops	(	Operation &	a,
		Operation &	b )

Returns the number of surrounding loops common to both A and B.

Returns the number of surrounding loops common to 'loopsA' and 'loopsB', where each lists loops from outer-most to inner-most in loop nest.

Definition at line 2108 of file Utils.cpp.

References b, and getAffineIVs().

Referenced by canFuseLoops(), getMaxLoopDepth(), hasCyclicDependence(), hasNoInterveningEffect(), mayHaveEffect(), and mustReachAtInnermost().

◆ getPerfectlyNestedLoops()

void mlir::affine::getPerfectlyNestedLoops	(	SmallVectorImpl< AffineForOp > &	nestedLoops,
		AffineForOp	root )

Get perfectly nested sequence of loops starting at root of loop nest (the first op being another AffineFor, and the second op - a terminator).

A loop is perfectly nested iff: the first op in the loop's body is another AffineForOp, and the second op is a terminator).

Definition at line 856 of file LoopUtils.cpp.

References mlir::Block::begin(), mlir::Block::end(), and mlir::Block::front().

Referenced by getTopLevelTileableBands(), and sinkSequentialLoops().

◆ getReachableAffineApplyOps()

void mlir::affine::getReachableAffineApplyOps	(	ArrayRef< Value >	operands,
		SmallVectorImpl< Operation * > &	affineApplyOps )

Returns in affineApplyOps, the sequence of those AffineApplyOp Operations that are reachable via a search starting from operands and ending at those operands that are not the result of an AffineApplyOp.

◆ getRelationFromMap() [1/2]

LogicalResult mlir::affine::getRelationFromMap	(	AffineMap &	map,
		presburger::IntegerRelation &	rel )

Builds a relation from the given AffineMap/AffineValueMap map, containing all pairs of the form operands -> result that satisfy map.

rel is set to the relation built. For example, give the AffineMap:

(d0, d1)[s0] -> (d0 + s0, d0 - s0)

the resulting relation formed is:

(d0, d1) -> (r1, r2) [d0 d1 r1 r2 s0 const] 1 0 -1 0 1 0 = 0 0 1 0 -1 -1 0 = 0

For AffineValueMap, the domain and symbols have Value set corresponding to the Value in map. Returns failure if the AffineMap could not be flattened (i.e., semi-affine is not yet handled).

Definition at line 498 of file AffineStructures.cpp.

References mlir::presburger::IntegerRelation::addEquality(), mlir::FlatLinearValueConstraints::appendDimVar(), mlir::getFlattenedAffineExprs(), mlir::presburger::IntegerRelation::getNumCols(), mlir::presburger::IntegerRelation::getNumDimAndSymbolVars(), mlir::AffineMap::getNumDims(), mlir::presburger::IntegerRelation::getNumDimVars(), mlir::AffineMap::getNumResults(), mlir::FlatLinearValueConstraints::setValue(), and success().

Referenced by mlir::affine::MemRefAccess::getAccessRelation(), and getRelationFromMap().

◆ getRelationFromMap() [2/2]

LogicalResult mlir::affine::getRelationFromMap	(	const AffineValueMap &	map,
		presburger::IntegerRelation &	rel )

Definition at line 538 of file AffineStructures.cpp.

References mlir::affine::AffineValueMap::getAffineMap(), mlir::AffineMap::getNumDims(), mlir::presburger::IntegerRelation::getNumDimVars(), mlir::AffineMap::getNumResults(), mlir::presburger::IntegerRelation::getNumSymbolVars(), mlir::affine::AffineValueMap::getOperand(), getRelationFromMap(), mlir::presburger::IntegerRelation::setId(), and success().

◆ getSequentialLoops()

void mlir::affine::getSequentialLoops	(	AffineForOp	forOp,
		llvm::SmallDenseSet< Value, 8 > *	sequentialLoops )

Returns in 'sequentialLoops' all sequential loops in loop nest rooted at 'forOp'.

Definition at line 2191 of file Utils.cpp.

Referenced by getComputationSliceState().

◆ getSliceIterationCount()

uint64_t mlir::affine::getSliceIterationCount ( const llvm::SmallDenseMap< Operation *, uint64_t, 8 > & sliceTripCountMap )

Return the number of iterations for the slicetripCountMap provided.

Definition at line 1856 of file Utils.cpp.

Referenced by fuseLoops(), getComputationSliceState(), and getFusionComputeCost().

◆ getSupportedReductions()

void mlir::affine::getSupportedReductions	(	AffineForOp	forOp,
		SmallVectorImpl< LoopReduction > &	supportedReductions )

Populate supportedReductions with descriptors of the supported reductions.

Definition at line 83 of file AffineAnalysis.cpp.

References getSupportedReduction().

Referenced by loopUnrollJamByFactor().

◆ getTripCountMapAndOperands()

void mlir::affine::getTripCountMapAndOperands	(	AffineForOp	forOp,
		AffineMap *	tripCountMap,
		SmallVectorImpl< Value > *	tripCountOperands )

Returns the trip count of the loop as an affine map with its corresponding operands if the latter is expressible as an affine expression, and nullptr otherwise.

Returns the trip count of the loop as an affine expression if the latter is expressible as an affine expression, and nullptr otherwise.

This method always succeeds as long as the lower bound is not a multi-result map. The trip count expression is simplified before returning. This method only utilizes map composition to construct lower and upper bounds before computing the trip count expressions

The trip count expression is simplified before returning. This method only utilizes map composition to construct lower and upper bounds before computing the trip count expressions.

Definition at line 166 of file LoopAnalysis.cpp.

References mlir::AffineExpr::ceilDiv(), mlir::affine::AffineValueMap::difference(), mlir::AffineMap::get(), mlir::affine::AffineValueMap::getAffineMap(), mlir::AffineMap::getConstantMap(), mlir::affine::AffineValueMap::getNumResults(), mlir::affine::AffineValueMap::getOperands(), mlir::affine::AffineValueMap::getResult(), and mlir::affine::AffineValueMap::setResult().

Referenced by getCleanupLoopLowerBound(), getConstantTripCount(), and getLargestDivisorOfTripCount().

◆ hasCyclicDependence()

bool mlir::affine::hasCyclicDependence ( AffineForOp root )

Returns true if the affine nest rooted at root has a cyclic dependence among its affine memory accesses.

The dependence could be through any dependences carried by loops contained in root (inclusive of root) and those carried by loop bodies (blocks) contained. Dependences carried by loops outer to root aren't relevant. This method doesn't consider/account for aliases.

Definition at line 575 of file LoopAnalysis.cpp.

References checkMemrefAccessDependence(), getNestingDepth(), getNumCommonSurroundingLoops(), and noDependence().

◆ hasDependence()

bool mlir::affine::hasDependence ( DependenceResult result )

inline

Utility function that returns true if the provided DependenceResult corresponds to a dependence result.

Definition at line 181 of file AffineAnalysis.h.

References mlir::affine::DependenceResult::HasDependence, and result.

Referenced by getDependenceComponents(), getMaxLoopDepth(), isTilingValid(), and mustReachAtInnermost().

◆ hasNoInterveningEffect()

template<typename EffectType, typename T>

bool mlir::affine::hasNoInterveningEffect	(	Operation *	start,
		T	memOp,
		llvm::function_ref< bool(Value, Value)>	mayAlias )

Hoists out affine.if/else to as high as possible, i.e., past all invariant affine.fors/parallel's.

Returns success if any hoisting happened; folded` is / set to true if the op was folded or erased. This hoisting could lead to / significant code expansion in some cases. LogicalResult hoistAffineIfOp(AffineIfOp ifOp, bool *folded = nullptr);

/ Holds parameters to perform n-D vectorization on a single loop nest. / For example, for the following loop nest: / / func @vec2d(in: memref<64x128x512xf32>, out: memref<64x128x512xf32>) { / affine.for i0 = 0 to 64 { / affine.for i1 = 0 to 128 { / affine.for i2 = 0 to 512 { / ld = affine.load in[i0, i1, i2] : memref<64x128x512xf32> / affine.store ld, out[i0, i1, i2] : memref<64x128x512xf32> / } / } / } / return / } / / and VectorizationStrategy = 'vectorSizes = {8, 4}', 'loopToVectorDim = / {{i1->0}, {i2->1}}', SuperVectorizer will generate: / / func @vec2d(arg0: memref<64x128x512xf32>, arg1: memref<64x128x512xf32>) { / affine.for arg2 = 0 to 64 { / affine.for arg3 = 0 to 128 step 8 { / affine.for arg4 = 0 to 512 step 4 { / cst = arith.constant 0.000000e+00 : f32 / %0 = vector.transfer_read arg0[arg2, arg3, arg4], cst : ... / vector.transfer_write %0, arg1[arg2, arg3, arg4] : ... / } / } / } / return / } TODO: Hoist to a VectorizationStrategy.cpp when appropriate. struct VectorizationStrategy { Vectorization factors to apply to each target vector dimension. Each factor will be applied to a different loop. SmallVector<int64_t, 8> vectorSizes; Maps each AffineForOp vectorization candidate with its vector dimension. The candidate will be vectorized using the vectorization factor in 'vectorSizes' for that dimension. DenseMap<Operation *, unsigned> loopToVectorDim; Maps loops that implement vectorizable reductions to the corresponding reduction descriptors. ReductionLoopMap reductionLoops; };

/ Replace affine store and load accesses by scalars by forwarding stores to / loads and eliminate invariant affine loads; consequently, eliminate dead / allocs. void affineScalarReplace(func::FuncOp f, DominanceInfo &domInfo, PostDominanceInfo &postDomInfo, AliasAnalysis &analysis);

/ Vectorizes affine loops in 'loops' using the n-D vectorization factors in / 'vectorSizes'. By default, each vectorization factor is applied / inner-to-outer to the loops of each loop nest. 'fastestVaryingPattern' can / be optionally used to provide a different loop vectorization order. / If reductionLoops is not empty, the given reduction loops may be / vectorized along the reduction dimension. / TODO: Vectorizing reductions is supported only for 1-D vectorization. void vectorizeAffineLoops( Operation *parentOp, llvm::DenseSet<Operation *, DenseMapInfo<Operation *>> &loops, ArrayRef<int64_t> vectorSizes, ArrayRef<int64_t> fastestVaryingPattern, const ReductionLoopMap &reductionLoops = ReductionLoopMap());

/ External utility to vectorize affine loops from a single loop nest using an / n-D vectorization strategy (see doc in VectorizationStrategy definition). / Loops are provided in a 2D vector container. The first dimension represents / the nesting level relative to the loops to be vectorized. The second / dimension contains the loops. This means that: / a) every loop in 'loops[i]' must have a parent loop in 'loops[i-1]', / b) a loop in 'loops[i]' may or may not have a child loop in 'loops[i+1]'. / / For example, for the following loop nest: / / func @vec2d(in0: memref<64x128x512xf32>, in1: memref<64x128x128xf32>, / out0: memref<64x128x512xf32>, / out1: memref<64x128x128xf32>) { / affine.for i0 = 0 to 64 { / affine.for i1 = 0 to 128 { / affine.for i2 = 0 to 512 { / ld = affine.load in0[i0, i1, i2] : memref<64x128x512xf32> / affine.store ld, out0[i0, i1, i2] : memref<64x128x512xf32> / } / affine.for i3 = 0 to 128 { / ld = affine.load in1[i0, i1, i3] : memref<64x128x128xf32> / affine.store ld, out1[i0, i1, i3] : memref<64x128x128xf32> / } / } / } / return / } / / loops = {{i0}, {i2, i3}}, to vectorize the outermost and the two / innermost loops; / loops = {{i1}, {i2, i3}}, to vectorize the middle and the two innermost / loops; / loops = {{i2}}, to vectorize only the first innermost loop; / loops = {{i3}}, to vectorize only the second innermost loop; / loops = {{i1}}, to vectorize only the middle loop. LogicalResult vectorizeAffineLoopNest(std::vector<SmallVector<AffineForOp, 2>> &loops, const VectorizationStrategy &strategy);

/ Normalize a affine.parallel op so that lower bounds are 0 and steps are 1. / As currently implemented, this transformation cannot fail and will return / early if the op is already in a normalized form. void normalizeAffineParallel(AffineParallelOp op);

/ Normalize an affine.for op. An affine.for op is normalized by converting the / lower bound to zero and loop step to one. The upper bound is set to the trip / count of the loop. Original loops must have a lower bound with only a single / result. There is no such restriction on upper bounds. Returns success if the / loop has been normalized (or is already in the normal form). If / promoteSingleIter is true, the loop is simply promoted if it has a single / iteration. LogicalResult normalizeAffineFor(AffineForOp op, bool promoteSingleIter = false);

/ Traverse e and return an AffineExpr where all occurrences of dim have / been replaced by either: / - min if positivePath is true when we reach an occurrence of dim / - max if positivePath is true when we reach an occurrence of dim / positivePath is negated each time we hit a multiplicative or divisive / binary op with a constant negative coefficient. AffineExpr substWithMin(AffineExpr e, AffineExpr dim, AffineExpr min, AffineExpr max, bool positivePath = true);

/ Replaces all "dereferencing" uses of oldMemRef with newMemRef while / optionally remapping the old memref's indices using the supplied affine map, / indexRemap. The new memref could be of a different shape or rank. / extraIndices provides any additional access indices to be added to the / start. / / indexRemap remaps indices of the old memref access to a new set of indices / that are used to index the memref. Additional input operands to indexRemap / can be optionally provided in extraOperands, and they occupy the start / of its input list. indexRemap's dimensional inputs are expected to / correspond to memref's indices, and its symbolic inputs if any should be / provided in symbolOperands.

/ If userFilterFn is specified, restrict replacement to only those users / that pass the specified filter (i.e., the filter returns true). / / 'allowNonDereferencingOps', if set, allows replacement of non-dereferencing / uses of a memref without any requirement for access index rewrites as long / as the user operation has the MemRefsNormalizable trait. The default value / of this flag is false. / / 'replaceInDeallocOp', if set, lets DeallocOp, a non-dereferencing user, to / also be a candidate for replacement. The default value of this flag is / false. / / Returns true on success and false if the replacement is not possible, / whenever a memref is used as an operand in a non-dereferencing context and / 'allowNonDereferencingOps' is false, except for dealloc's on the memref / which are left untouched. See comments at function definition for an / example.

Ex: to replace load A[i, j] with load Abuf[t mod 2, ii - i, j]: The SSA value corresponding to 't mod 2' should be in 'extraIndices', and index remap will perform (i, j) -> (ii - i, j), i.e., indexRemap = (d0, d1, d2) -> (d0 - d1, d2), and ii will be the extra operand. Without any extra operands, note that 'indexRemap' would just be applied to existing indices (i, j).

TODO: allow extraIndices to be added at any position. LogicalResult replaceAllMemRefUsesWith( Value oldMemRef, Value newMemRef, ArrayRef<Value> extraIndices = {}, AffineMap indexRemap = AffineMap(), ArrayRef<Value> extraOperands = {}, ArrayRef<Value> symbolOperands = {}, llvm::function_ref<bool(Operation *)> userFilterFn = nullptr, bool allowNonDereferencingOps = false, bool replaceInDeallocOp = false);

/ Performs the same replacement as the other version above but only for the / dereferencing uses of oldMemRef in op, except in cases where / 'allowNonDereferencingOps' is set to true where we replace the / non-dereferencing uses as well. LogicalResult replaceAllMemRefUsesWith(Value oldMemRef, Value newMemRef, Operation *op, ArrayRef<Value> extraIndices = {}, AffineMap indexRemap = AffineMap(), ArrayRef<Value> extraOperands = {}, ArrayRef<Value> symbolOperands = {}, bool allowNonDereferencingOps = false);

/ Rewrites the memref defined by alloc or reinterpret_cast op to have an / identity layout map and updates all its indexing uses. Returns failure if / any of its uses escape (while leaving the IR in a valid state). template <typename AllocLikeOp> LogicalResult normalizeMemRef(AllocLikeOp op); extern template LogicalResult normalizeMemRef<memref::AllocaOp>(memref::AllocaOp op); extern template LogicalResult normalizeMemRef<memref::AllocOp>(memref::AllocOp op); LogicalResult normalizeMemRef(memref::ReinterpretCastOp op);

/ Normalizes memrefType so that the affine layout map of the memref is / transformed to an identity map with a new shape being computed for the / normalized memref type and returns it. The old memref type is simplify / returned if the normalization failed. MemRefType normalizeMemRefType(MemRefType memrefType);

/ Given an operation, inserts one or more single result affine apply / operations, results of which are exclusively used by this operation. / The operands of these newly created affine apply ops are / guaranteed to be loop iterators or terminal symbols of a function. / / Before / / affine.for i = 0 to #map(N) / idx = affine.apply (d0) -> (d0 mod 2) (i) / send A[idx], ... / v = "compute"(idx, ...) / / After / / affine.for i = 0 to #map(N) / idx = affine.apply (d0) -> (d0 mod 2) (i) / send A[idx], ... / idx_ = affine.apply (d0) -> (d0 mod 2) (i) / v = "compute"(idx_, ...)

/ This allows the application of different transformations on send and / compute (for eg. different shifts/delays) / / Fills sliceOps with the list of affine.apply operations. / In the following cases, sliceOps remains empty: / 1. If none of opInst's operands were the result of an affine.apply / (i.e., there was no affine computation slice to create). / 2. If all the affine.apply op's supplying operands to this opInst did not / have any uses other than those in this opInst. void createAffineComputationSlice(Operation *opInst, SmallVectorImpl<AffineApplyOp> *sliceOps);

/ Emit code that computes the given affine expression using standard / arithmetic operations applied to the provided dimension and symbol values. Value expandAffineExpr(OpBuilder &builder, Location loc, AffineExpr expr, ValueRange dimValues, ValueRange symbolValues);

/ Create a sequence of operations that implement the affineMap applied to / the given operands (as it it were an AffineApplyOp). std::optional<SmallVector<Value, 8>> expandAffineMap(OpBuilder &builder, Location loc, AffineMap affineMap, ValueRange operands);

/ Holds the result of (div a, b) and (mod a, b). struct DivModValue { Value quotient; Value remainder; };

/ Create IR to calculate (div lhs, rhs) and (mod lhs, rhs). DivModValue getDivMod(OpBuilder &b, Location loc, Value lhs, Value rhs);

/ Generate the IR to delinearize linearIndex given the basis and return / the multi-index. hasOuterBound indicates whether basis has an entry / given the size of the first multi-index result - if it is true, the function / will return basis.size() values, otherwise, it will return basis.size() + 1. FailureOr<SmallVector<Value>> delinearizeIndex(OpBuilder &b, Location loc, Value linearIndex, ArrayRef<Value> basis, bool hasOuterBound = true);

FailureOr<SmallVector<Value>> delinearizeIndex(OpBuilder &b, Location loc, Value linearIndex, ArrayRef<OpFoldResult> basis, bool hasOuterBound = true);

Generate IR that extracts the linear index from a multi-index according to a basis/shape. The basis may contain either multiIndex.size() or multiIndex.size() - 1 elements. OpFoldResult linearizeIndex(ArrayRef<OpFoldResult> multiIndex, ArrayRef<OpFoldResult> basis, ImplicitLocOpBuilder &builder);

OpFoldResult linearizeIndex(OpBuilder &builder, Location loc, ArrayRef<OpFoldResult> multiIndex, ArrayRef<OpFoldResult> basis);

/ Ensure that all operations that could be executed after start / (noninclusive) and prior to memOp (e.g. on a control flow/op path / between the operations) do not have the potential memory effect / EffectType on memOp. memOp is an operation that reads or writes to / a memref. For example, if EffectType is MemoryEffects::Write, this method / will check if there is no write to the memory between start and memOp / that would change the read within memOp.

Definition at line 687 of file Utils.cpp.

References mlir::Block::end(), mlir::Operation::getBlock(), getNumCommonSurroundingLoops(), mlir::Operation::getParentOp(), mlir::Operation::getParentRegion(), mlir::Region::getParentRegion(), mlir::Block::getSuccessors(), mlir::Block::getTerminator(), mlir::Operation::isAncestor(), mlir::Region::isAncestor(), mayAlias(), and mayHaveEffect().

Referenced by findUnusedStore(), forwardStoreToLoad(), loadCSE(), and mlir::affine::hasNoInterveningEffect< mlir::MemoryEffects::Read, affine::AffineReadOpInterface >().

◆ insertBackwardComputationSlice()

AffineForOp mlir::affine::insertBackwardComputationSlice	(	Operation *	srcOpInst,
		Operation *	dstOpInst,
		unsigned	dstLoopDepth,
		ComputationSliceState *	sliceState )

Creates a clone of the computation contained in the loop nest surrounding 'srcOpInst', slices the iteration space of src loop based on slice bounds in 'sliceState', and inserts the computation slice at the beginning of the operation block of the loop at 'dstLoopDepth' in the loop nest surrounding 'dstOpInst'.

Creates a computation slice of the loop nest surrounding 'srcOpInst', updates the slice loop bounds with any non-null bound maps specified in 'sliceState', and inserts this slice into the loop nest surrounding 'dstOpInst' at loop depth 'dstLoopDepth'.

Returns the top-level loop of the computation slice on success, returns nullptr otherwise.

Definition at line 1982 of file Utils.cpp.

References b, mlir::Operation::emitError(), findInstPosition(), getAffineForIVs(), getInstAtPosition(), mlir::affine::ComputationSliceState::lbOperands, mlir::affine::ComputationSliceState::lbs, mlir::affine::ComputationSliceState::ubOperands, and mlir::affine::ComputationSliceState::ubs.

◆ interchangeLoops()

void mlir::affine::interchangeLoops	(	AffineForOp	forOpA,
		AffineForOp	forOpB )

Performs loop interchange on 'forOpA' and 'forOpB'.

Performs loop interchange on 'forOpA' and 'forOpB', where 'forOpB' is nested within 'forOpA' as the only non-terminator operation in its block.

Requires that 'forOpA' and 'forOpB' are part of a perfectly nested sequence of loops.

Definition at line 1279 of file LoopUtils.cpp.

◆ isAffineForInductionVar()

bool mlir::affine::isAffineForInductionVar ( Value val )

Returns true if the provided value is the induction variable of an AffineForOp.

Returns true if the provided value is the induction variable of a AffineForOp.

Definition at line 2926 of file AffineOps.cpp.

References getForInductionVarOwner().

Referenced by addMissingLoopIVBounds(), getNumCommonLoops(), isAccessIndexInvariant(), and isAffineInductionVar().

◆ isAffineInductionVar()

bool mlir::affine::isAffineInductionVar ( Value val )

Returns true if the provided value is the induction variable of an AffineForOp or AffineParallelOp.

Definition at line 2934 of file AffineOps.cpp.

References isAffineForInductionVar(), and isAffineParallelInductionVar().

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), mlir::affine::MemRefRegion::compute(), isValidDim(), and isValidDim().

◆ isAffineParallelInductionVar()

bool mlir::affine::isAffineParallelInductionVar ( Value val )

Returns true if val is the induction variable of an AffineParallelOp.

Definition at line 2930 of file AffineOps.cpp.

References getAffineParallelInductionVarOwner().

Referenced by getNumCommonLoops(), and isAffineInductionVar().

◆ isContiguousAccess()

template<typename LoadOrStoreOp>

bool mlir::affine::isContiguousAccess	(	Value	iv,
		LoadOrStoreOp	memoryOp,
		int *	memRefDim )

Given:

an induction variable iv of type AffineForOp;
a memoryOp of type const LoadOp& or const StoreOp&; determines whether memoryOp has a contiguous access along iv. Contiguous is defined as either invariant or varying only along a unique MemRef dim. Upon success, the unique MemRef dim is written in memRefDim (or -1 to convey the memRef access is invariant along iv).

Prerequisites:

memRefDim ~= nullptr;
iv of the proper type;
the MemRef accessed by memoryOp has no layout map or at most an identity layout map.

Currently only supports no layout map or identity layout map in the memref. Returns false if the memref has a non-identity layoutMap. This behavior is conservative.

Definition at line 320 of file LoopAnalysis.cpp.

References isAccessIndexInvariant().

Referenced by isVectorizableLoopBody().

◆ isInvariantAccess()

template<typename LoadOrStoreOp>

bool mlir::affine::isInvariantAccess	(	LoadOrStoreOp	memOp,
		AffineForOp	forOp )

Checks if an affine read or write operation depends on forOp's IV, i.e., if the memory access is invariant on forOp.

Definition at line 294 of file LoopAnalysis.cpp.

References mlir::affine::AffineValueMap::composeSimplifyAndCanonicalize(), and mlir::affine::AffineValueMap::getOperands().

◆ isLoopMemoryParallel()

bool mlir::affine::isLoopMemoryParallel ( AffineForOp forOp )

Returns true if ‘forOp’ is a parallel loop.

If parallelReductions is / provided, populates it with descriptors of the parallelizable reductions and / treats them as not preventing parallelization. bool isLoopParallel( AffineForOp forOp, SmallVectorImpl<LoopReduction> *parallelReductions = nullptr);

/ Returns true if ‘forOp’ doesn't have memory dependences preventing parallelization. Memrefs that are allocated inside forOp do not impact its dependences and parallelism. This function does not check iter_args (for values other than memref types) and should be used only as a building block for complete parallelism-checking functions.

◆ isLoopParallelAndContainsReduction()

bool mlir::affine::isLoopParallelAndContainsReduction ( AffineForOp forOp )

Returns whether a loop is a parallel loop and contains a reduction loop.

Returns whether a loop is parallel and contains a reduction loop.

Definition at line 2182 of file Utils.cpp.

Referenced by fuseLoops(), and getComputationSliceState().

◆ isOpwiseShiftValid()

bool mlir::affine::isOpwiseShiftValid	(	AffineForOp	forOp,
		ArrayRef< uint64_t >	shifts )

Checks where SSA dominance would be violated if a for op's body operations are shifted by the specified shifts.

Checks whether SSA dominance would be violated if a for op's body operations are shifted by the specified shifts.

This method checks if a 'def' and all its uses have the same shift factor.

Definition at line 483 of file LoopAnalysis.cpp.

References result.

Referenced by affineForOpBodySkew().

◆ isPerfectlyNested()

bool mlir::affine::isPerfectlyNested ( ArrayRef< AffineForOp > loops )

Returns true if loops is a perfectly nested loop nest, where loops appear in it from outermost to innermost.

Definition at line 1361 of file LoopUtils.cpp.

Referenced by performPreTilingChecks(), and permuteLoops().

◆ isTilingValid()

bool mlir::affine::isTilingValid ( ArrayRef< AffineForOp > loops )

Checks whether hyper-rectangular loop tiling of the nest represented by loops is valid.

The validity condition is from Irigoin and Triolet, which states that two tiles cannot depend on each other. We simplify such condition to just checking whether there is any negative dependence direction, since we have the prior knowledge that the tiling results will be hyper-rectangles, which are scheduled in the lexicographically increasing order on the vector of loop indices. This function will return failure when any dependence component is negative along any of loops.

Definition at line 520 of file LoopAnalysis.cpp.

References checkMemrefAccessDependence(), hasDependence(), mlir::affine::MemRefAccess::opInst, and result.

Referenced by getTopLevelTileableBands().

◆ isTopLevelValue() [1/2]

bool mlir::affine::isTopLevelValue ( Value value )

A utility function to check if a value is defined at the top level of an op with trait AffineScope or is a region argument for such an op.

A utility function to check if a value is defined at the top level of an op with trait AffineScope.

A value of index type defined at the top level is always a valid symbol for all its uses.

If the value is defined in an unlinked region, conservatively assume it is not top-level. A value of index type defined at the top level is always a valid symbol.

Definition at line 250 of file AffineOps.cpp.

References mlir::Value::getDefiningOp(), mlir::Operation::getParentOp(), and mlir::Operation::hasTrait().

Referenced by isDimOpValidSymbol(), isValidDim(), isValidSymbol(), and remainsLegalAfterInline().

◆ isTopLevelValue() [2/2]

bool mlir::affine::isTopLevelValue	(	Value	value,
		Region *	region )

A utility function to check if a value is defined at the top level of region or is an argument of region.

A value of index type defined at the top level of a AffineScope region is always a valid symbol for all uses in that region.

Definition at line 50 of file AffineOps.cpp.

References mlir::Value::getDefiningOp(), and mlir::Operation::getParentRegion().

◆ isValidDim() [1/2]

bool mlir::affine::isValidDim ( Value value )

Returns true if the given Value can be used as a dimension id in the region of the closest surrounding op that has the trait AffineScope.

Definition at line 290 of file AffineOps.cpp.

References getAffineScope(), mlir::Value::getDefiningOp(), mlir::Value::getType(), isAffineInductionVar(), mlir::Type::isIndex(), and isValidDim().

Referenced by isValidAffineIndexOperand(), isValidDim(), legalizeDemotedDims(), remainsLegalAfterInline(), remainsLegalAfterInline(), and verifyDimAndSymbolIdentifiers().

◆ isValidDim() [2/2]

bool mlir::affine::isValidDim	(	Value	value,
		Region *	region )

Returns true if the given Value can be used as a dimension id in region, i.e., for all its uses in region.

Definition at line 314 of file AffineOps.cpp.

References mlir::Value::getDefiningOp(), mlir::Value::getType(), isAffineInductionVar(), mlir::Type::isIndex(), isTopLevelValue(), and isValidSymbol().

◆ isValidLoopInterchangePermutation()

bool mlir::affine::isValidLoopInterchangePermutation	(	ArrayRef< AffineForOp >	loops,
		ArrayRef< unsigned >	loopPermMap )

Checks if the loop interchange permutation 'loopPermMap', of the perfectly nested sequence of loops in 'loops', would violate dependences (loop 'i' in 'loops' is mapped to location 'j = 'loopPermMap[i]' in the interchange).

Checks if the loop interchange permutation 'loopPermMap' of the perfectly nested sequence of loops in 'loops' would violate dependences.

Definition at line 1336 of file LoopUtils.cpp.

References checkLoopInterchangeDependences(), and getDependenceComponents().

◆ isValidSymbol() [1/2]

bool mlir::affine::isValidSymbol ( Value value )

Returns true if the given value can be used as a symbol in the region of the closest surrounding op that has the trait AffineScope.

Definition at line 412 of file AffineOps.cpp.

References getAffineScope(), mlir::Value::getDefiningOp(), mlir::Value::getType(), mlir::Type::isIndex(), isTopLevelValue(), and isValidSymbol().

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), canonicalizePromotedSymbols(), mlir::affine::MemRefRegion::compute(), mlir::affine::ComputationSliceState::getAsConstraints(), isMemRefSizeValidSymbol(), isValidAffineIndexOperand(), isValidDim(), isValidSymbol(), legalizeDemotedDims(), remainsLegalAfterInline(), remainsLegalAfterInline(), and verifyDimAndSymbolIdentifiers().

◆ isValidSymbol() [2/2]

bool mlir::affine::isValidSymbol	(	Value	value,
		Region *	region )

Returns true if the given Value can be used as a symbol for region, i.e., for all its uses in region.

A value can be used as a symbol for region iff it meets one of the following conditions: *) It is a constant.

*) It is a result of a Pure operation whose operands are valid symbolic *) identifiers. *) It is a result of the dim op on a memref whose corresponding size is a valid symbol. *) It is defined at the top level of 'region' or is its argument. *) It dominates region's parent op. If region is null, conservatively assume the symbol definition scope does not exist and only accept the values that would be symbols regardless of the surrounding region structure, i.e. the first three cases above.

Definition at line 457 of file AffineOps.cpp.

References mlir::Value::getDefiningOp(), mlir::Value::getType(), isDimOpValidSymbol(), mlir::Type::isIndex(), mlir::isPure(), isTopLevelValueOrAbove(), mlir::m_Constant(), and mlir::matchPattern().

◆ isVectorizableLoopBody() [1/2]

bool mlir::affine::isVectorizableLoopBody	(	AffineForOp	loop,
		int *	memRefDim,
		NestedPattern &	vectorTransferMatcher )

Checks whether the loop is structurally vectorizable and that all the LoadOp and StoreOp matched have access indexing functions that are either:

invariant along the loop induction variable created by 'loop';
varying along at most one memory dimension. If such a unique dimension is found, it is written into memRefDim.

Definition at line 447 of file LoopAnalysis.cpp.

References isContiguousAccess(), isVectorizableLoopBodyWithOpCond(), and load.

◆ isVectorizableLoopBody() [2/2]

bool mlir::affine::isVectorizableLoopBody	(	AffineForOp	loop,
		NestedPattern &	vectorTransferMatcher )

Checks whether the loop is structurally vectorizable; i.e.:

no conditionals are nested under the loop;
all nested load/stores are to scalar MemRefs. TODO: relax the no-conditionals restriction

Definition at line 473 of file LoopAnalysis.cpp.

References isVectorizableLoopBodyWithOpCond().

Referenced by vectorizeLoopNest().

◆ loopUnrollByFactor()

LogicalResult mlir::affine::loopUnrollByFactor	(	AffineForOp	forOp,
		uint64_t	unrollFactor,
		function_ref< void(unsigned, Operation *, OpBuilder)>	annotateFn = nullptr,
		bool	cleanUpUnroll = false )

Unrolls this for operation by the specified unroll factor.

Unrolls this loop by the specified factor.

Returns failure if the loop cannot be unrolled either due to restrictions or due to invalid unroll factors. Requires positive loop bounds and step. If specified, annotates the Ops in each unrolled iteration by applying annotateFn. When cleanUpUnroll is true, we can ensure the cleanup loop is unrolled regardless of the unroll factor.

Returns success if the loop is successfully unrolled.

Definition at line 995 of file LoopUtils.cpp.

References b, generateCleanupLoopForUnroll(), mlir::generateUnrolledLoop(), getConstantTripCount(), getLargestDivisorOfTripCount(), loopUnrollFull(), promoteIfSingleIteration(), and success().

Referenced by loopUnrollFull(), and loopUnrollUpToFactor().

◆ loopUnrollFull()

LogicalResult mlir::affine::loopUnrollFull ( AffineForOp forOp )

Unrolls this for operation completely if the trip count is known to be constant.

Unrolls this loop completely.

Returns failure otherwise.

Definition at line 871 of file LoopUtils.cpp.

References getConstantTripCount(), loopUnrollByFactor(), promoteIfSingleIteration(), and success().

Referenced by affineForOpBodySkew(), and loopUnrollByFactor().

◆ loopUnrollJamByFactor()

LogicalResult mlir::affine::loopUnrollJamByFactor	(	AffineForOp	forOp,
		uint64_t	unrollJamFactor )

Unrolls and jams this loop by the specified factor.

forOp can be a loop with iteration arguments performing supported reductions and its inner loops can have iteration arguments. Returns success if the loop is successfully unroll-jammed.

Definition at line 1084 of file LoopUtils.cpp.

References areInnerBoundsInvariant(), b, mlir::OpBuilder::clone(), generateCleanupLoopForUnroll(), mlir::AffineMap::get(), mlir::Builder::getAffineDimExpr(), getConstantTripCount(), getLargestDivisorOfTripCount(), mlir::arith::getReductionOp(), getSupportedReductions(), lhs, promoteIfSingleIteration(), rhs, mlir::OpBuilder::setInsertionPointAfter(), mlir::JamBlockGatherer< OpTy >::subBlocks, success(), mlir::Value::use_empty(), and mlir::JamBlockGatherer< OpTy >::walk().

Referenced by loopUnrollJamUpToFactor().

◆ loopUnrollJamUpToFactor()

LogicalResult mlir::affine::loopUnrollJamUpToFactor	(	AffineForOp	forOp,
		uint64_t	unrollJamFactor )

Unrolls and jams this loop by the specified factor or by the trip count (if constant), whichever is lower.

Definition at line 1061 of file LoopUtils.cpp.

References getConstantTripCount(), and loopUnrollJamByFactor().

◆ loopUnrollUpToFactor()

LogicalResult mlir::affine::loopUnrollUpToFactor	(	AffineForOp	forOp,
		uint64_t	unrollFactor )

Unrolls this loop by the specified unroll factor or its trip count, whichever is lower.

Unrolls this loop by the specified factor or by the trip count (if constant) whichever is lower.

Definition at line 886 of file LoopUtils.cpp.

References getConstantTripCount(), and loopUnrollByFactor().

◆ lowerAffineDelinearizeIndexOp()

LogicalResult mlir::affine::lowerAffineDelinearizeIndexOp	(	RewriterBase &	rewriter,
		AffineDelinearizeIndexOp	op )

Lowers affine.delinearize_index into a sequence of division and remainder operations.

Definition at line 87 of file AffineExpandIndexOps.cpp.

References mlir::computeStrides(), mlir::OpBuilder::createOrFold(), mlir::RewriterBase::replaceOp(), and success().

◆ lowerAffineLinearizeIndexOp()

LogicalResult mlir::affine::lowerAffineLinearizeIndexOp	(	RewriterBase &	rewriter,
		AffineLinearizeIndexOp	op )

Lowers affine.linearize_index into a sequence of multiplications and additions.

Make a best effort to sort the input indices so that the most loop-invariant terms are at the left of the additions to enable loop-invariant code motion.

Definition at line 145 of file AffineExpandIndexOps.cpp.

◆ makeComposedAffineApply() [1/2]

AffineApplyOp mlir::affine::makeComposedAffineApply	(	OpBuilder &	b,
		Location	loc,
		AffineExpr	e,
		ArrayRef< OpFoldResult >	operands,
		bool	composeAffineMin = false )

Definition at line 1427 of file AffineOps.cpp.

References b, mlir::AffineMap::inferFromExprList(), and makeComposedAffineApply().

◆ makeComposedAffineApply() [2/2]

AffineApplyOp mlir::affine::makeComposedAffineApply	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands,
		bool	composeAffineMin = false )

Returns a composed AffineApplyOp by composing map and operands with other AffineApplyOps supplying those operands.

The operands of the resulting AffineApplyOp do not change the length of AffineApplyOp chains.

Definition at line 1416 of file AffineOps.cpp.

References b, composeAffineMapAndOperands(), and mlir::foldAttributesIntoMap().

Referenced by commonLinearIdBuilderFn(), mlir::linalg::computeContinuousTileSizes(), mlir::linalg::computeMultiTileSizes(), createMask(), mlir::gpu::WarpDistributionPattern::delinearizeLaneId(), getCollapsedOutputDimFromInputShape(), getXferIndices(), invertSliceIndexing(), laneIdBuilderFn(), makeComposedAffineApply(), makeComposedFoldedAffineApply(), mlir::memref::multiBuffer(), and mlir::linalg::updateBoundsForCyclicDistribution().

◆ makeComposedAffineMin()

AffineMinOp mlir::affine::makeComposedAffineMin	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands )

Returns an AffineMinOp obtained by composing map and operands with AffineApplyOps supplying those operands.

Definition at line 1534 of file AffineOps.cpp.

References b, and makeComposedMinMax().

◆ makeComposedFoldedAffineApply() [1/2]

OpFoldResult mlir::affine::makeComposedFoldedAffineApply	(	OpBuilder &	b,
		Location	loc,
		AffineExpr	expr,
		ArrayRef< OpFoldResult >	operands,
		bool	composeAffineMin = false )

Variant of makeComposedFoldedAffineApply that applies to an expression.

Definition at line 1503 of file AffineOps.cpp.

References b, mlir::AffineMap::inferFromExprList(), and makeComposedFoldedAffineApply().

◆ makeComposedFoldedAffineApply() [2/2]

OpFoldResult mlir::affine::makeComposedFoldedAffineApply	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands,
		bool	composeAffineMin = false )

Constructs an AffineApplyOp that applies map to operands after composing the map with the maps of any other AffineApplyOp supplying the operands, then immediately attempts to fold it.

If folding results in a constant value, no ops are actually created. The map must be a single-result affine map.

Definition at line 1469 of file AffineOps.cpp.

References b, mlir::AffineMap::getNumResults(), mlir::m_Constant(), makeComposedAffineApply(), mlir::matchPattern(), and mlir::OpBuilder::setInsertionPoint().

Referenced by mlir::affine::AffineBuilder::add(), mlir::tensor::bubbleUpPadSlice(), HopperBuilder::buildBarrierArriveTx(), buildLinearId(), HopperBuilder::buildTmaAsyncLoad(), calculateExpandedAccessIndices(), calculateTileOffsetsAndSizes(), mlir::affine::AffineBuilder::ceil(), common3DIdBuilderFn(), commonLinearIdBuilderFn(), composedAffineMultiply(), mlir::linalg::computeContinuousTileSizes(), mlir::linalg::computePaddedShape(), mlir::linalg::computeSliceParameters(), mlir::memref::computeSuffixProductIRBlockImpl(), mlir::linalg::computeTileSizes(), mlir::memref::createExpandStridedMetadataPass(), createInBoundsCond(), createInitialTensorsForTiling(), mlir::tensor::createPadHighOp(), denormalizeInductionVariableForIndexType(), denormalizeIndVar(), emitNormalizedLoopBoundsForIndexType(), mlir::affine::AffineBuilder::floor(), mlir::tensor::getCollapsedExtractSliceInfo(), getCollapsedIndices(), getCompressedMaskOp(), getIndicesForLoadOrStore(), mlir::memref::getLinearizedMemRefOffsetAndSize(), mlir::memref::getLinearizedMemRefOffsetAndSize(), getOffsetForBitwidth(), getProductOfIndexes(), getSplitReductionIvs(), getTileOffsetAndSizesWithForAllOp(), getUserTileSizesAndNumThreads(), laneIdBuilderFn(), mlir::linalg::lowerPack(), makeComposedFoldedAffineApply(), mergeOffsetsSizesAndStrides(), mlir::affine::AffineBuilder::mul(), normalizeUpperBounds(), mlir::linalg::offsetIndices(), mlir::linalg::packMatmulGreedily(), resolveIndicesIntoOpWithOffsetsAndStrides(), mlir::memref::resolveSourceIndicesCollapseShape(), CopyBuilder::rewrite(), mlir::linalg::splitOp(), mlir::affine::AffineBuilder::sub(), and updateExpandedGenericOpRegion().

◆ makeComposedFoldedAffineMax()

OpFoldResult mlir::affine::makeComposedFoldedAffineMax	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands )

Constructs an AffineMinOp that computes a maximum across the results of applying map to operands, then immediately attempts to fold it.

If folding results in a constant value, no ops are actually created.

Definition at line 1579 of file AffineOps.cpp.

References b, and makeComposedFoldedMinMax().

Referenced by mlir::tensor::bubbleUpPadSlice(), buildMax(), mlir::memref::getLinearizedMemRefOffsetAndSize(), getTileOffsetAndSizesWithForAllOp(), and mlir::affine::AffineBuilder::max().

◆ makeComposedFoldedAffineMin()

OpFoldResult mlir::affine::makeComposedFoldedAffineMin	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands )

Constructs an AffineMinOp that computes a minimum across the results of applying map to operands, then immediately attempts to fold it.

If folding results in a constant value, no ops are actually created.

Definition at line 1572 of file AffineOps.cpp.

References b, and makeComposedFoldedMinMax().

Referenced by mlir::tensor::bubbleUpPadSlice(), buildMin(), mlir::linalg::computeSliceParameters(), getBoundedTileSize(), getTileOffsetAndSizesWithForAllOp(), mlir::affine::AffineBuilder::min(), and mlir::linalg::splitOp().

◆ makeComposedFoldedMultiResultAffineApply()

SmallVector< OpFoldResult > mlir::affine::makeComposedFoldedMultiResultAffineApply	(	OpBuilder &	b,
		Location	loc,
		AffineMap	map,
		ArrayRef< OpFoldResult >	operands,
		bool	composeAffineMin = false )

Variant of makeComposedFoldedAffineApply suitable for multi-result maps.

Note that this may create as many affine.apply operations as the map has results given that affine.apply must be single-result.

Definition at line 1514 of file AffineOps.cpp.

References b, and mlir::AffineMap::getNumResults().

Referenced by mlir::linalg::computeMultiTileSizes(), getGenericOpLoopRange(), mlir::linalg::makeTiledLoopRanges(), and tileLinalgOpImpl().

◆ materializeComputedBound()

OpFoldResult mlir::affine::materializeComputedBound	(	OpBuilder &	b,
		Location	loc,
		AffineMap	boundMap,
		ArrayRef< std::pair< Value, std::optional< int64_t > > >	mapOperands )

Materialize an already computed bound with Affine dialect ops.

ValueBoundsOpInterface::computeBound computes bounds but does not create IR. It is dialect independent.
materializeComputedBound materializes computed bounds with Affine dialect ops.
reifyIndexValueBound/reifyShapedValueDimBound are a combination of the two functions mentioned above.

Definition at line 34 of file ReifyValueBounds.cpp.

References b, canonicalizeMapAndOperands(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getResult(), mlir::AffineMap::getSingleConstantResult(), mlir::Value::getType(), mlir::Type::isIndex(), and mlir::AffineMap::isSingleConstant().

Referenced by makeIndependent(), and reifyValueBound().

◆ mergeOffsetsSizesAndStrides() [1/2]

LogicalResult mlir::affine::mergeOffsetsSizesAndStrides	(	OpBuilder &	builder,
		Location	loc,
		ArrayRef< OpFoldResult >	producerOffsets,
		ArrayRef< OpFoldResult >	producerSizes,
		ArrayRef< OpFoldResult >	producerStrides,
		const llvm::SmallBitVector &	droppedProducerDims,
		ArrayRef< OpFoldResult >	consumerOffsets,
		ArrayRef< OpFoldResult >	consumerSizes,
		ArrayRef< OpFoldResult >	consumerStrides,
		SmallVector< OpFoldResult > &	combinedOffsets,
		SmallVector< OpFoldResult > &	combinedSizes,
		SmallVector< OpFoldResult > &	combinedStrides )

Fills the combinedOffsets, combinedSizes and combinedStrides to use when combining a producer slice into a consumer slice.

This function performs the following computation:

Combined offsets = producer_offsets * consumer_strides + consumer_offsets
Combined sizes = consumer_sizes
Combined strides = producer_strides * consumer_strides

Definition at line 17 of file ViewLikeInterfaceUtils.cpp.

References mlir::bindSymbols(), mlir::Builder::getContext(), makeComposedFoldedAffineApply(), and success().

Referenced by mergeOffsetsSizesAndStrides().

◆ mergeOffsetsSizesAndStrides() [2/2]

LogicalResult mlir::affine::mergeOffsetsSizesAndStrides	(	OpBuilder &	builder,
		Location	loc,
		OffsetSizeAndStrideOpInterface	producer,
		OffsetSizeAndStrideOpInterface	consumer,
		const llvm::SmallBitVector &	droppedProducerDims,
		SmallVector< OpFoldResult > &	combinedOffsets,
		SmallVector< OpFoldResult > &	combinedSizes,
		SmallVector< OpFoldResult > &	combinedStrides )

Fills the combinedOffsets, combinedSizes and combinedStrides to use when combining a producer slice op into a consumer slice op.

Definition at line 61 of file ViewLikeInterfaceUtils.cpp.

References mergeOffsetsSizesAndStrides().

◆ noDependence()

bool mlir::affine::noDependence ( DependenceResult result )

inline

Returns true if the provided DependenceResult corresponds to the absence of a dependence.

Definition at line 187 of file AffineAnalysis.h.

References mlir::affine::DependenceResult::NoDependence, and result.

Referenced by hasCyclicDependence(), mayDependence(), and mayHaveEffect().

◆ numEnclosingInvariantLoops()

int64_t mlir::affine::numEnclosingInvariantLoops ( OpOperand & operand )

Performs explicit copying for the contiguous sequence of operations in the block iterator range [‘begin’, ‘end’), where ‘end’ can't be past the / terminator of the block (since additional operations are potentially / inserted right before end.

copyOptions provides various parameters, and / the output argument copyNests is the set of all copy nests inserted, each / represented by its root affine.for. Since we generate alloc's and dealloc's / for all fast buffers (before and after the range of operations resp. or at a / hoisted position), all of the fast memory capacity is assumed to be / available for processing this block range. When 'filterMemRef' is specified, / copies are only generated for the provided MemRef. Returns success if the / explicit copying succeeded for all memrefs on which affine load/stores were / encountered. For memrefs for whose element types a size in bytes can't be / computed (index type), their capacity is not accounted for and the / fastMemCapacityBytes copy option would be non-functional in such cases. LogicalResult affineDataCopyGenerate(Block::iterator begin, Block::iterator end, const AffineCopyOptions &copyOptions, std::optional<Value> filterMemRef, DenseSet<Operation *> &copyNests);

/ A convenience version of affineDataCopyGenerate for all ops in the body of / an AffineForOp. LogicalResult affineDataCopyGenerate(AffineForOp forOp, const AffineCopyOptions &copyOptions, std::optional<Value> filterMemRef, DenseSet<Operation *> &copyNests);

/ Result for calling generateCopyForMemRegion. struct CopyGenerateResult { Number of bytes used by alloc. uint64_t sizeInBytes;

The newly created buffer allocation. Operation *alloc;

Generated loop nest for copying data between the allocated buffer and the original memref. Operation *copyNest; };

/ generateCopyForMemRegion is similar to affineDataCopyGenerate, but works / with a single memref region. memrefRegion is supposed to contain analysis / information within analyzedOp. The generated prologue and epilogue always / surround analyzedOp. / / Note that analyzedOp is a single op for API convenience, and the / [begin, end) version can be added as needed. / / Also note that certain options in copyOptions aren't looked at anymore, / like slowMemorySpace. LogicalResult generateCopyForMemRegion(const MemRefRegion &memrefRegion, Operation *analyzedOp, const AffineCopyOptions &copyOptions, CopyGenerateResult &result);

/ Replace a perfect nest of "for" loops with a single linearized loop. Assumes / loops contains a list of perfectly nested loops outermost to innermost / that are normalized (step one and lower bound of zero) and with bounds and / steps independent of any loop induction variable involved in the nest. / Coalescing affine.for loops is not always possible, i.e., the result may not / be representable using affine.for. LogicalResult coalesceLoops(MutableArrayRef<AffineForOp> loops);

/ Maps forOp for execution on a parallel grid of virtual processorIds of / size given by numProcessors. This is achieved by embedding the SSA values / corresponding to processorIds and numProcessors into the bounds and step / of the forOp. No check is performed on the legality of the rewrite, it is / the caller's responsibility to ensure legality. / / Requires that processorIds and numProcessors have the same size and that / for each idx, processorIds[idx] takes, at runtime, all values between 0 / and numProcessors[idx] - 1. This corresponds to traditional use cases for: / 1. GPU (threadIdx, get_local_id(), ...) / 2. MPI (MPI_Comm_rank) / 3. OpenMP (omp_get_thread_num) / / Example: / Assuming a 2-d grid with processorIds = [blockIdx.x, threadIdx.x] and / numProcessors = [gridDim.x, blockDim.x], the loop: / / / scf.for i = lb to ub step step { / ... / } / / / is rewritten into a version resembling the following pseudo-IR: / / / scf.for i = lb + step * (threadIdx.x + blockIdx.x * blockDim.x) / to ub step gridDim.x * blockDim.x * step { / ... / } / void mapLoopToProcessorIds(scf::ForOp forOp, ArrayRef<Value> processorId, ArrayRef<Value> numProcessors);

/ Gathers all AffineForOps in 'func.func' grouped by loop depth. void gatherLoops(func::FuncOp func, std::vector<SmallVector<AffineForOp, 2>> &depthToLoops);

/ Creates an AffineForOp while ensuring that the lower and upper bounds are / canonicalized, i.e., unused and duplicate operands are removed, any constant / operands propagated/folded in, and duplicate bound maps dropped. AffineForOp createCanonicalizedAffineForOp(OpBuilder b, Location loc, ValueRange lbOperands, AffineMap lbMap, ValueRange ubOperands, AffineMap ubMap, int64_t step = 1);

/ Separates full tiles from partial tiles for a perfect nest nest by / generating a conditional guard that selects between the full tile version / and the partial tile version using an AffineIfOp. The original loop nest / is replaced by this guarded two version form. / / affine.if (cond) / // full_tile / else / // partial tile / LogicalResult separateFullTiles(MutableArrayRef<AffineForOp> nest, SmallVectorImpl<AffineForOp> *fullTileNest = nullptr);

/ Walk an affine.for to find a band to coalesce. LogicalResult coalescePerfectlyNestedAffineLoops(AffineForOp op);

/ Count the number of loops surrounding operand such that operand could be hoisted above. Stop counting at the first loop over which the operand cannot be hoisted. This counts any LoopLikeOpInterface, not just affine.for.

Definition at line 2816 of file LoopUtils.cpp.

References mlir::IROperand< DerivedT, IRValueT >::get(), mlir::detail::IROperandBase::getOwner(), and mlir::Operation::getParentOfType().

◆ parseDimAndSymbolList()

ParseResult mlir::affine::parseDimAndSymbolList	(	OpAsmParser &	parser,
		SmallVectorImpl< Value > &	operands,
		unsigned &	numDims )

Parses dimension and symbol list.

Parses dimension and symbol list and returns true if parsing failed.

numDims is set to the number of dimensions in the list parsed.

Definition at line 507 of file AffineOps.cpp.

References mlir::AsmParser::getBuilder(), mlir::Builder::getIndexType(), mlir::AsmParser::OptionalSquare, mlir::AsmParser::Paren, mlir::OpAsmParser::parseOperandList(), and mlir::OpAsmParser::resolveOperands().

Referenced by parseBound().

◆ permuteLoops()

unsigned mlir::affine::permuteLoops	(	ArrayRef< AffineForOp >	inputNest,
		ArrayRef< unsigned >	permMap )

Performs a loop permutation on a perfectly nested loop nest inputNest (where the contained loops appear from outer to inner) as specified by the permutation permMap: loop 'i' in inputNest is mapped to location 'loopPermMap[i]', where positions 0, 1, ... are from the outermost position to inner.

Returns the position in inputNest of the AffineForOp that becomes the new outermost loop of this nest. This method always succeeds, asserts out on invalid input / specifications.

Definition at line 1383 of file LoopUtils.cpp.

References mlir::Block::begin(), mlir::Block::end(), mlir::Block::getOperations(), and isPerfectlyNested().

Referenced by sinkSequentialLoops().

◆ populateAffineExpandIndexOpsAsAffinePatterns()

void mlir::affine::populateAffineExpandIndexOpsAsAffinePatterns ( RewritePatternSet & patterns )

Populate patterns that expand affine index operations into their equivalent affine.apply representations.

Definition at line 89 of file AffineExpandIndexOpsAsAffine.cpp.

References mlir::patterns.

◆ populateAffineExpandIndexOpsPatterns()

void mlir::affine::populateAffineExpandIndexOpsPatterns ( RewritePatternSet & patterns )

Populate patterns that expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect).

Definition at line 229 of file AffineExpandIndexOps.cpp.

References mlir::patterns.

◆ promoteIfSingleIteration()

LogicalResult mlir::affine::promoteIfSingleIteration ( AffineForOp forOp )

Promotes the loop body of a AffineForOp to its containing block if the loop was known to have a single iteration.

Promotes the loop body of a forOp to its containing block if the forOp was known to have a single iteration.

Definition at line 119 of file LoopUtils.cpp.

References mlir::arith::ConstantIndexOp::create(), getConstantTripCount(), mlir::Builder::getDimIdentityMap(), replaceIterArgsAndYieldResults(), and success().

Referenced by fuseLoops(), generateCleanupLoopForUnroll(), generateShiftedLoop(), loopUnrollByFactor(), loopUnrollFull(), and loopUnrollJamByFactor().

◆ promoteSingleIterationLoops()

void mlir::affine::promoteSingleIterationLoops ( func::FuncOp f )

Promotes all single iteration AffineForOp's in the Function, i.e., moves their body into the containing Block.

References target.

◆ registerAffineDataCopyGeneration()

void mlir::affine::registerAffineDataCopyGeneration ( )

inline

Definition at line 1287 of file Passes.h.

◆ registerAffineDataCopyGenerationPass()

void mlir::affine::registerAffineDataCopyGenerationPass ( )

inline

Definition at line 1294 of file Passes.h.

◆ registerAffineExpandIndexOps()

void mlir::affine::registerAffineExpandIndexOps ( )

inline

Definition at line 1308 of file Passes.h.

◆ registerAffineExpandIndexOpsAsAffine()

void mlir::affine::registerAffineExpandIndexOpsAsAffine ( )

inline

Definition at line 1329 of file Passes.h.

◆ registerAffineExpandIndexOpsAsAffinePass()

void mlir::affine::registerAffineExpandIndexOpsAsAffinePass ( )

inline

Definition at line 1336 of file Passes.h.

◆ registerAffineExpandIndexOpsPass()

void mlir::affine::registerAffineExpandIndexOpsPass ( )

inline

Definition at line 1315 of file Passes.h.

◆ registerAffineLoopFusion()

void mlir::affine::registerAffineLoopFusion ( )

inline

Definition at line 1350 of file Passes.h.

◆ registerAffineLoopFusionPass()

void mlir::affine::registerAffineLoopFusionPass ( )

inline

Definition at line 1357 of file Passes.h.

◆ registerAffineLoopInvariantCodeMotion()

void mlir::affine::registerAffineLoopInvariantCodeMotion ( )

inline

Definition at line 1371 of file Passes.h.

◆ registerAffineLoopInvariantCodeMotionPass()

void mlir::affine::registerAffineLoopInvariantCodeMotionPass ( )

inline

Definition at line 1378 of file Passes.h.

◆ registerAffineLoopNormalize()

void mlir::affine::registerAffineLoopNormalize ( )

inline

Definition at line 1392 of file Passes.h.

◆ registerAffineLoopNormalizePass()

void mlir::affine::registerAffineLoopNormalizePass ( )

inline

Definition at line 1399 of file Passes.h.

◆ registerAffineLoopTiling()

void mlir::affine::registerAffineLoopTiling ( )

inline

Definition at line 1413 of file Passes.h.

◆ registerAffineLoopTilingPass()

void mlir::affine::registerAffineLoopTilingPass ( )

inline

Definition at line 1420 of file Passes.h.

◆ registerAffineLoopUnroll()

void mlir::affine::registerAffineLoopUnroll ( )

inline

Definition at line 1434 of file Passes.h.

◆ registerAffineLoopUnrollAndJam()

void mlir::affine::registerAffineLoopUnrollAndJam ( )

inline

Definition at line 1455 of file Passes.h.

◆ registerAffineLoopUnrollAndJamPass()

void mlir::affine::registerAffineLoopUnrollAndJamPass ( )

inline

Definition at line 1462 of file Passes.h.

◆ registerAffineLoopUnrollPass()

void mlir::affine::registerAffineLoopUnrollPass ( )

inline

Definition at line 1441 of file Passes.h.

◆ registerAffineParallelize()

void mlir::affine::registerAffineParallelize ( )

inline

Definition at line 1476 of file Passes.h.

◆ registerAffineParallelizePass()

void mlir::affine::registerAffineParallelizePass ( )

inline

Definition at line 1483 of file Passes.h.

◆ registerAffinePasses()

void mlir::affine::registerAffinePasses ( )

inline

Definition at line 1644 of file Passes.h.

Referenced by mlir::registerAllPasses().

◆ registerAffinePipelineDataTransfer()

void mlir::affine::registerAffinePipelineDataTransfer ( )

inline

Definition at line 1497 of file Passes.h.

◆ registerAffinePipelineDataTransferPass()

void mlir::affine::registerAffinePipelineDataTransferPass ( )

inline

Definition at line 1504 of file Passes.h.

◆ registerAffineScalarReplacement()

void mlir::affine::registerAffineScalarReplacement ( )

inline

Definition at line 1518 of file Passes.h.

◆ registerAffineScalarReplacementPass()

void mlir::affine::registerAffineScalarReplacementPass ( )

inline

Definition at line 1525 of file Passes.h.

◆ registerAffineVectorize()

void mlir::affine::registerAffineVectorize ( )

inline

Definition at line 1539 of file Passes.h.

◆ registerAffineVectorizePass()

void mlir::affine::registerAffineVectorizePass ( )

inline

Definition at line 1546 of file Passes.h.

◆ registerLoopCoalescing()

void mlir::affine::registerLoopCoalescing ( )

inline

Definition at line 1560 of file Passes.h.

◆ registerLoopCoalescingPass()

void mlir::affine::registerLoopCoalescingPass ( )

inline

Definition at line 1567 of file Passes.h.

◆ registerRaiseMemrefDialect()

void mlir::affine::registerRaiseMemrefDialect ( )

inline

Definition at line 1581 of file Passes.h.

◆ registerRaiseMemrefDialectPass()

void mlir::affine::registerRaiseMemrefDialectPass ( )

inline

Definition at line 1588 of file Passes.h.

◆ registerSimplifyAffineMinMaxPass()

void mlir::affine::registerSimplifyAffineMinMaxPass ( )

inline

Definition at line 1602 of file Passes.h.

◆ registerSimplifyAffineMinMaxPassPass()

void mlir::affine::registerSimplifyAffineMinMaxPassPass ( )

inline

Definition at line 1609 of file Passes.h.

◆ registerSimplifyAffineStructures()

void mlir::affine::registerSimplifyAffineStructures ( )

inline

Definition at line 1623 of file Passes.h.

◆ registerSimplifyAffineStructuresPass()

void mlir::affine::registerSimplifyAffineStructuresPass ( )

inline

Definition at line 1630 of file Passes.h.

◆ registerTransformDialectExtension()

void mlir::scf::registerTransformDialectExtension ( DialectRegistry & registry )

Definition at line 215 of file AffineTransformOps.cpp.

References mlir::DialectRegistry::addExtensions().

Referenced by mlir::registerAllExtensions().

◆ registerValueBoundsOpInterfaceExternalModels()

void mlir::affine::registerValueBoundsOpInterfaceExternalModels ( DialectRegistry & registry )

Definition at line 155 of file ValueBoundsOpInterfaceImpl.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllDialects().

◆ reifyIndexValueBound()

FailureOr< OpFoldResult > mlir::affine::reifyIndexValueBound	(	OpBuilder &	b,
		Location	loc,
		presburger::BoundType	type,
		Value	value,
		ValueBoundsConstraintSet::StopConditionFn	stopCondition = nullptr,
		bool	closedUB = false )

Reify a bound for the given index-typed value in terms of SSA values for which stopCondition is met.

If no stop condition is specified, reify in terms of the operands of the owner op.

By default, lower/equal bounds are closed and upper bounds are open. If closedUB is set to "true", upper bounds are also closed.

Example: %0 = arith.addi a, b : index %1 = arith.addi %0, c : index

If stopCondition evaluates to "true" for %0 and c, "%0 + %c" is an EQ bound for %1.
If stopCondition evaluates to "true" for a, b and c, "%a + %b + %c" is an EQ bound for %1.
Otherwise, if the owners of a, b or c do not implement the ValueBoundsOpInterface, no bound can be computed.

Definition at line 100 of file ReifyValueBounds.cpp.

References b, and reifyValueBound().

◆ reifyShapedValueDimBound()

FailureOr< OpFoldResult > mlir::affine::reifyShapedValueDimBound	(	OpBuilder &	b,
		Location	loc,
		presburger::BoundType	type,
		Value	value,
		int64_t	dim,
		ValueBoundsConstraintSet::StopConditionFn	stopCondition = nullptr,
		bool	closedUB = false )

Reify a bound for the specified dimension of the given shaped value in terms of SSA values for which stopCondition is met.

If no stop condition is specified, reify in terms of the operands of the owner op.

By default, lower/equal bounds are closed and upper bounds are open. If closedUB is set to "true", upper bounds are also closed.

Definition at line 82 of file ReifyValueBounds.cpp.

References b, and reifyValueBound().

◆ reifyValueBound()

FailureOr< OpFoldResult > mlir::affine::reifyValueBound	(	OpBuilder &	b,
		Location	loc,
		presburger::BoundType	type,
		const ValueBoundsConstraintSet::Variable &	var,
		ValueBoundsConstraintSet::StopConditionFn	stopCondition,
		bool	closedUB = false )

Reify a bound for the given variable in terms of SSA values for which stopCondition is met.

By default, lower/equal bounds are closed and upper bounds are open. If closedUB is set to "true", upper bounds are also closed.

Definition at line 19 of file ReifyValueBounds.cpp.

References b, mlir::ValueBoundsConstraintSet::computeBound(), and materializeComputedBound().

Referenced by reifyIndexValueBound(), and reifyShapedValueDimBound().

◆ reorderOperandsByHoistability()

void mlir::affine::reorderOperandsByHoistability	(	RewriterBase &	rewriter,
		AffineApplyOp	op )

Helper function to rewrite op's affine map and reorder its operands such that they are in increasing order of hoistability (i.e.

the least hoistable) operands come first in the operand list.

Definition at line 41 of file DecomposeAffineOps.cpp.

References canonicalizeMapAndOperands(), mlir::RewriterBase::finalizeOpModification(), mlir::AffineMap::get(), mlir::getAffineSymbolExpr(), mlir::AffineExpr::getContext(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getResult(), mlir::AffineMap::replaceDimsAndSymbols(), mlir::simplifyAffineExpr(), and mlir::RewriterBase::startOpModification().

◆ resolveIndicesIntoOpWithOffsetsAndStrides() [1/2]

void mlir::affine::resolveIndicesIntoOpWithOffsetsAndStrides	(	RewriterBase &	rewriter,
		Location	loc,
		ArrayRef< OpFoldResult >	mixedSourceOffsets,
		ArrayRef< OpFoldResult >	mixedSourceStrides,
		const llvm::SmallBitVector &	rankReducedDims,
		ArrayRef< OpFoldResult >	consumerIndices,
		SmallVectorImpl< Value > &	resolvedIndices )

Given the 'consumerIndices' of a load/store operation operating on an op with offsets and strides, return the combined indices.

For example, using memref.load and memref.subview as an illustration:

%0 = ... : memref<12x42xf32>
%1 = memref.subview %0[%arg0, %arg1][...][%stride1, %stride2] :
  memref<12x42xf32> to memref<4x4xf32, offset=?, strides=[?, ?]>
%2 = load %1[%i1, %i2] : memref<4x4xf32, offset=?, strides=[?, ?]>

could be folded into:

%2 = load %0[%arg0 + %i1 * %stride1][%arg1 + %i2 * %stride2] :

memref<12x42xf32>å

Definition at line 80 of file ViewLikeInterfaceUtils.cpp.

References mlir::bindSymbols(), mlir::AffineMap::get(), mlir::Builder::getContext(), mlir::Builder::getIndexAttr(), mlir::getValueOrCreateConstantIndexOp(), indices, and makeComposedFoldedAffineApply().

Referenced by mlir::amdgpu::foldMemrefViewOp(), InsertSliceOfInsertSliceFolder< OpTy >::matchAndRewrite(), and resolveIndicesIntoOpWithOffsetsAndStrides().

◆ resolveIndicesIntoOpWithOffsetsAndStrides() [2/2]

void mlir::affine::resolveIndicesIntoOpWithOffsetsAndStrides	(	RewriterBase &	rewriter,
		Location	loc,
		ArrayRef< OpFoldResult >	mixedSourceOffsets,
		ArrayRef< OpFoldResult >	mixedSourceStrides,
		const llvm::SmallBitVector &	rankReducedDims,
		ValueRange	consumerIndices,
		SmallVectorImpl< Value > &	resolvedIndices )

inline

Definition at line 82 of file ViewLikeInterfaceUtils.h.

References mlir::getAsOpFoldResult(), and resolveIndicesIntoOpWithOffsetsAndStrides().

◆ resolveSizesIntoOpWithSizes()

void mlir::affine::resolveSizesIntoOpWithSizes	(	ArrayRef< OpFoldResult >	sourceSizes,
		ArrayRef< OpFoldResult >	destSizes,
		const llvm::SmallBitVector &	rankReducedSourceDims,
		SmallVectorImpl< OpFoldResult > &	resolvedSizes )

Given sourceSizes, destSizes and information about which dimensions are dropped by the source: rankReducedSourceDims, compute the resolved sizes that correspond to dest_op(source_op).

In practice, this amounts to filtering by rankReducedSourceDims and taking from sourceSizes if a dimension is dropped, otherwise taking from destSizes.

Definition at line 112 of file ViewLikeInterfaceUtils.cpp.

Referenced by InsertSliceOfInsertSliceFolder< OpTy >::matchAndRewrite().

◆ simplifyAffineMaxOp()

bool mlir::affine::simplifyAffineMaxOp	(	RewriterBase &	rewriter,
		AffineMaxOp	op )

This transform tries to simplify the affine max operation op, by finding a common upper bound for a set of expressions in the affine map results.

It returns whether the transform updated op's affine map.

In concrete terms, given an operation like: affine.max affine_map<(d0)[s0, s1] -> (d0, s1, s0, 128)>(i)[s0, s1] If d0 > 128 and s0 > s1 > 128, the transform will update op to: affine.max affine_map<(d0)[s0, s1] -> (d0, s0)>(i)[s0, s1].

Definition at line 145 of file SimplifyAffineMinMax.cpp.

References simplifyAffineMinMaxOp().

Referenced by simplifyAffineMinMaxOps().

◆ simplifyAffineMinMaxOps()

LogicalResult mlir::affine::simplifyAffineMinMaxOps	(	RewriterBase &	rewriter,
		ArrayRef< Operation * >	ops,
		bool *	modified = nullptr )

This transform applies simplifyAffineMinOp and simplifyAffineMaxOp to all the affine.min or affine.max operations in ops.

After simplification, it invokes the affine.min/max canonicalization patterns on ops.

This transform returns failure if the greedy pattern rewriter failed to converge during canonicalization, otherwise it returns success. If provided, modified is set to true if the IR was modified in any way.

Definition at line 149 of file SimplifyAffineMinMax.cpp.

References mlir::applyOpPatternsGreedily(), mlir::changed, mlir::ExistingAndNewOps, mlir::Builder::getContext(), mlir::OpBuilder::getListener(), mlir::patterns, mlir::GreedyRewriteConfig::setStrictness(), simplifyAffineMaxOp(), simplifyAffineMinOp(), and success().

◆ simplifyAffineMinOp()

bool mlir::affine::simplifyAffineMinOp	(	RewriterBase &	rewriter,
		AffineMinOp	op )

This transform tries to simplify the affine min operation op, by finding a common lower bound for a set of expressions in the affine map results.

It returns whether the transform updated op's affine map.

In concrete terms, given an operation like: affine.min affine_map<(d0)[s0, s1] -> (d0, s1, s0, 128)>(i)[s0, s1] If d0 < 128 and 128 < s1 < s0, the transform will update op to: affine.min affine_map<(d0)[s0, s1] -> (d0, 128)>(i)[s0, s1].

Definition at line 141 of file SimplifyAffineMinMax.cpp.

References simplifyAffineMinMaxOp().

Referenced by simplifyAffineMinMaxOps().

◆ simplifyConstrainedMinMaxOp()

FailureOr< AffineValueMap > mlir::affine::simplifyConstrainedMinMaxOp	(	Operation *	op,
		FlatAffineValueConstraints	constraints )

Try to simplify the given affine.min or affine.max op to an affine map with a single result and operands, taking into account the specified constraint set.

Return failure if no simplified version could be found.

Definition at line 2286 of file Utils.cpp.

References mlir::affine::FlatAffineValueConstraints::addBound(), addConstToResults(), mlir::presburger::IntegerRelation::addInequality(), alignAndAddBound(), mlir::FlatLinearValueConstraints::appendDimVar(), canonicalizeMapAndOperands(), mlir::Builder::getAffineConstantExpr(), mlir::Builder::getAffineDimExpr(), mlir::Builder::getAffineSymbolExpr(), mlir::presburger::IntegerRelation::getConstantBound64(), mlir::getConstantIntValue(), mlir::Operation::getContext(), mlir::FlatLinearValueConstraints::getMaybeValues(), mlir::presburger::IntegerRelation::getNumCols(), mlir::presburger::IntegerRelation::getNumDimAndSymbolVars(), mlir::AffineMap::getNumDims(), mlir::AffineMap::getNumResults(), mlir::AffineMap::getNumSymbols(), mlir::Operation::getOperands(), mlir::FlatLinearConstraints::getSliceBounds(), mlir::AffineMap::getSubMap(), mlir::presburger::IntegerRelation::isEmpty(), mlir::AffineMap::replace(), mlir::AffineMap::shiftDims(), and unpackOptionalValues().

Referenced by canonicalizeMinMaxOp().

◆ simplifyIntegerSet()

IntegerSet mlir::affine::simplifyIntegerSet ( IntegerSet set )

Simplify the integer set by simplifying the underlying affine expressions by flattening and some simple inference.

Also, drop any duplicate constraints. Returns the simplified integer set. This method runs in time linear in the number of constraints.

Definition at line 2200 of file Utils.cpp.

References mlir::FlatLinearConstraints::getAsIntegerSet(), mlir::IntegerSet::getContext(), mlir::IntegerSet::getEmptySet(), mlir::IntegerSet::getNumDims(), mlir::IntegerSet::getNumSymbols(), mlir::presburger::IntegerRelation::isEmpty(), and mlir::presburger::IntegerRelation::removeTrivialRedundancy().

◆ sinkSequentialLoops()

AffineForOp mlir::affine::sinkSequentialLoops ( AffineForOp forOp )

Definition at line 1453 of file LoopUtils.cpp.

References checkLoopInterchangeDependences(), getDependenceComponents(), getPerfectlyNestedLoops(), mlir::affine::DependenceComponent::lb, permuteLoops(), and mlir::affine::DependenceComponent::ub.

Referenced by sinkSequentialLoops().

◆ tile() [1/2]

SmallVector< AffineForOp, 8 > mlir::affine::tile	(	ArrayRef< AffineForOp >	forOps,
		ArrayRef< uint64_t >	sizes,
		AffineForOp	target )

Performs tiling (with interchange) by strip-mining the forOps by sizes and sinking them, in their order of occurrence in forOps, under target.

Returns the new AffineForOps, one per forOps, nested immediately under target.

Definition at line 1596 of file LoopUtils.cpp.

References target, and tile().

◆ tile() [2/2]

SmallVector< SmallVector< AffineForOp, 8 >, 8 > mlir::affine::tile	(	ArrayRef< AffineForOp >	forOps,
		ArrayRef< uint64_t >	sizes,
		ArrayRef< AffineForOp >	targets )

Performs tiling fo imperfectly nested loops (with interchange) by strip-mining the forOps by sizes and sinking them, in their order of occurrence in forOps, under each of the targets.

Returns the new AffineForOps, one per each of (forOps, targets) pair, nested immediately under each of targets.

Definition at line 1584 of file LoopUtils.cpp.

References stripmineSink().

Referenced by tile().

◆ tilePerfectlyNested()

LogicalResult mlir::affine::tilePerfectlyNested	(	MutableArrayRef< AffineForOp >	input,
		ArrayRef< unsigned >	tileSizes,
		SmallVectorImpl< AffineForOp > *	tiledNest = nullptr )

Tiles the specified band of perfectly nested loops creating tile-space loops and intra-tile loops.

A band is a contiguous set of loops. This utility doesn't check for the validity of tiling itself, but just performs it.

Definition at line 772 of file LoopUtils.cpp.

References constructTiledIndexSetHyperRect(), constructTiledLoopNest(), extractForInductionVars(), performPreTilingChecks(), and success().

◆ tilePerfectlyNestedParametric()

LogicalResult mlir::affine::tilePerfectlyNestedParametric	(	MutableArrayRef< AffineForOp >	input,
		ArrayRef< Value >	tileSizes,
		SmallVectorImpl< AffineForOp > *	tiledNest = nullptr )

Tiles the specified band of perfectly nested loops creating tile-space loops and intra-tile loops, using SSA values as tiling parameters.

A band is a contiguous set of loops.

Definition at line 814 of file LoopUtils.cpp.

References constructParametricallyTiledIndexSetHyperRect(), constructTiledLoopNest(), extractForInductionVars(), performPreTilingChecks(), and success().

Namespaces

Classes

Typedefs

Enumerations

Functions

Typedef Documentation

◆ FilterFunctionType

◆ ReductionLoopMap

◆ VectorizableLoopFun

Enumeration Type Documentation

◆ FusionMode

Function Documentation

◆ affineForOpBodySkew()

◆ affineParallelize()

◆ boundCheckLoadOrStoreOp()

◆ buildAffineLoopNest() [1/2]

◆ buildAffineLoopNest() [2/2]

◆ buildSliceTripCountMap()

◆ canFuseLoops()

◆ canonicalizeMapAndOperands()

◆ canonicalizeSetAndOperands()

◆ checkMemrefAccessDependence()

◆ computeProduct()

◆ computeSliceUnion()

◆ createAffineDataCopyGenerationPass() [1/2]

◆ createAffineDataCopyGenerationPass() [2/2]

◆ createAffineExpandIndexOpsAsAffinePass()

◆ createAffineExpandIndexOpsPass()

◆ createAffineLoopInvariantCodeMotionPass()

◆ createAffineLoopNormalizePass()

◆ createAffineParallelize() [1/2]

◆ createAffineParallelize() [2/2]

◆ createAffineParallelizePass()

◆ createAffineScalarReplacementPass()

◆ createAffineVectorize() [1/2]

◆ createAffineVectorize() [2/2]

◆ createLoopCoalescingPass()

◆ createLoopFusionPass()

◆ createLoopTilingPass() [1/2]

◆ createLoopTilingPass() [2/2]

◆ createLoopUnrollAndJamPass()

◆ createLoopUnrollPass()

◆ createPipelineDataTransferPass()

◆ createRaiseMemrefToAffine()

◆ createSimplifyAffineMinMaxPass()

◆ createSimplifyAffineStructuresPass()

◆ decompose()

◆ defaultFilterFunction()

◆ extractForInductionVars()

◆ extractInductionVars()

◆ findInnermostCommonBlockInScope()

◆ fullyComposeAffineMapAndOperands()

◆ fullyComposeAndComputeConstantDelta()

◆ fuseLoops()

◆ gatherProducerConsumerMemrefs()

◆ getAffineAnalysisScope()

◆ getAffineForIVs()

◆ getAffineIVs()

◆ getAffineParallelInductionVarOwner()

◆ getAffineScope()

◆ getComputationSliceState()

◆ getComputeCost()

◆ getConstantTripCount()

◆ getDependenceComponents()

◆ getEnclosingAffineOps()

◆ getForInductionVarOwner()

◆ getFusionComputeCost()

◆ getIndexSet()

◆ getInnermostCommonLoopDepth()

◆ getIntOrFloatMemRefSizeInBytes()

◆ getInvariantAccesses()

◆ getLargestDivisorOfTripCount()

◆ getLoopNestStats()

◆ getMemoryFootprintBytes()

◆ getMemRefIntOrFloatEltSizeInBytes()

◆ getNestingDepth()

◆ getNumCommonSurroundingLoops()

◆ getPerfectlyNestedLoops()

◆ getReachableAffineApplyOps()

◆ getRelationFromMap() [1/2]