Commit Graph

324 Commits

Author SHA1 Message Date
Ahmed S. Taei 39389587d2 {mhlo.is_finite, lmhlo.is_finite} -> linalg.generic conversion
PiperOrigin-RevId: 334414295
2020-09-29 10:50:42 -07:00
Benjamin Kramer 6459f12235 Lower mhlo.not to a xor with all ones
PiperOrigin-RevId: 334361499
2020-09-29 05:59:40 -07:00
Ahmed S. Taei 9c6640cbb6 Only apply GeneralDotOpLoweringPatterns for static shaped inputs
PiperOrigin-RevId: 333439680
2020-09-23 21:42:01 -07:00
A. Unique TensorFlower 08e0d09463 [MLIR][KernelGen] Rename `legalize-tanh-to-approximation` to `legalize-trigonometric-to-approximation`
To add more approximation lowerings in the future, generalize the pass name.

PiperOrigin-RevId: 333340075
2020-09-23 11:53:45 -07:00
Benjamin Kramer dd92c8ef61 Integrate LLVM at llvm/llvm-project@7e78d89052
Updates LLVM usage to match
[7e78d89052b1](https://github.com/llvm/llvm-project/commit/7e78d89052b1)

PiperOrigin-RevId: 333090785
2020-09-22 09:08:23 -07:00
Mehdi Amini 7abd557a61 Add a header for table-gen generated pass for MHLO and use it in SinkConstantsToControlFlowPass
A non globally registered pass should define `getName()` in order to generate correct crash reproducers.
This is something we get "for free" when using the TableGen generated base class.
We should also migrate the other passes to the same mechanism and remove the static
global registration.

PiperOrigin-RevId: 332976907
2020-09-21 18:01:57 -07:00
A. Unique TensorFlower 4002077261 [MLIR][KernelGen] Lower `tf.Sinh` to MLHLO
PiperOrigin-RevId: 332425724
2020-09-18 04:27:07 -07:00
A. Unique TensorFlower 2fbbbe9cf1 [MLIR][KernelGen] Lower `tf.Acos` to LMHLO.
- Add ranked code generation for `mhlo.compare/select`
- Add bufferization for `tensor_cast`
- Add lowerings for `Atan2Op`

PiperOrigin-RevId: 332407734
2020-09-18 01:40:18 -07:00
Hanhan Wang b29dd5ef8f Fix a bug in the case check of reshape op lowering.
PiperOrigin-RevId: 332044191
2020-09-16 11:06:16 -07:00
A. Unique TensorFlower 69b80d8deb [MLIR] Extend unranked transformation to CHLO dialect
PiperOrigin-RevId: 332026604
2020-09-16 09:49:18 -07:00
Stephan Herhut 2aa07b0091 Insert explicit casts to model extra shape knowledge for unranked chlo transform
When transforming unranked binary operations from CHLO to HLO, we insert `shape.broadcast` operations. Due to context, we know that the result of the `shape.broadcast` operation has a static shape. Instead of modelling this in the type of the broadcast operation itself, which is illegal,  we now use an explicit cast.

PiperOrigin-RevId: 331989879
2020-09-16 06:14:49 -07:00
A. Unique TensorFlower a6fdebdc6c [MLIR] Lower `chlo.constant_like` to MHLO
Lower `chlo.constant_like` to a constant and, if needed, a broadcast.

PiperOrigin-RevId: 331964137
2020-09-16 02:41:44 -07:00
A. Unique TensorFlower da43c8596b [MLIR] Simplify and generalize `transform-unranked-hlo`
This refactoring allows to support a wider range of n-ary operations in future
changes.

PiperOrigin-RevId: 331953362
2020-09-16 01:13:23 -07:00
A. Unique TensorFlower 48022987ce [XLA][MLIR] Lower `tf.Tan` and `tf.Sin` to MLHLO
Add `tan` op and lowering to CHLO dialect, move CHLO lowerings to
`chlo_legalize_to_hlo_patterns` and extend missing patterns.

PiperOrigin-RevId: 331506094
2020-09-14 02:34:52 -07:00
Ehsan Toosi ce1c8a1ebc [MLIR][LHLO] Replace lhlo-copy-removal pass with mlir-copy-removal pass
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/43137

This PR removes lhlo-copy-removal pass entirely and replace its usages with ```mlir::createCopyRemovalPass()```.

--
7ce1a06f507c8db46c6d7b43c7870cf56002e18e by Ehsan Toosi <ehsan.nadjaran_toosi@dfki.de>:

[mlir][lhlo] Replace lhlo-copy-removal pass with mlir-copy-removal pass

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/43137 from dfki-ehna:using_mlir_copy_removal 7ce1a06f507c8db46c6d7b43c7870cf56002e18e
PiperOrigin-RevId: 331498501
2020-09-14 01:22:19 -07:00
A. Unique TensorFlower a7a7184eb6 [XLA][MLIR] Lower `tf.Tan` and `tf.Sin` to MLHLO
Add `tan` op and lowering to CHLO dialect, move CHLO lowerings to
`chlo_legalize_to_hlo_patterns` and extend missing patterns.

PiperOrigin-RevId: 331128170
2020-09-11 05:05:58 -07:00
A. Unique TensorFlower 90927f6b53 [XLA][MLIR] Lower `tf.Tan` and `tf.Sin` to MLHLO
Add `tan` op and lowering to CHLO dialect, move CHLO lowerings to
`chlo_legalize_to_hlo_patterns` and extend missing patterns.

PiperOrigin-RevId: 331125286
2020-09-11 04:39:28 -07:00
A. Unique TensorFlower b22f2f0eea Integrate LLVM at llvm/llvm-project@52f0837778
Updates LLVM usage to match
[52f0837778b6](https://github.com/llvm/llvm-project/commit/52f0837778b6)

PiperOrigin-RevId: 330939173
2020-09-10 08:14:23 -07:00
Ehsan Toosi d599485e06 PR #43069: [hlo] Unbreak hlo-legalize-to-lhlo test
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/43069

The legalization of mlhlo.ReturnOp to lhlo.TerminatorOp by using BufferAssignmentReturnOpConverter fails since the Memref typed results (or the Memref typed operands of Return operation) are set to stay as results after legalization but lhlo.TerminatorOp doesn't accept any operands. Therefore, BufferAssignmentReturnOpConverter must be replaced with a manual conversion that removes all operands of mlhlo.ReturnOp and inserts copy operations in their places.
Copybara import of the project:

--
8be0435b0147263c3872bedec58fd215f784b450 by Ehsan Toosi <ehsan.nadjaran_toosi@dfki.de>:

[hlo] Unbreak hlo-legalize-to-lhlo test

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/43069 from dfki-ehna:fix_hlo_legalize_to_lhlo_test 8be0435b0147263c3872bedec58fd215f784b450
PiperOrigin-RevId: 330907602
2020-09-10 04:09:03 -07:00
A. Unique TensorFlower f46ba09653 Integrate LLVM at llvm/llvm-project@4964d75d70
Updates LLVM usage to match
[4964d75d7078](https://github.com/llvm/llvm-project/commit/4964d75d7078)

PiperOrigin-RevId: 330713009
2020-09-09 06:50:26 -07:00
Alexander Belyaev ebc7992d31 [MLIR][KERNEL_GEN] Add a library to lower kernels with the host side.
* Unified TF->Cubin and TF->Kernel_with_host side lowering in `kernel_creator.h|cc`
* Added a pass that attaches GPU binary blob to GPUModuleOp
* Refactored most of the code.
* Added tf_to_kernel binary that emits obj file

PiperOrigin-RevId: 330494488
2020-09-08 06:06:29 -07:00
Eugene Burmako dde1ed56cc Add support for legalizing mhlo.slice to lmhlo.slice
PiperOrigin-RevId: 330153599
2020-09-04 21:43:01 -07:00
Eugene Burmako 2a9d7ac084 Add support for legalizing lmhlo.transpose to linalg.generic
PiperOrigin-RevId: 330130704
2020-09-04 15:37:42 -07:00
Eugene Burmako f5d12604ed Add support for legalizing mhlo.transpose to lmhlo.transpose
PiperOrigin-RevId: 330127758
2020-09-04 14:59:08 -07:00
A. Unique TensorFlower 801262688c Integrate LLVM at llvm/llvm-project@1426ac0482
Updates LLVM usage to match
[1426ac048295](https://github.com/llvm/llvm-project/commit/1426ac048295)

PiperOrigin-RevId: 329839640
2020-09-02 20:03:41 -07:00
A. Unique TensorFlower 65b0613491 Integrate LLVM at llvm/llvm-project@202766947e
Updates LLVM usage to match
[202766947edb](https://github.com/llvm/llvm-project/commit/202766947edb)

PiperOrigin-RevId: 329673065
2020-09-02 02:27:52 -07:00
Benjamin Kramer 049f034116 Lower mhlo.floor to lmhlo to linalg
PiperOrigin-RevId: 329304882
2020-08-31 08:16:36 -07:00
Jacques Pienaar 344c500fca [mhlo] Add legalize to SCF pass
Start of pass to legalize MHLO control flow to SCF for further optimization in common form. The current version just matches a very simple instance (which also happens to occur a few times). Exposes some further canonicalization opportunities that aren't yet addressed.

PiperOrigin-RevId: 329017723
2020-08-28 15:11:58 -07:00
Mehdi Amini 36ddbeb6b2 Remove the dependency on global dialect registry from mlir-hlo
PiperOrigin-RevId: 328457105
2020-08-25 20:30:42 -07:00
Hanhan Wang bfd629ecb0 Enhance lowering reshape op to Linalg.
Handle non-expansion and non-collapsion cases by rewriting it to two reshape
ops.

PiperOrigin-RevId: 327926863
2020-08-21 23:27:34 -07:00
Robert Suderman 5252aeae8f Lowering for mhlo.ceil to std.ceil
PiperOrigin-RevId: 326335301
2020-08-12 16:15:35 -07:00
Alexander Belyaev 843af36e05 [MLIR] Add e2e test for unranked unary TF op, lowered and run with CPU runner.
PiperOrigin-RevId: 325665428
2020-08-09 02:38:00 -07:00
A. Unique TensorFlower e00565b544 Integrate LLVM at llvm/llvm-project@9dbdaea9a0
Updates LLVM usage to match
[9dbdaea9a0e6](https://github.com/llvm/llvm-project/commit/9dbdaea9a0e6)

PiperOrigin-RevId: 325354353
2020-08-07 22:20:06 -07:00
Rahul Joshi a6978cf4ab [MLIR][NFC] Adopt FuncOp/Region argument API's.
- Use FuncOp::getArguments() and Region::getArguments() and friends where possible
  instead of going through the front() block.

PiperOrigin-RevId: 325352975
2020-08-07 22:19:21 -07:00
A. Unique TensorFlower a68a16cdc7 [MLIR][XLA] Allow for choice of safe/unsafe variant in broadcast utils
Create safe or unsafe variants of `shape.broadcast` depending on the context.
The representation by means of an extent tensor is only legal if the operands
are known to be broadcastable. Currently, there is no use in a safe context in
the codebase but it will be used for shape inference eventually.

PiperOrigin-RevId: 325228073
2020-08-07 22:16:11 -07:00
Alexander Belyaev bc3293a05f [MLIR] Remove Affine->STD and SCF->STD patterns from lhlo->llvm pass.
PiperOrigin-RevId: 325219360
2020-08-07 22:15:24 -07:00
Mehdi Amini 701312720c Add CMake files and lit configurations, enough for `ninja check-mlir-hlo` to pass on all the tests
PiperOrigin-RevId: 325172984
2020-08-07 22:14:34 -07:00
A. Unique TensorFlower 4372124362 [MLIR][XLA] Allow for choice of safe/unsafe variant in broadcast utils
Create safe or unsafe variants of `shape.broadcast` depending on the context.
The representation by means of an extent tensor is only legal if the operands
are known to be broadcastable. Currently, there is no use in a safe context in
the codebase but it will be used for shape inference eventually.

PiperOrigin-RevId: 325079842
2020-08-05 12:43:29 -07:00
A. Unique TensorFlower 5d3cc2105e [MLIR][HLO] Remove redundant casts from unranked to ranked transformation
The transformation of unranked to ranked operations no longer generates cast
operations for shapes and sizes. Instead, we use the newly introduced support
for extent tensor and index types directly.

PiperOrigin-RevId: 325057440
2020-08-05 11:11:43 -07:00
A. Unique TensorFlower 37c36a4389 [MLIR][XLA] Allow for choice of safe/unsafe variant in broadcast utils
Create safe or unsafe variants of `shape.broadcast` depending on the context.
The representation by means of an extent tensor is only legal if the operands
are known to be broadcastable. Currently, there is no use in a safe context in
the codebase but it will be used for shape inference eventually.

PiperOrigin-RevId: 325056915
2020-08-05 11:09:23 -07:00
Stephan Herhut 6584c2ab1f Remove optional static registration for hlo dialects again.
Instead, we invoke multiple test tools in a row in end to end tests now. For hlo dialects and passes, we use mlir-hlo-opt explicitly.

PiperOrigin-RevId: 324989884
2020-08-05 03:53:17 -07:00
Mehdi Amini 6c7d1a7c7e Internal change
PiperOrigin-RevId: 324753700
2020-08-03 22:59:44 -07:00
Smit Hinsu 3fe9a7d2db Legalize TensorFlow NonMaxSuppressionV4 and SelfAdjointEigV2Op ops to HLO
Added support for HLO ops bitcast-convert, sort and while in MlirHloBuilder and enabled tests for NonMaxSuppressionV4 and SelfAdjointEigV2Op using these ops.

PiperOrigin-RevId: 324360651
2020-08-03 19:30:54 -07:00
A. Unique TensorFlower 7809320a5e Integrate LLVM at llvm/llvm-project@b7cfa6ca92
Updates LLVM usage to match
[b7cfa6ca9283](https://github.com/llvm/llvm-project/commit/b7cfa6ca9283)

PiperOrigin-RevId: 324331764
2020-08-03 19:30:43 -07:00
Smit Hinsu 577a81a66d Sink standard dialect constants in sink_constants_to_control_flow pass
This is required before exporting HLO dialect ops with standard dialect constant to XLA.

Also, sink constants for sort op as well. Added a TODO to generalize this pass to handle more ops and non-const values defined outside.

PiperOrigin-RevId: 324301911
2020-08-03 19:30:29 -07:00
Stephan Herhut 734b9b25fd Add optional static registration for mhlo/lmhlo passes.
PiperOrigin-RevId: 324190465
2020-08-03 19:29:40 -07:00
Thomas Joerg 735ae2838c Integrate LLVM at llvm/llvm-project@cd4e8d7f6f
Updates LLVM usage to match
[cd4e8d7f6f5e](https://github.com/llvm/llvm-project/commit/cd4e8d7f6f5e)

PiperOrigin-RevId: 324173542
2020-08-03 19:29:29 -07:00
Tres Popp cce4bddf4b Remove unnecessary conversions between Shape and ExtentTensor.
PiperOrigin-RevId: 323981215
2020-08-03 19:29:04 -07:00
Tres Popp ffef8d6593 Support CHLO->LHLO lowering for broadcasting operations with both inputs unranked.
PiperOrigin-RevId: 323960733
2020-08-03 19:28:52 -07:00
Mehdi Amini cd01bb4c4e More cleanup in mlir-hlo to prepare for the standalone build
Shuffle files around, use TableGen to register passes, and introduce
a `mlir-hlo-opt.cpp` file to hold the main entry point of the -opt tool
and stop relying on static registration for dialect/passes.

PiperOrigin-RevId: 323674455
2020-08-03 19:28:00 -07:00
Stephan Herhut effd3fb4f9 Extend unranked to ranked pattern for hlo operations to all unary and binary ops.
As this is essentially always the same pattern, only one operation is tested.

PiperOrigin-RevId: 323525418
2020-08-03 19:27:49 -07:00
Thomas Joerg 739758f9cc Integrate LLVM at llvm/llvm-project@eed333149d
Updates LLVM usage to match
[eed333149d17](https://github.com/llvm/llvm-project/commit/eed333149d17)

PiperOrigin-RevId: 323354988
2020-08-03 19:27:25 -07:00
Robert Suderman 8023baa959 Modified HLOAbsOp lowering for differing types.
PiperOrigin-RevId: 323082107
2020-08-03 19:27:12 -07:00
Hanhan Wang 8f262ae8f5 Add support for lowering mhlo.iota to Linalg.
PiperOrigin-RevId: 322799853
2020-07-30 22:34:45 +00:00
Tres Popp 4251630426 Support CHLO broadcasting operations between scalar and unranked tensors.
This is done through reshaping the unranked tensor into a 1D ranked tensor which will result in a safe broadcast/indexing logic when the other operand is a scalar.

PiperOrigin-RevId: 322553661
2020-07-30 22:34:40 +00:00
Robert Suderman c23ad602c8 Add a transform for Gathers to torch_index_select.
Some gathers can be interpreted as torch index selects. Transforming these
cases allow torch_index_select lowerings to be used for certain gathers.

PiperOrigin-RevId: 322255835
2020-07-30 22:34:32 +00:00
Robert Suderman cc776071fe Fix namespace for complex lowerings
PiperOrigin-RevId: 322180317
2020-07-30 22:34:27 +00:00
A. Unique TensorFlower 30fa3db949 Update comments to reflect the new names of `GenericOp` constructor parameters
PiperOrigin-RevId: 321795872
2020-07-30 22:34:23 +00:00
Stephan Herhut c44e08351d Fix mhlo to lmhlo conversion for ReduceOp.
The existing conversion no longer worked and was not save to undo. Furthermore, the pattern for mhlo.return had been removed.

Also adds some tests to ensure this does not degrade again.

PiperOrigin-RevId: 321542071
2020-07-30 22:34:19 +00:00
Robert Suderman 98a1e3b108 Add an optimization that converts some Gathers to Slices.
Some Gathers can be represented as slices. This lowering transforms
these gathers into slices.

PiperOrigin-RevId: 321394868
2020-07-30 22:34:10 +00:00
Stephan Herhut 86f290896d Implement lowering of lmhlo.reshape_memref_cast to LLVM for unknown length shape operand.
PiperOrigin-RevId: 320959625
2020-07-30 22:34:02 +00:00
Uday Bondhugula d166b66cba PR #40925: [MLIR] Update lhlo.const to linalg lowering to use affine.store inste…
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/40925

…ad of std.store

The xla_lhlo.const lowering uses std.store to store a constant to
0-d memrefs. Update it to affine.store since such an access is trivially
affine (no indices). An affine.store can always be lowered to std.store.
Copybara import of the project:

--
9e18ede72fbbca107177bd742921e4cbf77adc82 by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Update lhlo.const to linalg lowering to use affine.store instead of std.store

The xla_lhlo.const lowering uses std.store to store a constant to
0-d memrefs. Update it to affine.store since such an access is trivially
affine (no indices). An affine.store can always be lowered to std.store.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/40925 from polymage-labs:lhlo_to_linalg_affine_store 9e18ede72fbbca107177bd742921e4cbf77adc82
PiperOrigin-RevId: 320623152
2020-07-30 22:33:51 +00:00
Alexander Belyaev 6eaccefdab [MLIR][LHLO] Lower ReshapeMemRefCastOp to LLVM.
PiperOrigin-RevId: 320572751
2020-07-30 22:33:47 +00:00
Mehdi Amini 506ddd9c4a Cleanup build rule names in compiler/mlir/hlo to remove the redundant/obsolete xla_ prefix
PiperOrigin-RevId: 320320140
2020-07-30 22:33:29 +00:00
Alexander Belyaev f4303855c4 Add mlir_cpu_runner tests infra for CHLO->LHLO->LLVM lowering.
PiperOrigin-RevId: 320218897
2020-07-30 22:33:25 +00:00
Mehdi Amini a575636862 Rename XlaHloDialect class into MhloDialect following the recent dialect namespace renaming
PiperOrigin-RevId: 320213526
2020-07-30 22:33:20 +00:00
Mehdi Amini 94dcb90d38 Rename xla_chlo dialect into chlo
Following on the plan of isolating the compiler/mlir/hlo directory.

PiperOrigin-RevId: 320212018
2020-07-30 22:33:16 +00:00
Mehdi Amini 7c4a5d62b5 Rename xla_lhlo dialect into lmhlo
Following on the plan of isolating the compiler/mlir/hlo directory.
Another xla_lhlo dialect will be created under compiler/mlir/xla/ later.

PiperOrigin-RevId: 320210326
2020-07-30 22:33:11 +00:00
Alexander Belyaev b076e018a8 [MLIR][LHLO] Legalize CallOp that call funcs with tensor args/results.
PiperOrigin-RevId: 320172723
2020-07-30 22:33:07 +00:00
Alexander Belyaev e8cfdee592 [MLIR][LHLO] Convert mhlo.dynamic_reshape -> lhlo.reshape_memref_cast.
PiperOrigin-RevId: 320149593
2020-07-30 22:33:02 +00:00
Alexander Belyaev 8692fde3f9 [MLIR] Convert FuncOp signature with unranked types in HLO->LHLO conversion.
PiperOrigin-RevId: 320146856
2020-07-30 22:32:58 +00:00
Mehdi Amini 8900222fed Rename `xla_hlo` dialect to `mhlo`
This is part of the current refactoring of the HLO related dialect.
`xla_hlo` will be reintroduced in a new form later.

PiperOrigin-RevId: 319916753
2020-07-30 22:32:50 +00:00
Mehdi Amini 31dc1b21eb Move XLA-independent transforms to the new MLIR-HLO directory
This is as straighforward as possible, more cleanup/rewrite to come.

PiperOrigin-RevId: 319849713
2020-07-30 22:32:40 +00:00
Mehdi Amini fcf3df1541 Move the HLO/LHLO dialects to a new directory: tensorflow/compiler/mlir/hlo
We're preparing to restructure the MLIR HLO ecosystem with 5 dialects:

- chlo: client dialect with explicit broadcast and multiple composite operations
- mhlo: hlo with dynamic shape, decouple from XLA for evolution purpose
- lmhlo: same as above, but after buffer assignment.
- xla_hlo: mapping 1:1 to the XLA HloInstruction class.
- xla_lhlo: same as above, but after buffer assignment.

The first three dialects are intended to live in the new tensorflow/compiler/mlir/hlo
path, the latter two will be created in tensorflow/compiler/mlir/xla.

This patch only moves the directory, will followup with other transformations and tests.

The structure of the new directory follows: https://llvm.discourse.group/t/rfc-canonical-file-paths-to-dialects/621 as we intend to make it a standalone buildable component (see also https://github.com/google/mlir-npcomp as another example).

PiperOrigin-RevId: 319273229
2020-07-30 22:32:32 +00:00