summaryrefslogtreecommitdiffstats
Commit message (Collapse)AuthorAgeFilesLines
...
| * Add @below and @above directives to verify-diagnostics.River Riddle2019-10-235-314/+357
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | This simplifies defining expected-* directives when there are multiple that apply to the next or previous line. @below applies the directive to the next non-designator line, i.e. the next line that does not contain an expected-* designator. @above applies to the previous non designator line. Examples: // Expect an error on the next line that does not contain a designator. // expected-remark@below {{remark on function below}} // expected-remark@below {{another remark on function below}} func @bar(%a : f32) // Expect an error on the previous line that does not contain a designator. func @baz(%a : f32) // expected-remark@above {{remark on function above}} // expected-remark@above {{another remark on function above}} PiperOrigin-RevId: 276369085
| * Fix "set-but-unused" warning in DialectConversionAlex Zinenko2019-10-231-0/+1
| | | | | | | | | | | | | | The variable in question is only used in an assertion, leading to a warning in opt builds. PiperOrigin-RevId: 276352259
| * NFC: Remove references to the toy.generic attribute.River Riddle2019-10-236-41/+8
| | | | | | | | | | | | This was used for shape inference in the previous tutorial flow. PiperOrigin-RevId: 276351916
| * Drop MemRefUtils from the ExecutionEngineAlex Zinenko2019-10-234-163/+0
| | | | | | | | | | | | | | | | | | | | The ExecutionEngine was updated recently to only take the LLVM dialect as input. Memrefs are no longer expected in the signature of the entry point function by the executor so there is no need to allocate and free them. The code in MemRefUtils is therefore dead and furthermore out of sync with the recent evolution of memref type to support strides. Drop it. PiperOrigin-RevId: 276272302
| * Update chapter 3 code snippet to match the actual output of the codeMLIR Team2019-10-221-2/+1
| | | | | | | | PiperOrigin-RevId: 276117540
| * Update loop.for verifier messageUday Bondhugula2019-10-222-3/+3
| | | | | | | | | | | | | | | | | | 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
| * Expose optimizations flags in Python bindingsAlex Zinenko2019-10-221-3/+30
| | | | | | | | | | | | | | | | | | ExecutionEngine currently supports additional parameters that can be used to run LLVM transformations during JIT compilation. Expose this to Python bindings. When the ExecutionEngine functionality is moved to LLVM, the bindings-specific code can be updated to interact with LLVM. PiperOrigin-RevId: 276060475
| * NFC: Remove a right parenthesis from comment.Hanhan Wang2019-10-211-1/+1
| | | | | | | | PiperOrigin-RevId: 275998781
| * NFC: Add support for parsing attributes programmatically via ↵River Riddle2019-10-212-9/+60
| | | | | | | | | | | | | | | | mlir::parseAttribute. This matches the behavior of the public mlir::parseType, and even uses the internal implementation. PiperOrigin-RevId: 275989777
| * [DRR] Allow interleaved operands and attributesLei Zhang2019-10-215-38/+101
| | | | | | | | | | | | | | | | 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-215-36/+106
| | | | | | | | | | | | | | | | | | | | | | | | 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
| * Use LLVM_Type instead of AnyType in the definition of LLVM_CallOpAlex Zinenko2019-10-212-5/+3
| | | | | | | | | | | | | | | | | | | | The type constraint had to be relaxed due to the order of lowering passes in the examples, that since has been fixed. The relaxed version was still used by the CUDA lowering for launch sizes of `index` type. This is not necessary since the GPU dialect does not restrict the type of the launch size operands. Use an LLVM type instead and restore the check in the LLVM_CallOp definition. PiperOrigin-RevId: 275920109
| * Cleanup and rewrite Ch-4.md.River Riddle2019-10-2110-90/+370
| | | | | | | | | | | | 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-2119-214/+177
| | | | | | | | | | | | MulOp now represents an element-wise multiplication instead of a matrix multiplication. PiperOrigin-RevId: 275886774
| * Unify GPU op definition names with other dialects.Christian Sigg2019-10-218-42/+49
| | | | | | | | | | | | 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-212-3/+12
| | | | | | | | | | | | 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-2110-44/+79
| | | | | | | | | | | | 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
| * Update Ch1 to reflect new changes in the tutorial.River Riddle2019-10-201-21/+21
| | | | | | | | | | | | The chapter list is out of date, as well as mentions of matrix multiplication(now element-wise multiplication). PiperOrigin-RevId: 275744911
| * [DRR] Address GCC warning by wrapping for statement body with {}Lei Zhang2019-10-201-1/+2
| | | | | | | | | | | | | | Otherwise, we'll see the following warning when compiling with GCC 8: warning: this ?for? clause does not guard... [-Wmisleading-indentation] PiperOrigin-RevId: 275735925
| * Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-2040-62/+62
| | | | | | | | | | | | Closes tensorflow/mlir#175 PiperOrigin-RevId: 275726876
| * Fix minor spelling tweaks (NFC)Kazuaki Ishizaki2019-10-2038-106/+107
| | | | | | | | | | | | Closes tensorflow/mlir#177 PiperOrigin-RevId: 275692653
| * Add SourceMgrDiagnosticHandler to toyJacques Pienaar2019-10-193-9/+15
| | | | | | | | PiperOrigin-RevId: 275659433
| * Add missing include to StringMap in Verifier and DialectConversion.Jacques Pienaar2019-10-192-0/+3
| | | | | | | | PiperOrigin-RevId: 275656416
| * Add missing include to llvm Allocator.hMehdi Amini2019-10-192-2/+2
| | | | | | | | | | | | This header is not self-contained otherwise. PiperOrigin-RevId: 275651582
| * Slight rewording in toy ch2 to make persistence of name clearerJacques Pienaar2019-10-191-5/+7
| | | | | | | | PiperOrigin-RevId: 275650756
| * Use new eraseOp instead of replaceOp with empty valuesGeoffrey Martin-Noble2019-10-194-5/+5
| | | | | | | | PiperOrigin-RevId: 275631166
| * Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-199-33/+117
| | | | | | | | | | | | | | | | 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-182-68/+45
| | | | | | | | | | | | 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: Add missing include for StringMap.River Riddle2019-10-181-0/+1
| | | | | | | | PiperOrigin-RevId: 275588019
| * NFC: Rename SPIR-V serializer find*ID() to get*ID() to be consistentLei Zhang2019-10-182-30/+30
| | | | | | | | | | | | We use get*() in deserizer and other places across the codebase. PiperOrigin-RevId: 275582390
| * Add support for function result attributes.Sean Silva2019-10-1813-25/+394
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | 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-1810-41/+52
| | | | | | | | | | | | 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-184-17/+150
| | | | | | | | | | | | | | | | | | | | 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
| * Slightly rephrase a difficult-to-parse sentence.reinerp2019-10-181-3/+3
| | | | | | | | PiperOrigin-RevId: 275499524
| * [DRR] Use eraseOp() to replace no-result opsLei Zhang2019-10-181-10/+18
| | | | | | | | PiperOrigin-RevId: 275475229
| * Automated rollback of commit 575405f4d6762830c1c4520569de4e4ed3c8eed5Nicolas Vasilache2019-10-187-40/+138
| | | | | | | | 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-183-7/+13
| | | | | | | | | | | | constraints. PiperOrigin-RevId: 275448372
| * Add documentation on restrictions to dialect conversion rewritesGeoffrey Martin-Noble2019-10-181-4/+23
| | | | | | | | PiperOrigin-RevId: 275435593
| * Add gpu.barrier op to synchronize invocations of a local workgroup.Christian Sigg2019-10-187-5/+143
| | | | | | | | | | | | | | | | 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-1738-223/+128
| | | | | | | | | | | | 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-176-135/+39
| | | | | | | | PiperOrigin-RevId: 275370861
| * Introduce a wrapper around ConversionPattern that operates on the derived classGeoffrey Martin-Noble2019-10-172-6/+80
| | | | | | | | | | | | Analogous to OpRewritePattern, this makes writing conversion patterns more convenient. PiperOrigin-RevId: 275349854
| * Add EDSC support for loop.for operationsNicolas Vasilache2019-10-176-39/+135
| | | | | | | | | | | | | | 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-1773-6883/+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-1733-2/+4254
| | | | | | | | | | | | 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-177-282/+332
| | | | | | | | | | | | 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-174-0/+406
| | | | | | | | | | | | | | | | | | | | | | 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-174-11/+16
| | | | | | | | | | | | | | | | | | '_' 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
OpenPOWER on IntegriCloud