Commit Graph

1004 Commits

Author SHA1 Message Date
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
A. Unique TensorFlower af5343ff86 Integrate LLVM at llvm/llvm-project@2dbe88db58
Updates LLVM usage to match
[2dbe88db5804](https://github.com/llvm/llvm-project/commit/2dbe88db5804)

PiperOrigin-RevId: 357587981
2021-02-15 09:58:25 -08:00
Adrian Kuegel b594254c79 [mhlo] Lower int->bool to a comparison with zero
This matches what TF (and C++) do in this case.

PiperOrigin-RevId: 357566262
2021-02-15 06:38:09 -08:00
Benjamin Kramer 240a44de82 [mhlo] Lower int->int cast to sign extension instead of zero extension
Signless does not mean unsigned here. Currently mhlo only has signed types.

PiperOrigin-RevId: 357561712
2021-02-15 05:58:47 -08:00
Adrian Kuegel 8672735e9a [mhlo] Lower float->bool to a comparison with zero
This matches what TF (and C++) do in this case.

PiperOrigin-RevId: 357553098
2021-02-15 04:36:36 -08:00
A. Unique TensorFlower 2fe0c33083 Integrate LLVM at llvm/llvm-project@16428a8d91
Updates LLVM usage to match
[16428a8d91a9](https://github.com/llvm/llvm-project/commit/16428a8d91a9)

PiperOrigin-RevId: 357550807
2021-02-15 04:17:58 -08:00
A. Unique TensorFlower 89d81adf6d [mhlo] Lower float->bool to a comparison with zero
This matches what TF (and C++) do in this case.

PiperOrigin-RevId: 357541594
2021-02-15 03:11:56 -08:00
Benjamin Kramer 3e80d91e73 [mhlo] Lower float->bool to a comparison with zero
This matches what TF (and C++) do in this case.

PiperOrigin-RevId: 357534118
2021-02-15 02:17:19 -08:00
Adrian Kuegel 824bc9c425 Improve broadcast transformation to treat dynamic shapes with 1 element as scalar.
A shape that contains exactly one element is effectively a scalar. This leads
to a speedup in cases where we have a binary op with one operand that is
effectively a scalar, because we can use the fast path.

PiperOrigin-RevId: 357515552
2021-02-14 23:25:41 -08:00
A. Unique TensorFlower 4060a86fe2 Integrate LLVM at llvm/llvm-project@2bfe27da17
Updates LLVM usage to match
[2bfe27da171e](https://github.com/llvm/llvm-project/commit/2bfe27da171e)

PiperOrigin-RevId: 357196336
2021-02-12 08:32:03 -08:00
A. Unique TensorFlower e993082b97 Integrate LLVM at llvm/llvm-project@6f04addc8b
Updates LLVM usage to match
[6f04addc8b2e](https://github.com/llvm/llvm-project/commit/6f04addc8b2e)

PiperOrigin-RevId: 357170209
2021-02-12 05:21:12 -08:00
A. Unique TensorFlower f7ac575dd9 Integrate LLVM at llvm/llvm-project@0c118831a3
Updates LLVM usage to match
[0c118831a37a](https://github.com/llvm/llvm-project/commit/0c118831a37a)

PiperOrigin-RevId: 357158589
2021-02-12 03:32:55 -08:00
A. Unique TensorFlower a7365aadbe Integrate LLVM at llvm/llvm-project@8151c1b442
Updates LLVM usage to match
[8151c1b44211](https://github.com/llvm/llvm-project/commit/8151c1b44211)

PiperOrigin-RevId: 357080940
2021-02-11 16:17:46 -08:00
A. Unique TensorFlower 077df4f8cc Integrate LLVM at llvm/llvm-project@418c218efa
Updates LLVM usage to match
[418c218efa95](https://github.com/llvm/llvm-project/commit/418c218efa95)

PiperOrigin-RevId: 357010454
2021-02-11 10:49:27 -08:00
Tim Shen 6fa6974e8d [XLA/GPU] Plumb through Bitcast op for LMHLO.
Also remove BitcastOp. XLA bitcast requires the input buffer to alias the output buffer, which makes bitcast always a no-op.

PiperOrigin-RevId: 356884383
2021-02-10 19:45:40 -08:00
A. Unique TensorFlower 4e202135d1 Integrate LLVM at llvm/llvm-project@9db6e97a86
Updates LLVM usage to match
[9db6e97a8605](https://github.com/llvm/llvm-project/commit/9db6e97a8605)

PiperOrigin-RevId: 356753027
2021-02-10 09:32:14 -08:00
Alexander Belyaev 36e04d92c0 [KERNEL_GEN] Add a pattern to bufferize `mhlo.reshape(<unranked_tensor>)`.
PiperOrigin-RevId: 356720899
2021-02-10 06:32:21 -08:00
A. Unique TensorFlower 54c2a49866 Integrate LLVM at llvm/llvm-project@a5222aa085
Updates LLVM usage to match
[a5222aa0858a](https://github.com/llvm/llvm-project/commit/a5222aa0858a)

PiperOrigin-RevId: 356515740
2021-02-09 09:13:45 -08:00
A. Unique TensorFlower 7446b6022d Integrate LLVM at llvm/llvm-project@2fa4186d4e
Updates LLVM usage to match
[2fa4186d4e1c](https://github.com/llvm/llvm-project/commit/2fa4186d4e1c)

PiperOrigin-RevId: 356383544
2021-02-08 16:53:07 -08:00
A. Unique TensorFlower 4a29ca3b1d Add layout to mhlo::InfeedOp td.
PiperOrigin-RevId: 356286875
2021-02-08 09:48:14 -08:00
A. Unique TensorFlower 80d753c1fe Integrate LLVM at llvm/llvm-project@f89f6d1e5d
Updates LLVM usage to match
[f89f6d1e5d7d](https://github.com/llvm/llvm-project/commit/f89f6d1e5d7d)

PiperOrigin-RevId: 356265374
2021-02-08 09:47:00 -08:00
Tres Popp d086b8a0ec Correct HLO atan2 lowering in cases of -inf and -0 inputs.
This is being done by just removing the approximation and lowering to atan2 lib calls later to make the implementation the same as XLA. Note that if the approximation is brought back later, it can be fixed by changing the IR checking `less-than(X, 0)` to `less-than(copysign(X, 1), 0)`

PiperOrigin-RevId: 356253941
2021-02-08 06:58:04 -08:00
A. Unique TensorFlower bd0856578f Integrate LLVM at llvm/llvm-project@a4fa667dee
Updates LLVM usage to match
[a4fa667dee60](https://github.com/llvm/llvm-project/commit/a4fa667dee60)

PiperOrigin-RevId: 355945749
2021-02-05 15:53:55 -08:00
A. Unique TensorFlower 2aa8a90c69 Integrate LLVM at llvm/llvm-project@a1a1d338e9
Updates LLVM usage to match
[a1a1d338e99d](https://github.com/llvm/llvm-project/commit/a1a1d338e99d)

PiperOrigin-RevId: 355927079
2021-02-05 14:20:29 -08:00
Rahul Joshi b251712b1d [XLA:GPU] Add conversion from HLO -> MLIR LMHLO for TriangularSolve
- Also add layout attributes for inputs and output for error checking.

PiperOrigin-RevId: 355863625
2021-02-05 09:18:02 -08:00
Adrian Kuegel 1c4521cc42 Integrate LLVM at llvm/llvm-project@d1978fa4bf
Updates LLVM usage to match
[d1978fa4bf0d](https://github.com/llvm/llvm-project/commit/d1978fa4bf0d)

PiperOrigin-RevId: 355848094
2021-02-05 07:42:06 -08:00
A. Unique TensorFlower 99bc05f2e4 Integrate LLVM at llvm/llvm-project@91e7a17133
Updates LLVM usage to match
[91e7a1713332](https://github.com/llvm/llvm-project/commit/91e7a1713332)

PiperOrigin-RevId: 355702100
2021-02-04 13:42:31 -08:00
Stephan Herhut 60e1b6882c Add kernel definition for zeta operation.
PiperOrigin-RevId: 355575619
2021-02-04 01:27:43 -08:00
A. Unique TensorFlower 945bf8768d Integrate LLVM at llvm/llvm-project@56fcd4ea8d
Updates LLVM usage to match
[56fcd4ea8daf](https://github.com/llvm/llvm-project/commit/56fcd4ea8daf)

PiperOrigin-RevId: 355529409
2021-02-03 18:48:22 -08:00
A. Unique TensorFlower 3380d0a761 Integrate LLVM at llvm/llvm-project@ea5b75de49
Updates LLVM usage to match
[ea5b75de4999](https://github.com/llvm/llvm-project/commit/ea5b75de4999)

PiperOrigin-RevId: 355493884
2021-02-03 15:20:56 -08:00
Mahesh Ravishankar 44d0464d16 Use linalg.fill on tensors instead of tensor.generate in MHLO -> Linalg conversion.
linalg.fill on tensors is a structured op that allows use tile + fuse
to reduce the fill overhead.

PiperOrigin-RevId: 355490400
2021-02-03 15:03:49 -08:00
A. Unique TensorFlower 10cd797d6d Integrate LLVM at llvm/llvm-project@bbb7555403
Updates LLVM usage to match
[bbb755540385](https://github.com/llvm/llvm-project/commit/bbb755540385)

PiperOrigin-RevId: 355462338
2021-02-03 12:58:16 -08:00