summaryrefslogtreecommitdiffstats
path: root/mlir/lib/Dialect
Commit message (Collapse)AuthorAgeFilesLines
...
* Add support for nested symbol references.River Riddle2019-11-117-20/+23
| | | | | | | | | | | | | | | | | | This change allows for adding additional nested references to a SymbolRefAttr to allow for further resolving a symbol if that symbol also defines a SymbolTable. If a referenced symbol also defines a symbol table, a nested reference can be used to refer to a symbol within that table. Nested references are printed after the main reference in the following form: symbol-ref-attribute ::= symbol-ref-id (`::` symbol-ref-id)* Example: module @reference { func @nested_reference() } my_reference_op @reference::@nested_reference Given that SymbolRefAttr is now more general, the existing functionality centered around a single reference is moved to a derived class FlatSymbolRefAttr. Followup commits will add support to lookups, rauw, etc. for scoped references. PiperOrigin-RevId: 279860501
* Adds std.subview operation which takes dynamic offsets, sizes and strides ↵Andy Davis2019-11-117-50/+152
| | | | | | | | and returns a memref type which represents sub/reduced-size view of its memref argument. This operation is a companion operation to the std.view operation added as proposed in "Updates to the MLIR MemRefType" RFC. PiperOrigin-RevId: 279766410
* Look for SymbolRefAttr in KernelOutlining instead of hard-coding CallOpMLIR Team2019-11-081-15/+17
| | | | | | | | This code should be exercised using the existing kernel outlining unit test, but let me know if I should add a dedicated unit test using a fake call instruction as well. PiperOrigin-RevId: 279436321
* [spirv] Add bit opsDenis Khalikov2019-11-081-0/+34
| | | | | | | | | | | | | | | | | | | This CL added op definitions for a few bit operations: * OpShiftLeftLogical * OpShiftRightArithmetic * OpShiftRightLogical * OpBitCount * OpBitReverse * OpNot Also moved the definition of spv.BitwiseAnd to follow the lexicographical order. Closes tensorflow/mlir#215 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/215 from denis0x0D:sandbox/bit_ops d9b0852b689ac6c4879a9740b1740a2357f44d24 PiperOrigin-RevId: 279350470
* mlir-translate: support -verify-diagnosticsAlex Zinenko2019-11-071-13/+10
| | | | | | | | | | | MLIR translation tools can emit diagnostics and we want to be able to check if it is indeed the case in tests. Reuse the source manager error handlers provided for mlir-opt to support the verification in mlir-translate. This requires us to change the signature of the functions that are registered to translate sources to MLIR: it now takes a source manager instead of a memory buffer. PiperOrigin-RevId: 279132972
* Swap operand order in std.view operation so that offset appears before ↵Andy Davis2019-11-071-30/+42
| | | | | | dynamic sizes in the operand list. PiperOrigin-RevId: 279114236
* Add canonicalizer for ViewOp which folds constants into the ViewOp memref ↵Andy Davis2019-11-071-0/+112
| | | | | | shape and layout map strides and offset. PiperOrigin-RevId: 279088023
* Update Linalg to use std.viewNicolas Vasilache2019-11-075-386/+51
| | | | | | Now that a view op has graduated to the std dialect, we can update Linalg to use it and remove ops that have become obsolete. As a byproduct, the linalg buffer and associated ops can also disappear. PiperOrigin-RevId: 279073591
* Add IndexedGenericOp to Linalg.Alexander Belyaev2019-11-062-27/+99
| | | | PiperOrigin-RevId: 279013404
* Adding an m_NonZero constant integer matcher.Ben Vanik2019-11-061-20/+13
| | | | | | This is useful for making matching cases where a non-zero value is required more readable, such as the results of a constant comparison that are expected to be equal. PiperOrigin-RevId: 278932874
* Add ViewOp verification for dynamic strides, and address some comments from ↵Andy Davis2019-11-061-12/+20
| | | | | | previous change. PiperOrigin-RevId: 278903187
* Add ViewOp to the StandardOps dialect, which casts a 1D/i8 element type ↵Andy Davis2019-11-063-4/+82
| | | | | | | | | | | memref type to an N-D memref type. Proposed in RFC: https://groups.google.com/a/tensorflow.org/forum/#!searchin/mlir/std.view%7Csort:date/mlir/-wKHANzDNTg/4K6nUAp8AAAJ Supports creating the N-D memref type with dynamic sizes and at a dynamic offset within the 1D base memref. This change contains op definition/parsing/printing and tests. Follow up changes will handle constant shape/layout map folding and llvm lowering. PiperOrigin-RevId: 278869990
* Add (parse|print)OptionalAttrDictWithKeyword hooks to simplify parsing ↵River Riddle2019-11-051-17/+3
| | | | | | | | attribute dictionaries with regions. Many operations with regions add an additional 'attributes' prefix when printing the attribute dictionary to differentiate it from the region body. This leads to duplicated logic for detecting when to actually print the attribute dictionary. PiperOrigin-RevId: 278747681
* [llvm] Allow GlobalOp to take a region for complex initializersJames Molloy2019-11-051-1/+23
| | | | | | | | | | | | | | | This allows GlobalOp to either take a value attribute (for simple constants) or a region that can contain IR instructions (that must be constant-foldable) to create a ConstantExpr initializer. Example: // A complex initializer is constructed with an initializer region. llvm.mlir.global constant @int_gep() : !llvm<"i32*"> { %0 = llvm.mlir.addressof @g2 : !llvm<"i32*"> %1 = llvm.mlir.constant(2 : i32) : !llvm.i32 %2 = llvm.getelementptr %0[%1] : (!llvm<"i32*">, !llvm.i32) -> !llvm<"i32*"> llvm.return %2 : !llvm<"i32*"> } PiperOrigin-RevId: 278717836
* NFC: Rename parseOptionalAttributeDict -> parseOptionalAttrDict to match the ↵River Riddle2019-11-0510-74/+71
| | | | | | name of the print method. PiperOrigin-RevId: 278696668
* Add a PatternRewriter hook to merge blocks, and use it to support for ↵River Riddle2019-11-051-0/+27
| | | | | | | | | | | | | | | | | | | 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
* [NVVM] Add mma.sync operation.MLIR Team2019-11-041-0/+83
| | | | PiperOrigin-RevId: 278440547
* Update the SPV dialect type parser to use the methods on DialectAsmParser ↵River Riddle2019-11-011-250/+158
| | | | | | | | 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-011-39/+31
| | | | | | 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-011-499/+143
| | | | | | 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-014-7/+9
| | | | | | | | 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-014-23/+34
| | | | | | | | | | | | | | 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
* Move BitEnumAttr from SPIRVBase.td to OpBase.tdLei Zhang2019-11-012-4/+2
| | | | | | | | | | | | | | | | | 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
* Add Linalg pattern for producer-consumer fusionNicolas Vasilache2019-11-013-22/+126
| | | | | | | | | | | | | | | 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
* LinalgDependenceGraph: add const modifiers to accessorsAlex Zinenko2019-10-312-24/+24
| | | | | | | | | | | | | 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
* [spirv] Add cast operationsDenis Khalikov2019-10-301-22/+34
| | | | | | | | | | | | | | | | | | | 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
* Add a utility accessor 'has_single_element' for ranges.River Riddle2019-10-302-3/+2
| | | | | | This provides an easy way to check if a range has a single element. PiperOrigin-RevId: 277544647
* Add basic support for declarative Linalg transformationsNicolas Vasilache2019-10-302-0/+69
| | | | | | | | | 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-301-0/+20
| | | | | | | | | 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] 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-281-0/+5
| | | | | | | | 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-282-113/+406
| | | | | | | | | | 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
* Standardize Linalg transformations to take an OpBuilder and an ↵Nicolas Vasilache2019-10-285-68/+75
| | | | | | | | 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
* [spirv] AccessChainOp canonicalization.Denis Khalikov2019-10-241-3/+39
| | | | | | | | | | 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
* Wrap ODS to 80 lines and remove const qualifier for local `int` variable (NFC)Alexander Belyaev2019-10-231-3/+3
| | | | | | This addresses post-submit comments on 00d2a37e32 PiperOrigin-RevId: 276419770
* Update loop.for verifier messageUday Bondhugula2019-10-221-1/+1
| | | | | | | | | fix: nonnegative -> positive Closes tensorflow/mlir#206 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/206 from bondhugula:bondhugula-patch-1 9a47ca7dfd230180a9df33e9a64b33d02252d30a PiperOrigin-RevId: 276060885
* [spirv] Allow block arguments on spv.Branch(Conditional)Lei Zhang2019-10-211-2/+4
| | | | | | | | | | | | We will use block arguments as the way to model SPIR-V OpPhi in the SPIR-V dialect. This CL also adds a few useful helper methods to both ops to get the block arguments. Also added tests for branch weight (de)serialization. PiperOrigin-RevId: 275960797
* Unify GPU op definition names with other dialects.Christian Sigg2019-10-212-8/+13
| | | | | | Rename GPU op names from gpu_Foo to GPU_FooOp. PiperOrigin-RevId: 275882232
* Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-2011-36/+37
| | | | | | Closes tensorflow/mlir#177 PiperOrigin-RevId: 275692653
* Use new eraseOp instead of replaceOp with empty valuesGeoffrey Martin-Noble2019-10-192-2/+2
| | | | PiperOrigin-RevId: 275631166
* Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-192-11/+24
| | | | | | | | nvvm.shfl.sync.bfly optionally returns a predicate whether source lane was active. Support for this was added to clang in https://reviews.llvm.org/D68892. Add an optional 'pred' unit attribute to the instruction to return this predicate. Specify this attribute in the partial warp reduction so we don't need to manually compute the predicate. PiperOrigin-RevId: 275616564
* NFC: Rename SPIR-V serializer find*ID() to get*ID() to be consistentLei Zhang2019-10-181-29/+29
| | | | | | We use get*() in deserizer and other places across the codebase. PiperOrigin-RevId: 275582390
* Add support for function result attributes.Sean Silva2019-10-181-0/+15
| | | | | | | | | | | | | | | | | | | | | | This allows dialect-specific attributes to be attached to func results. (or more specifically, FunctionLike ops). For example: ``` func @f() -> (i32 {my_dialect.some_attr = 3}) ``` This attaches my_dialect.some_attr with value 3 to the first result of func @f. Another more complex example: ``` func @g() -> (i32, f32 {my_dialect.some_attr = "foo", other_dialect.some_other_attr = [1,2,3]}, i1) ``` Here, the second result has two attributes attached. PiperOrigin-RevId: 275564165
* Lower vector transfer ops to loop.for operations.Nicolas Vasilache2019-10-181-1/+1
| | | | | | This allows mixing linalg operations with vector transfer operations (with additional modifications to affine ops) and is a step towards solving tensorflow/mlir#189. PiperOrigin-RevId: 275543361
* Use StrEnumAttr for gpu.allreduce op instead of StringAttr to better encode ↵Stephan Herhut2019-10-181-5/+0
| | | | | | constraints. PiperOrigin-RevId: 275448372
* NFC: Remove trivial builder get methods.River Riddle2019-10-177-30/+28
| | | | | | These don't add any value, and some are even more restrictive than the respective static 'get' method. PiperOrigin-RevId: 275391240
* Decouple Linalg promotion from Linalg tiling - NFCNicolas Vasilache2019-10-173-217/+248
| | | | | | This CL creates a new Linalg promotion pass that operates on SubViewOp and decouples it from Linalg tiling. This is mostly moving code around. PiperOrigin-RevId: 275329213
* [spirv] Add a canonicalization pattern for spv.selection.Denis Khalikov2019-10-171-0/+164
| | | | | | | | | | | Add a canonicalization pattern for spv.selection operation. Convert spv.selection operation to spv.Select based on simple pattern. Closes tensorflow/mlir#183 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/183 from denis0x0D:sandbox/canon_select 43d04d923272dd60b9da39f70bdbc51a5168db62 PiperOrigin-RevId: 275312748
* Use a SmallVector instead of an ArrayRef to materialize a temporary local arrayMehdi Amini2019-10-171-1/+1
| | | | | | | | | | | | This pattern is error prone and unfortunately none of the sanitizer is catching it at the moment. Fixes tensorflow/mlir#192 Closes tensorflow/mlir#193 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/193 from joker-eph:fix_array_ref 8092252e64c426c6a8a790b7638f847bea4818b1 PiperOrigin-RevId: 275280201
OpenPOWER on IntegriCloud