summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAgeFilesLines
...
| * NFC: Rename parseOptionalAttributeDict -> parseOptionalAttrDict to match the ↵River Riddle2019-11-0515-82/+79
| | | | | | | | | | | | name of the print method. PiperOrigin-RevId: 278696668
| * Add a PatternRewriter hook to merge blocks, and use it to support for ↵River Riddle2019-11-057-7/+112
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | folding branches. A pattern rewriter hook, mergeBlock, is added that allows for merging the operations of one block into the end of another. This is used to support a canonicalization pattern for branch operations that folds the branch when the successor has a single predecessor(the branch block). Example: ^bb0: %c0_i32 = constant 0 : i32 br ^bb1(%c0_i32 : i32) ^bb1(%x : i32): return %x : i32 becomes: ^bb0: %c0_i32 = constant 0 : i32 return %c0_i32 : i32 PiperOrigin-RevId: 278677825
| * Emit empty lines after headers when generating op docsLei Zhang2019-11-051-7/+7
| | | | | | | | | | | | | | | | | | This makes the generated doc easier to read and it is also more friendly to certain markdown parsers like kramdown. Fixes tensorflow/mlir#221 PiperOrigin-RevId: 278643469
| * Rename Region::RegionType to Region::BlockListTypeSean Silva2019-11-052-10/+10
| | | | | | | | | | | | Region::RegionType doesn't make much sense when reading it. It's just a list of blocks. So call it that. PiperOrigin-RevId: 278632500
| * [NVVM] Add mma.sync operation.MLIR Team2019-11-046-3/+213
| | | | | | | | PiperOrigin-RevId: 278440547
| * Update the SPV dialect type parser to use the methods on DialectAsmParser ↵River Riddle2019-11-017-315/+260
| | | | | | | | | | | | | | | | directly. This simplifies the implementation quite a bit, and removes the need for explicit string munging. One change is made to some of the enum elements of SPV_DimAttr to ensure that they are proper identifiers; The string form is now prefixed with 'Dim'. PiperOrigin-RevId: 278027132
| * Drop spurious debug spew.Nicolas Vasilache2019-11-011-2/+0
| | | | | | | | PiperOrigin-RevId: 278023371
| * Refactor LinalgDialect::parseType to use the DialectAsmParser methods directly.River Riddle2019-11-014-39/+69
| | | | | | | | | | | | This simplifies the implementation, and removes the need to do explicit string manipulation. A utility method 'parseDimensionList' is added to the DialectAsmParser to simplify defining types and attributes that contain shapes. PiperOrigin-RevId: 278020604
| * Refactor QuantOps TypeParser to use the DialectAsmParser methods directly.River Riddle2019-11-015-524/+516
| | | | | | | | | | | | This greatly simplifies the implementation and removes custom parser functionality. The necessary methods are added to the DialectAsmParser. PiperOrigin-RevId: 278015983
| * Remove the need for passing a location to parseAttribute/parseType.River Riddle2019-11-0114-56/+120
| | | | | | | | | | | | | | | | Now that a proper parser is passed to these methods, there isn't a need to explicitly pass a source location. The source location can be recovered from the parser as necessary. This removes the need to explicitly decode an SMLoc in the case where we don't need to, which can be expensive. This requires adding some basic nesting support to the parser for supporting nested parsers to allow for remapping source locations of the nested parsers to the top level parser for accurate diagnostics. This is due to the fact that the attribute and type parsers use different source buffers than the top level parser, as they may be represented in string form. PiperOrigin-RevId: 278014858
| * Add DialectAsmParser/Printer classes to simplify dialect attribute and type ↵River Riddle2019-11-0113-111/+413
| | | | | | | | | | | | | | | | | | | | | | | | | | | | parsing. These classes are functionally similar to the OpAsmParser/Printer classes and provide hooks for parsing attributes/tokens/types/etc. This change merely sets up the base infrastructure and updates the parser hooks, followups will add hooks as needed to simplify existing handrolled dialect parsers. This has various different benefits: *) Attribute/Type parsing is much simpler to define. *) Dialect attributes/types that contain other attributes/types can now use aliases. *) It provides a 'spec' with which we may use in the future to auto-generate parsers/printers. *) Error messages emitted by attribute/type parsers can provide character exact locations rather than "beginning of the string" PiperOrigin-RevId: 278005322
| * Add ReferToOp attribute constraint for SymbolRefAttrLei Zhang2019-11-015-1/+37
| | | | | | | | | | | | | | This constraint can be used to limit a SymbolRefAttr to point to a specific kind of op in the closest parent with a symbol table. PiperOrigin-RevId: 278001364
| * Delete spurious fileNicolas Vasilache2019-11-011-281/+0
| | | | | | | | PiperOrigin-RevId: 277967079
| * Move BitEnumAttr from SPIRVBase.td to OpBase.tdLei Zhang2019-11-0114-346/+420
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | BitEnumAttr is a mechanism for modelling attributes whose value is a bitfield. It should not be scoped to the SPIR-V dialect and can be used by other dialects too. This CL is mostly shuffling code around and adding tests and docs. Functionality changes are: * Fixed to use `getZExtValue()` instead of `getSExtValue()` when getting the value from the underlying IntegerAttr for a case. * Changed to auto-detect whether there is a case whose value is all bits unset (i.e., zero). If so handle it specially in all helper methods. PiperOrigin-RevId: 277964926
| * Support lowering of imperfectly nested loops into GPU dialect.Mahesh Ravishankar2019-11-0112-37/+691
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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
| * Add Linalg pattern for producer-consumer fusionNicolas Vasilache2019-11-019-32/+543
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This CL adds a simple pattern for specifying producer-consumer fusion on Linalg operations. Implementing such an extension reveals some interesting properties. Since Linalg operates on a buffer abstraction, the output buffers are specified as in/out parameters to the ops. As a consequence, there are no SSA use-def chains and one cannot specify complex dag input patterns with the current infrastructure. Instead this CL uses constraints based on the existing linalg dependence analysis to focus the pattern and refine patterns based on the type of op that last wrote in a buffer. This is a very local property and is less powerful than the generic dag specification based on SSA use-def chains. This will be generalized in the future. PiperOrigin-RevId: 277931503
| * [mlir][llvm] Add missing cast opsJames Molloy2019-11-013-1/+40
| | | | | | | | | | | | Also adds a builder method for fcmp, identical to that for icmp. PiperOrigin-RevId: 277923158
| * NFC: Use #ifndef in various .td files instead of #ifdef and #elseLei Zhang2019-10-3142-186/+93
| | | | | | | | | | | | | | | | | | | | Upstream LLVM gained support for #ifndef with https://reviews.llvm.org/D61888 This is changed mechanically via the following command: find . -name "*.td" -exec sed -i -e ':a' -e 'N' -e '$!ba' -e 's/#ifdef \([A-Z_]*\)\n#else/#ifndef \1/g' {} \; PiperOrigin-RevId: 277789427
| * Add a test for lowering GPU ops that cover cases where the symbol table ↵Mehdi Amini2019-10-312-0/+45
| | | | | | | | | | | | isn't held by a ModuleOp (NFC) PiperOrigin-RevId: 277752004
| * Add a test.symbol_scope operation that has the SymbolTable Traits to the ↵Mehdi Amini2019-10-314-0/+23
| | | | | | | | | | | | Test dialect PiperOrigin-RevId: 277741687
| * LinalgDependenceGraph: add const modifiers to accessorsAlex Zinenko2019-10-314-37/+39
| | | | | | | | | | | | | | | | | | | | | | | | | | MLIR const-correctness policy is to avoid having `const` on IR objects. LinalgDependenceGraph is not an IR object but an auxiliary data structure. Furthermore, it is not updated once constructed unlike IR objects. Add const qualifiers to get* and find* methods of LinalgDependenceGraph since they are not modifying the graph. This allows transformation functions that require the dependence graph to take it by const-reference, clearly indicating that they are not modifying it (and that the graph may have to be recomputed after the transformation). PiperOrigin-RevId: 277731608
| * NFC: Simplify UseRange::iterator to just be a std::vector::const_iterator.River Riddle2019-10-301-20/+4
| | | | | | | | | | | | At some point the implementation of UseRange was more complex, but now it is just a simple wrapper around a std::vector<SymbolUse>. PiperOrigin-RevId: 277597294
| * [spirv] Add cast operationsDenis Khalikov2019-10-307-98/+564
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This CL added op definitions for a few cast operations: * OpConvertFToU * OpConvertFToS * OpConvertSToF * OpConvertUToF * OpUConvert * OpSConvert * OpFConvert Also moved the definition of spv.Bitcast to the new file. Closes tensorflow/mlir#208 and tensorflow/mlir#174 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/208 from denis0x0D:sandbox/cast_ops 79bc9b37398aafddee6cf6beb301807988fe67f9 PiperOrigin-RevId: 277587891
| * Use `not` to invert return code in expected to fail testsLei Zhang2019-10-301-1/+1
| | | | | | | | | | | | Windows does not like the RUN command of `(... || true) | ...`. PiperOrigin-RevId: 277587031
| * Dump op location in createPrintOpGraphPass for easier debugging.Jing Pu2019-10-301-0/+4
| | | | | | | | PiperOrigin-RevId: 277546527
| * Add support to GreedyPatternRewriter for erasing unreachable blocks.River Riddle2019-10-302-0/+71
| | | | | | | | | | | | Rewrite patterns may make modifications to the CFG, including dropping edges between blocks. This change adds a simple unreachable block elimination run at the end of each iteration to ensure that the CFG remains valid. PiperOrigin-RevId: 277545805
| * Add a utility accessor 'has_single_element' for ranges.River Riddle2019-10-304-4/+8
| | | | | | | | | | | | This provides an easy way to check if a range has a single element. PiperOrigin-RevId: 277544647
| * Fix segfault when no symbol is given to an constraint operandLei Zhang2019-10-302-1/+25
| | | | | | | | | | | | | | This fixed the segfault when we see the following pattern: Pat<(...), (...), [(... 1, 2, 3), ...]> PiperOrigin-RevId: 277544300
| * Add basic support for declarative Linalg transformationsNicolas Vasilache2019-10-309-0/+237
| | | | | | | | | | | | | | | | | | Linalg ops provide a good anchor for pattern matching/rewriting transformations. This CL adds a simple example of how multi-level tiling may be specified by attaching a simple StringAttr to ops as they are transformed so we can easily specify partial lowering to control transformation application. This is a first stab at taking advantage of higher-level information contained in Linalg ops and will evolve in the future. PiperOrigin-RevId: 277497958
| * [spirv] Fix gen_spirv_dialect.py and add spv.UnreachableLei Zhang2019-10-306-7/+137
| | | | | | | | | | | | | | | | | | This CL fixed gen_spirv_dialect.py to support nested delimiters when chunking existing ODS entries in .td files and to allow ops without correspondence in the spec. This is needed to pull in the definition of OpUnreachable. PiperOrigin-RevId: 277486465
| * [spirv] Mark control flow ops as InFunctionScopeLei Zhang2019-10-291-4/+5
| | | | | | | | PiperOrigin-RevId: 277373473
| * Add "[TOC]" to generated documentationMLIR Team2019-10-291-0/+3
| | | | | | | | PiperOrigin-RevId: 277354482
| * Bugfix: Keep worklistMap in sync with worklist in GreedyPatternRewriterDiego Caballero2019-10-291-0/+1
| | | | | | | | | | | | | | | | | | | | | | When we removed a pattern, we removed it from worklist but not from worklistMap. Then, when we tried to add a new pattern on the same Operation again, the pattern wasn't added since it already existed in the worklistMap (but not in the worklist). Closes tensorflow/mlir#211 PiperOrigin-RevId: 277319669
| * [spirv] Use LLVM graph traversal utility for PrettyBlockOrderVisitorLei Zhang2019-10-291-72/+33
| | | | | | | | | | | | | | This removes a bunch of special tailored DFS code in favor of the common LLVM utility. Besides, we avoid recursion with system stack given that llvm::depth_first_ext is iterator based and maintains its own stack. PiperOrigin-RevId: 277272961
| * Add a convenient operation build method for spirv::SelectOpMahesh Ravishankar2019-10-282-0/+9
| | | | | | | | | | | | | | | | The SelectOp always has the same result type as its true/false value. Add a builder method that uses the operand type to get the result type. PiperOrigin-RevId: 277217978
| * [spirv] Support OpPhi using block argumentsLei Zhang2019-10-285-116/+657
| | | | | | | | | | | | | | | | | | | | This CL adds another control flow instruction in SPIR-V: OpPhi. It is modelled as block arguments to be idiomatic with MLIR. See the rationale.md doc for "Block Arguments vs PHI nodes". Serialization and deserialization is updated to convert between block arguments and SPIR-V OpPhi instructions. PiperOrigin-RevId: 277161545
| * Parse locations in parseGenericOperationSean Silva2019-10-283-14/+21
| | | | | | | | | | | | | | | | | | For ops that recursively re-enter the parser to parse an operation (such as ops with a "wraps" pretty form), this ensures that the wrapped op will parse its location, which can then be used for the locations of the wrapping op and any other implicit ops. PiperOrigin-RevId: 277152636
| * Standardize Linalg transformations to take an OpBuilder and an ↵Nicolas Vasilache2019-10-287-81/+114
| | | | | | | | | | | | | | | | OperationFolder - NFC This will be used to specify declarative Linalg transformations in a followup CL. In particular, the PatternRewrite mechanism does not allow folding and has its own way of tracking erasure. PiperOrigin-RevId: 277149158
| * Add support for marking an operation as recursively legal.River Riddle2019-10-285-82/+257
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | In some cases, it may be desirable to mark entire regions of operations as legal. This provides an additional granularity of context to the concept of "legal". The `ConversionTarget` supports marking operations, that were previously added as `Legal` or `Dynamic`, as `recursively` legal. Recursive legality means that if an operation instance is legal, either statically or dynamically, all of the operations nested within are also considered legal. An operation can be marked via `markOpRecursivelyLegal<>`: ```c++ ConversionTarget &target = ...; /// The operation must first be marked as `Legal` or `Dynamic`. target.addLegalOp<MyOp>(...); target.addDynamicallyLegalOp<MySecondOp>(...); /// Mark the operation as always recursively legal. target.markOpRecursivelyLegal<MyOp>(); /// Mark optionally with a callback to allow selective marking. target.markOpRecursivelyLegal<MyOp, MySecondOp>([](Operation *op) { ... }); /// Mark optionally with a callback to allow selective marking. target.markOpRecursivelyLegal<MyOp>([](MyOp op) { ... }); ``` PiperOrigin-RevId: 277086382
| * Print reason why dynamic library could not be loaded during execution.Christian Sigg2019-10-281-1/+2
| | | | | | | | PiperOrigin-RevId: 277037138
| * Lookup function declaration in SymbolTable not ModuleOp.Alexander Belyaev2019-10-281-4/+2
| | | | | | | | PiperOrigin-RevId: 277033167
| * Fix include guards and add tests for OpToFuncCallLowering.Alexander Belyaev2019-10-263-7/+49
| | | | | | | | PiperOrigin-RevId: 276859463
| * Define AnyRankedTensor Type in TableGenSmit Hinsu2019-10-253-0/+33
| | | | | | | | PiperOrigin-RevId: 276714649
| * Add support for parsing multiple result name groups.River Riddle2019-10-252-48/+47
| | | | | | | | | | | | | | | | | | | | | | | | This allows for parsing things like: %name_1, %name_2:5, %name_3:2 = "my.op" ... This is useful for operations that have groups of variadic result values. The total number of results is expected to match the number of results defined by the operation. PiperOrigin-RevId: 276703280
| * [spirv] AccessChainOp canonicalization.Denis Khalikov2019-10-243-3/+99
| | | | | | | | | | | | | | | | | | | | Combine chained `spirv::AccessChainOp` operations into one `spirv::AccessChainOp` operation. Closes tensorflow/mlir#198 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/198 from denis0x0D:sandbox/canon_access_chain 0cb87955a85511071143d62637ff939d0dabc2bd PiperOrigin-RevId: 276609345
| * Convert the Canonicalize and CSE passes to generic Operation Passes.River Riddle2019-10-2425-89/+83
| | | | | | | | | | | | This allows for them to be used on other non-function, or even other function-like, operations. The algorithms are already generic, so this is simply changing the derived pass type. The majority of this change is just ensuring that the nesting of these passes remains the same, as the pass manager won't auto-nest them anymore. PiperOrigin-RevId: 276573038
| * Add support for replacing all uses of a symbol.River Riddle2019-10-244-45/+259
| | | | | | | | | | | | This requires reconstructing the attribute dictionary of each operation containing a use. PiperOrigin-RevId: 276520544
| * Add missing dependency on MLIRIR on MLIREDSCInterfaceMehdi Amini2019-10-241-0/+1
| | | | | | | | | | | | | | | | MLIRIR includes generated header for interfaces, including these headers require an extra dependency to ensure these headers are generated before we attempt to build MLIREDSCInterface. PiperOrigin-RevId: 276518255
| * Add custom lowering of ExpOp for NVVM and ROCM.Alexander Belyaev2019-10-244-1/+112
| | | | | | | | PiperOrigin-RevId: 276440911
| * Wrap ODS to 80 lines and remove const qualifier for local `int` variable (NFC)Alexander Belyaev2019-10-232-4/+5
| | | | | | | | | | | | This addresses post-submit comments on 00d2a37e32 PiperOrigin-RevId: 276419770
OpenPOWER on IntegriCloud