MLIR  19.0.0git
Namespaces | Classes | Typedefs | Enumerations | Functions
mlir::affine Namespace Reference

Namespaces

 matcher
 

Classes

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  DependenceComponent
 
struct  DependenceResult
 Checks whether two accesses to the same memref access the same element. More...
 
class  FlatAffineValueConstraints
 FlatAffineValueConstraints is an extension of FlatLinearValueConstraints with helper functions for Affine dialect ops. More...
 
class  FlatAffineRelation
 A FlatAffineRelation represents a set of ordered pairs (domain -> range) where "domain" and "range" are tuples of variables. 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  LoopNestStateCollector
 
struct  MemRefDependenceGraph
 
struct  SliceComputationResult
 Enumerates different result statuses of slice computation by computeSliceUnion More...
 
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  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  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...
 
class  AffineBound
 AffineBound represents a lower or upper bound in the for operation. More...
 
class  AffineValueMap
 An AffineValueMap is an affine map plus its ML value operands and results for analysis purposes. More...
 
struct  FusionResult
 
class  FusionStrategy
 Describes the fusion strategy to be used in the Affine loop fusion utilities. More...
 
struct  LoopNestStats
 LoopNestStats aggregates various per-loop statistics (eg. More...
 
struct  AffineCopyOptions
 Explicit copy / DMA generation options for mlir::affineDataCopyGenerate. More...
 
struct  CopyGenerateResult
 Result for calling generateCopyForMemRegion. More...
 
struct  VectorizationStrategy
 Holds parameters to perform n-D vectorization on a single loop nest. More...
 
struct  DivModValue
 Holds the result of (div a, b) and (mod a, b). More...
 
struct  AffineValueExpr
 
struct  AffineBuilder
 Helper struct to build simple AffineValueExprs with minimal type inference support. More...
 

Typedefs

using VectorizableLoopFun = std::function< bool(AffineForOp)>
 
using FilterFunctionType = std::function< bool(Operation &)>
 A NestedPattern is a nested operation walker that: More...
 
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. More...
 
bool isLoopParallel (AffineForOp forOp, SmallVectorImpl< LoopReduction > *parallelReductions=nullptr)
 Returns true if ‘forOp’ is a parallel loop. More...
 
bool isLoopMemoryParallel (AffineForOp forOp)
 Returns true if ‘forOp’ doesn't have memory dependences preventing parallelization. More...
 
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. More...
 
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. More...
 
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. More...
 
bool noDependence (DependenceResult result)
 Returns true if the provided DependenceResult corresponds to the absence of a dependence. More...
 
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]. More...
 
LogicalResult getRelationFromMap (AffineMap &map, FlatAffineRelation &rel)
 Builds a relation from the given AffineMap/AffineValueMap map, containing all pairs of the form operands -> result that satisfy map. More...
 
LogicalResult getRelationFromMap (const AffineValueMap &map, FlatAffineRelation &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. More...
 
std::optional< uint64_t > getConstantTripCount (AffineForOp forOp)
 Returns the trip count of the loop if it's a constant, std::nullopt otherwise. More...
 
uint64_t getLargestDivisorOfTripCount (AffineForOp forOp)
 Returns the greatest known integral divisor of the trip count. More...
 
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. More...
 
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. More...
 
template<typename LoadOrStoreOp >
bool isContiguousAccess (Value iv, LoadOrStoreOp memoryOp, int *memRefDim)
 Given: More...
 
bool isVectorizableLoopBody (AffineForOp loop, NestedPattern &vectorTransferMatcher)
 Checks whether the loop is structurally vectorizable; i.e. More...
 
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: More...
 
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. More...
 
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. More...
 
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. More...
 
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. More...
 
unsigned getNestingDepth (Operation *op)
 Returns the nesting depth of this operation, i.e., the number of loops surrounding this operation. More...
 
bool isLoopParallelAndContainsReduction (AffineForOp forOp)
 Returns whether a loop is a parallel loop and contains a reduction loop. More...
 
void getSequentialLoops (AffineForOp forOp, llvm::SmallDenseSet< Value, 8 > *sequentialLoops)
 Returns in 'sequentialLoops' all sequential loops in loop nest rooted at 'forOp'. More...
 
void getComputationSliceState (Operation *depSourceOp, Operation *depSinkOp, 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'. More...
 
uint64_t getSliceIterationCount (const llvm::SmallDenseMap< Operation *, uint64_t, 8 > &sliceTripCountMap)
 Return the number of iterations for the slicetripCountMap provided. More...
 
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'. More...
 
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. More...
 
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'. More...
 
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. More...
 
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. More...
 
unsigned getNumCommonSurroundingLoops (Operation &a, Operation &b)
 Returns the number of surrounding loops common to both A and B. More...
 
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. More...
 
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. More...
 
IntegerSet simplifyIntegerSet (IntegerSet set)
 Simplify the integer set by simplifying the underlying affine expressions by flattening and some simple inference. More...
 
unsigned getInnermostCommonLoopDepth (ArrayRef< Operation * > ops, SmallVectorImpl< AffineForOp > *surroundingLoops=nullptr)
 Returns the innermost common loop depth for the set of operations in 'ops'. More...
 
FailureOr< AffineValueMapsimplifyConstrainedMinMaxOp (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. More...
 
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. More...
 
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. More...
 
RegiongetAffineScope (Operation *op)
 Returns the closest region enclosing op that is held by an operation with trait AffineScope; nullptr if there is no such region. More...
 
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. More...
 
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. More...
 
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. More...
 
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. More...
 
ParseResult parseDimAndSymbolList (OpAsmParser &parser, SmallVectorImpl< Value > &operands, unsigned &numDims)
 Parses dimension and symbol list. More...
 
void canonicalizeMapAndOperands (AffineMap *map, SmallVectorImpl< Value > *operands)
 Modifies both map and operands in-place so as to: More...
 
void canonicalizeSetAndOperands (IntegerSet *set, SmallVectorImpl< Value > *operands)
 Canonicalizes an integer set the same way canonicalizeMapAndOperands does for affine maps. More...
 
AffineApplyOp makeComposedAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
 Returns a composed AffineApplyOp by composing map and operands with other AffineApplyOps supplying those operands. More...
 
AffineApplyOp makeComposedAffineApply (OpBuilder &b, Location loc, AffineExpr e, ArrayRef< OpFoldResult > operands)
 
OpFoldResult makeComposedFoldedAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
 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. More...
 
OpFoldResult makeComposedFoldedAffineApply (OpBuilder &b, Location loc, AffineExpr expr, ArrayRef< OpFoldResult > operands)
 Variant of makeComposedFoldedAffineApply that applies to an expression. More...
 
SmallVector< OpFoldResultmakeComposedFoldedMultiResultAffineApply (OpBuilder &b, Location loc, AffineMap map, ArrayRef< OpFoldResult > operands)
 Variant of makeComposedFoldedAffineApply suitable for multi-result maps. More...
 
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. More...
 
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. More...
 
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. More...
 
void fullyComposeAffineMapAndOperands (AffineMap *map, SmallVectorImpl< Value > *operands)
 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. More...
 
bool isAffineForInductionVar (Value val)
 Returns true if the provided value is the induction variable of an AffineForOp. More...
 
bool isAffineParallelInductionVar (Value val)
 Returns true if val is the induction variable of an AffineParallelOp. More...
 
bool isAffineInductionVar (Value val)
 Returns true if the provided value is the induction variable of an AffineForOp or AffineParallelOp. More...
 
AffineForOp getForInductionVarOwner (Value val)
 Returns the loop parent of an induction variable. More...
 
AffineParallelOp getAffineParallelInductionVarOwner (Value val)
 Returns true if the provided value is among the induction variables of an AffineParallelOp. More...
 
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. More...
 
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. More...
 
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. More...
 
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. More...
 
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'. More...
 
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'. More...
 
bool getLoopNestStats (AffineForOp forOp, LoopNestStats *stats)
 Collect loop nest statistics (eg. More...
 
int64_t getComputeCost (AffineForOp forOp, LoopNestStats &stats)
 Computes the total cost of the loop nest rooted at 'forOp' using 'stats'. More...
 
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'. More...
 
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'. More...
 
LogicalResult loopUnrollFull (AffineForOp forOp)
 Unrolls this for operation completely if the trip count is known to be constant. More...
 
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. More...
 
LogicalResult loopUnrollUpToFactor (AffineForOp forOp, uint64_t unrollFactor)
 Unrolls this loop by the specified unroll factor or its trip count, whichever is lower. More...
 
bool LLVM_ATTRIBUTE_UNUSED isPerfectlyNested (ArrayRef< AffineForOp > loops)
 Returns true if loops is a perfectly nested loop nest, where loops appear in it from outermost to innermost. More...
 
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). More...
 
LogicalResult loopUnrollJamByFactor (AffineForOp forOp, uint64_t unrollJamFactor)
 Unrolls and jams this loop by the specified factor. More...
 
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. More...
 
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. More...
 
void promoteSingleIterationLoops (func::FuncOp f)
 Promotes all single iteration AffineForOp's in the Function, i.e., moves their body into the containing Block. More...
 
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. More...
 
void getTileableBands (func::FuncOp f, std::vector< SmallVector< AffineForOp, 6 >> *bands)
 Identify valid and profitable bands of loops to tile. More...
 
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. More...
 
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. More...
 
void interchangeLoops (AffineForOp forOpA, AffineForOp forOpB)
 Performs loop interchange on 'forOpA' and 'forOpB'. More...
 
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). More...
 
unsigned permuteLoops (MutableArrayRef< 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, ... More...
 
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. More...
 
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. More...
 
LogicalResult affineDataCopyGenerate (Block::iterator begin, Block::iterator end, const AffineCopyOptions &copyOptions, std::optional< Value > filterMemRef, DenseSet< Operation * > &copyNests)
 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. More...
 
LogicalResult affineDataCopyGenerate (AffineForOp forOp, const AffineCopyOptions &copyOptions, std::optional< Value > filterMemRef, DenseSet< Operation * > &copyNests)
 A convenience version of affineDataCopyGenerate for all ops in the body of an AffineForOp. More...
 
LogicalResult generateCopyForMemRegion (const MemRefRegion &memrefRegion, Operation *analyzedOp, const AffineCopyOptions &copyOptions, CopyGenerateResult &result)
 generateCopyForMemRegion is similar to affineDataCopyGenerate, but works with a single memref region. More...
 
LogicalResult coalesceLoops (MutableArrayRef< AffineForOp > loops)
 Replace a perfect nest of "for" loops with a single linearized loop. More...
 
void mapLoopToProcessorIds (scf::ForOp forOp, ArrayRef< Value > processorId, ArrayRef< Value > numProcessors)
 Maps forOp for execution on a parallel grid of virtual processorIds of size given by numProcessors. More...
 
void gatherLoops (func::FuncOp func, std::vector< SmallVector< AffineForOp, 2 >> &depthToLoops)
 Gathers all AffineForOps in 'func.func' grouped by loop depth. More...
 
AffineForOp createCanonicalizedAffineForOp (OpBuilder b, Location loc, ValueRange lbOperands, AffineMap lbMap, ValueRange ubOperands, AffineMap ubMap, int64_t step=1)
 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. More...
 
LogicalResult separateFullTiles (MutableArrayRef< AffineForOp > nest, SmallVectorImpl< AffineForOp > *fullTileNest=nullptr)
 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. More...
 
template<typename LoopOpTy >
LogicalResult coalescePerfectlyNestedLoops (LoopOpTy op)
 Walk either an scf.for or an affine.for to find a band to coalesce. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createSimplifyAffineStructuresPass ()
 Creates a simplification pass for affine structures (maps and sets). More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createAffineLoopInvariantCodeMotionPass ()
 Creates a loop invariant code motion pass that hoists loop invariant operations out of affine loops. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createAffineParallelizePass ()
 Creates a pass to convert all parallel affine.for's into 1-d affine.parallel ops. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createAffineLoopNormalizePass (bool promoteSingleIter=false)
 Apply normalization transformations to affine loop-like ops. More...
 
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. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createAffineDataCopyGenerationPass ()
 Overload relying on pass options for initialization. More...
 
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. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createLoopCoalescingPass ()
 Creates a pass that transforms perfectly nested loops with independent bounds into a single loop. More...
 
std::unique_ptr< PasscreateLoopFusionPass (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. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createLoopTilingPass (uint64_t cacheSizeBytes)
 Creates a pass to perform tiling on loop nests. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createLoopTilingPass ()
 Overload relying on pass options for initialization. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createLoopUnrollPass (int unrollFactor=-1, bool unrollUpToFactor=false, bool unrollFull=false, const std::function< unsigned(AffineForOp)> &getUnrollFactor=nullptr)
 Creates a loop unrolling pass with the provided parameters. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createLoopUnrollAndJamPass (int unrollJamFactor=-1)
 Creates a loop unroll jam pass to unroll jam by the specified factor. More...
 
std::unique_ptr< OperationPass< func::FuncOp > > createPipelineDataTransferPass ()
 Creates a pass to pipeline explicit movement of data across levels of the memory hierarchy. More...
 
std::unique_ptr< PasscreateAffineExpandIndexOpsPass ()
 Creates a pass to expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect). More...
 
void registerTransformDialectExtension (DialectRegistry &registry)
 
void populateAffineExpandIndexOpsPatterns (RewritePatternSet &patterns)
 Populate patterns that expand affine index operations into more fundamental operations (not necessarily restricted to Affine dialect). More...
 
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. More...
 
FailureOr< AffineApplyOp > decompose (RewriterBase &rewriter, AffineApplyOp op)
 Split an "affine.apply" operation into smaller ops. More...
 
FailureOr< OpFoldResultreifyIndexValueBound (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. More...
 
FailureOr< OpFoldResultreifyShapedValueDimBound (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. More...
 
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. More...
 
LogicalResult affineParallelize (AffineForOp forOp, ArrayRef< LoopReduction > parallelReductions={}, AffineParallelOp *resOp=nullptr)
 Replaces a parallel affine.for op with a 1-d affine.parallel op. More...
 
LogicalResult hoistAffineIfOp (AffineIfOp ifOp, bool *folded=nullptr)
 Hoists out affine.if/else to as high as possible, i.e., past all invariant affine.fors/parallel's. More...
 
void affineScalarReplace (func::FuncOp f, DominanceInfo &domInfo, PostDominanceInfo &postDomInfo)
 Replace affine store and load accesses by scalars by forwarding stores to loads and eliminate invariant affine loads; consequently, eliminate dead allocs. More...
 
void vectorizeAffineLoops (Operation *parentOp, llvm::DenseSet< Operation *, DenseMapInfo< Operation * >> &loops, ArrayRef< int64_t > vectorSizes, ArrayRef< int64_t > fastestVaryingPattern, const ReductionLoopMap &reductionLoops=ReductionLoopMap())
 Vectorizes affine loops in 'loops' using the n-D vectorization factors in 'vectorSizes'. More...
 
LogicalResult vectorizeAffineLoopNest (std::vector< SmallVector< AffineForOp, 2 >> &loops, const VectorizationStrategy &strategy)
 External utility to vectorize affine loops from a single loop nest using an n-D vectorization strategy (see doc in VectorizationStrategy definition). More...
 
void normalizeAffineParallel (AffineParallelOp op)
 Normalize a affine.parallel op so that lower bounds are 0 and steps are 1. More...
 
LogicalResult normalizeAffineFor (AffineForOp op, bool promoteSingleIter=false)
 Normalize an affine.for op. More...
 
AffineExpr substWithMin (AffineExpr e, AffineExpr dim, AffineExpr min, AffineExpr max, bool positivePath=true)
 Traverse e and return an AffineExpr where all occurrences of dim have been replaced by either: More...
 
LogicalResult replaceAllMemRefUsesWith (Value oldMemRef, Value newMemRef, ArrayRef< Value > extraIndices={}, AffineMap indexRemap=AffineMap(), ArrayRef< Value > extraOperands={}, ArrayRef< Value > symbolOperands={}, Operation *domOpFilter=nullptr, Operation *postDomOpFilter=nullptr, bool allowNonDereferencingOps=false, bool replaceInDeallocOp=false)
 Replaces all "dereferencing" uses of oldMemRef with newMemRef while optionally remapping the old memref's indices using the supplied affine map, indexRemap. More...
 
LogicalResult replaceAllMemRefUsesWith (Value oldMemRef, Value newMemRef, Operation *op, ArrayRef< Value > extraIndices={}, AffineMap indexRemap=AffineMap(), ArrayRef< Value > extraOperands={}, ArrayRef< Value > symbolOperands={}, bool allowNonDereferencingOps=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. More...
 
LogicalResult normalizeMemRef (memref::AllocOp *op)
 Rewrites the memref defined by this alloc op to have an identity layout map and updates all its indexing uses. More...
 
MemRefType normalizeMemRefType (MemRefType memrefType)
 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. More...
 
void createAffineComputationSlice (Operation *opInst, SmallVectorImpl< AffineApplyOp > *sliceOps)
 Given an operation, inserts one or more single result affine apply operations, results of which are exclusively used by this operation. More...
 
Value expandAffineExpr (OpBuilder &builder, Location loc, AffineExpr expr, ValueRange dimValues, ValueRange symbolValues)
 Emit code that computes the given affine expression using standard arithmetic operations applied to the provided dimension and symbol values. More...
 
std::optional< SmallVector< Value, 8 > > expandAffineMap (OpBuilder &builder, Location loc, AffineMap affineMap, ValueRange operands)
 Create a sequence of operations that implement the affineMap applied to the given operands (as it it were an AffineApplyOp). More...
 
DivModValue getDivMod (OpBuilder &b, Location loc, Value lhs, Value rhs)
 Create IR to calculate (div lhs, rhs) and (mod lhs, rhs). More...
 
FailureOr< SmallVector< Value > > delinearizeIndex (OpBuilder &b, Location loc, Value linearIndex, ArrayRef< Value > basis)
 Generate the IR to delinearize linearIndex given the basis and return the multi-index. More...
 
OpFoldResult linearizeIndex (ArrayRef< OpFoldResult > multiIndex, ArrayRef< OpFoldResult > basis, ImplicitLocOpBuilder &builder)
 
template<typename EffectType , typename T >
bool hasNoInterveningEffect (Operation *start, T memOp)
 Ensure that all operations that could be executed after start (noninclusive) and prior to memOp (e.g. More...
 
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. More...
 
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. More...
 
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. More...
 
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). More...
 

Typedef Documentation

◆ FilterFunctionType

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

A NestedPattern is a nested operation walker that:

  1. recursively matches a substructure in the tree;
  2. uses a filter function to refine matches with extra semantic constraints (passed via a lambda of type FilterFunctionType);
  3. 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

Definition at line 42 of file Utils.h.

◆ VectorizableLoopFun

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

Definition at line 88 of file LoopAnalysis.h.

Enumeration Type Documentation

◆ FusionMode

Fusion mode to attempt.

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

Enumerator
Greedy 
ProducerConsumer 
Sibling 

Definition at line 31 of file Passes.h.

Function Documentation

◆ affineDataCopyGenerate() [1/2]

LogicalResult mlir::affine::affineDataCopyGenerate ( AffineForOp  forOp,
const AffineCopyOptions copyOptions,
std::optional< Value filterMemRef,
DenseSet< Operation * > &  copyNests 
)

A convenience version of affineDataCopyGenerate for all ops in the body of an AffineForOp.

Definition at line 2494 of file LoopUtils.cpp.

References affineDataCopyGenerate().

◆ affineDataCopyGenerate() [2/2]

LogicalResult mlir::affine::affineDataCopyGenerate ( Block::iterator  begin,
Block::iterator  end,
const AffineCopyOptions copyOptions,
std::optional< Value filterMemRef,
DenseSet< Operation * > &  copyNests 
)

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.

Definition at line 2305 of file LoopUtils.cpp.

References mlir::detail::divideCeil(), mlir::Operation::emitError(), mlir::Operation::emitWarning(), mlir::failed(), mlir::failure(), mlir::affine::AffineCopyOptions::fastMemCapacityBytes, findHighestBlockForPlacement(), generateCopy(), getFullMemRefAsRegion(), mlir::Operation::getLoc(), getNestingDepth(), mlir::Block::getParentOp(), mlir::affine::AffineCopyOptions::slowMemorySpace, mlir::succeeded(), mlir::success(), and mlir::Block::walk().

Referenced by affineDataCopyGenerate().

◆ 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 237 of file LoopUtils.cpp.

References getConstantTripCount(), isOpwiseShiftValid(), and mlir::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 348 of file Utils.cpp.

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

◆ affineScalarReplace()

void mlir::affine::affineScalarReplace ( func::FuncOp  f,
DominanceInfo domInfo,
PostDominanceInfo postDomInfo 
)

Replace affine store and load accesses by scalars by forwarding stores to loads and eliminate invariant affine loads; consequently, eliminate dead allocs.

Definition at line 1036 of file Utils.cpp.

References forwardStoreToLoad().

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

◆ 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 2667 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 2675 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 1601 of file Utils.cpp.

References getForInductionVarOwner(), and mlir::affine::ComputationSliceState::ivs.

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 249 of file LoopFusionUtils.cpp.

References mlir::affine::FusionResult::FailBlockDependence, mlir::affine::FusionResult::FailFusionDependence, mlir::affine::FusionResult::FailPrecondition, gatherLoadsAndStores(), mlir::affine::FusionStrategy::Generic, getFusedLoopNestInsertionPoint(), getMaxLoopDepth(), getNumCommonSurroundingLoops(), mlir::affine::FusionStrategy::getStrategy(), and mlir::affine::FusionStrategy::ProducerConsumer.

◆ canonicalizeMapAndOperands()

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

◆ 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 1434 of file AffineOps.cpp.

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 
)

◆ coalesceLoops()

LogicalResult mlir::affine::coalesceLoops ( MutableArrayRef< AffineForOp >  loops)

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.

Definition at line 1647 of file LoopUtils.cpp.

References mlir::OpBuilder::create(), mlir::failure(), mlir::AffineExpr::floorDiv(), mlir::AffineMap::get(), mlir::Builder::getAffineDimExpr(), mlir::Builder::getAffineSymbolExpr(), mlir::Builder::getContext(), mlir::affine::AffineBound::getMap(), mlir::affine::AffineBound::getOperands(), mlir::AffineMap::getResults(), mlir::replaceAllUsesInRegionWith(), mlir::OpBuilder::setInsertionPointToStart(), and mlir::success().

◆ coalescePerfectlyNestedLoops()

template<typename LoopOpTy >
LogicalResult mlir::affine::coalescePerfectlyNestedLoops ( LoopOpTy  op)

Walk either an scf.for or an affine.for to find a band to coalesce.

Definition at line 304 of file LoopUtils.h.

◆ 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 1417 of file Utils.cpp.

References addMissingLoopIVBounds(), mlir::FlatLinearValueConstraints::areVarsAlignedWithOther(), checkMemrefAccessDependence(), mlir::affine::ComputationSliceState::clearBounds(), mlir::affine::FlatAffineValueConstraints::convertLoopIVSymbolsToDims(), mlir::failed(), mlir::affine::DependenceResult::Failure, mlir::affine::SliceComputationResult::GenericFailure, mlir::affine::ComputationSliceState::getAsConstraints(), getComputationSliceState(), getContext(), getInnermostCommonLoopDepth(), getNestingDepth(), mlir::presburger::IntegerRelation::getNumDimAndSymbolVars(), mlir::presburger::IntegerRelation::getNumDimVars(), mlir::presburger::IntegerRelation::getNumLocalVars(), mlir::FlatLinearConstraints::getSliceBounds(), mlir::FlatLinearValueConstraints::getValue(), mlir::FlatLinearValueConstraints::getValues(), mlir::affine::SliceComputationResult::IncorrectSliceFailure, mlir::affine::ComputationSliceState::insertPoint, mlir::affine::ComputationSliceState::isSliceValid(), mlir::affine::ComputationSliceState::ivs, mlir::affine::ComputationSliceState::lbOperands, mlir::affine::ComputationSliceState::lbs, mlir::affine::MemRefAccess::memref, mlir::FlatLinearValueConstraints::mergeAndAlignVarsWithOther(), mlir::affine::DependenceResult::NoDependence, mlir::affine::MemRefAccess::opInst, mlir::affine::SliceComputationResult::Success, mlir::affine::ComputationSliceState::ubOperands, mlir::affine::ComputationSliceState::ubs, mlir::FlatLinearValueConstraints::unionBoundingBox(), and mlir::affine::DependenceResult::value.

◆ createAffineComputationSlice()

void mlir::affine::createAffineComputationSlice ( Operation opInst,
SmallVectorImpl< AffineApplyOp > *  sliceOps 
)

Given an operation, inserts one or more single result affine apply operations, results of which are exclusively used by this operation.

Given an operation, inserts one or more single result affine apply operations, results of which are exclusively used by this operation 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.

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"(idx, A, ...) "compute"(idx)

After

affine.for i = 0 to #map(N) idx = affine.apply (d0) -> (d0 mod 2) (i) "send"(idx, A, ...) idx_ = affine.apply (d0) -> (d0 mod 2) (i) "compute"(idx_)

This allows applying different transformations on send and compute (for eg. different shifts/delays).

Returns nullptr either if none of opInst's operands were the result of an affine.apply and thus there was no affine computation slice to create, or if all the affine.apply op's supplying operands to this opInst did not have any uses besides this opInst; otherwise returns the list of affine.apply operations created in output argument sliceOps.

Definition at line 1383 of file Utils.cpp.

References mlir::Operation::getNumOperands(), mlir::Operation::getOperands(), and getReachableAffineApplyOps().

◆ createAffineDataCopyGenerationPass() [1/2]

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

Overload relying on pass options for initialization.

Definition at line 99 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.

Only load op's handled for now.

Definition at line 91 of file AffineDataCopyGeneration.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 70 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 249 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.

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

Definition at line 95 of file AffineParallelize.cpp.

◆ 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 44 of file AffineScalarReplacement.cpp.

◆ createCanonicalizedAffineForOp()

AffineForOp mlir::affine::createCanonicalizedAffineForOp ( OpBuilder  b,
Location  loc,
ValueRange  lbOperands,
AffineMap  lbMap,
ValueRange  ubOperands,
AffineMap  ubMap,
int64_t  step = 1 
)

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.

Definition at line 2558 of file LoopUtils.cpp.

References canonicalizeMapAndOperands(), mlir::OpBuilder::create(), fullyComposeAffineMapAndOperands(), and mlir::removeDuplicateExprs().

Referenced by createFullTiles(), and generatePointWiseCopy().

◆ 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 52 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 1423 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 72 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 68 of file LoopTiling.cpp.

◆ createLoopUnrollAndJamPass()

std::unique_ptr< OperationPass< func::FuncOp > > 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 79 of file LoopUnrollAndJam.cpp.

◆ createLoopUnrollPass()

std::unique_ptr< OperationPass< func::FuncOp > > mlir::affine::createLoopUnrollPass ( int  unrollFactor = -1,
bool  unrollUpToFactor = false,
bool  unrollFull = 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 148 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 56 of file PipelineDataTransfer.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 96 of file DecomposeAffineOps.cpp.

◆ defaultFilterFunction()

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

Definition at line 92 of file NestedMatcher.h.

◆ delinearizeIndex()

FailureOr< SmallVector< Value > > mlir::affine::delinearizeIndex ( OpBuilder b,
Location  loc,
Value  linearIndex,
ArrayRef< Value basis 
)

Generate the IR to delinearize linearIndex given the basis and return the multi-index.

Definition at line 1849 of file Utils.cpp.

References mlir::failed(), mlir::failure(), getDivMod(), getIndexProduct(), mlir::getValueOrCreateConstantIndexOp(), mlir::affine::DivModValue::quotient, and mlir::affine::DivModValue::remainder.

Referenced by mlir::linalg::unrollIndex().

◆ expandAffineExpr()

mlir::Value mlir::affine::expandAffineExpr ( OpBuilder builder,
Location  loc,
AffineExpr  expr,
ValueRange  dimValues,
ValueRange  symbolValues 
)

Emit code that computes the given affine expression using standard arithmetic operations applied to the provided dimension and symbol values.

Create a sequence of operations that implement the expr applied to the given dimension and symbol values.

Definition at line 215 of file Utils.cpp.

◆ expandAffineMap()

std::optional< SmallVector< Value, 8 > > mlir::affine::expandAffineMap ( OpBuilder builder,
Location  loc,
AffineMap  affineMap,
ValueRange  operands 
)

Create a sequence of operations that implement the affineMap applied to the given operands (as it it were an AffineApplyOp).

Definition at line 225 of file Utils.cpp.

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

Referenced by lowerAffineMapMax(), and lowerAffineMapMin().

◆ 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 2581 of file AffineOps.cpp.

Referenced by 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 2588 of file AffineOps.cpp.

◆ fullyComposeAffineMapAndOperands()

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

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 1128 of file AffineOps.cpp.

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

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

◆ 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 107 of file ValueBoundsOpInterfaceImpl.cpp.

References mlir::ValueBoundsConstraintSet::computeConstantBound(), mlir::presburger::EQ, fullyComposeAffineMapAndOperands(), mlir::AffineMap::get(), mlir::Builder::getAffineDimExpr(), 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 
)

◆ gatherLoops()

void mlir::affine::gatherLoops ( func::FuncOp  func,
std::vector< SmallVector< AffineForOp, 2 >> &  depthToLoops 
)

Gathers all AffineForOps in 'func.func' grouped by loop depth.

Definition at line 2546 of file LoopUtils.cpp.

References gatherLoopsInBlock().

◆ 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 643 of file LoopFusionUtils.cpp.

Referenced by gatherProducerConsumerMemrefs(), and getMaxLoopDepth().

◆ generateCopyForMemRegion()

LogicalResult mlir::affine::generateCopyForMemRegion ( const MemRefRegion memrefRegion,
Operation analyzedOp,
const AffineCopyOptions copyOptions,
CopyGenerateResult result 
)

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.

Definition at line 2502 of file LoopUtils.cpp.

References mlir::affine::CopyGenerateResult::alloc, mlir::affine::CopyGenerateResult::copyNest, mlir::failed(), mlir::failure(), generateCopy(), mlir::Operation::getBlock(), mlir::affine::MemRefRegion::memref, mlir::affine::CopyGenerateResult::sizeInBytes, and mlir::success().

◆ 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 633 of file Utils.cpp.

Referenced by findHighestBlockForPlacement(), getComputationSliceState(), getFusionComputeCost(), getInnermostCommonLoopDepth(), insertBackwardComputationSlice(), and isFusionProfitable().

◆ 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 1876 of file Utils.cpp.

Referenced by 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 2568 of file AffineOps.cpp.

References mlir::Operation::getParentOp().

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), 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 257 of file AffineOps.cpp.

Referenced by checkMemrefAccessDependence(), isValidDim(), isValidSymbol(), mayHaveEffect(), mustReachAtInnermost(), mlir::affine::AffineDmaStartOp::verifyInvariantsImpl(), and mlir::affine::AffineDmaWaitOp::verifyInvariantsImpl().

◆ getComputationSliceState()

void mlir::affine::getComputationSliceState ( Operation depSourceOp,
Operation depSinkOp,
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 1653 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 564 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 87 of file LoopAnalysis.cpp.

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

Referenced by adjustToDivisorsOfTripCounts(), affineForOpBodySkew(), constructTiledIndexSetHyperRect(), 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 676 of file AffineAnalysis.cpp.

Referenced by isValidLoopInterchangePermutation(), and sinkSequentialLoops().

◆ getDivMod()

DivModValue mlir::affine::getDivMod ( OpBuilder b,
Location  loc,
Value  lhs,
Value  rhs 
)

Create IR to calculate (div lhs, rhs) and (mod lhs, rhs).

Definition at line 1823 of file Utils.cpp.

References mlir::bindDims(), mlir::AffineExpr::floorDiv(), mlir::Builder::getContext(), makeComposedAffineApply(), mlir::affine::DivModValue::quotient, and mlir::affine::DivModValue::remainder.

Referenced by delinearizeIndex().

◆ 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 647 of file Utils.cpp.

◆ getForInductionVarOwner()

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

◆ 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 575 of file LoopFusionUtils.cpp.

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

Referenced by isFusionProfitable().

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

Definition at line 241 of file AffineAnalysis.cpp.

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

◆ 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 1385 of file Utils.cpp.

References getAffineForIVs(), max(), and min().

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 1246 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):

  1. iv and indices of the proper type;
  2. 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 178 of file LoopAnalysis.cpp.

References isAccessIndexInvariant().

Referenced by 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 113 of file LoopAnalysis.cpp.

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

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

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

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 1953 of file Utils.cpp.

References getMemoryFootprintBytes().

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 1196 of file Utils.cpp.

References mlir::detail::divideCeil().

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 1848 of file Utils.cpp.

Referenced by affineDataCopyGenerate(), computeSliceUnion(), and getMemoryFootprintBytes().

◆ 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 1893 of file Utils.cpp.

References getAffineIVs(), and min().

Referenced by canFuseLoops(), 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 864 of file LoopUtils.cpp.

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

Referenced by getTileableBands(), 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.

Returns the sequence of AffineApplyOp Operations operation in 'affineApplyOps', which are reachable via a search starting from 'operands', and ending at operands which are not defined by AffineApplyOps.

Definition at line 191 of file AffineAnalysis.cpp.

Referenced by createAffineComputationSlice().

◆ getRelationFromMap() [1/2]

LogicalResult mlir::affine::getRelationFromMap ( AffineMap map,
FlatAffineRelation 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 493 of file AffineStructures.cpp.

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

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

◆ getRelationFromMap() [2/2]

LogicalResult mlir::affine::getRelationFromMap ( const AffineValueMap map,
FlatAffineRelation rel 
)

◆ 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 1971 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 1639 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.

Referenced by isLoopParallel(), and loopUnrollJamByFactor().

◆ getTileableBands()

void mlir::affine::getTileableBands ( func::FuncOp  f,
std::vector< SmallVector< AffineForOp, 6 >> *  bands 
)

Identify valid and profitable bands of loops to tile.

This is currently just a temporary placeholder to test the mechanics of tiled code generation. Returns all maximal outermost perfect loop nests to tile.

Definition at line 881 of file LoopUtils.cpp.

References getPerfectlyNestedLoops().

◆ 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 38 of file LoopAnalysis.cpp.

References mlir::ceilDiv(), 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().

◆ 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 178 of file AffineAnalysis.h.

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

Referenced by mustReachAtInnermost().

◆ hasNoInterveningEffect()

template<typename EffectType , typename T >
bool mlir::affine::hasNoInterveningEffect ( Operation start,
memOp 
)

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 681 of file Utils.cpp.

◆ hoistAffineIfOp()

LogicalResult mlir::affine::hoistAffineIfOp ( AffineIfOp  ifOp,
bool *  folded = nullptr 
)

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.

Definition at line 410 of file Utils.cpp.

References mlir::applyOpPatternsAndFold(), mlir::applyPatternsAndFoldGreedily(), mlir::ExistingOps, mlir::failure(), mlir::RewritePatternSet::getContext(), getOutermostInvariantForOp(), mlir::GreedyRewriteConfig::strictMode, and mlir::success().

◆ 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 1767 of file Utils.cpp.

References mlir::OpBuilder::clone(), 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 1330 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 2545 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 2553 of file AffineOps.cpp.

References isAffineForInductionVar(), and isAffineParallelInductionVar().

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol().

◆ isAffineParallelInductionVar()

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

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

Definition at line 2549 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:

  1. an induction variable iv of type AffineForOp;
  2. 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:

  1. memRefDim ~= nullptr;
  2. iv of the proper type;
  3. 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 191 of file LoopAnalysis.cpp.

References isAccessIndexInvariant().

◆ 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 164 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’ 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.

Definition at line 137 of file AffineAnalysis.cpp.

Referenced by isLoopParallel().

◆ isLoopParallel()

bool mlir::affine::isLoopParallel ( AffineForOp  forOp,
SmallVectorImpl< LoopReduction > *  parallelReductions = nullptr 
)

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.

Definition at line 99 of file AffineAnalysis.cpp.

References getSupportedReductions(), and isLoopMemoryParallel().

Referenced by isLoopParallelAndContainsReduction().

◆ 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 1962 of file Utils.cpp.

References isLoopParallel().

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 356 of file LoopAnalysis.cpp.

References mlir::detail::enumerate().

Referenced by affineForOpBodySkew().

◆ isPerfectlyNested()

bool LLVM_ATTRIBUTE_UNUSED 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 1401 of file LoopUtils.cpp.

Referenced by permuteLoops().

◆ 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 242 of file AffineOps.cpp.

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

Referenced by mlir::affine::FlatAffineValueConstraints::addInductionVarOrTerminalSymbol(), isDimOpValidSymbol(), 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 42 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 272 of file AffineOps.cpp.

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

Referenced by isValidAffineIndexOperand().

◆ 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 292 of file AffineOps.cpp.

References mlir::Value::getType(), mlir::Type::isIndex(), 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 1387 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 386 of file AffineOps.cpp.

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

Referenced by canonicalizePromotedSymbols(), mlir::affine::ComputationSliceState::getAsConstraints(), isMemRefSizeValidSymbol(), isValidAffineIndexOperand(), isValidDim(), and isValidSymbol().

◆ 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 the result of an affine apply operation with symbol arguments. *) 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 415 of file AffineOps.cpp.

References mlir::Value::getDefiningOp(), mlir::Region::getParentOp(), mlir::Operation::getParentRegion(), mlir::Value::getType(), mlir::Operation::hasTrait(), isDimOpValidSymbol(), mlir::Type::isIndex(), isTopLevelValue(), isValidSymbol(), 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:

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

Definition at line 320 of file LoopAnalysis.cpp.

◆ isVectorizableLoopBody() [2/2]

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

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

:

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

Definition at line 346 of file LoopAnalysis.cpp.

References isVectorizableLoopBodyWithOpCond().

Referenced by isVectorizableLoopPtrFactory().

◆ linearizeIndex()

OpFoldResult mlir::affine::linearizeIndex ( ArrayRef< OpFoldResult multiIndex,
ArrayRef< OpFoldResult basis,
ImplicitLocOpBuilder builder 
)

◆ 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 1016 of file LoopUtils.cpp.

References mlir::failed(), mlir::failure(), generateCleanupLoopForUnroll(), generateUnrolledLoop(), getConstantTripCount(), getLargestDivisorOfTripCount(), loopUnrollFull(), promoteIfSingleIteration(), and mlir::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 893 of file LoopUtils.cpp.

References mlir::failure(), getConstantTripCount(), loopUnrollByFactor(), promoteIfSingleIteration(), and mlir::success().

Referenced by 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 1134 of file LoopUtils.cpp.

References areInnerBoundsInvariant(), mlir::OpBuilder::clone(), mlir::OpBuilder::create(), mlir::failed(), mlir::failure(), generateCleanupLoopForUnroll(), mlir::AffineMap::get(), mlir::Builder::getAffineDimExpr(), getConstantTripCount(), getLargestDivisorOfTripCount(), mlir::arith::getReductionOp(), getSupportedReductions(), promoteIfSingleIteration(), mlir::OpBuilder::setInsertionPointAfter(), JamBlockGatherer::subBlocks, mlir::success(), mlir::Value::use_empty(), and JamBlockGatherer::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 1083 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 908 of file LoopUtils.cpp.

References getConstantTripCount(), and loopUnrollByFactor().

◆ makeComposedAffineApply() [1/2]

AffineApplyOp mlir::affine::makeComposedAffineApply ( OpBuilder b,
Location  loc,
AffineExpr  e,
ArrayRef< OpFoldResult operands 
)

◆ makeComposedAffineApply() [2/2]

AffineApplyOp mlir::affine::makeComposedAffineApply ( OpBuilder b,
Location  loc,
AffineMap  map,
ArrayRef< OpFoldResult operands 
)

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 1138 of file AffineOps.cpp.

References composeAffineMapAndOperands(), mlir::OpBuilder::create(), and mlir::foldAttributesIntoMap().

Referenced by commonLinearIdBuilderFn(), getCollapsedOutputDimFromInputShape(), mlir::linalg::getConvolvedIndex(), getDivMod(), getExpandedOutputDimFromInputShape(), getXferIndices(), invertSliceIndexing(), 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 1255 of file AffineOps.cpp.

◆ makeComposedFoldedAffineApply() [1/2]

OpFoldResult mlir::affine::makeComposedFoldedAffineApply ( OpBuilder b,
Location  loc,
AffineExpr  expr,
ArrayRef< OpFoldResult operands 
)

Variant of makeComposedFoldedAffineApply that applies to an expression.

Definition at line 1224 of file AffineOps.cpp.

References mlir::Builder::getContext(), mlir::AffineMap::inferFromExprList(), and makeComposedFoldedAffineApply().

◆ makeComposedFoldedAffineApply() [2/2]

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

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 1188 of file AffineOps.cpp.

References mlir::failed(), mlir::Builder::getContext(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::OpBuilder::getListener(), 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(), mlir::linalg::computeSliceParameters(), mlir::linalg::computeTileSizes(), createInBoundsCond(), mlir::tensor::createPadHighOp(), mlir::affine::AffineBuilder::floor(), getCompressedMaskOp(), getFlatOffsetAndStrides(), getIndexProduct(), getIndicesForLoadOrStore(), mlir::memref::getLinearizedMemRefOffsetAndSize(), getOffsetForBitwidth(), linearizeIndex(), mlir::linalg::lowerPack(), makeComposedFoldedAffineApply(), mergeOffsetsSizesAndStrides(), mlir::affine::AffineBuilder::mul(), mlir::linalg::offsetIndices(), mlir::linalg::packMatmulGreedily(), resolveIndicesIntoOpWithOffsetsAndStrides(), resolveSourceIndicesCollapseShape(), resolveSourceIndicesExpandShape(), CopyBuilder::rewrite(), and mlir::affine::AffineBuilder::sub().

◆ 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 1301 of file AffineOps.cpp.

Referenced by mlir::tensor::bubbleUpPadSlice(), buildMax(), 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 1294 of file AffineOps.cpp.

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

◆ makeComposedFoldedMultiResultAffineApply()

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

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 1235 of file AffineOps.cpp.

References mlir::AffineMap::getNumResults().

Referenced by mlir::linalg::makeTiledLoopRanges().

◆ mapLoopToProcessorIds()

void mlir::affine::mapLoopToProcessorIds ( scf::ForOp  forOp,
ArrayRef< Value processorId,
ArrayRef< Value numProcessors 
)

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

Definition at line 1761 of file LoopUtils.cpp.

References mlir::bindSymbols(), mlir::OpBuilder::create(), and mlir::AffineMap::get().

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

◆ 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 35 of file ReifyValueBounds.cpp.

References canonicalizeMapAndOperands(), mlir::OpBuilder::create(), mlir::Builder::getIndexAttr(), 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 mlir::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 62 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 184 of file AffineAnalysis.h.

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

Referenced by mayHaveEffect().

◆ normalizeAffineFor()

LogicalResult mlir::affine::normalizeAffineFor ( AffineForOp  op,
bool  promoteSingleIter = false 
)

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.

Definition at line 555 of file Utils.cpp.

◆ normalizeAffineParallel()

void mlir::affine::normalizeAffineParallel ( AffineParallelOp  op)

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.

Definition at line 491 of file Utils.cpp.

◆ normalizeMemRef()

LogicalResult mlir::affine::normalizeMemRef ( memref::AllocOp *  op)

Rewrites the memref defined by this alloc 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).

Definition at line 1687 of file Utils.cpp.

References mlir::OpBuilder::create(), createNewDynamicSizes(), mlir::Operation::erase(), mlir::failed(), mlir::failure(), getTileSizePos(), mlir::Value::getType(), mlir::Value::getUsers(), normalizeMemRefType(), and replaceAllMemRefUsesWith().

◆ normalizeMemRefType()

MemRefType mlir::affine::normalizeMemRefType ( MemRefType  memrefType)

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

References mlir::failure(), 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 ( MutableArrayRef< 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 1423 of file LoopUtils.cpp.

References mlir::detail::enumerate(), and isPerfectlyNested().

Referenced by sinkSequentialLoops().

◆ 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 65 of file AffineExpandIndexOps.cpp.

References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().

◆ 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 131 of file LoopUtils.cpp.

References mlir::OpBuilder::create(), mlir::failure(), getConstantTripCount(), mlir::Builder::getDimIdentityMap(), mlir::Operation::replaceAllUsesWith(), replaceIterArgsAndYieldResults(), and mlir::success().

Referenced by fuseLoops(), generateCleanupLoopForUnroll(), 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.

◆ registerTransformDialectExtension()

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

◆ registerValueBoundsOpInterfaceExternalModels()

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

◆ 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 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 83 of file ReifyValueBounds.cpp.

References reifyValueBound().

◆ 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 42 of file DecomposeAffineOps.cpp.

◆ replaceAllMemRefUsesWith() [1/2]

LogicalResult mlir::affine::replaceAllMemRefUsesWith ( Value  oldMemRef,
Value  newMemRef,
ArrayRef< Value extraIndices = {},
AffineMap  indexRemap = AffineMap(),
ArrayRef< Value extraOperands = {},
ArrayRef< Value symbolOperands = {},
Operation domOpFilter = nullptr,
Operation postDomOpFilter = nullptr,
bool  allowNonDereferencingOps = false,
bool  replaceInDeallocOp = false 
)

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.

domOpFilter, if non-null, restricts the replacement to only those operations that are dominated by the former; similarly, postDomOpFilter restricts replacement to only those operations that are postdominated by it.

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

Definition at line 1267 of file Utils.cpp.

References mlir::AffineMap::getNumInputs(), mlir::AffineMap::getNumResults(), mlir::AffineMap::getNumSymbols(), mlir::Operation::getParentOfType(), and mlir::Value::getType().

Referenced by createPrivateMemRef(), doubleBuffer(), and normalizeMemRef().

◆ replaceAllMemRefUsesWith() [2/2]

LogicalResult mlir::affine::replaceAllMemRefUsesWith ( Value  oldMemRef,
Value  newMemRef,
Operation op,
ArrayRef< Value extraIndices = {},
AffineMap  indexRemap = AffineMap(),
ArrayRef< Value extraOperands = {},
ArrayRef< Value symbolOperands = {},
bool  allowNonDereferencingOps = 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.

Definition at line 1094 of file Utils.cpp.

◆ 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 81 of file ViewLikeInterfaceUtils.cpp.

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

Referenced by 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

◆ 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 113 of file ViewLikeInterfaceUtils.cpp.

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

◆ separateFullTiles()

LogicalResult mlir::affine::separateFullTiles ( MutableArrayRef< AffineForOp >  nest,
SmallVectorImpl< AffineForOp > *  fullTileNest = nullptr 
)

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

Definition at line 2715 of file LoopUtils.cpp.

References createFullTiles(), createSeparationCondition(), mlir::Block::end(), mlir::failed(), mlir::failure(), mlir::Block::getOperations(), and mlir::success().

◆ 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 2066 of file Utils.cpp.

◆ 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 1980 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)

◆ substWithMin()

AffineExpr mlir::affine::substWithMin ( AffineExpr  e,
AffineExpr  dim,
AffineExpr  min,
AffineExpr  max,
bool  positivePath = true 
)

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.

Definition at line 464 of file Utils.cpp.

References mlir::Add, mlir::getAffineBinaryOpExpr(), max(), and min().

◆ 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 1636 of file LoopUtils.cpp.

References 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 1624 of file LoopUtils.cpp.

References stripmineSink().

Referenced by tile(), and mlir::tilePerfectlyNested().

◆ 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 780 of file LoopUtils.cpp.

References constructTiledIndexSetHyperRect(), constructTiledLoopNest(), extractForInductionVars(), mlir::failed(), mlir::failure(), performPreTilingChecks(), and mlir::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 822 of file LoopUtils.cpp.

References constructParametricallyTiledIndexSetHyperRect(), constructTiledLoopNest(), extractForInductionVars(), mlir::failed(), mlir::failure(), performPreTilingChecks(), and mlir::success().

◆ vectorizeAffineLoopNest()

LogicalResult mlir::affine::vectorizeAffineLoopNest ( std::vector< SmallVector< AffineForOp, 2 >> &  loops,
const VectorizationStrategy strategy 
)

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.

Definition at line 1886 of file SuperVectorize.cpp.

References mlir::failed(), mlir::failure(), vectorizeLoopNest(), and verifyLoopNesting().

◆ vectorizeAffineLoops()

void mlir::affine::vectorizeAffineLoops ( Operation parentOp,
llvm::DenseSet< Operation *, DenseMapInfo< Operation * >> &  loops,
ArrayRef< int64_t >  vectorSizes,
ArrayRef< int64_t >  fastestVaryingPattern,
const ReductionLoopMap reductionLoops = ReductionLoopMap() 
)

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.