Commit Graph

1074 Commits

Author SHA1 Message Date
Chris Jones 968226b9d7 [XLA:GPU] Add AllReduce{Start,Done} to MLIR LHLO dialect.
PiperOrigin-RevId: 378640706
2021-06-10 06:54:42 -07:00
Adrian Kuegel 6088eb697c Fix Cosh approximation for F16.
We should upcast F16 to F32 to prevent precision loss.
E.g. cosh(-9) would evaluate to 4042 previously instead of 4052.
This allows to enable the MLIR generated kernel for F16 type.
Also move template instantiation for Sinh to inside the #ifdef block.
This was missed in a previous commit.

PiperOrigin-RevId: 378635042
2021-06-10 06:16:44 -07:00
A. Unique TensorFlower 837a1de7c5 Integrate LLVM at llvm/llvm-project@e11b5b87be
Updates LLVM usage to match
[e11b5b87bebf](https://github.com/llvm/llvm-project/commit/e11b5b87bebf)

PiperOrigin-RevId: 378589304
2021-06-10 00:18:25 -07:00
A. Unique TensorFlower 9f67417b41 [MLIR][HLO] Avoid duplicate cluster operands when merging
When merging rank specialization clusters, avoid duplicating operands. A fewer
number of operands usually allows better rank specialization.

PiperOrigin-RevId: 378445946
2021-06-09 10:54:55 -07:00
A. Unique TensorFlower b580722041 [MLIR][KernelGen] Merge rank specialization clusters
Merge adjacent rank specialization clusters. Combine their operands, bodies, and
results.

PiperOrigin-RevId: 378433222
2021-06-09 10:07:47 -07:00
Adrian Kuegel b6d8160611 Add Broadcasting and BroadcastingElementwise traits to ConstantLikeOp.
This allows to include such ops in rank specialization clusters.

PiperOrigin-RevId: 378380915
2021-06-09 05:09:26 -07:00
A. Unique TensorFlower b9e45007d5 [MLIR][HLO] Extend broadcast propagation pass to enable more fusion
Move element-wise operations into assuming regions. This enables fusion
opportunities within the region.

PiperOrigin-RevId: 378362725
2021-06-09 03:03:37 -07:00
A. Unique TensorFlower d828b457b3 Handle empty tensors in SimplifyConcatSlice.
If the result of the slice is an empty tensor, do nothing.
This fixes a crash: we can't create a `concat` with an
empty operand range.

PiperOrigin-RevId: 378354956
2021-06-09 02:15:47 -07:00
Mehdi Amini 1770ed455f Remove unnecessary duplicated source from "lhlo" (NFC)
PiperOrigin-RevId: 378291564
2021-06-08 18:07:26 -07:00
A. Unique TensorFlower 4134923d4f Integrate LLVM at llvm/llvm-project@f96b5e801d
Updates LLVM usage to match
[f96b5e801d67](https://github.com/llvm/llvm-project/commit/f96b5e801d67)

PiperOrigin-RevId: 378139137
2021-06-08 06:24:26 -07:00
Adrian Kuegel 9a8c254526 Support complex types for Sinh.
Because mhlo::ConstantLike doesn't support complex types, we need to use
GetScalarOfType and broadcast it to the needed shape.
Disable the tf2xla fallback, now that MLIR fully supports Sinh.

PiperOrigin-RevId: 378123151
2021-06-08 04:23:19 -07:00
A. Unique TensorFlower c47869f931 [MLIR][HLO] Rename `move-up-dynamic-broadcasts-for-fusion` to `broadcast-propagation`
PiperOrigin-RevId: 378102608
2021-06-08 01:51:10 -07:00
A. Unique TensorFlower b2839c735b Integrate LLVM at llvm/llvm-project@8344e215ec
Updates LLVM usage to match
[8344e215ec6c](https://github.com/llvm/llvm-project/commit/8344e215ec6c)

PiperOrigin-RevId: 378043689
2021-06-07 17:32:52 -07:00
A. Unique TensorFlower c11de49300 Integrate LLVM at llvm/llvm-project@7ed7d4ccb8
Updates LLVM usage to match
[7ed7d4ccb899](https://github.com/llvm/llvm-project/commit/7ed7d4ccb899)

PiperOrigin-RevId: 377972571
2021-06-07 12:05:09 -07:00
Benjamin Kramer d1c60df2fe [MHLO:linalg] Be more aggressive about turning mhlo.const into std.constant
On tensors the only difference between these ops is that mhlo.const supports unsigned types.

PiperOrigin-RevId: 377970948
2021-06-07 11:58:23 -07:00
Hanhan Wang 25b93c8d66 Add support for lowering mhlo.iota/dynamic_iota to Linalg on unsigned types.
PiperOrigin-RevId: 377956338
2021-06-07 10:59:33 -07:00
Adrian Kuegel 5315997402 Fix Sinh approximation for F16.
We should upcast F16 to F32 to prevent precision loss.
E.g. sinh(-9) would evaluate to -4042 previously instead of -4052.
This allows to enable the MLIR generated kernel for F16 type.

PiperOrigin-RevId: 377901896
2021-06-07 06:38:42 -07:00
Tobias Gysi fc723380e6 Update lhlo to use the new structured op interface.
Replace deprecated methods in lhlo_fuse_linalg.cc. The new structured op interface has been introduced in https://reviews.llvm.org/D103394.

PiperOrigin-RevId: 377875452
2021-06-07 03:11:03 -07:00
Wenyi Zhao ade873a5e0 PR #49970: [MLIR][DISC] bufferize DynamicReshape and DynamicBroadcastInDim
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/49970

1, add hlo-to-lhlo support for DynamicReshape and DynamicBroadcastInDim

2, add a flag `convert-to-lmhlo-only` to seperate following two case:
   - hlo-to-lhlo only. Simply lowers all mhlo ops to their lmhlo
     counterparts, do not apply any optimization (e.g. elide any
     buffer copy). Buffer optimization is not easy in dynamic
     shape world especially when involving control flow, thus we
     leave this to another dedicated pass.

   - hlo-to-lhlo-or-memref-directly. Lowers some metadata-only mhlo
     ops (e.g. reshape) to memref dialect directly and Lowers others
     to their lmhlo counterparts.
Copybara import of the project:

--
562bd65a368f6194405c4ae6900e3b4388a5ec03 by Wenyi Zhao <reyizero@gmail.com>:

[MLIR][DISC] bufferize DynamicReshape and DynamicBroadcastInDim

1, add hlo-to-lhlo support for DynamicReshape and DynamicBroadcastInDim

2, add a flag `convert-to-lmhlo-only` to seperate following two case:
   - hlo-to-lhlo only. Simply lowers all mhlo ops to their lmhlo
     counterparts, do not apply any optimization (e.g. elide any
     buffer copy). Buffer optimization is not easy in dynamic
     shape world especially when involving control flow, thus we
     leave this to another dedicated pass.

   - hlo-to-lhlo-or-memref-directly. Lowers some metadata-only mhlo
     ops (e.g. reshape) to memref dialect directly and Lowers others
     to their lmhlo counterparts.

PiperOrigin-RevId: 377603395
2021-06-04 15:36:03 -07:00
A. Unique TensorFlower 8b3a75ea25 Integrate LLVM at llvm/llvm-project@b109172d99
Updates LLVM usage to match
[b109172d993e](https://github.com/llvm/llvm-project/commit/b109172d993e)

PiperOrigin-RevId: 377549160
2021-06-04 11:11:34 -07:00
A. Unique TensorFlower 9c895af2f1 Integrate LLVM at llvm/llvm-project@23a116c8c4
Updates LLVM usage to match
[23a116c8c446](https://github.com/llvm/llvm-project/commit/23a116c8c446)

PiperOrigin-RevId: 377501435
2021-06-04 06:46:18 -07:00
A. Unique TensorFlower f1f4c903df Integrate LLVM at llvm/llvm-project@fcf8827a98
Updates LLVM usage to match
[fcf8827a98be](https://github.com/llvm/llvm-project/commit/fcf8827a98be)

PiperOrigin-RevId: 377485560
2021-06-04 04:31:28 -07:00
A. Unique TensorFlower db05388a3c Integrate LLVM at llvm/llvm-project@da3ed58b97
Updates LLVM usage to match
[da3ed58b97c1](https://github.com/llvm/llvm-project/commit/da3ed58b97c1)

PiperOrigin-RevId: 377432380
2021-06-03 20:45:18 -07:00
A. Unique TensorFlower aba16adfa5 Add `mhlo.all_gather` op to MHLO dialect.
Adds import/export/verifier support as well.
Also makes `channel_handle` uniform across mhlo.all_reduce and mhlo.all-gather.

PiperOrigin-RevId: 377323468
2021-06-03 10:45:29 -07:00
Jacques Pienaar 4fc2e87a42 Add mhlo python binding generator target
This just invokes the generator backend & creates a filegroup.

PiperOrigin-RevId: 377318653
2021-06-03 10:26:30 -07:00
A. Unique TensorFlower fe42a08fc9 Use channel_handle for ChannelHandles in MHLO ops. This makes the naming of these properties consistent across these ops.
PiperOrigin-RevId: 377309518
2021-06-03 09:49:47 -07:00
A. Unique TensorFlower 063086dd78 Integrate LLVM at llvm/llvm-project@c89dff5855
Updates LLVM usage to match
[c89dff5855bb](https://github.com/llvm/llvm-project/commit/c89dff5855bb)

PiperOrigin-RevId: 377199586
2021-06-02 19:36:24 -07:00
A. Unique TensorFlower 4620410f18 Integrate LLVM at llvm/llvm-project@b25546a4b4
Updates LLVM usage to match
[b25546a4b406](https://github.com/llvm/llvm-project/commit/b25546a4b406)

PiperOrigin-RevId: 377077163
2021-06-02 09:32:59 -07:00
A. Unique TensorFlower 75a1c450ea [MLIR][KernelGen] Fix Windows build failure
Fix usage of default constructor. Instead, always use the parameterized
constructor and make the maximum supported rank explicit.

PiperOrigin-RevId: 377037155
2021-06-02 05:34:44 -07:00
A. Unique TensorFlower 557e56362e [MLIR][KernelGen] Simplify rank specialization tests with smaller target rank
For the tests rank specialize only up to rank 3. The remaining cases for higher
ranks are analogous.

PiperOrigin-RevId: 377024370
2021-06-02 03:48:07 -07:00
wyzhao 968d4b8709 PR #49598: [MLIR][DISC] legalize tensor_load inserted during hlo-to-lhlo conversion
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/49598

This PR implements logic for lowering memref.tensor_load ops that are
inserted during `mhlo-legalize-to-lmhlo`
Copybara import of the project:

--
80eb377af4e02182e1aecc943a41ca5d7d1c2100 by Wenyi Zhao <reyizero@gmail.com>:

[MLIR][DISC] legalize tensor_load inserted during hlo-to-lhlo conversion

This PR implements logic for lowering memref.tensor_load ops that are
inserted during `mhlo-legalize-to-lmhlo`.

--
ac452fe3dcd591211cd5c59be9189fe2f7153b41 by Wenyi Zhao <reyizero@gmail.com>:

minor fix

--
6b36017f8632a06adbc3e05a62975fa641d0260f by Wenyi Zhao <reyizero@gmail.com>:

minor refine

--
846005cc76d0033112e47825c2e9a97790b6925f by Wenyi Zhao <reyizero@gmail.com>:

minor fix

--
f6a4becaa287d5ca323b2d152a4d0ae053730fd9 by Wenyi Zhao <reyizero@gmail.com>:

fix

--
5555749f60f7fce8f57962860ef65efccf0362ba by Wenyi Zhao <reyizero@gmail.com>:

fix

--
8873b9b6d9315c1199ca9f7c133ecf377ecd2fa6 by Wenyi Zhao <reyizero@gmail.com>:

fix

PiperOrigin-RevId: 376942547
2021-06-01 16:27:56 -07:00
A. Unique TensorFlower 5baf6e7709 Integrate LLVM at llvm/llvm-project@97d234935f
Updates LLVM usage to match
[97d234935f15](https://github.com/llvm/llvm-project/commit/97d234935f15)

PiperOrigin-RevId: 376840022
2021-06-01 08:38:32 -07:00
A. Unique TensorFlower d1828625ab [MLIR][KernelGen] Make maximum supported rank in rank specialization configurable
The maximum supported target rank of 5 is sufficient for all operations but
`select`. Make the maximum target rank configurable in the rank specialization.
This reduces the number of generated kernels for operations that don't require
it.

PiperOrigin-RevId: 376822496
2021-06-01 06:54:31 -07:00
A. Unique TensorFlower c7c245eaf1 [MLIR][KernelGen] Add MLIR-generated Xlogy kernel
Add the first MLIR-generated kernel that relies on an in-TF lowering. Fusion for
this kernel relies on the generalized rank specialization for operation groups.

PiperOrigin-RevId: 376805435
2021-06-01 04:48:18 -07:00
A. Unique TensorFlower 6a570502b6 Integrate LLVM at llvm/llvm-project@f000c4cfb6
Updates LLVM usage to match
[f000c4cfb66c](https://github.com/llvm/llvm-project/commit/f000c4cfb66c)

PiperOrigin-RevId: 376798033
2021-06-01 03:51:22 -07:00
A. Unique TensorFlower f16e5a3a67 [MLIR][HLO] Use canonicalization patterns in broadcast propagation pass
Replace local canonicalization patterns with those from upstream.

PiperOrigin-RevId: 376794178
2021-06-01 03:14:31 -07:00
A. Unique TensorFlower adff130b12 Integrate LLVM at llvm/llvm-project@41d7909368
Updates LLVM usage to match
[41d7909368be](https://github.com/llvm/llvm-project/commit/41d7909368be)

PiperOrigin-RevId: 376785495
2021-06-01 02:10:18 -07:00
A. Unique TensorFlower 31536431e0 [MLIR][HLO] Eliminate duplicate broadcastable constraints
PiperOrigin-RevId: 376718433
2021-05-31 13:50:23 -07:00
A. Unique TensorFlower 0f341012c6 [MLIR][HLO] Eliminate duplicate broadcastable constraints
PiperOrigin-RevId: 376715240
2021-05-31 13:08:02 -07:00
A. Unique TensorFlower 511a1db4f3 [MLIR][HLO] Use canonicalization patterns in broadcast propagation pass
Replace local canonicalization patterns with those from upstream.

PiperOrigin-RevId: 376708719
2021-05-31 12:01:26 -07:00
A. Unique TensorFlower 5f5db13715 [MLIR][HLO] Use canonicalization patterns in broadcast propagation pass
Replace local canonicalization patterns with those from upstream.

PiperOrigin-RevId: 376707588
2021-05-31 11:44:50 -07:00
A. Unique TensorFlower 1f786eb934 Integrate LLVM at llvm/llvm-project@c36ff6424f
Updates LLVM usage to match
[c36ff6424f24](https://github.com/llvm/llvm-project/commit/c36ff6424f24)

PiperOrigin-RevId: 376687137
2021-05-31 08:16:39 -07:00
A. Unique TensorFlower cc1b22e618 [HLO][Linalg] Support scalar broadcasts in point-wise converter
This is needed for operations that support this limited form of broadcasting,
namely `mhlo.select`.

PiperOrigin-RevId: 376655844
2021-05-31 03:50:23 -07:00
A. Unique TensorFlower 832f39b871 Integrate LLVM at llvm/llvm-project@577fea4e1a
Updates LLVM usage to match
[577fea4e1a13](https://github.com/llvm/llvm-project/commit/577fea4e1a13)

PiperOrigin-RevId: 376640971
2021-05-31 01:54:46 -07:00
A. Unique TensorFlower 8110683c0b Integrate LLVM at llvm/llvm-project@9968896cd6
Updates LLVM usage to match
[9968896cd62a](https://github.com/llvm/llvm-project/commit/9968896cd62a)

PiperOrigin-RevId: 376382711
2021-05-28 12:09:55 -07:00
A. Unique TensorFlower e422ce5a9a Integrate LLVM at llvm/llvm-project@8086f9d87e
Updates LLVM usage to match
[8086f9d87ee8](https://github.com/llvm/llvm-project/commit/8086f9d87ee8)

PiperOrigin-RevId: 376289371
2021-05-27 18:20:22 -07:00
Hanhan Wang 402b74ed7f Fix type bug in mhlo.dynamic-update-slice lowering.
The operand type can be f32. We should not use operand type to do clamp
operations.

PiperOrigin-RevId: 376286524
2021-05-27 17:53:49 -07:00
A. Unique TensorFlower b5f444232f Integrate LLVM at llvm/llvm-project@d47dd11071
Updates LLVM usage to match
[d47dd1107132](https://github.com/llvm/llvm-project/commit/d47dd1107132)

PiperOrigin-RevId: 376269354
2021-05-27 16:06:54 -07:00
Benjamin Kramer 6ea8f1ddf8 Integrate LLVM at llvm/llvm-project@eba69b59d1
Updates LLVM usage to match
[eba69b59d1a3](https://github.com/llvm/llvm-project/commit/eba69b59d1a3)

PiperOrigin-RevId: 376175310
2021-05-27 08:37:41 -07:00
Benjamin Kramer e319f87857 Integrate LLVM at llvm/llvm-project@c467585682
Updates LLVM usage to match
[c467585682dc](https://github.com/llvm/llvm-project/commit/c467585682dc)

PiperOrigin-RevId: 376155243
2021-05-27 06:25:17 -07:00