Commit Graph

2984 Commits

Author SHA1 Message Date
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
Nicolas Vasilache 10039d04e2 Rename LoopNestBuilder to AffineLoopNestBuilder - NFC
PiperOrigin-RevId: 275310747
2019-10-17 12:13:59 -07:00
Mehdi Amini 6ebc7318b0 Use a SmallVector instead of an ArrayRef to materialize a temporary local array
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
2019-10-17 10:10:46 -07:00
Lei Zhang 23d21af65c [DRR] Allow capturing and referencing no-result ops
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
2019-10-17 09:02:31 -07:00
Lei Zhang 603117b2d6 Fix RewriterGen to support using NativeCodeCall as auxiliary pattern
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
2019-10-17 08:39:59 -07:00
Lei Zhang 1358df19ca Add LLVM_DEBUG in RewritersGen.cpp and Pattern.cpp
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
2019-10-17 07:26:22 -07:00
River Riddle bdc250c5a7 Fix invalid transpose in example and add proper verification.
The transpose in the example had the same result type as its input, which is incorrect.

PiperOrigin-RevId: 275186568
2019-10-16 22:37:00 -07:00
Lei Zhang 0e3efb32c6 [spirv] Implement inliner interface
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
2019-10-16 17:46:19 -07:00
River Riddle 1ba9bb0507 Add Ch.5 of the toy tutorial.
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
2019-10-16 17:45:09 -07:00
River Riddle 7045471913 Add support for inlining toy call operations.
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
2019-10-16 17:32:57 -07:00
reinerp 7053a30f4b Fix typo in tutorial.
PiperOrigin-RevId: 275147795
2019-10-16 17:15:33 -07:00
Rob Suderman a245023c1c Add ComplexType to TableGen with Tensor support
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
2019-10-16 16:59:08 -07:00
River Riddle ab79c25d64 Code cleanups on Ch.4
This change performs general cleanups of the implementation of ch.4 and fixes some bugs. For example, the operations currently don't inherit from the shape inference interface.

PiperOrigin-RevId: 275089914
2019-10-16 12:34:26 -07:00
Sana Damani 3940b90d84 Update Chapter 4 of the Toy tutorial
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
2019-10-16 12:19:39 -07:00
Jacques Pienaar e88dbc8c95 Update comments in ast.toy
PiperOrigin-RevId: 275084969
2019-10-16 12:08:24 -07:00
Geoffrey Martin-Noble a3726a13f7 NFC: Update VectorOrTensor -> Shaped
This was missed when the type was renamed.

PiperOrigin-RevId: 275082588
2019-10-16 11:58:26 -07:00