Commit Graph

106 Commits

Author SHA1 Message Date
Jacques Pienaar 8cc6cddf9f Use location instead of unknown loc
Follow up on TODO as placeholder has been added to rewriter.

PiperOrigin-RevId: 331488583
2020-09-13 23:52:40 -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 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
Robert Suderman 81d51d810b Bitwise and/or/xor folders
Includes both and/or/xor on same inputs, constant all ones/zeros single arg
folder, and constant input folders.

PiperOrigin-RevId: 330610858
2020-09-08 16:27:13 -07:00
Mehdi Amini 9d4273b5a7 Delete mhlo static dialect registration
This is obsolete now.

PiperOrigin-RevId: 330605210
2020-09-08 15:57:36 -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
A. Unique TensorFlower 73b4861f2c Canonicalize mhlo.gather to mhlo.slice if it has a single set of constant indices
PiperOrigin-RevId: 330380755
2020-09-07 07:28:46 -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 f5d12604ed Add support for legalizing mhlo.transpose to lmhlo.transpose
PiperOrigin-RevId: 330127758
2020-09-04 14:59:08 -07:00
Prakalp Srivastava b7248424ae Fix return type of replica-id to unsigned int 32 tensor.
PiperOrigin-RevId: 330012181
2020-09-03 16:08:38 -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
Robert Suderman 7c93352a40 Scalar / Trivial folding for mhlo.select
This covers the case where the predicate is a splat or the on_true/on_false
values are the same.

PiperOrigin-RevId: 329622785
2020-09-01 18:34:12 -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
Uday Bondhugula 94296bb7ec PR #42509: [MLIR] Add folder for mhlo get_dimension_size
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/42509

Add folder for mhlo GetDimensionSizeOp (`mhlo.get_dimension_size`).
`get_dimension_size` folds to a constant when the corresponding tensor
dimension size is statically known / constant.
Copybara import of the project:

--
5994915525ec2e932125aa1f133ce2260ba100af by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Add folder for mhlo get_dimension_size

Add folder for mhlo GetDimensionSizeOp. get_dimension_size folds to a
constant when the corresponding tensor dimension size is statically
known / constant.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/42509 from polymage-labs:get_dimension_size_fold 5994915525ec2e932125aa1f133ce2260ba100af
PiperOrigin-RevId: 328222517
2020-08-24 15:36:30 -07:00
Mehdi Amini 73b5a44f33 Update third_party/tensorflow/compiler/mlir/tensorflow/utils/... to not depend on the global Dialect Registry (NFC)
PiperOrigin-RevId: 328171679
2020-08-24 10:59:22 -07:00
Uday Bondhugula 282dba6d38 PR #42508: [MLIR] Erase dead lmhlo.constant ops
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/42508

An lmhlo.constant op on an memref that is locally allocated and with
no users other than dealloc's can be deleted. Add a canonicalization
pattern for this.
Copybara import of the project:

--
8758c409a15f567e7cb8e1077faa020f5705c85a by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Erase dead lmhlo.constant ops

An lmhlo.constant op on an memref that is locally allocated and with
no other users (other than dealloc's) can be deleted. Add a
canonicalization patter for this.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/42508 from polymage-labs:lhlo_constant_erase 8758c409a15f567e7cb8e1077faa020f5705c85a
PiperOrigin-RevId: 328042416
2020-08-23 12:28:54 -07:00
A. Unique TensorFlower 75d50980f5 Integrate LLVM at llvm/llvm-project@e75bc5c791
Updates LLVM usage to match
[e75bc5c791e0](https://github.com/llvm/llvm-project/commit/e75bc5c791e0)

PiperOrigin-RevId: 327449456
2020-08-19 09:39:06 -07:00
Richard Uhler 9a232c7012 Fix failing segment_reduction_ops_mlir_bridge_test
By adding support for complex types to GetScalarOfType and using appropriate
choice of limits for initial values in the unsorted segment reduction ops.

PiperOrigin-RevId: 327061577
2020-08-17 11:28:20 -07:00
River Riddle 7acc765ed7 Remove kindof definitions on derived mlir Attribute and Type classes
This method no longer needs to be explicitly defined by the user, so all of the existing definitions are dead and can be removed.

PiperOrigin-RevId: 326350650
2020-08-12 17:42:07 -07:00
Jacques Pienaar 5dac76f4af Add chlo.constant_like op which splats a constant to shape of operand
This allows specifying a constant whose shape is only known when operand shape is. Also use it to update tf.Acos legalization.

PiperOrigin-RevId: 325860604
2020-08-11 14:54:48 -07:00
Jacques Pienaar 17ccca7f4b Add HLO RngBitGenerator
This adds the XlaBuilder RngBitGenerator to the MHLO dialect. The op is currently represented very directly using int attribute for random algorithm and direct import/export.

PiperOrigin-RevId: 325814134
2020-08-10 09:01:11 -07:00
A. Unique TensorFlower e6fa003bf2 Integrate LLVM at llvm/llvm-project@b6d9add71b
Updates LLVM usage to match
[b6d9add71b1a](https://github.com/llvm/llvm-project/commit/b6d9add71b1a)

PiperOrigin-RevId: 325589103
2020-08-08 04:35:25 -07:00
Andy Ly 53fdda7f3e Update mhlo.constant to use a custom assembly format instead of a custom printer and parser (NFC).
PiperOrigin-RevId: 325560779
2020-08-07 22:23:06 -07:00
Mehdi Amini 5f8da992f2 Fix mlir-hlo CMakefiles after internal change
PiperOrigin-RevId: 325366142
2020-08-07 22:20:50 -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
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
Andy Ly c340367702 Add canonicalization for unpacking and repacking the same tuple (e.g. tuple -> get_tuple_element -> tuple).
These unpacking and repacking of tuples may be generated when modifying tuple arguments or results.

PiperOrigin-RevId: 325162694
2020-08-05 21:38:02 -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 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 7bff244b8b Allow index typed memrefs in reshape_memref_cast.
With the recent change to allow memref of index in MLIR core, we should also
allow this in the mhlo dialect.

PiperOrigin-RevId: 324879354
2020-08-04 13:34:11 -07:00
A. Unique TensorFlower bf8836effb Integrate LLVM at llvm/llvm-project@b5059b7140
Updates LLVM usage to match
[b5059b714023](https://github.com/llvm/llvm-project/commit/b5059b714023)

PiperOrigin-RevId: 324754243
2020-08-03 23:05:03 -07:00
Andy Ly 4c8fead3e0 Add support for token operands to mhlo.tuple.
mhlo.get_tuple_element supports extracting a mhlo.token type from a tuple. This updates the creation of tuples to allow for mhlo.token typed operands.

PiperOrigin-RevId: 324628663
2020-08-03 19:31:15 -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 1c535f1718 Restrict GetDimensionSize HLO op result type to 32 bit integer
XLA implementation has this limitation and always uses 32 bit result for this instruction. This will cause mismatch between the result type in MLIR and XLA at the time of export.

This should be resolved once we have a special dialect mapping directly to HLOInstructionProto. Another option until then could be to introduce a pass to legalize mhlo itself to match XLA semantics.

PiperOrigin-RevId: 324286936
2020-08-03 19:30:15 -07:00
Tim Shen fcb91fb0b9 [MLIR] Add conversion between XLA Fusion and MHLO FusionOp.
PiperOrigin-RevId: 324279065
2020-08-03 19:30:03 -07:00
Stephan Herhut 9cbe5f2285 Constrain mhlo.const to static shaped tensors.
Constants of unknown shape cannot be materialized. In most cases, one likely wants to use a scalar constant and rely on broadcasting instead.

PiperOrigin-RevId: 324252475
2020-08-03 19:29:51 -07:00
Jacques Pienaar d83d437588 Add chlo.acos and legalization
Add client HLO op for arc cosine and legalize TF op to it & legalization from it to HLO.

PiperOrigin-RevId: 324053167
2020-08-03 19:29:17 -07:00
Robert Suderman de5ddaf7c9 HLO Random operations should match shape constraints.
PiperOrigin-RevId: 323844002
2020-08-03 19:28:23 -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
Kazuaki Ishizaki 882468da13 PR #41662: NFC - minor spelling tweaks
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/41662

This PR addresses minor spelling tweaks in documents
Copybara import of the project:

--
b806191a117990a479944b40ec7a4b79843287a2 by Kazuaki Ishizaki <ishizaki@jp.ibm.com>:

fix trivial typo

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/41662 from kiszk:spelling_tweaks_docs b806191a117990a479944b40ec7a4b79843287a2
PiperOrigin-RevId: 322955351
2020-08-03 19:23:45 -07: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 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 7a6adc6a84 Add canonicalization patterns for dynamic_broadcast_in_dim where the target shape is the shape of the operand.
PiperOrigin-RevId: 321312182
2020-07-30 22:34:06 +00:00
Robert Suderman 06ae59074f Fold xla iota across a 1-length dimension into a zero value
Iota across length-1 is just a constant. Fold into it.

PiperOrigin-RevId: 320443468
2020-07-30 22:33:43 +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
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 e8cfdee592 [MLIR][LHLO] Convert mhlo.dynamic_reshape -> lhlo.reshape_memref_cast.
PiperOrigin-RevId: 320149593
2020-07-30 22:33:02 +00:00
Robert Suderman e1651b6090 Canonicalize multidimensional iota to use broadcast
There is no reason to have a multidimensional iota for codegen.
This should be canonicalized to a single dimensional iota followed
by a broadcast. Changing iota to on a single dimension  and a broadcast
substantially simplifies implementing iota operations.

PiperOrigin-RevId: 320095470
2020-07-30 22:32:54 +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
Alexander Belyaev 72010faaa7 [MLIR][LHLO] Add ReshapeMemrefCastOp to LHLO.
PiperOrigin-RevId: 319799171
2020-07-30 22:32:36 +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