summaryrefslogtreecommitdiffstats
path: root/mlir/lib/Transforms/Utils/LoopUtils.cpp
Commit message (Collapse)AuthorAgeFilesLines
* Make helper functions static or move them into anonymous namespaces. NFC.Benjamin Kramer2020-01-141-1/+1
|
* [mlir] NFC: Remove Value::operator* and Value::operator-> now that Value is ↵River Riddle2020-01-111-19/+18
| | | | | | | | | | properly value-typed. Summary: These were temporary methods used to simplify the transition. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D72548
* NFC: Replace ValuePtr with Value and remove it now that Value is value-typed.River Riddle2019-12-231-77/+74
| | | | | | ValuePtr was a temporary typedef during the transition to a value-typed Value. PiperOrigin-RevId: 286945714
* Adjust License.txt file to use the LLVM licenseMehdi Amini2019-12-231-13/+4
| | | | PiperOrigin-RevId: 286906740
* NFC: Introduce new ValuePtr/ValueRef typedefs to simplify the transition to ↵River Riddle2019-12-221-84/+85
| | | | | | | | | | Value being value-typed. This is an initial step to refactoring the representation of OpResult as proposed in: https://groups.google.com/a/tensorflow.org/g/mlir/c/XXzzKhqqF_0/m/v6bKb08WCgAJ This change will make it much simpler to incrementally transition all of the existing code to use value-typed semantics. PiperOrigin-RevId: 286844725
* Add integer bit-shift operations to the standard dialect.Manuel Freiberger2019-12-221-6/+6
| | | | | | | | | | | | | | | | | | | Rename the 'shlis' operation in the standard dialect to 'shift_left'. Add tests for this operation (these have been missing so far) and add a lowering to the 'shl' operation in the LLVM dialect. Add also 'shift_right_signed' (lowered to LLVM's 'ashr') and 'shift_right_unsigned' (lowered to 'lshr'). The original plan was to name these operations 'shift.left', 'shift.right.signed' and 'shift.right.unsigned'. This works if the operations are prefixed with 'std.' in MLIR assembly. Unfortunately during import the short form is ambigous with operations from a hypothetical 'shift' dialect. The best solution seems to omit dots in standard operations for now. Closes tensorflow/mlir#226 PiperOrigin-RevId: 286803388
* NFC: Remove unnecessary 'llvm::' prefix from uses of llvm symbols declared ↵River Riddle2019-12-181-8/+8
| | | | | | | | in `mlir` namespace. Aside from being cleaner, this also makes the codebase more consistent. PiperOrigin-RevId: 286206974
* Minor spelling tweaksKazuaki Ishizaki2019-12-091-1/+1
| | | | | | Closes tensorflow/mlir#304 PiperOrigin-RevId: 284568358
* Replace spurious SmallVector constructions with ValueRangeUday Bondhugula2019-12-091-2/+1
| | | | | | | | | Signed-off-by: Uday Bondhugula <uday@polymagelabs.com> Closes tensorflow/mlir#305 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/305 from bondhugula:value_range 21d1fae73f549e3c8e72b60876eff1b864cea39c PiperOrigin-RevId: 284541027
* Update the builder API to take ValueRange instead of ArrayRef<Value *>River Riddle2019-12-071-3/+2
| | | | | | This allows for users to provide operand_range and result_range in builder.create<> calls, instead of requiring an explicit copy into a separate data structure like SmallVector/std::vector. PiperOrigin-RevId: 284360710
* Loop coalescing: fix pointer chainsing in use-chain traversalAlex Zinenko2019-12-041-1/+1
| | | | | | | | | | | | | | | In the replaceAllUsesExcept utility function called from loop coalescing the iteration over the use-chain is incorrect. The use list nodes (IROperands) have next/prev links, and bluntly resetting the use would make the loop to continue on uses of the value that was replaced instead of the original one. As a result, it could miss the existing uses and update the wrong ones. Make sure we increment the iterator before updating the use in the loop body. Reported-by: Uday Bondhugula <uday@polymagelabs.com> Closes tensorflow/mlir#291. PiperOrigin-RevId: 283754195
* NFC: Convert CmpIPredicate in StandardOps to use EnumAttrLei Zhang2019-11-151-1/+1
| | | | | | This turns several hand-written functions to auto-generated ones. PiperOrigin-RevId: 280684326
* Support lowering of imperfectly nested loops into GPU dialect.Mahesh Ravishankar2019-11-011-3/+4
| | | | | | | | | | | | | | | | | | | The current lowering of loops to GPU only supports lowering of loop nests where the loops mapped to workgroups and workitems are perfectly nested. Here a new lowering is added to handle lowering of imperfectly nested loop body with the following properties 1) The loops partitioned to workgroups are perfectly nested. 2) The loop body of the inner most loop partitioned to workgroups can contain one or more loop nests that are to be partitioned across workitems. Each individual loops nests partitioned to workitems should also be perfectly nested. 3) The number of workgroups and workitems are not deduced from the loop bounds but are passed in by the caller of the lowering as values. 4) For statements within the perfectly nested loop nest partitioned across workgroups that are not loops, it is valid to have all threads execute that statement. This is NOT verified. PiperOrigin-RevId: 277958868
* Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-201-5/+5
| | | | | | Closes tensorflow/mlir#177 PiperOrigin-RevId: 275692653
* NFC: Remove trivial builder get methods.River Riddle2019-10-171-12/+12
| | | | | | These don't add any value, and some are even more restrictive than the respective static 'get' method. PiperOrigin-RevId: 275391240
* NFC - clean up op accessor usage, std.load/store op verify, other stale infoUday Bondhugula2019-09-271-4/+1
| | | | | | | | | | | - also remove stale terminology/references in docs Signed-off-by: Uday Bondhugula <uday@polymagelabs.com> Closes tensorflow/mlir#148 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/148 from bondhugula:cleanup e846b641a3c2936e874138aff480a23cdbf66591 PiperOrigin-RevId: 271618279
* Fix a number of Clang-Tidy warnings.Christian Sigg2019-09-231-1/+0
| | | | PiperOrigin-RevId: 270632324
* Support symbolic operands for memref replacement; fix memrefNormalizeUday Bondhugula2019-09-181-0/+1
| | | | | | | | | | | | | - allow symbols in index remapping provided for memref replacement - fix memref normalize crash on cases with layout maps with symbols Signed-off-by: Uday Bondhugula <uday@polymagelabs.com> Reported by: Alex Zinenko Closes tensorflow/mlir#139 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/139 from bondhugula:memref-rep-symbols 2f48c1fdb5d4c58915bbddbd9f07b18541819233 PiperOrigin-RevId: 269851182
* Unify error messages to start with lower-case.MLIR Team2019-09-181-2/+2
| | | | PiperOrigin-RevId: 269803466
* NFC - Move explicit copy/dma generation utility out of pass and into LoopUtilsUday Bondhugula2019-09-141-6/+668
| | | | | | | | | | | | | | | | - turn copy/dma generation method into a utility in LoopUtils, allowing it to be reused elsewhere. - no functional/logic change to the pass/utility - trim down header includes in files affected Signed-off-by: Uday Bondhugula <uday@polymagelabs.com> Closes tensorflow/mlir#124 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/124 from bondhugula:datacopy 9f346e62e5bd9dd1986720a30a35f302eb4d3252 PiperOrigin-RevId: 269106088
* Refactor the 'walk' methods for operations.River Riddle2019-08-291-2/+1
| | | | | | | | | | | | This change refactors and cleans up the implementation of the operation walk methods. After this refactoring is that the explicit template parameter for the operation type is no longer needed for the explicit op walks. For example: op->walk<AffineForOp>([](AffineForOp op) { ... }); is now accomplished via: op->walk([](AffineForOp op) { ... }); PiperOrigin-RevId: 266209552
* NFC: Move AffineOps dialect to the Dialect sub-directory.River Riddle2019-08-201-1/+1
| | | | PiperOrigin-RevId: 264482571
* NFC: Move LLVMIR, SDBM, and StandardOps to the Dialect/ directory.River Riddle2019-08-191-1/+1
| | | | PiperOrigin-RevId: 264193915
* 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
* 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-161-9/+11
| | | | | | | 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
* 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-121-0/+1
| | | | | | 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-101-3/+3
| | | | PiperOrigin-RevId: 257293379
* Fix a test broken on some systems due to a mis-rebase.Alex Zinenko2019-07-091-2/+0
| | | | 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-081-1/+1
| | | | | | 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: Refactor Function to be value typed.River Riddle2019-07-011-1/+1
| | | | | | 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
* 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
* NFC: Rename FuncBuilder to OpBuilder and refactor to take a top level region ↵River Riddle2019-06-091-13/+12
| | | | | | instead of a function. PiperOrigin-RevId: 251563898
* Remove "size" property of affine maps.MLIR Team2019-06-011-7/+5
| | | | | | -- PiperOrigin-RevId: 250572818
* 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-201-3/+2
| | | | | | -- 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-201-1/+1
| | | | | | | | replace usages of Operation::dyn_cast with llvm::dyn_cast. -- PiperOrigin-RevId: 247780086
* Automated rollback of changelist 247778391.MLIR Team2019-05-201-1/+1
| | | | PiperOrigin-RevId: 247778691
* Add support for using llvm::dyn_cast/cast/isa for operation casts and ↵River Riddle2019-05-201-1/+1
| | | | | | | | replace usages of Operation::dyn_cast with llvm::dyn_cast. -- PiperOrigin-RevId: 247778391
* Add support for basic remark diagnostics. This is the minimal ↵River Riddle2019-05-061-1/+1
| | | | | | | | functionality needed to separate notes from remarks. It also provides a starting point to start building out better remark infrastructure. -- PiperOrigin-RevId: 246175216
* Create a LoopUtil function to return perfectly nested loop setMLIR Team2019-04-051-0/+16
| | | | | | -- PiperOrigin-RevId: 242019230
* Remove the non-postorder walk functions from Function/Block/Instruction ↵River Riddle2019-04-051-1/+1
| | | | | | | | and rename walkPostOrder to walk. -- PiperOrigin-RevId: 241965239
* Replace usages of Instruction with Operation in the Transforms/ directory.River Riddle2019-03-291-52/+52
| | | | PiperOrigin-RevId: 240636130
OpenPOWER on IntegriCloud