summaryrefslogtreecommitdiffstats
path: root/mlir/test/Dialect
Commit message (Collapse)AuthorAgeFilesLines
...
* [spirv] Add bit opsDenis Khalikov2019-11-082-0/+114
| | | | | | | | | | | | | | | | | | | 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
* Swap operand order in std.view operation so that offset appears before ↵Andy Davis2019-11-071-6/+6
| | | | | | dynamic sizes in the operand list. PiperOrigin-RevId: 279114236
* Update Linalg to use std.viewNicolas Vasilache2019-11-075-276/+57
| | | | | | 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-2/+57
| | | | PiperOrigin-RevId: 279013404
* Add support for the LLVM FNeg instructionEric Schweitz2019-11-061-0/+3
| | | | | | | Closes tensorflow/mlir#216 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/216 from schweitzpgi:llvmir-fneg-op f9b5f185845d671b745ab6fc213d5d9aff044b34 PiperOrigin-RevId: 278795325
* [llvm] Allow GlobalOp to take a region for complex initializersJames Molloy2019-11-051-0/+33
| | | | | | | | | | | | | | | 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
* [NVVM] Add mma.sync operation.MLIR Team2019-11-042-0/+105
| | | | PiperOrigin-RevId: 278440547
* Update the SPV dialect type parser to use the methods on DialectAsmParser ↵River Riddle2019-11-011-50/+45
| | | | | | | | 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
* Refactor LinalgDialect::parseType to use the DialectAsmParser methods directly.River Riddle2019-11-011-0/+20
| | | | | | 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-012-25/+25
| | | | | | This greatly simplifies the implementation and removes custom parser functionality. The necessary methods are added to the DialectAsmParser. PiperOrigin-RevId: 278015983
* Delete spurious fileNicolas Vasilache2019-11-011-281/+0
| | | | PiperOrigin-RevId: 277967079
* Add Linalg pattern for producer-consumer fusionNicolas Vasilache2019-11-012-0/+363
| | | | | | | | | | | | | | | 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-011-1/+11
| | | | | | Also adds a builder method for fcmp, identical to that for icmp. PiperOrigin-RevId: 277923158
* [spirv] Add cast operationsDenis Khalikov2019-10-302-21/+177
| | | | | | | | | | | | | | | | | | | 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 basic support for declarative Linalg transformationsNicolas Vasilache2019-10-301-0/+72
| | | | | | | | | 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-302-0/+44
| | | | | | | | | 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] Support OpPhi using block argumentsLei Zhang2019-10-281-0/+165
| | | | | | | | | | 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
* [spirv] AccessChainOp canonicalization.Denis Khalikov2019-10-241-0/+58
| | | | | | | | | | 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-246-6/+6
| | | | | | 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
* Update loop.for verifier messageUday Bondhugula2019-10-221-2/+2
| | | | | | | | | 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-213-26/+32
| | | | | | | | | | | | 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
* Add a Symbol trait to simplify defining operations that represent symbols.River Riddle2019-10-211-1/+1
| | | | | | This trait provides accessors for the name, symbol use list methods, verification, with more to be added. PiperOrigin-RevId: 275864554
* Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-206-7/+7
| | | | | | Closes tensorflow/mlir#175 PiperOrigin-RevId: 275726876
* Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-201-1/+1
| | | | | | Closes tensorflow/mlir#177 PiperOrigin-RevId: 275692653
* Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-192-0/+33
| | | | | | | | 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
* Lower vector transfer ops to loop.for operations.Nicolas Vasilache2019-10-183-4/+4
| | | | | | 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-1/+1
| | | | | | constraints. PiperOrigin-RevId: 275448372
* Add gpu.barrier op to synchronize invocations of a local workgroup.Christian Sigg2019-10-181-0/+2
| | | | | | | | Adding gen table for rewrite patterns from GPU to NVVM dialect. Copy missing op documentation from GPUOps.td to GPU.md. PiperOrigin-RevId: 275419588
* Decouple Linalg promotion from Linalg tiling - NFCNicolas Vasilache2019-10-171-41/+70
| | | | | | 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/+231
| | | | | | | | | | | 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
* Allow '_' when pretty printing dialect symbolsLei Zhang2019-10-171-2/+2
| | | | | | | | | '_' is used frequently enough as the separator of words in symbols. We should allow it in dialect symbols when considering pretty printing. Also updated LangRef.md regarding pretty form. PiperOrigin-RevId: 275312494
* [spirv] Implement inliner interfaceLei Zhang2019-10-161-0/+182
| | | | | | | | | | | | | | We just need to implement a few interface hooks to DialectInlinerInterface and CallOpInterface to gain the benefits of an inliner. :) Right now only supports some trivial cases: * Inlining single block with spv.Return/spv.ReturnValue * Inlining multi block with spv.Return * Inlining spv.selection/spv.loop without return ops More advanced cases will require block argument and Phi support. PiperOrigin-RevId: 275151132
* Support custom accumulator provided as region to gpu.all_reduce.Christian Sigg2019-10-162-1/+79
| | | | | | | | | | In addition to specifying the type of accumulation through the 'op' attribute, the accumulation can now also be specified as arbitrary code region. Adds a gpu.yield op to specify the result of the accumulation. Also support more types (integers) and accumulations (mul). PiperOrigin-RevId: 275065447
* Add support for OpBitwiseOr, OpBitwiseXor, and OpBitwiseAnd in SPIR-V dialect.Hanhan Wang2019-10-151-0/+77
| | | | PiperOrigin-RevId: 274935374
* [spirv] Add support for SpecId decoration on spv.specConstantLei Zhang2019-10-152-5/+15
| | | | | | | | The SpecId decoration is the handle for providing external specialization. Similar to descriptor set and binding on global variables, we directly bake it into assembly parsing and printing. PiperOrigin-RevId: 274893879
* Fix linalg.subview behavior in (partially) static cases.Nicolas Vasilache2019-10-142-20/+52
| | | | | | | | | | | | | | | | | | | | | When the implementation of the strided memref [RFC](https://groups.google.com/a/tensorflow.org/forum/#!msg/mlir/MaL8m2nXuio/1scRqZa6AQAJ) landed, linalg started using this type instead of the now retired !linalg.view. As static and partially static cases appear, the stride information needs to be maintained properly. In particular, the result type of the subview op was generally incorrect. This CL fixes the issue by computing a return type that: 1. always has dynamic sizes, which is generally the only correct way to construct a subview in the absence of data padding and/or code versioning. 2. has the same strides as the base strided memref. Point 1. above can be further refined but will needs further analysis and canonicalization to optimize the particular case where: 1. The base memref has static size along a given dimension. 2. The subview size can be statically derived (e.g. after canonicalization). 3. *And* the subview size is an even divisor of the base memref. This 3rd constraint is well-known in the case of tiled layouts that don't assume implicit padding: the boundary tile may be only partial and has size given by `problem_size % tile_size`. Tests are updated as appropriate. PiperOrigin-RevId: 274578624
* Add lowering of VectorOps dialect to LLVM to the Linalg LLVM lowering passNicolas Vasilache2019-10-141-0/+42
| | | | | | | | | This fixes an omission that prevents Linalg to lower generic ops regions operating on ops in the VectorOps dialect. To achieve this we simply need to `populateVectorToLLVMConversionPatterns` in the conversion. Relevant tests are added. PiperOrigin-RevId: 274577325
* Add LLVM IR dialect hooks for FP128 and X86_FP80 typesEric Schweitz2019-10-111-0/+7
| | | | | | | Closes tensorflow/mlir#184 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/184 from schweitzpgi:more-float-types ca27d00510a86ffc9c79c65fb3a0193b5ea097a0 PiperOrigin-RevId: 274288813
* LLVM Dialect: introduce llvm.mlir.null operationAlex Zinenko2019-10-112-0/+17
| | | | | | | | Similarly to `llvm.mlir.undef`, this auxiliary operation creates an SSA value that corresponds to `null` in LLVM IR. This operation is necessary to model sizeof(<...>) behavior when allocating memory. PiperOrigin-RevId: 274158760
* Add lowering of constant ops to SPIR-V.Mahesh Ravishankar2019-10-101-46/+0
| | | | | | | | | | | The lowering is specified as a pattern and is done only if the result is a SPIR-V scalar type or vector type. Handling ConstantOp with index return type needs special handling since SPIR-V dialect does not have index types. Based on the bitwidth of the attribute value, either i32 or i64 is chosen. Other constant lowerings are left as a TODO. PiperOrigin-RevId: 274056805
* [spirv] Add a pass to decorate the composite types with layout info.Denis Khalikov2019-10-081-0/+99
| | | | | | | | | | | Add a pass to decorate the composite types used by composite objects in the StorageBuffer, PhysicalStorageBuffer, Uniform, and PushConstant storage classes with layout information. Closes tensorflow/mlir#156 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/156 from denis0x0D:sandbox/layout_info_decoration 7c50840fd38ca169a2da7ce9886b52b50c868b84 PiperOrigin-RevId: 273634140
* Use named modules for gpu.launch_funcAlex Zinenko2019-10-083-123/+215
| | | | | | | | | | | | | | | | | | The kernel function called by gpu.launch_func is now placed into an isolated nested module during the outlining stage to simplify separate compilation. Until recently, modules did not have names and could not be referenced. This limitation was circumvented by introducing a stub kernel at the same name at the same nesting level as the module containing the actual kernel. This relation is only effective in one direction: from actual kernel function to its launch_func "caller". Leverage the recently introduced symbol name attributes on modules to refer to a specific nested module from `gpu.launch_func`. This removes the implicit connection between the identically named stub and kernel functions. It also enables support for `gpu.launch_func`s to call different kernels located in the same module. PiperOrigin-RevId: 273491891
* [spirv] Disable a crashing spv.loop testLei Zhang2019-10-071-98/+99
| | | | PiperOrigin-RevId: 273379318
* Update UndefOp (de)serialization to generate OpUndef at module level.Mahesh Ravishankar2019-10-072-14/+32
| | | | | | | | | | | | | The SPIR-V spec recommends all OpUndef instructions be generated at module level. For the SPIR-V dialect its better for UndefOp to produce an SSA value for use with other instructions. If UndefOp is to be used at module level, it cannot produce an SSA value (use of this SSA value within FuncOp would need implicit capture). To satisfy needs of the SPIR-V spec while making it simpler to represent UndefOp in the SPIR-V dialect, the serialization is updated to create OpUndef instruction at module scope. PiperOrigin-RevId: 273355526
* [spirv] Fix function entry block erase after moving to spv.selectionLei Zhang2019-10-072-5/+157
| | | | | | | | | | | | | | | | The structured selection/loop's entry block does not have arguments. If the function's header block is also part of the structured control flow, we cannot just simply erase it because it may contain arguments matching the function signature and used by the cloned blocks. Instead, turn it into a block only containing a spv.Branch op. Also, we can directly emit instructions for the spv.selection header block to the block containing the spv.selection op. This eliminates unnecessary branches in the SPIR-V blob. Added a test for nested spv.loop. PiperOrigin-RevId: 273351424
* Support AllocOp terminal in Linalg::AliasAnalysis.Nicolas Vasilache2019-10-071-0/+43
| | | | | | | | Now that linalg.view and strided memrefs are unified, there is no reason to disallow AllocOp in alias analysis. This CLs adds support for AllocOp which allows writing shorter tests that do not require explicitly creating a view for each operation. PiperOrigin-RevId: 273303060
* [spirv] Allow return ops to be in control flow opsLei Zhang2019-10-041-0/+66
| | | | | | | Use `getParentOfType<FunctionOp>()` instead of `cast<FuncOp>(getParentOp())` to avoid crash when return ops are used inside spv.selection/spv.loop. PiperOrigin-RevId: 273006041
* Add spv.Undef op to support OpUndef instruction in SPIR-V.Mahesh Ravishankar2019-10-042-0/+44
| | | | | | | | Adding support for OpUndef instruction. Updating the dialect generation script to fix a few bugs in the instruction spec generation. PiperOrigin-RevId: 272975685
* Add missing Linalg lowerings to allow roundtrip.mlir to lower to LLVMNicolas Vasilache2019-10-043-7/+12
| | | | | | | | | | | | Certain lowering patterns were reported as [missing](https://groups.google.com/a/tensorflow.org/forum/#!topic/mlir/dkdmHa77sSQ). This CL adds them and allows Linalg/roundtrip.mlir and Linalg/loops.mlir to lower to LLVM directly. Those 2 tests are updated to additionally check that the direct lowering to LLVM does not crash. The following points, left as TODOs still need to be addressed for correct end-to-end execution: 1. the lowering for ConvOp needs to pass attributes such as strides and dilations; the external library call needs to support it. 2. the lowering for GenericOp needs to support lowering to loops as a DialectConversion pattern. This is blocked on the DialectConversion infrastructure accepting an OperationFolder. PiperOrigin-RevId: 272878131
* Add `axis` attribute to the quant.stats opFeng Liu2019-10-032-7/+23
| | | | | | | The first dim length of the axisStats attribute should equals to the slice size of the input argument when splitted by the axis dimension. PiperOrigin-RevId: 272798042
OpenPOWER on IntegriCloud