Commit Graph

107 Commits

Author SHA1 Message Date
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
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
Mehdi Amini 56689ce908 Temporarily mark hlo-legalize-to-lhlo.mlir as failing, pending fixes
Upstream changed the buffer assignment API and this needs adjustment on our side.

PiperOrigin-RevId: 329959756
2020-09-03 11:36:20 -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
A. Unique TensorFlower a622bf479b Fix mhlo::SliceOp::fold to not crash on unknown shapes
PiperOrigin-RevId: 329504383
2020-09-01 07:37:36 -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
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
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
A. Unique TensorFlower 1bcdcc3b3b Integrate LLVM at llvm/llvm-project@9c5e25a696
Updates LLVM usage to match
[9c5e25a696b2](https://github.com/llvm/llvm-project/commit/9c5e25a696b2)

PiperOrigin-RevId: 327140496
2020-08-17 18:22:29 -07:00
Robert Suderman 5252aeae8f Lowering for mhlo.ceil to std.ceil
PiperOrigin-RevId: 326335301
2020-08-12 16:15:35 -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
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
Geoffrey Martin-Noble 2b0a244a6b Add license header to lit.site.cfg.py.in
PiperOrigin-RevId: 325515949
2020-08-07 22:22:24 -07:00
Lucy Fox d742477c02 Verify that MHLO DynamicUpdateSlice start indices have matching element types.
HLO requires that the element types match for all start index parameters. Right now we don't catch this invalid case until export, so adding a check in the verifier so that we catch this sooner.

This also requires a small tweak to the TF InplaceUpdate op lowering.

PiperOrigin-RevId: 325463796
2020-08-07 22:21:40 -07:00
Lucy Fox cd22ecd136 Relax DynamicBroadcastInDim verifier when dimensions are dynamic.
For input and output dimensions which must match, we shouldn't fail in the case where one dim is dynamic and the other is static. This is insufficient information to conclude a dimension mismatch.

PiperOrigin-RevId: 325344738
2020-08-07 22:18:38 -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
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 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
A. Unique TensorFlower 9a97cffbb5 Integrate LLVM at llvm/llvm-project@28e322ea93
Updates LLVM usage to match
[28e322ea9393](https://github.com/llvm/llvm-project/commit/28e322ea9393)

PiperOrigin-RevId: 324957510
2020-08-04 22:31:45 -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
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 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
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
Stephan Herhut b09bf2a4dc Move hlo end to end tests to the hlo directory tree.
PiperOrigin-RevId: 323955773
2020-08-03 19:28:37 -07:00
Stephan Herhut 1b0eb4baa7 Do not mandate the result type of shape computations but have it be inferred from context.
The computation of a broadcasted shape forced the use of the shape type unnecessarily, which blocked further canonicalizations.

PiperOrigin-RevId: 323783998
2020-08-03 19:28:12 -07:00
Benjamin Kramer b7c4314e7f Allow running tests with FILECHECK_OPTS=-enable-var-scope
PiperOrigin-RevId: 323420636
2020-08-03 19:27:38 -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
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
A. Unique TensorFlower c8bb0ff54d Integrate LLVM at https://github.com/llvm/llvm-project/commit/f233b92f92a6
PiperOrigin-RevId: 321454533
2020-07-30 22:34:14 +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
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
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
Stephan Herhut c3be2474dd Integrate LLVM at https://github.com/llvm/llvm-project/commit/305b500eaf8c
PiperOrigin-RevId: 320395881
2020-07-30 22:33:38 +00:00
Uday Bondhugula de0578b4f9 PR #40745: [MLIR] Add constant folder for xla_hlo.broadcast_in_dim op
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/40745

Fold broadcast_in_dim op if the operand is the result of a tensor splat.
Copybara import of the project:

--
26c9f631448b8d6ffd20ece39ea8d4132b5550c7 by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Add constant folder for xla_hlo.broadcast_in_dim op

Fold broadcast_in_dim op if the operand is the result of a tensor
splat.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/40745 from polymage-labs:broadcast_in_dim_fold 26c9f631448b8d6ffd20ece39ea8d4132b5550c7
PiperOrigin-RevId: 320365164
2020-07-30 22:33:34 +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 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
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 fa057cc0bc Move HLO tests to `third_party/tensorflow/compiler/mlir/hlo/`
Also add a localized `mlir-hlo-opt` binary for the testing of
tensorflow/compiler/mlir/hlo/... ; this directory is intended to be self-contained
and depend only on MLIR.

PiperOrigin-RevId: 319878984
2020-07-30 22:32:45 +00:00