Commit Graph

2999 Commits

Author SHA1 Message Date
River Riddle 21ee4e987f Add @below and @above directives to verify-diagnostics.
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
2019-10-23 15:56:29 -07:00
Alex Zinenko edffbbcdae Fix "set-but-unused" warning in DialectConversion
The variable in question is only used in an assertion,
leading to a warning in opt builds.

PiperOrigin-RevId: 276352259
2019-10-23 14:32:13 -07:00
River Riddle 5ee610a091 NFC: Remove references to the toy.generic attribute.
This was used for shape inference in the previous tutorial flow.

PiperOrigin-RevId: 276351916
2019-10-23 14:30:35 -07:00
Alex Zinenko 0d33703f2a Drop MemRefUtils from the ExecutionEngine
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
2019-10-23 07:43:06 -07:00
MLIR Team d499976098 Update chapter 3 code snippet to match the actual output of the code
PiperOrigin-RevId: 276117540
2019-10-22 12:18:27 -07:00
Uday Bondhugula ad6925f479 Update loop.for verifier message
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
2019-10-22 07:34:56 -07:00
Alex Zinenko 43de1c4303 Expose optimizations flags in Python bindings
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
2019-10-22 07:31:48 -07:00
Hanhan Wang 0237e52dde NFC: Remove a right parenthesis from comment.
PiperOrigin-RevId: 275998781
2019-10-21 22:44:48 -07:00
River Riddle 057ee97c73 NFC: Add support for parsing attributes programmatically via mlir::parseAttribute.
This matches the behavior of the public mlir::parseType, and even uses the internal implementation.

PiperOrigin-RevId: 275989777
2019-10-21 21:34:51 -07:00
Lei Zhang 020f9eb68c [DRR] Allow interleaved operands and attributes
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
2019-10-21 20:48:17 -07:00
Lei Zhang d9fe892e42 [spirv] Allow block arguments on spv.Branch(Conditional)
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
2019-10-21 17:32:00 -07:00
Alex Zinenko 5f867d26b4 Use LLVM_Type instead of AnyType in the definition of LLVM_CallOp
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
2019-10-21 14:12:19 -07:00
River Riddle 4514cdd5eb Cleanup and rewrite Ch-4.md.
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
2019-10-21 11:32:39 -07:00
River Riddle 941a1c4332 NFC: Fix remaining usages of MulOp as matrix multiplication.
MulOp now represents an element-wise multiplication instead of a matrix multiplication.

PiperOrigin-RevId: 275886774
2019-10-21 11:31:32 -07:00
Christian Sigg b74af4aa5c Unify GPU op definition names with other dialects.
Rename GPU op names from gpu_Foo to GPU_FooOp.

PiperOrigin-RevId: 275882232
2019-10-21 11:10:56 -07:00
River Riddle 03d7be2aca NFC: Elide the value of a UnitAttr within nested attribute dictionaries.
This matches the behavior of the top level attribute dictionary.

PiperOrigin-RevId: 275879828
2019-10-21 11:02:07 -07:00
River Riddle 9ac459e871 Add a Symbol trait to simplify defining operations that represent symbols.
This trait provides accessors for the name, symbol use list methods, verification, with more to be added.

PiperOrigin-RevId: 275864554
2019-10-21 09:58:59 -07:00
River Riddle 1bdfc9e74d NFC: Fix typo : Retur -> Return
PiperOrigin-RevId: 275745931
2019-10-20 15:13:20 -07:00
River Riddle 0bebd06f9a Update Ch1 to reflect new changes in the tutorial.
The chapter list is out of date, as well as mentions of matrix multiplication(now element-wise multiplication).

PiperOrigin-RevId: 275744911
2019-10-20 15:00:31 -07:00
Lei Zhang aad15d812e [DRR] Address GCC warning by wrapping for statement body with {}
Otherwise, we'll see the following warning when compiling with GCC 8:

warning: this ?for? clause does not guard... [-Wmisleading-indentation]
PiperOrigin-RevId: 275735925
2019-10-20 12:24:07 -07:00
Kazuaki Ishizaki f28c5aca17 Fix minor spelling tweaks (NFC)
Closes tensorflow/mlir#175

PiperOrigin-RevId: 275726876
2019-10-20 09:44:36 -07:00
Kazuaki Ishizaki 8bfedb3ca5 Fix minor spelling tweaks (NFC)
Closes tensorflow/mlir#177

PiperOrigin-RevId: 275692653
2019-10-20 00:11:34 -07:00
Jacques Pienaar 8317bd85e5 Add SourceMgrDiagnosticHandler to toy
PiperOrigin-RevId: 275659433
2019-10-19 14:36:36 -07:00
Jacques Pienaar 305dafd3b1 Add missing include to StringMap in Verifier and DialectConversion.
PiperOrigin-RevId: 275656416
2019-10-19 13:39:02 -07:00
Mehdi Amini 5b1345ff76 Add missing include to llvm Allocator.h
This header is not self-contained otherwise.

PiperOrigin-RevId: 275651582
2019-10-19 12:11:01 -07:00
Jacques Pienaar f9462da569 Slight rewording in toy ch2 to make persistence of name clearer
PiperOrigin-RevId: 275650756
2019-10-19 12:00:08 -07:00
Geoffrey Martin-Noble bc577eaf44 Use new eraseOp instead of replaceOp with empty values
PiperOrigin-RevId: 275631166
2019-10-19 06:04:18 -07:00
Christian Sigg c3e56cd12c Get active source lane predicate from shuffle instruction.
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
2019-10-19 01:53:25 -07:00
River Riddle 5f6bdd144a NFC: Cleanup the implementation of walkSymbolUses.
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
2019-10-18 21:29:15 -07:00
River Riddle d9db842e68 NFC: Add missing include for StringMap.
PiperOrigin-RevId: 275588019
2019-10-18 19:20:58 -07:00
Lei Zhang c5b9fefddc NFC: Rename SPIR-V serializer find*ID() to get*ID() to be consistent
We use get*() in deserizer and other places across the codebase.

PiperOrigin-RevId: 275582390
2019-10-18 18:16:05 -07:00
Sean Silva 9c9a7e9268 Add support for function result attributes.
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
2019-10-18 16:03:28 -07:00
Nicolas Vasilache 9e7e297da3 Lower vector transfer ops to loop.for operations.
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
2019-10-18 14:10:10 -07:00
Nicolas Vasilache 2823b68580 Implement lowering of VectorTypeCastOp to LLVM
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
2019-10-18 14:00:06 -07:00
reinerp 02b3ea6038 Slightly rephrase a difficult-to-parse sentence.
PiperOrigin-RevId: 275499524
2019-10-18 10:28:33 -07:00
Lei Zhang 3aae473658 [DRR] Use eraseOp() to replace no-result ops
PiperOrigin-RevId: 275475229
2019-10-18 08:20:25 -07:00
Nicolas Vasilache 151e7e61e8 Automated rollback of commit 575405f4d6
PiperOrigin-RevId: 275461067
2019-10-18 06:45:06 -07:00
Nicolas Vasilache 3e3ab38021 Fix OSS target name GPUtoNVVMTransforms -> MLIRGPUtoNVVMTransforms
This unbreaks the `cmake -G Ninja ../llvm -DLLVM_BUILD_EXAMPLES=ON -DLLVM_TARGETS_TO_BUILD="host"`
 in my local OSS build

PiperOrigin-RevId: 275452330
2019-10-18 05:22:38 -07:00
Stephan Herhut 3622e1833f Use StrEnumAttr for gpu.allreduce op instead of StringAttr to better encode constraints.
PiperOrigin-RevId: 275448372
2019-10-18 04:44:48 -07:00
Geoffrey Martin-Noble 234b8e85ba Add documentation on restrictions to dialect conversion rewrites
PiperOrigin-RevId: 275435593
2019-10-18 02:40:38 -07:00
Christian Sigg fe0ee32da5 Add gpu.barrier op to synchronize invocations of a local workgroup.
Adding gen table for rewrite patterns from GPU to NVVM dialect.

Copy missing op documentation from GPUOps.td to GPU.md.

PiperOrigin-RevId: 275419588
2019-10-18 00:30:44 -07:00
River Riddle 2acc220f17 NFC: Remove trivial builder get methods.
These don't add any value, and some are even more restrictive than the respective static 'get' method.

PiperOrigin-RevId: 275391240
2019-10-17 20:08:34 -07:00
River Riddle 575405f4d6 Automated rollback of commit b65c8bb5d6
PiperOrigin-RevId: 275370861
2019-10-17 17:11:39 -07:00
Geoffrey Martin-Noble 6090643877 Introduce a wrapper around ConversionPattern that operates on the derived class
Analogous to OpRewritePattern, this makes writing conversion patterns more convenient.

PiperOrigin-RevId: 275349854
2019-10-17 15:30:38 -07:00
Nicolas Vasilache b65c8bb5d6 Add EDSC support for loop.for operations
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
2019-10-17 15:18:34 -07:00
River Riddle dae0ae6879 NFC: Delete the Linalg tutorial.
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
2019-10-17 14:27:37 -07:00
River Riddle 0372eb413f Add Ch.6 of the Toy tutorial.
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
2019-10-17 14:22:13 -07:00
Nicolas Vasilache 5b03e692f6 Decouple Linalg promotion from Linalg tiling - NFC
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
2019-10-17 13:41:17 -07:00
Denis Khalikov a560505d1a [spirv] Add a canonicalization pattern for spv.selection.
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
2019-10-17 12:36:47 -07:00
Lei Zhang 057dc41bf6 Allow '_' when pretty printing dialect symbols
'_' 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
2019-10-17 12:24:18 -07:00