summaryrefslogtreecommitdiffstats
path: root/mlir/lib/Transforms/Utils
Commit message (Collapse)AuthorAgeFilesLines
...
* NFC: Standardize the terminology used for parent ops/regions/etc.River Riddle2019-08-092-6/+6
| | | | | | There are currently several different terms used to refer to a parent IR unit in 'get' methods: getParent/getEnclosing/getContaining. This cl standardizes all of these methods to use 'getParent*'. PiperOrigin-RevId: 262680287
* NFC: Update usages of OwningRewritePatternList to pass by & instead of &&.River Riddle2019-08-091-1/+1
| | | | | | This will allow for reusing the same pattern list, which may be costly to continually reconstruct, on multiple invocations. PiperOrigin-RevId: 262664599
* Add utility 'replaceAllUsesWith' methods to Operation.River Riddle2019-08-071-5/+2
| | | | | | These methods will allow replacing the uses of results with an existing operation, with the same number of results, or a range of values. This removes a number of hand-rolled result replacement loops and simplifies replacement for operations with multiple results. PiperOrigin-RevId: 262206600
* Remove ops in regions/blocks from worklist when parent op is being removed ↵Andy Ly2019-08-061-0/+4
| | | | | | | | via GreedyPatternRewriteDriver::replaceOp. This fixes a bug where ops inside the parent op are visited even though the parent op has been removed. PiperOrigin-RevId: 261953580
* NFC: Implement OwningRewritePatternList as a class instead of a using directive.River Riddle2019-08-051-3/+3
| | | | | | This allows for proper forward declaration, as opposed to leaking the internal implementation via a using directive. This also allows for all pattern building to go through 'insert' methods on the OwningRewritePatternList, replacing uses of 'push_back' and 'RewriteListBuilder'. PiperOrigin-RevId: 261816316
* Refactor LoopParametricTiling as a test pass - NFCNicolas Vasilache2019-07-221-4/+76
| | | | | | This CL moves LoopParametricTiling into test/lib as a pass for purely testing purposes. PiperOrigin-RevId: 259300264
* Refactor stripmineSink for AffineForOp - NFCNicolas Vasilache2019-07-191-34/+9
| | | | | | More moving less cloning. PiperOrigin-RevId: 258947575
* Utility function to map a loop on a parametric grid of virtual processorsNicolas Vasilache2019-07-191-0/+22
| | | | | | | | | | | | | | | | | | | | | | This CL introduces a simple loop utility function which rewrites the bounds and step of a loop so as to become mappable on a regular grid of processors whose identifiers are given by SSA values. A corresponding unit test is added. For example, using CUDA terminology, and assuming a 2-d grid with processorIds = [blockIdx.x, threadIdx.x] and numProcessors = [gridDim.x, blockDim.x], the loop: ``` loop.for %i = %lb to %ub step %step { ... } ``` is rewritten into a version resembling the following pseudo-IR: ``` loop.for %i = %lb + threadIdx.x + blockIdx.x * blockDim.x to %ub step %gridDim.x * blockDim.x { ... } ``` PiperOrigin-RevId: 258945942
* Uniformize the API for the mlir::tile functions on AffineForOp and loop::ForOpNicolas Vasilache2019-07-191-84/+88
| | | | | | | | | | | This CL adapts the recently introduced parametric tiling to have an API matching the tiling of AffineForOp. The transformation using stripmineSink is more general and produces imperfectly nested loops. Perfect nesting invariants of the tiled version are obtained by selectively applying hoisting of ops to isolate perfectly nested bands. Such hoisting may fail to produce a perfect loop nest in cases where ForOp transitively depend on enclosing induction variables. In such cases, the API provides a LogicalResult return but the SimpleParametricLoopTilingPass does not currently use this result. A new unit test is added with a triangular loop for which the perfect nesting property does not hold. For this example, the old behavior was to produce IR that did not verify (some use was not dominated by its def). PiperOrigin-RevId: 258928309
* Move affine.for and affine.if to ODSNicolas Vasilache2019-07-161-9/+9
| | | | | | As the move to ODS is made, body and region names across affine and loop dialects are uniformized. PiperOrigin-RevId: 258416590
* Update 'applyPatternsGreedily' to work on the regions of any operations.River Riddle2019-07-161-20/+24
| | | | | | 'applyPatternsGreedily' is a useful utility outside of just function regions. PiperOrigin-RevId: 258182937
* Introduce loop coalescing utility and a simple passAlex Zinenko2019-07-161-0/+144
| | | | | | | | | | | Multiple (perfectly) nested loops with independent bounds can be combined into a single loop and than subdivided into blocks of arbitrary size for load balancing or more efficient parallelism exploitation. However, MLIR wants to preserve the multi-dimensional multi-loop structure at higher levels of abstraction. Introduce a transformation that coalesces nested loops with independent bounds so that they can be further subdivided by tiling. PiperOrigin-RevId: 258151016
* Extract std.for std.if and std.terminator in their own dialectNicolas Vasilache2019-07-162-9/+12
| | | | | | | These ops should not belong to the std dialect. This CL extracts them in their own dialect and updates the corresponding conversions and tests. PiperOrigin-RevId: 258123853
* Update the PatternRewriter constructor to take a context instead of a region.River Riddle2019-07-121-3/+3
| | | | | | This will allow for cleanly using a rewriter for multiple different regions. PiperOrigin-RevId: 257845371
* Remove the 'region' field from OpBuilder.River Riddle2019-07-121-5/+5
| | | | | | This field wasn't updated as the insertion point changed, making it potentially dangerous given the multi-level of MLIR(e.g. 'createBlock' would always insert the new block in 'region'). This also allows for building an OpBuilder with just a context. PiperOrigin-RevId: 257829135
* Fix a bug in the canonicalizer when replacing constants via patterns.River Riddle2019-07-121-3/+6
| | | | | | The GreedyPatternRewriteDriver currently does not notify the OperationFolder when constants are removed as part of a pattern match. This materializes in a nasty bug where a different operation may be allocated to the same address. This causes an assertion in the OperationFolder when it gets notified of the new operations removal. PiperOrigin-RevId: 257817627
* Lower affine control flow to std control flow to LLVM dialectNicolas Vasilache2019-07-121-6/+0
| | | | | | | | | | | | | This CL splits the lowering of affine to LLVM into 2 parts: 1. affine -> std 2. std -> LLVM The conversions mostly consists of splitting concerns between the affine and non-affine worlds from existing conversions. Short-circuiting of affine `if` conditions was never tested or exercised and is removed in the process, it can be reintroduced later if needed. LoopParametricTiling.cpp is updated to reflect the newly added ForOp::build. PiperOrigin-RevId: 257794436
* Rename FunctionAttr to SymbolRefAttr.River Riddle2019-07-124-0/+4
| | | | | | This allows for the attribute to hold symbolic references to other operations than FuncOp. This also allows for removing the dependence on FuncOp from the base Builder. PiperOrigin-RevId: 257650017
* NFC: Rename Function to FuncOp.River Riddle2019-07-102-5/+5
| | | | PiperOrigin-RevId: 257293379
* Fix a test broken on some systems due to a mis-rebase.Alex Zinenko2019-07-092-2/+1
| | | | PiperOrigin-RevId: 257190161
* Implement parametric tiling on standard for loopsAlex Zinenko2019-07-091-8/+165
| | | | | | | | | | | | | | | | | | Parametric tiling can be used to extract outer loops with fixed number of iterations. This in turn enables mapping to GPU kernels on a fixed grid independently of the range of the original loops, which may be unknown statically, making the kernel adaptable to different sizes. Provide a utility function that also computes the parametric tile size given the range of the loop. Exercise the utility function through a simple pass that applies it to all top-level loop nests. Permutability or parallelism checks must be performed before calling this utility function in actual passes. Note that parametric tiling cannot be implemented in a purely affine way, although it can be encoded using semi-affine maps. The choice to implement it on standard loops is guided by them being the common representation between Affine loops, Linalg and GPU kernels. PiperOrigin-RevId: 257180251
* NFC: Remove the various "::getFunction" methods.River Riddle2019-07-082-4/+5
| | | | | | These methods assume that a function is a valid builtin top-level operation, and removing these methods allows for decoupling FuncOp and IR/. Utility "getParentOfType" methods have been added to Operation/OpState to allow for querying the first parent operation of a given type. PiperOrigin-RevId: 257018913
* NFC: Remove Region::getContainingFunction as Functions are now Operations.River Riddle2019-07-042-4/+3
| | | | PiperOrigin-RevId: 256579717
* Globally change load/store/dma_start/dma_wait operations over to ↵Andy Davis2019-07-032-36/+104
| | | | | | | | | | | affine.load/store/dma_start/dma_wait. In most places, this is just a name change (with the exception of affine.dma_start swapping the operand positions of its tag memref and num_elements operands). Significant code changes occur here: *) Vectorization: LoopAnalysis.cpp, Vectorize.cpp *) Affine Transforms: Transforms/Utils/Utils.cpp PiperOrigin-RevId: 256395088
* NFC: Refactor Function to be value typed.River Riddle2019-07-012-3/+3
| | | | | | Move the data members out of Function and into a new impl storage class 'FunctionStorage'. This allows for Function to become value typed, which will greatly simplify the transition of Function to FuncOp(given that FuncOp is also value typed). PiperOrigin-RevId: 255983022
* Add a folder-based EDSC intrinsics constructor (NFC)Nicolas Vasilache2019-07-011-0/+18
| | | | PiperOrigin-RevId: 255908660
* Update the OperationFolder to find a valid insertion point when ↵River Riddle2019-06-252-44/+59
| | | | | | | | materializing constants. The OperationFolder currently just inserts into the entry block of a Function, but regions may be isolated above, i.e. explicit capture only, and blindly inserting constants may break the invariants of these regions. PiperOrigin-RevId: 254987796
* Add a new dialect hook 'materializeConstant' to create a constant operation ↵River Riddle2019-06-222-80/+124
| | | | | | that materializes an attribute value with the given type. This effectively adds support for dialect specific constant values that have different invariants than std.constant. 'OperationFolder' is updated to use this new hook, or attempt to default to std.constant when legal. PiperOrigin-RevId: 254570153
* NFC: Remove the 'context' parameter from OperationState.River Riddle2019-06-221-2/+1
| | | | | | Now that Locations are Attributes they contain a direct reference to the MLIRContext, i.e. the context can be directly accessed from the given location instead of being explicitly passed in. PiperOrigin-RevId: 254568329
* Uniformize usage of OpBuilder& (NFC)Nicolas Vasilache2019-06-221-19/+19
| | | | | | | Historically the pointer-based version of builders was used. This CL uniformizes to OpBuilder & PiperOrigin-RevId: 254280885
* Factor fusion compute cost calculation out of LoopFusion and into ↵Andy Davis2019-06-191-0/+234
| | | | | | LoopFusionUtils (NFC). PiperOrigin-RevId: 253797886
* Factor Region::getUsedValuesDefinedAbove into Transforms/RegionUtilsAlex Zinenko2019-06-191-0/+24
| | | | | | | | Arguably, this function is only useful for transformations and should not pollute the main IR. Also make sure it accepts a the resulting container by-reference instead of returning it. PiperOrigin-RevId: 253622981
* LoopFusion: adds support for computing forward computation slices, which ↵Andy Davis2019-06-191-15/+22
| | | | | | will enable fusion of consumer loop nests into their producers in subsequent CLs. PiperOrigin-RevId: 253601994
* Convert a nest affine loops to a GPU kernelAlex Zinenko2019-06-191-0/+31
| | | | | | | | | This converts entire loops into threads/blocks. No check on the size of the block or grid, or on the validity of parallelization is performed, it is under the responsibility of the caller to strip-mine the loops and to perform the dependence analysis before calling the conversion. PiperOrigin-RevId: 253189268
* Add utility 'create' methods to OperationFolder that will create an ↵River Riddle2019-06-091-1/+1
| | | | | | operation with a given OpBuilder and automatically try to fold it, similarly to OpBuilder::createOrFold. The difference here is that these methods enable folding to constants in addition to existing values. This functionality is then used to replace linalg::FunctionConstants. PiperOrigin-RevId: 251716247
* NFC: Rename FuncBuilder to OpBuilder and refactor to take a top level region ↵River Riddle2019-06-094-21/+22
| | | | | | instead of a function. PiperOrigin-RevId: 251563898
* NFC: Rename FoldHelper to OperationFolder and split a large function in two.River Riddle2019-06-092-30/+51
| | | | PiperOrigin-RevId: 251485843
* Remove "size" property of affine maps.MLIR Team2019-06-012-11/+8
| | | | | | -- PiperOrigin-RevId: 250572818
* LoopFusionUtils CL 2/n: Factor out and generalize slice union computation.Andy Davis2019-06-011-7/+50
| | | | | | | | | | | *) Factors slice union computation out of LoopFusion into Analysis/Utils (where other iteration slice utilities exist). *) Generalizes slice union computation to take the union of slices computed on all loads/stores pairs between source and destination loop nests. *) Fixes a bug in FlatAffineConstraints::addSliceBounds where redundant constraints were added. *) Takes care of a TODO to expose FlatAffineConstraints::mergeAndAlignIds as a public method. -- PiperOrigin-RevId: 250561529
* Affine Loop Fusion Utility Module (1/n).Andy Davis2019-06-011-0/+202
| | | | | | | | | | *) Adds LoopFusionUtils which will expose a set of loop fusion utilities (e.g. dependence checks, fusion cost/storage reduction, loop fusion transformation) for use by loop fusion algorithms. Support for checking block-level fusion-preventing dependences is added in this CL (additional loop fusion utilities will be added in subsequent CLs). *) Adds TestLoopFusion test pass for testing LoopFusionUtils at a fine granularity. *) Adds unit test for testing dependence check for block-level fusion-preventing dependences. -- PiperOrigin-RevId: 249861071
* Refactor FunctionAttr to hold the internal function reference by name ↵River Riddle2019-06-011-42/+0
| | | | | | | | | | | | | | | | instead of pointer. The one downside to this is that the function reference held by a FunctionAttr needs to be explicitly looked up from the parent module. This provides several benefits though: * There is no longer a need to explicitly remap function attrs. - This removes a potentially expensive call from the destructor of Function. - This will enable some interprocedural transformations to now run intraprocedurally. - This wasn't scalable and forces dialect defined attributes to override a virtual function. * Replacing a function is now a trivial operation. * This is a necessary first step to representing functions as operations. -- PiperOrigin-RevId: 249510802
* Add user iterators to IRObjects, i.e. Values.River Riddle2019-05-202-13/+8
| | | | | | -- PiperOrigin-RevId: 248877752
* Rewrite the DialectOpConversion patterns to inherit from RewritePattern ↵River Riddle2019-05-201-2/+2
| | | | | | | | instead of Pattern. This simplifies the infrastructure a bit by being able to reuse PatternRewriter and the RewritePatternMatcher, but also starts to lay the groundwork for a more generalized legalization framework that can operate on DialectOpConversions as well as normal RewritePatterns. -- PiperOrigin-RevId: 248836492
* Refactor PatternRewriter to inherit from FuncBuilder instead of Builder. ↵River Riddle2019-05-201-8/+4
| | | | | | | | This is necessary for allowing more complicated rewrites in the future that may do things like update the insertion point (e.g. for rewrites involving regions). -- PiperOrigin-RevId: 248803153
* Unify the 'constantFold' and 'fold' hooks on an operation into just ↵River Riddle2019-05-202-70/+53
| | | | | | | | 'fold'. This new unified fold hook will take constant attributes as operands, and may return an existing 'Value *' or a constant 'Attribute' when folding. This removes the awkward situation where a simple canonicalization like "sub(x,x)->0" had to be written as a canonicalization pattern as opposed to a fold. -- PiperOrigin-RevId: 248582024
* Factor out loop interchange code from LoopFusion into LoopUtils (NFC).Andy Davis2019-05-201-0/+123
| | | | | | -- PiperOrigin-RevId: 247926512
* Replace Operation::isa with llvm::isa.River Riddle2019-05-203-12/+8
| | | | | | -- PiperOrigin-RevId: 247789235
* Replace Operation::cast with llvm::cast.River Riddle2019-05-201-2/+2
| | | | | | -- PiperOrigin-RevId: 247785983
* Add support for using llvm::dyn_cast/cast/isa for operation casts and ↵River Riddle2019-05-202-2/+2
| | | | | | | | replace usages of Operation::dyn_cast with llvm::dyn_cast. -- PiperOrigin-RevId: 247780086
* Automated rollback of changelist 247778391.MLIR Team2019-05-202-2/+2
| | | | PiperOrigin-RevId: 247778691
OpenPOWER on IntegriCloud