Commit Graph

835 Commits

Author SHA1 Message Date
Adrian Kuegel 0683db3b24 Legalize MinimumBroadcastShapes op.
Use it in TransformUnrankedHloPass, which allows to reduce the maximum
rank for rank specialized broadcast from 6 to 5.

PiperOrigin-RevId: 360415743
2021-03-02 06:39:01 -08:00
Jacques Pienaar 329b1fd071 Verify compatible shapes in unpack verification rather than exact
Previously this would be too strict and fail if dynamic and static dims were
compared. Dynamic/unknown are treated as "maybe equal" to a static value without further info, so at this layer don't flag as invalid unless truly are.

PiperOrigin-RevId: 360189086
2021-03-01 08:00:16 -08:00
Christian Sigg 70ee9369d5 Use mlir::OpState::operator->() to get to Operation::getAttrs().
This is a preparation step to remove getAttrs() from OpState.

PiperOrigin-RevId: 360159716
2021-03-01 04:53:00 -08:00
Benjamin Kramer 7c071e8ee6 Integrate LLVM at llvm/llvm-project@99c24f7aa8
Updates LLVM usage to match
[99c24f7aa8cc](https://github.com/llvm/llvm-project/commit/99c24f7aa8cc)

PiperOrigin-RevId: 360150476
2021-03-01 03:44:30 -08:00
Benjamin Kramer e19ccf975e Filter static dimensions from dynamic_broadcast_in_dim's init_tensor
Otherwise we'd generate invalid IR for those cases.

PiperOrigin-RevId: 360144122
2021-03-01 03:03:54 -08:00
Adrian Kuegel e6a1f5f0f9 Add MinimumBroadcastShapesOp to chlo dialect.
This op is useful for rank specialization of broadcasts. Kernel Generator
needs to generate one kernel for each rank, so if we can minimize the rank
of the broadcast shape, we can support more cases with the same number of
special-cased kernels.

PiperOrigin-RevId: 360137827
2021-03-01 02:23:52 -08:00
Christian Sigg 2d818c4fd9 Use mlir::OpState::operator->() to get to methods of mlir::Operation.
This is a preparation step to remove those methods from OpState.

PiperOrigin-RevId: 360043992
2021-02-28 09:02:33 -08:00
A. Unique TensorFlower 006b58c476 Integrate LLVM at llvm/llvm-project@5077d42cfa
Updates LLVM usage to match
[5077d42cfa42](https://github.com/llvm/llvm-project/commit/5077d42cfa42)

PiperOrigin-RevId: 359899172
2021-02-26 22:19:53 -08:00
Hanhan Wang a8f99ee0f5 Fix the shape of linalg.init_tensor in conv op lowering.
The output spatial dims are not as same as the input spatial dims. Only supports
static output spatial dims for now.

PiperOrigin-RevId: 359775479
2021-02-26 09:34:11 -08:00
A. Unique TensorFlower c616963501 Integrate LLVM at llvm/llvm-project@72b18a86e1
Updates LLVM usage to match
[72b18a86e11e](https://github.com/llvm/llvm-project/commit/72b18a86e11e)

PiperOrigin-RevId: 359762921
2021-02-26 08:22:16 -08:00
A. Unique TensorFlower c68b71b5dc Integrate LLVM at llvm/llvm-project@cb81135f94
Updates LLVM usage to match
[cb81135f94e5](https://github.com/llvm/llvm-project/commit/cb81135f94e5)

PiperOrigin-RevId: 359746866
2021-02-26 06:47:30 -08:00
A. Unique TensorFlower c63ac91206 Integrate LLVM at llvm/llvm-project@bf6380c096
Updates LLVM usage to match
[bf6380c0966b](https://github.com/llvm/llvm-project/commit/bf6380c0966b)

PiperOrigin-RevId: 359728311
2021-02-26 04:05:03 -08:00
A. Unique TensorFlower fe4234bcff Integrate LLVM at llvm/llvm-project@5d64dd8e3c
Updates LLVM usage to match
[5d64dd8e3c22](https://github.com/llvm/llvm-project/commit/5d64dd8e3c22)

PiperOrigin-RevId: 359653225
2021-02-25 17:37:37 -08:00
A. Unique TensorFlower e7a3ec18ad Integrate LLVM at llvm/llvm-project@7f6e331645
Updates LLVM usage to match
[7f6e3316456f](https://github.com/llvm/llvm-project/commit/7f6e3316456f)

PiperOrigin-RevId: 359595260
2021-02-25 13:12:35 -08:00
A. Unique TensorFlower bb22553219 Integrate LLVM at llvm/llvm-project@f4d78a5e3a
Updates LLVM usage to match
[f4d78a5e3aee](https://github.com/llvm/llvm-project/commit/f4d78a5e3aee)

PiperOrigin-RevId: 359543192
2021-02-25 09:40:18 -08:00
Hanhan Wang 90f0d7f935 Add support for lowering mhlo.conv to Linalg on tensors.
This pattern only works for normal convolutions. It does not work for depthwise
convolutions. The Linalg conv ops are defined with static rank, so it only
supports 1d/2d/3d cases, which are the most typical cases.

This also refactors out the same check in lmhlo.conv lowering.

PiperOrigin-RevId: 359503527
2021-02-25 05:59:08 -08:00
Rahul Joshi c5f5d13930 [MLIR] Add dependency from "lhlo_ops_structs_inc_gen" to "lhlo" target.
PiperOrigin-RevId: 359488020
2021-02-25 03:51:36 -08:00
Hanhan Wang 45a1249fe2 Add support for lowering mhlo.pad to linalg.pad_tensor
The change upstreams the pattern from IREE repo to MHLO repo.

PiperOrigin-RevId: 359481543
2021-02-25 03:00:39 -08:00
A. Unique TensorFlower 459362b206 Integrate LLVM at llvm/llvm-project@d8a8e5d624
Updates LLVM usage to match
[d8a8e5d6240a](https://github.com/llvm/llvm-project/commit/d8a8e5d6240a)

PiperOrigin-RevId: 359475769
2021-02-25 02:20:12 -08:00
A. Unique TensorFlower 7dc1543bb2 Integrate LLVM at llvm/llvm-project@96a3dfeb93
Updates LLVM usage to match
[96a3dfeb9303](https://github.com/llvm/llvm-project/commit/96a3dfeb9303)

PiperOrigin-RevId: 359394503
2021-02-24 16:17:23 -08:00
A. Unique TensorFlower d9d6ab105e Integrate LLVM at llvm/llvm-project@5c74c6be3c
Updates LLVM usage to match
[5c74c6be3c29](https://github.com/llvm/llvm-project/commit/5c74c6be3c29)

PiperOrigin-RevId: 359354498
2021-02-24 13:19:53 -08:00
Geoffrey Martin-Noble 89f7f2bd65 Lower integer matmuls to linalg
PiperOrigin-RevId: 359306495
2021-02-24 09:45:07 -08:00
Hanhan Wang 475b4a06a5 Add support for lowering mhlo.slice to subtensor.
PiperOrigin-RevId: 359297978
2021-02-24 09:06:09 -08:00
A. Unique TensorFlower b478bdf00e Integrate LLVM at llvm/llvm-project@de40423c85
Updates LLVM usage to match
[de40423c8512](https://github.com/llvm/llvm-project/commit/de40423c8512)

PiperOrigin-RevId: 359280337
2021-02-24 07:32:25 -08:00
A. Unique TensorFlower 3e2ff6e253 Integrate LLVM at llvm/llvm-project@861dbe1a02
Updates LLVM usage to match
[861dbe1a021e](https://github.com/llvm/llvm-project/commit/861dbe1a021e)

PiperOrigin-RevId: 359091800
2021-02-23 11:28:36 -08:00
A. Unique TensorFlower ac0552f127 [MLIR][HLO] Remove duplicate `PopulateTransformUnrankedHloPatterns`
PiperOrigin-RevId: 359046173
2021-02-23 07:50:47 -08:00
A. Unique TensorFlower 2df1bb3d6b Integrate LLVM at llvm/llvm-project@6c9541d4dd
Updates LLVM usage to match
[6c9541d4ddfd](https://github.com/llvm/llvm-project/commit/6c9541d4ddfd)

PiperOrigin-RevId: 359020255
2021-02-23 04:38:26 -08:00
A. Unique TensorFlower 855cd484fc Integrate LLVM at llvm/llvm-project@557d2ade01
Updates LLVM usage to match
[557d2ade016f](https://github.com/llvm/llvm-project/commit/557d2ade016f)

PiperOrigin-RevId: 358925409
2021-02-22 16:05:25 -08:00
A. Unique TensorFlower b45eaf8d33 Integrate LLVM at llvm/llvm-project@12edddafac
Updates LLVM usage to match
[12edddafac45](https://github.com/llvm/llvm-project/commit/12edddafac45)

PiperOrigin-RevId: 358889182
2021-02-22 13:21:11 -08:00
Rahul Joshi 0da7ea2545 [MLIR][HLO] Cleanup CMakeLists and comments.
- Cleanup CMakeLists file to remove unused argument and use a new function for
  setting up lmhlo and lmhlo_gpu dialect targets.
- Fix inconsistently formatted copyright comment and fix header include guards.

PiperOrigin-RevId: 358865838
2021-02-22 11:36:29 -08:00
Rahul Joshi 5adb7c6e12 [MLIR:LHLO] Add optional call target arg mapping to LMHLO CustomCall operations.
- XLA:HLO -> LMHLO conversion drops all token arguments and return values, however
  custom calls that users write still expect to get buffer pointers for these token types.
- To be able to support this, add an optional call target argument mapping attribute to
  LMHLO custom calls. When this attribute is present, it indicates the number of
  arguments and returns that the custom call expects and also indicates which LMHLO
  arg() or output() maps to which arg or result number of the custom call.

PiperOrigin-RevId: 358826664
2021-02-22 08:43:00 -08:00
Benjamin Kramer a9cc1dcfa0 [mlir][hlo] Add basic rank-specialization for select
This just blows up everything to ranked (up to 6) and is probably quite slow.
This is sufficient to make kernelgen compile SelectV2.

PiperOrigin-RevId: 358777728
2021-02-22 02:41:12 -08:00
Prakalp Srivastava 909574e393 Pass result element type to XlaBuilder for `mhlo.dot_general` and `mhlo.convolution` ops.
`mhlo.dot_general` and `mhlo.convolution` result element type might be different from operand element type. See `preferred_element_type` attribute that allows i8xi8 to i32 dot computation. `mhlo` to HLO exporter should pass the result element type to Xla builder to override the shape inference of XLA.

PiperOrigin-RevId: 358580718
2021-02-20 07:07:17 -08:00
A. Unique TensorFlower f63c93399a Integrate LLVM at llvm/llvm-project@7e54d7304b
Updates LLVM usage to match
[7e54d7304be2](https://github.com/llvm/llvm-project/commit/7e54d7304be2)

PiperOrigin-RevId: 358516349
2021-02-19 17:49:00 -08:00
A. Unique TensorFlower 5cc9fa9d9d Integrate LLVM at llvm/llvm-project@cd8cc00b9e
Updates LLVM usage to match
[cd8cc00b9e2b](https://github.com/llvm/llvm-project/commit/cd8cc00b9e2b)

PiperOrigin-RevId: 358329421
2021-02-18 21:09:36 -08:00
A. Unique TensorFlower 5400287d69 Integrate LLVM at llvm/llvm-project@b7e05c874b
Updates LLVM usage to match
[b7e05c874b5b](https://github.com/llvm/llvm-project/commit/b7e05c874b5b)

PiperOrigin-RevId: 358289004
2021-02-18 16:19:04 -08:00
A. Unique TensorFlower 27a8b18526 Integrate LLVM at llvm/llvm-project@f70cdc5b5c
Updates LLVM usage to match
[f70cdc5b5c7c](https://github.com/llvm/llvm-project/commit/f70cdc5b5c7c)

PiperOrigin-RevId: 358240325
2021-02-18 12:40:08 -08:00
A. Unique TensorFlower 72d3eee118 Integrate LLVM at llvm/llvm-project@5b094bfeb3
Updates LLVM usage to match
[5b094bfeb3cc](https://github.com/llvm/llvm-project/commit/5b094bfeb3cc)

PiperOrigin-RevId: 358214975
2021-02-18 10:52:25 -08:00
Benjamin Kramer b42def4612 [mlir][hlo] Refactor rank specialization to allow an arbitrary number of inputs
This actually simplifies the code a bit.

PiperOrigin-RevId: 358201038
2021-02-18 09:53:03 -08:00
Benjamin Kramer ca4034b56e [mlir][hlo] Make select ready for dynamic shapes (ranked only for now)
Move tf.SelectV2 broadcast lowering to a chlo.broadcast_select op, and lower it
to broadcasts on mhlo from there.

PiperOrigin-RevId: 358179975
2021-02-18 08:08:40 -08:00
A. Unique TensorFlower a6f03aecfb Integrate LLVM at llvm/llvm-project@624fccba87
Updates LLVM usage to match
[624fccba87cc](https://github.com/llvm/llvm-project/commit/624fccba87cc)

PiperOrigin-RevId: 358167564
2021-02-18 06:50:29 -08:00
Adrian Kuegel 37e31f8b26 Lower Expm1 kernel to math.ExpM1.
PiperOrigin-RevId: 358152908
2021-02-18 04:54:23 -08:00
A. Unique TensorFlower 90b222cf6e Integrate LLVM at llvm/llvm-project@55756f32f7
Updates LLVM usage to match
[55756f32f735](https://github.com/llvm/llvm-project/commit/55756f32f735)

PiperOrigin-RevId: 358093073
2021-02-17 20:33:14 -08:00
Richard Uhler b579bd5d9e Support dynamic-shaped operand in verification of BroadcastInDim.
Verification of HLO_BroadcastInDimOp was previously failing or crashing if the
operand had a dynamic shape or was unranked. Update the verification code to
allow the operand to be unranked or have dynamic shape.

PiperOrigin-RevId: 358056793
2021-02-17 16:18:09 -08:00
A. Unique TensorFlower dd237d4267 Integrate LLVM at llvm/llvm-project@892d2822b6
Updates LLVM usage to match
[892d2822b62e](https://github.com/llvm/llvm-project/commit/892d2822b62e)

PiperOrigin-RevId: 358038509
2021-02-17 14:52:43 -08:00
A. Unique TensorFlower 220deb3709 [MLIR][CHLO] Add legalization for `chlo.polygamma` to MHLO
PiperOrigin-RevId: 357954624
2021-02-17 08:33:01 -08:00
A. Unique TensorFlower c06de24f6c [MLIR][CHLO] Generalize lowering with upcast to n-ary operation
Allows reuse for zeta lowering now and for the polygamma lowering soon.

PiperOrigin-RevId: 357739910
2021-02-16 09:47:24 -08:00
A. Unique TensorFlower 81abaf364d [MLIR][MHLO] Add polygamma op to the CHLO dialect
PiperOrigin-RevId: 357724465
2021-02-16 08:32:33 -08:00
A. Unique TensorFlower 02dab94054 Integrate LLVM at llvm/llvm-project@9f581815ae
Updates LLVM usage to match
[9f581815ae4d](https://github.com/llvm/llvm-project/commit/9f581815ae4d)

PiperOrigin-RevId: 357683119
2021-02-16 03:33:12 -08:00
Benjamin Kramer 9447374676 Integrate LLVM at llvm/llvm-project@715dc556b7
Updates LLVM usage to match
[715dc556b782](https://github.com/llvm/llvm-project/commit/715dc556b782)

PiperOrigin-RevId: 357607282
2021-02-15 13:35:33 -08:00