summaryrefslogtreecommitdiffstats
path: root/mlir/lib
Commit message (Collapse)AuthorAgeFilesLines
...
* [DRR] Allow interleaved operands and attributesLei Zhang2019-10-211-0/+12
| | | | | | | | Previously DRR assumes attributes to appear after operands. This was the previous requirements on ODS, but that has changed some time ago. Fix DRR to also support interleaved operands and attributes. PiperOrigin-RevId: 275983485
* [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-217-28/+35
| | | | | | Rename GPU op names from gpu_Foo to GPU_FooOp. PiperOrigin-RevId: 275882232
* NFC: Elide the value of a UnitAttr within nested attribute dictionaries.River Riddle2019-10-211-1/+7
| | | | | | 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-1/+9
| | | | | | 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-2037-105/+106
| | | | | | Closes tensorflow/mlir#177 PiperOrigin-RevId: 275692653
* Add missing include to StringMap in Verifier and DialectConversion.Jacques Pienaar2019-10-191-0/+2
| | | | PiperOrigin-RevId: 275656416
* Add missing include to llvm Allocator.hMehdi Amini2019-10-191-2/+1
| | | | | | This header is not self-contained otherwise. PiperOrigin-RevId: 275651582
* Use new eraseOp instead of replaceOp with empty valuesGeoffrey Martin-Noble2019-10-193-4/+4
| | | | PiperOrigin-RevId: 275631166
* Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-194-29/+55
| | | | | | | | 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: Cleanup the implementation of walkSymbolUses.River Riddle2019-10-181-66/+35
| | | | | | Refactor the implementation to be much cleaner by adding a `make_second_range` utility to walk the `second` value of a range of pairs. PiperOrigin-RevId: 275598985
* 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-183-15/+120
| | | | | | | | | | | | | | | | | | | | | | 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-184-16/+19
| | | | | | 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-182-15/+126
| | | | | | | | | | 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-17/+57
| | | | PiperOrigin-RevId: 275461067
* Fix OSS target name GPUtoNVVMTransforms -> MLIRGPUtoNVVMTransformsNicolas Vasilache2019-10-181-1/+1
| | | | | | | This unbreaks the `cmake -G Ninja ../llvm -DLLVM_BUILD_EXAMPLES=ON -DLLVM_TARGETS_TO_BUILD="host"` in my local OSS build PiperOrigin-RevId: 275452330
* Use StrEnumAttr for gpu.allreduce op instead of StringAttr to better encode ↵Stephan Herhut2019-10-181-5/+0
| | | | | | constraints. PiperOrigin-RevId: 275448372
* Add gpu.barrier op to synchronize invocations of a local workgroup.Christian Sigg2019-10-183-0/+50
| | | | | | | | 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-1721-145/+77
| | | | | | 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-172-57/+17
| | | | PiperOrigin-RevId: 275370861
* Introduce a wrapper around ConversionPattern that operates on the derived classGeoffrey Martin-Noble2019-10-171-6/+4
| | | | | | Analogous to OpRewritePattern, this makes writing conversion patterns more convenient. PiperOrigin-RevId: 275349854
* Add EDSC support for loop.for operationsNicolas Vasilache2019-10-172-17/+57
| | | | | | | 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
* 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
* 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
* Rename LoopNestBuilder to AffineLoopNestBuilder - NFCNicolas Vasilache2019-10-172-9/+8
| | | | PiperOrigin-RevId: 275310747
* 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
* [DRR] Allow capturing and referencing no-result opsLei Zhang2019-10-171-0/+7
| | | | | | | | | | 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
* Add LLVM_DEBUG in RewritersGen.cpp and Pattern.cppLei Zhang2019-10-171-10/+53
| | | | | | | | | It's usually hard to understand what went wrong if mlir-tblgen crashes on some input. This CL adds a few useful LLVM_DEBUG statements so that we can use mlir-tblegn -debug to figure out the culprit for a crash. PiperOrigin-RevId: 275253532
* [spirv] Implement inliner interfaceLei Zhang2019-10-162-0/+73
| | | | | | | | | | | | | | 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
* Update Chapter 4 of the Toy tutorialSana Damani2019-10-161-0/+4
| | | | | | | | | | 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
* NFC: Update VectorOrTensor -> ShapedGeoffrey Martin-Noble2019-10-161-1/+1
| | | | | | This was missed when the type was renamed. PiperOrigin-RevId: 275082588
* Makes spv.module generated by GPU->SPIRV conversion spec compliantMahesh Ravishankar2019-10-165-196/+186
| | | | | | | | | | | | | | | | 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-162-8/+137
| | | | | | | | | | 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
* Allow for remapping argument to a Value in SignatureConversion.Mahesh Ravishankar2019-10-162-4/+23
| | | | | | | | | | | | | | | | The current SignatureConversion framework (part of DialectConversion) allows remapping input arguments to a function from 1->0, 1->1 or 1->many arguments during conversion. Another case is where the argument itself is dropped, but it's use are remapped to another Value*. An example of this is: The Vulkan/SPIR-V spec requires entry functions to be of type void(void). The GPU -> SPIR-V conversion implemented this without having the DialectConversion framework track the remapping that lead to some undefined behavior. The changes here addresses that. PiperOrigin-RevId: 275059656
* Add support for PatternRewriter::eraseOp.River Riddle2019-10-169-27/+42
| | | | | | This hook is useful when an operation is known to be dead, and no replacement values make sense. PiperOrigin-RevId: 275052756
* Fix CMake configuration after introduction of LICM and LoopLikeInterfaceMehdi Amini2019-10-163-3/+10
| | | | | | | b843cc5d5a introduced a new op LICM transformation and a LoopLike interface, but missed the CMake aspects of it. This should fix the build. PiperOrigin-RevId: 275038533
* Implement simple loop-invariant-code-motion based on dialect interfaces.Stephan Herhut2019-10-164-185/+396
| | | | PiperOrigin-RevId: 275004258
* [spirv] Add support for SpecId decoration on spv.specConstantLei Zhang2019-10-153-35/+88
| | | | | | | | 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
* Add conversion for splat of vectors of 2+DNicolas Vasilache2019-10-151-48/+142
| | | | | | | 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
* Expose mlir::parseType to bindingsAlex Zinenko2019-10-152-26/+33
| | | | | | | | | | | | | | Python bindings currently currently provide a makeScalarType function that constructs one of the predefined types. It was implemented in the bindings directly to circumvent the absence of standalone type parsing function. Now that mlir::parseType has been made available, rely on the core parsing procedure to construct types from strings in the bindings. This changes includes a library reshuffling that splits out "CoreAPIs" implementing the binding helper APIs into a separate library and makes that dependent on the Parser library. PiperOrigin-RevId: 274794516
* AsmPrinter: avoid unused-variable warningAlex Zinenko2019-10-151-1/+1
| | | | | | | | The value defined in a loop was not being used and the function producing it re-evaluated instead. Use the value to avoid both the warning and the re-evaluation. PiperOrigin-RevId: 274794459
* NFC: Replace usages of Value::getKind with explicit isa/casts.River Riddle2019-10-142-29/+23
| | | | | | It is more idiomatic to use the llvm::cast infrastructure for checking the type of a value. PiperOrigin-RevId: 274684945
* Allowing replacing non-root operations in DialectConversion.River Riddle2019-10-141-1/+16
| | | | | | 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
* Fix linalg.subview behavior in (partially) static cases.Nicolas Vasilache2019-10-143-21/+44
| | | | | | | | | | | | | | | | | | | | | 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/+2
| | | | | | | | | 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-1/+10
| | | | | | | 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-111-12/+15
| | | | | | | | | | | | | | | | | 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-111-0/+27
| | | | | | | | 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
OpenPOWER on IntegriCloud