summaryrefslogtreecommitdiffstats
path: root/mlir/test
Commit message (Collapse)AuthorAgeFilesLines
...
* Cleanup and rewrite Ch-4.md.River Riddle2019-10-213-15/+9
| | | | | | This change rewrites Ch-4.md to introduced interfaces in a detailed step-by-step manner, adds examples, and fixes some errors. PiperOrigin-RevId: 275887017
* NFC: Fix remaining usages of MulOp as matrix multiplication.River Riddle2019-10-2114-131/+149
| | | | | | MulOp now represents an element-wise multiplication instead of a matrix multiplication. PiperOrigin-RevId: 275886774
* NFC: Elide the value of a UnitAttr within nested attribute dictionaries.River Riddle2019-10-211-2/+5
| | | | | | This matches the behavior of the top level attribute dictionary. PiperOrigin-RevId: 275879828
* Add a Symbol trait to simplify defining operations that represent symbols.River Riddle2019-10-212-3/+3
| | | | | | This trait provides accessors for the name, symbol use list methods, verification, with more to be added. PiperOrigin-RevId: 275864554
* NFC: Fix typo : Retur -> ReturnRiver Riddle2019-10-206-6/+6
| | | | PiperOrigin-RevId: 275745931
* Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-2017-24/+24
| | | | | | 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-193-0/+43
| | | | | | | | 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
* Add support for function result attributes.Sean Silva2019-10-185-0/+60
| | | | | | | | | | | | | | | | | | | | | | 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-186-25/+33
| | | | | | 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
* Implement lowering of VectorTypeCastOp to LLVMNicolas Vasilache2019-10-181-2/+19
| | | | | | | | | | A VectorTypeCastOp can only be used to lower between statically sized contiguous memrefs of scalar and matching vector type. The sizes and strides are thus fully static and easy to determine. A relevant test is added. This is a step towards solving tensorflow/mlir#189. PiperOrigin-RevId: 275538981
* Automated rollback of commit 575405f4d6762830c1c4520569de4e4ed3c8eed5Nicolas Vasilache2019-10-182-9/+35
| | | | PiperOrigin-RevId: 275461067
* 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-182-0/+13
| | | | | | | | Adding gen table for rewrite patterns from GPU to NVVM dialect. Copy missing op documentation from GPUOps.td to GPU.md. PiperOrigin-RevId: 275419588
* NFC: Remove trivial builder get methods.River Riddle2019-10-171-1/+1
| | | | | | These don't add any value, and some are even more restrictive than the respective static 'get' method. PiperOrigin-RevId: 275391240
* Automated rollback of commit b65c8bb5d6ab418bb3fcd0302aee19d3615d90f1River Riddle2019-10-171-32/+8
| | | | PiperOrigin-RevId: 275370861
* Add EDSC support for loop.for operationsNicolas Vasilache2019-10-171-8/+32
| | | | | | | This CL adds support for loop.for operations in EDSC and adds a test. This will be used in a followup commit to implement lowering of vector_transfer ops so that it works more generally and is not subject to affine constraints. PiperOrigin-RevId: 275349796
* NFC: Delete the Linalg tutorial.River Riddle2019-10-172-170/+0
| | | | | | This part of the tutorial is now covered by a new flow in Toy. This also removes a point of confusion as there is also a proper Linalg dialect. PiperOrigin-RevId: 275338933
* Add Ch.6 of the Toy tutorial.River Riddle2019-10-178-0/+247
| | | | | | This chapters introduces the notion of a full conversion, and adds support for lowering down to the LLVM dialect, LLVM IR, and thus code generation. PiperOrigin-RevId: 275337786
* Decouple Linalg promotion from Linalg tiling - NFCNicolas Vasilache2019-10-172-43/+72
| | | | | | 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-172-7/+8
| | | | | | | | | '_' 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
* Rename LoopNestBuilder to AffineLoopNestBuilder - NFCNicolas Vasilache2019-10-172-10/+10
| | | | PiperOrigin-RevId: 275310747
* [DRR] Allow capturing and referencing no-result opsLei Zhang2019-10-173-0/+21
| | | | | | | | | | Previously when we bind a symbol to an op in DRR, it means to capture the op's result(s) and later references will be expanded to result(s). This means for ops without result, we are replacing the symbol with nothing. This CL treats non-result op capturing and referencing as a special case to mean the op itself. PiperOrigin-RevId: 275269702
* Fix RewriterGen to support using NativeCodeCall as auxiliary patternLei Zhang2019-10-173-0/+21
| | | | | | | | | NativeCodeCall is handled differently than normal op creation in RewriterGen (because its flexibility). It will only be materialized to output stream if it is used. But when using it for auxiliary patterns, we still want the side effect even if it is not replacing matched root op's results. PiperOrigin-RevId: 275265467
* Fix invalid transpose in example and add proper verification.River Riddle2019-10-161-36/+36
| | | | | | The transpose in the example had the same result type as its input, which is incorrect. PiperOrigin-RevId: 275186568
* [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
* Add Ch.5 of the toy tutorial.River Riddle2019-10-1610-139/+174
| | | | | | This chapter adds a partial lowering of toy operations, all but PrintOp, to a combination of the Affine and Std dialects. This chapter focuses on introducing the conversion framework, the benefits of partial lowering, and how easily dialects may co-exist in the IR. PiperOrigin-RevId: 275150649
* Add support for inlining toy call operations.River Riddle2019-10-161-0/+30
| | | | | | The GenericCallOp needed to have the CallOpInterface to be picked up by the inliner. This also adds a CastOp to perform shape casts that are generated during inlining. The casts generated by the inliner will be folded away after shape inference. PiperOrigin-RevId: 275150438
* Add ComplexType to TableGen with Tensor supportRob Suderman2019-10-162-0/+33
| | | | | | | | Create a ComplexType for table gen references. Include an AnyComplex type to check whether the resulting tensor can be complex. Expand tensors to allow complex types. PiperOrigin-RevId: 275144804
* Update Chapter 4 of the Toy tutorialSana Damani2019-10-166-67/+21
| | | | | | | | | | This Chapter now introduces and makes use of the Interface concept in MLIR to demonstrate ShapeInference. END_PUBLIC Closes tensorflow/mlir#191 PiperOrigin-RevId: 275085151
* Update comments in ast.toyJacques Pienaar2019-10-163-84/+87
| | | | PiperOrigin-RevId: 275084969
* Makes spv.module generated by GPU->SPIRV conversion spec compliantMahesh Ravishankar2019-10-162-9/+9
| | | | | | | | | | | | | | | | Makes the spv.module generated by the GPU to SPIR-V conversion SPIR-V spec compliant (validated using spirv-val from Vulkan tools). 1) Separate out the VulkanLayoutUtils from DecorateSPIRVCompositeTypeLayoutPass to make it reusable within the Type converter in SPIR-V lowering infrastructure. This is used to compute the layout of the !spv.struct used in global variable type description. 2) Set the capabilities of the spv.module to Shader (needed for use of Logical Memory Model, and the extensions to SPV_KHR_storage_buffer_storage_class for use of Storage Buffer) PiperOrigin-RevId: 275081486
* Support custom accumulator provided as region to gpu.all_reduce.Christian Sigg2019-10-165-5/+142
| | | | | | | | | | 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 PatternRewriter::eraseOp.River Riddle2019-10-162-4/+4
| | | | | | This hook is useful when an operation is known to be dead, and no replacement values make sense. PiperOrigin-RevId: 275052756
* Implement simple loop-invariant-code-motion based on dialect interfaces.Stephan Herhut2019-10-162-324/+568
| | | | PiperOrigin-RevId: 275004258
* 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
* Update Chapter 3 to demonstrate pattern match and rewrite optimizationsSana Damani2019-10-153-23/+20
| | | | | | | | | This is using Table-driven Declarative Rewrite Rules (DRR), the previous version of the tutorial only showed the C++ patterns. Closes tensorflow/mlir#187 PiperOrigin-RevId: 274852321
* Consistent use of int in mlir_runner_utils.cppNicolas Vasilache2019-10-151-1/+1
| | | | | | This should fix the OSS build by only using int in template types. PiperOrigin-RevId: 274843584
* Add conversion for splat of vectors of 2+DNicolas Vasilache2019-10-153-39/+143
| | | | | | | This CL adds a missing lowering for splat of multi-dimensional vectors. Additional support is also added to the runtime utils library to allow printing memrefs with such vectors. PiperOrigin-RevId: 274794723
* Merge Ch3 of the Toy tutorial into chapter 2.River Riddle2019-10-142-5/+17
| | | | | | This effectively rewrites Ch.2 to introduce dialects, operations, and registration instead of deferring to Ch.3. This allows for introducing the best practices up front(using ODS, registering operations, etc.), and limits the opaque API to the chapter document instead of the code. PiperOrigin-RevId: 274724289
* Allowing replacing non-root operations in DialectConversion.River Riddle2019-10-143-3/+31
| | | | | | When dealing with regions, or other patterns that need to generate temporary operations, it is useful to be able to replace other operations than the root op being matched. Before this PR, these operations would still be considered for legalization meaning that the conversion would either fail, erroneously need to mark these ops as legal, or add unnecessary patterns. PiperOrigin-RevId: 274598513
* Use single quotes to wrap '@HOST_LDFLAGS@' in LIT config fileMehdi Amini2019-10-141-1/+2
| | | | | | ldflags can contain double-quoted paths, so must use single quotes here. PiperOrigin-RevId: 274581983
* 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
* Emit LLVM IR equivalent of sizeof when lowering alloc operationsAlex Zinenko2019-10-112-14/+26
| | | | | | | | | | | | | | | | | Originally, the lowering of `alloc` operations has been computing the number of bytes to allocate when lowering based on the properties of MLIR type. This does not take into account type legalization that happens when compiling LLVM IR down to target assembly. This legalization can widen the type, potentially leading to out-of-bounds accesses to `alloc`ed data due to mismatches between address computation that takes the widening into account and allocation that does not. Use the LLVM IR's equivalent of `sizeof` to compute the number of bytes to be allocated: %0 = getelementptr %type* null, %indexType 0 %1 = ptrtoint %type* %0 to %indexType adapted from http://nondot.org/sabre/LLVMNotes/SizeOf-OffsetOf-VariableSizedStructs.txt PiperOrigin-RevId: 274159900
* LLVM Dialect: introduce llvm.mlir.null operationAlex Zinenko2019-10-113-0/+25
| | | | | | | | 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 unary ops and ExpOp to Standard Dialect.Alexander Belyaev2019-10-112-0/+10
| | | | PiperOrigin-RevId: 274152154
OpenPOWER on IntegriCloud