81e8d778c4
Fix bug using std.rank instead of shape.rank
...
PiperOrigin-RevId: 339890070
2020-10-30 09:59:24 -07:00
76b30fd426
Move unranked chlo lowering to transform_unranked_hlo.
...
Additionally:
- Forward listeners through new if/else op builders.
This corrects an error that led to incomplete legalization of broadcasted op
lowering.
- Use OpConversionPattern to ensure up to date operand values are used.
PiperOrigin-RevId: 339838833
2020-10-30 02:56:44 -07:00
880f603239
Drop OperationFolder usage with Linalg fusion.
...
PiperOrigin-RevId: 339653466
2020-10-29 06:10:14 -07:00
3a6580bf75
Integrate LLVM at llvm/llvm-project@4d11daa659
...
Updates LLVM usage to match
[4d11daa659a1](https://github.com/llvm/llvm-project/commit/4d11daa659a1 )
PiperOrigin-RevId: 339631220
2020-10-29 03:02:38 -07:00
3bf4277ea4
[MLIR] Add a lmhlo.reduce -> linalg.generic converter
...
Doesn't support tensors right now, as it's somewhat hairy to support both at
the same time. Since we use a generic lowering the result is messy
and needs a mem2reg pass to eliminate extra load/store/allocas.
PiperOrigin-RevId: 339562971
2020-10-28 16:38:19 -07:00
7363748bae
Integrate LLVM at llvm/llvm-project@0fc1aa22ee
...
Updates LLVM usage to match
[0fc1aa22ee6a](https://github.com/llvm/llvm-project/commit/0fc1aa22ee6a )
PiperOrigin-RevId: 339239851
2020-10-27 06:56:16 -07:00
b7aa01dbe0
Integrate LLVM at llvm/llvm-project@26750a1264
...
Updates LLVM usage to match
[26750a1264b3](https://github.com/llvm/llvm-project/commit/26750a1264b3 )
PiperOrigin-RevId: 339127601
2020-10-26 15:11:52 -07:00
6eda9ed273
Add compare_type optional attribute to CompareOp in HLO dialects
...
If unspecified, `compare_type` is FLOAT for float element types, SIGNED for signed element types and UNSIGNED for unsigned element types. compare_type can be TOTALORDER for float element types.
- Added import and export support the attribute.
- Restricted legalization from HLO to TF to the default compare types.
- Updated existing usage of the CompareOp
PiperOrigin-RevId: 339099219
2020-10-26 12:58:29 -07:00
f9843fabe1
Use InferTypeOpInterface for HLO_SliceOp.
...
Instead of having a custom builder to construct a slice op without an explicit
return type.
PiperOrigin-RevId: 339058864
2020-10-26 09:54:13 -07:00
444fae9bac
[NFC] Make naming style consistent.
...
Use lowercase with underscores between words instead of camelStyle.
PiperOrigin-RevId: 338722328
2020-10-23 12:23:25 -07:00
c3f4ab06d3
Use OpConversionPattern instead of BufferizeOpConversionPattern
...
As described in mlir/Transforms/Bufferize.h, patterns that don't need the special methods on a BufferizeTypeConverter should use a regular OpConversionPattern.
PiperOrigin-RevId: 338424819
2020-10-22 01:20:43 -07:00
ae00ae487c
[mlir] Simplify DDR matching patterns with equal operands for operators.
...
This https://reviews.llvm.org/D89254 diff introduced implicit matching between same name arguments. Modify usages accordingly.
PiperOrigin-RevId: 338090110
2020-10-20 10:47:09 -07:00
7a983ea389
Add folder for mhlo::pad
...
PiperOrigin-RevId: 337827560
2020-10-19 04:21:44 -07:00
4a18aa41ee
Add folder to mhlo::round_nearest_afz
...
PiperOrigin-RevId: 337823786
2020-10-19 03:45:15 -07:00
27968619b7
Verify non-scalar inputs for HLO concat
...
XLA HLO concat does not accept scalars, so fail verification if this occurs. Avoids segfault when accessing an empty output shape.
PiperOrigin-RevId: 337618167
2020-10-16 19:39:31 -07:00
2e30b59ddc
Extend hlo-fuse-linalg slightly to support aliased returns.
...
The fusion heuristic identifies the root of a fusion by checking whether an
output of a linalg operation is a function result. It did not consider outputs
flowing through aliasing operations (like casts).
PiperOrigin-RevId: 337479910
2020-10-16 04:04:23 -07:00
51cd4200b6
Make LMHLO's Dot have the same power as MHLO's DotGeneral.
...
PiperOrigin-RevId: 337391565
2020-10-15 15:09:06 -07:00
05ee41baf8
Add folder for mhlo::scatter
...
PiperOrigin-RevId: 337274351
2020-10-15 03:26:05 -07:00
0f36979e2c
Integrate LLVM at llvm/llvm-project@220de1f32a
...
Updates LLVM usage to match
[220de1f32add](https://github.com/llvm/llvm-project/commit/220de1f32add )
PiperOrigin-RevId: 337270545
2020-10-15 02:54:24 -07:00
f6b4e6758a
Add GPU specific LMHLO level ops
...
- Introduce operations in a new lmhlo_gpu dialect that map to GPU library function calls
in the XLA:GPU backend.
- Add basic unit tests as well.
PiperOrigin-RevId: 337132166
2020-10-14 11:23:55 -07:00
15fe5ca737
Integrate LLVM at llvm/llvm-project@6713332fdd
...
Updates LLVM usage to match
[6713332fddb7](https://github.com/llvm/llvm-project/commit/6713332fddb7 )
PiperOrigin-RevId: 336988895
2020-10-13 17:21:14 -07:00
53d16d9f8e
Integrate LLVM at llvm/llvm-project@93377888ae
...
Updates LLVM usage to match
[93377888ae89](https://github.com/llvm/llvm-project/commit/93377888ae89 )
PiperOrigin-RevId: 336669842
2020-10-12 08:42:31 -07:00
f6af1fc134
Support hlo to lhlo buffer placement through shape.assuming ops.
...
PiperOrigin-RevId: 336287728
2020-10-09 07:14:12 -07:00
d986bd7ad7
Use tensor_cast instead of mhlo::reshape in the lowering of unranked binary operations.
...
We know that the value already is a scalar and we just want to update the type, so no need to reshape anything.
PiperOrigin-RevId: 336252315
2020-10-09 01:46:48 -07:00
3736c5542f
[MLIR][KernelGen] Fix unranked codegeneration in kernel generator
...
PiperOrigin-RevId: 335847086
2020-10-07 05:39:55 -07:00
84277ed784
Remove unused `using` statements.
...
PiperOrigin-RevId: 335517711
2020-10-05 15:53:59 -07:00
a33441907c
[MLIR][KernelGen] Add E2E test for `tf.Acos`
...
PiperOrigin-RevId: 335419310
2020-10-05 08:14:36 -07:00
bae0815ef0
[MLIR][KernelGen] Legalize `atan` to approximation
...
PiperOrigin-RevId: 335417836
2020-10-05 08:05:52 -07:00
7f84a86cf5
[MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
...
PiperOrigin-RevId: 335394668
2020-10-05 05:07:13 -07:00
d927e32451
[HLO] Clean-up dynamic allocation in hlo-legalize-to-lhlo pass.
...
PiperOrigin-RevId: 335385243
2020-10-05 03:55:24 -07:00
7367eac074
Add folder for mhlo::remainder
...
PiperOrigin-RevId: 335372628
2020-10-05 02:20:01 -07:00
3eb767b43d
Add missing lowering step for IsFiniteOp.
...
Also add a BUILD target for generating the GPU kernel.
PiperOrigin-RevId: 334993362
2020-10-02 03:09:00 -07:00
5f303440da
[MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
...
PiperOrigin-RevId: 334843070
2020-10-01 10:25:51 -07:00
458e861254
[MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
...
PiperOrigin-RevId: 334810730
2020-10-01 07:38:40 -07:00
049ca060a1
[MLIR][KernelGen] Legalize `atan2` to approximation
...
Legalize `atan2` analogously to XLA. `atan2` is first reduced to `atan` on the
interval [-1, 1] and subsequently approximated. This CL also adds e2e tests for
trigonometric approximations.
PiperOrigin-RevId: 334794336
2020-10-01 05:34:48 -07:00
4b1809784a
Support collapse_slice_dims in the mhlo.gather->mhlo.slice canonicalizer
...
PiperOrigin-RevId: 334774763
2020-10-01 02:46:49 -07:00
bce128b070
Introduce CustomCall operation in LHLO Dialect
...
- And add conversion from MHLO CustomCall to LHLO CustomCall
- According to XLA documentation, the called function should not be side effecting,
so marking the argument MemRefs as MemRead.
PiperOrigin-RevId: 334737196
2020-09-30 20:56:18 -07:00
dfe64d3958
Implement InferShapedTypeOpInterface for mhlo.complex
...
Binary companion for 8bcd33e4b7
2020-09-30 12:14:15 -07:00
c8919f8419
Implement InferShapedTypeOpInterface and use inferReturnTypes for mhlo.imag and mhlo.real
...
This makes the lhlo lowering work with dynamic shapes.
PiperOrigin-RevId: 334553472
2020-09-30 02:02:26 -07:00
39389587d2
{mhlo.is_finite, lmhlo.is_finite} -> linalg.generic conversion
...
PiperOrigin-RevId: 334414295
2020-09-29 10:50:42 -07:00
6459f12235
Lower mhlo.not to a xor with all ones
...
PiperOrigin-RevId: 334361499
2020-09-29 05:59:40 -07:00
26ac5baae4
Make mhlo.sort return variadic results instead of a tuple
...
Tuple is only used on XLA's sort to return multiple inputs. MLIR supports
multiple inputs, switch to a tuple return.
PiperOrigin-RevId: 334226937
2020-09-28 13:32:23 -07:00
9c6640cbb6
Only apply GeneralDotOpLoweringPatterns for static shaped inputs
...
PiperOrigin-RevId: 333439680
2020-09-23 21:42:01 -07:00
233f1a8a1a
Folders for mhlo.compare
...
Constant evaluation of compare for the case where inputs are either the same
variable or the values are constant.
PiperOrigin-RevId: 333342328
2020-09-23 12:03:48 -07:00
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
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
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
4002077261
[MLIR][KernelGen] Lower `tf.Sinh` to MLHLO
...
PiperOrigin-RevId: 332425724
2020-09-18 04:27:07 -07:00
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
b1fd4d27cf
[MLIR][KernelGen] Implement InferShapedTypeOpInterface for `mhlo.compare/select`
...
PiperOrigin-RevId: 332227340
2020-09-17 07:10:10 -07:00
Tres Popp
Nicolas Vasilache
Thomas Joerg
Benjamin Kramer
Smit Hinsu
Richard Uhler
Hanhan Wang
Sean Silva
Roman Dzhabarov
A. Unique TensorFlower
Jacques Pienaar
Stephan Herhut
Adrian Kuegel
Rahul Joshi
Lucy Fox
Alexander Belyaev
Chao Xie
Ahmed S. Taei
Robert Suderman
Mehdi Amini