Commit Graph

98 Commits

Author SHA1 Message Date
A. Unique TensorFlower 33c450e4cb Fix the MHLO to LMHLO lowering of 'gather'
The lowering assumes that the 'gather' op attributes are identical in both MHLO and LMHLO. But that's not true; some time ago the MHLO version was changed to pack 4 of its attributes into a struct. By doing the same for the LMHLO version we both fix the lowering for this op and resolve a longstanding TODO.

PiperOrigin-RevId: 337943946
2020-10-19 15:14:05 -07:00
A. Unique TensorFlower 7a983ea389 Add folder for mhlo::pad
PiperOrigin-RevId: 337827560
2020-10-19 04:21:44 -07:00
A. Unique TensorFlower 4a18aa41ee Add folder to mhlo::round_nearest_afz
PiperOrigin-RevId: 337823786
2020-10-19 03:45:15 -07:00
A. Unique TensorFlower 706718b4fb Permit vector types in lmhlo to std lowering.
PiperOrigin-RevId: 337523303
2020-10-16 09:47:02 -07:00
A. Unique TensorFlower 51cd4200b6 Make LMHLO's Dot have the same power as MHLO's DotGeneral.
PiperOrigin-RevId: 337391565
2020-10-15 15:09:06 -07:00
A. Unique TensorFlower 05ee41baf8 Add folder for mhlo::scatter
PiperOrigin-RevId: 337274351
2020-10-15 03:26:05 -07:00
Adrian Kuegel 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
Rahul Joshi 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
Adrian Kuegel 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
Jacques Pienaar ab4c03b84b Replace operands now injected in OpBuilder
These are now injected automatically.

PiperOrigin-RevId: 336356284
2020-10-09 13:39:33 -07:00
Rahul Joshi 41436ea0d9 [NFC] Extract common attributes of MHLO and LMHLO ConvOp
- Create a common class that hold ConvOp attributes that are common to MHLO and
  LHLO Dialects and use Tablegen DAG concat to append these common attributes to
  dialect specific arguments in ODS for ConvOp.

PiperOrigin-RevId: 336114003
2020-10-08 10:29:32 -07:00
Pavel Krajcevski ac483f4091 Remove superfluous quote from lhlo_ops.td
PiperOrigin-RevId: 335921873
2020-10-07 12:06:36 -07:00
A. Unique TensorFlower 3736c5542f [MLIR][KernelGen] Fix unranked codegeneration in kernel generator
PiperOrigin-RevId: 335847086
2020-10-07 05:39:55 -07:00
A. Unique TensorFlower bae0815ef0 [MLIR][KernelGen] Legalize `atan` to approximation
PiperOrigin-RevId: 335417836
2020-10-05 08:05:52 -07:00
A. Unique TensorFlower 7f84a86cf5 [MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
PiperOrigin-RevId: 335394668
2020-10-05 05:07:13 -07:00
A. Unique TensorFlower 7367eac074 Add folder for mhlo::remainder
PiperOrigin-RevId: 335372628
2020-10-05 02:20:01 -07:00
Tim Shen c708bfd6d0 [MLIR] Add cbrt, reduce-precision, and bitcast ops to MHLO.
PiperOrigin-RevId: 335109804
2020-10-02 15:13:18 -07:00
Chao Xie 5f303440da [MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
PiperOrigin-RevId: 334843070
2020-10-01 10:25:51 -07:00
A. Unique TensorFlower 458e861254 [MLIR][KerneGen] Lower `tf.Atan` all the way to LLVM
PiperOrigin-RevId: 334810730
2020-10-01 07:38:40 -07:00
Rahul Joshi 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
Benjamin Kramer dfe64d3958 Implement InferShapedTypeOpInterface for mhlo.complex
Binary companion for 8bcd33e4b7

PiperOrigin-RevId: 334651523
2020-09-30 12:14:15 -07:00
Benjamin Kramer 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
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 336ee14538 Make lmhlo.sign to linalg lowering work for more floating point and integer types
PiperOrigin-RevId: 334400341
2020-09-29 09:52:13 -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
Robert Suderman 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
Robert Suderman 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
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
Rahul Joshi 3ca740103d [NFC] Specify "reifyReturnTypeShapes" in DeclareOpInterfaceMethod overridenMethods
- This causes declaration of reifyReturnTypeShapes to be inserted in the TableGen
  generated code without explicitly adding it in extraClassDeclaration.

PiperOrigin-RevId: 333315601
2020-09-23 10:08:32 -07:00
Rahul Joshi 87d9e6951e [MLIR] Use MLIR provided 0DTensorOf/1DTensorOf instead of custom constraints.
- Use MLIR provided constraints for HLO_ScalarIntTensor and HLO_DimensionTensor.
- Update unit tests to expect new error messages.

PiperOrigin-RevId: 333313131
2020-09-23 09:57:12 -07:00
Rahul Joshi 7d01a60de8 [NFC] Fix typos in comments.
PiperOrigin-RevId: 333311070
2020-09-23 09:46:02 -07:00
A. Unique TensorFlower 08514eaa5f [MLIR][KernelGen] Add E2E tests for `tf.Tan` and `tf.Sinh`
PiperOrigin-RevId: 333275089
2020-09-23 05:50:54 -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
A. Unique TensorFlower b1fd4d27cf [MLIR][KernelGen] Implement InferShapedTypeOpInterface for `mhlo.compare/select`
PiperOrigin-RevId: 332227340
2020-09-17 07:10:10 -07:00
A. Unique TensorFlower 69b80d8deb [MLIR] Extend unranked transformation to CHLO dialect
PiperOrigin-RevId: 332026604
2020-09-16 09:49:18 -07:00
A. Unique TensorFlower 1880f87737 Integrate LLVM at llvm/llvm-project@e1669843f2
Updates LLVM usage to match
[e1669843f2aa](https://github.com/llvm/llvm-project/commit/e1669843f2aa)

PiperOrigin-RevId: 331987679
2020-09-16 05:58:08 -07:00
Hanhan Wang 1800f44a29 Add sqrt folder.
PiperOrigin-RevId: 331974344
2020-09-16 04:04:09 -07:00
Robert Suderman dff98f8cd5 Folder for mhlo.negate
PiperOrigin-RevId: 331592342
2020-09-14 11:38:44 -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
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