Commit Graph

766 Commits

Author SHA1 Message Date
A. Unique TensorFlower 446e25ed7d Integrate LLVM at llvm/llvm-project@b40b3196b3
Updates LLVM usage to match
[b40b3196b321](https://github.com/llvm/llvm-project/commit/b40b3196b321)

PiperOrigin-RevId: 345287702
2020-12-02 12:24:43 -08:00
Rahul Joshi dbbdfea95b [MLIR:HLO] Generate enum decls for HLO and LHLO GPU dialects.
- Split out enum definitions in hlo dialect into a separate .td file (similar to structs)
  and generate enum decl/defs for these enums.
- Also split out the LHLO GPU enums into a separate .td file and generate enum
  decl/defs for these enums as well.
- Remove unused dialect from ConvolutionAttributes and generate lhlo_gpu enums.
- Add appropriate namespace for all the enums.

PiperOrigin-RevId: 345277240
2020-12-02 11:39:23 -08:00
Rahul Joshi d7bd5233ab [XLA:GPU] Migrate GEMM Thunk emission to MLIR.
- Map Custom call for GEMM in XLA HLO to Gemm/Gemm bias operations in LHLO GPU
  dialect.
- Make 'algorithm' an optional attribute to better match with XLA HLO backend config.
- Replace 'alpha' with 'alpha_real' and 'alpha_complex' to support complex GEMM correctly.
- Generate GemmThunk off of LHLO GPU Gemm operations.

PiperOrigin-RevId: 345250840
2020-12-02 09:43:12 -08:00
A. Unique TensorFlower 1b711670bc Fix handling of negative seeds in random number generator op kernels for XLA
Casting negative s32 number to u64 directly will have leading 1s in the representation which is not what we want to get a single u64 out of two s32 seeds. Fixed this by first getting unsigned number of the same bit-width.

PiperOrigin-RevId: 345239817
2020-12-02 08:42:07 -08:00
Smit Hinsu 733fc6d032 Fix handling of negative seeds in random number generator op kernels for XLA
Casting negative s32 number to u64 directly will have leading 1s in the representation which is not what we want to get a single u64 out of two s32 seeds. Fixed this by first getting unsigned number of the same bit-width.

PiperOrigin-RevId: 345227848
2020-12-02 07:24:10 -08:00
A. Unique TensorFlower 2f6ced90fc Integrate LLVM at llvm/llvm-project@c266c56d54
Updates LLVM usage to match
[c266c56d545d](https://github.com/llvm/llvm-project/commit/c266c56d545d)

PiperOrigin-RevId: 345222350
2020-12-02 06:41:39 -08:00
A. Unique TensorFlower 27b10932cb Integrate LLVM at llvm/llvm-project@aafb366210
Updates LLVM usage to match
[aafb3662103f](https://github.com/llvm/llvm-project/commit/aafb3662103f)

PiperOrigin-RevId: 345121003
2020-12-01 16:10:48 -08:00
River Riddle f89244381d [mlir][NFC] Replace usages of Function.h and Module.h with BuiltinOps.h
This is part of a larger refactoring cleaning up the BuiltinDialect of MLIR.

PiperOrigin-RevId: 345085278
2020-12-01 13:18:06 -08:00
A. Unique TensorFlower eb9015a2b0 Integrate LLVM at llvm/llvm-project@bb993b1d9d
Updates LLVM usage to match
[bb993b1d9de3](https://github.com/llvm/llvm-project/commit/bb993b1d9de3)

PiperOrigin-RevId: 345083620
2020-12-01 13:09:14 -08:00
A. Unique TensorFlower 7138157ed9 Integrate LLVM at llvm/llvm-project@523775f967
Updates LLVM usage to match
[523775f96742](https://github.com/llvm/llvm-project/commit/523775f96742)

PiperOrigin-RevId: 345009972
2020-12-01 06:47:02 -08:00
A. Unique TensorFlower fd286b11f1 Integrate LLVM at llvm/llvm-project@40659cd2c6
Updates LLVM usage to match
[40659cd2c6f4](https://github.com/llvm/llvm-project/commit/40659cd2c6f4)

PiperOrigin-RevId: 344995763
2020-12-01 04:41:59 -08:00
A. Unique TensorFlower 00ab6722ac Integrate LLVM at llvm/llvm-project@d928dfc6f9
Updates LLVM usage to match
[d928dfc6f924](https://github.com/llvm/llvm-project/commit/d928dfc6f924)

PiperOrigin-RevId: 344948827
2020-11-30 21:45:24 -08:00
A. Unique TensorFlower 37b1110ae0 Integrate LLVM at llvm/llvm-project@8cdf4920c4
Updates LLVM usage to match
[8cdf4920c47d](https://github.com/llvm/llvm-project/commit/8cdf4920c47d)

PiperOrigin-RevId: 344933001
2020-11-30 19:06:43 -08:00
A. Unique TensorFlower 3f815138b6 Integrate LLVM at llvm/llvm-project@750049d78b
Updates LLVM usage to match
[750049d78b74](https://github.com/llvm/llvm-project/commit/750049d78b74)

PiperOrigin-RevId: 344889997
2020-11-30 14:32:11 -08:00
Rahul Joshi 5350b8080d [NFC] Add default value for attributes in lmhlo::FusionOp custom builder.
PiperOrigin-RevId: 344856882
2020-11-30 11:47:39 -08:00
A. Unique TensorFlower a4a3e9b73c Integrate LLVM at llvm/llvm-project@234a5297aa
Updates LLVM usage to match
[234a5297aa00](https://github.com/llvm/llvm-project/commit/234a5297aa00)

PiperOrigin-RevId: 344824858
2020-11-30 09:17:33 -08:00
A. Unique TensorFlower dd15c6cd84 [MLIR][KernelGen] Generate assertion message in `transform_unranked_hlo` pass
Use constant to generate the correct assertion message. This avoids
confusion when lowering the max rank specialization for debugging.

PiperOrigin-RevId: 344769021
2020-11-30 01:41:09 -08:00
Benjamin Kramer 1fe04c0aff Integrate LLVM at llvm/llvm-project@564628014c
Updates LLVM usage to match
[564628014c40](https://github.com/llvm/llvm-project/commit/564628014c40)

PiperOrigin-RevId: 344542990
2020-11-27 10:05:28 -08:00
Adrian Kuegel d14c63da54 Add a canonicalization pattern to remove redundant dynamic_reshapes.
PiperOrigin-RevId: 344517381
2020-11-27 04:46:50 -08:00
Stephan Herhut f183e95e69 Register the new branch op rewrite patterns in hlo to lhlo conversion
for good measure.

PiperOrigin-RevId: 344512196
2020-11-27 03:44:28 -08:00
Benjamin Kramer d617f3edb6 Integrate LLVM at llvm/llvm-project@b33fbbaa34
Updates LLVM usage to match
[b33fbbaa34f0](https://github.com/llvm/llvm-project/commit/b33fbbaa34f0)

PiperOrigin-RevId: 344510775
2020-11-27 03:28:09 -08:00
Benjamin Kramer df244927d3 Integrate LLVM at llvm/llvm-project@20c926e079
Updates LLVM usage to match
[20c926e0797e](https://github.com/llvm/llvm-project/commit/20c926e0797e)

PiperOrigin-RevId: 344437441
2020-11-26 10:12:18 -08:00
Adrian Kuegel 6a71a84302 Support different input/output type for TransformUnrankedHlo.
Also generate the tf.Equal kernel, now that it works.

PiperOrigin-RevId: 344402014
2020-11-26 04:20:34 -08:00
Benjamin Kramer 1b98bf5fab Integrate LLVM at llvm/llvm-project@54ec9bb551
Updates LLVM usage to match
[54ec9bb5510d](https://github.com/llvm/llvm-project/commit/54ec9bb5510d)

PiperOrigin-RevId: 344400393
2020-11-26 04:06:38 -08:00
A. Unique TensorFlower 651afab874 Integrate LLVM at llvm/llvm-project@dc35368ccf
Updates LLVM usage to match
[dc35368ccf17](https://github.com/llvm/llvm-project/commit/dc35368ccf17)

PiperOrigin-RevId: 344332738
2020-11-25 16:16:33 -08:00
Alexander Belyaev 5583c63cab [KERNEL_GEN] Add unranked Conj kernel.
PiperOrigin-RevId: 344243271
2020-11-25 06:37:26 -08:00
Benjamin Kramer a6948f6b41 Integrate LLVM at llvm/llvm-project@a38d13ed36
Updates LLVM usage to match
[a38d13ed3635](https://github.com/llvm/llvm-project/commit/a38d13ed3635)

PiperOrigin-RevId: 344235713
2020-11-25 05:33:34 -08:00
A. Unique TensorFlower 42345f41f7 Integrate LLVM at llvm/llvm-project@e0f4dea0d0
Updates LLVM usage to match
[e0f4dea0d0f1](https://github.com/llvm/llvm-project/commit/e0f4dea0d0f1)

PiperOrigin-RevId: 344202478
2020-11-25 00:27:24 -08:00
A. Unique TensorFlower e970a97efb Integrate LLVM at llvm/llvm-project@42eaf4fe0a
Updates LLVM usage to match
[42eaf4fe0ade](https://github.com/llvm/llvm-project/commit/42eaf4fe0ade)

PiperOrigin-RevId: 344158599
2020-11-24 17:12:52 -08:00
Tim Shen 3e01448481 [MLIR] Move some walk()ing functions to the lmhlo::FusionOp API.
PiperOrigin-RevId: 344109366
2020-11-24 12:23:57 -08:00
Lucy Fox 85f92a1651 [KernelGen] Lower tf.Erf and tf.Erfc ops to CHLO.
This does not include the lowerings from CHLO to LMHLO.

PiperOrigin-RevId: 344091604
2020-11-24 10:55:43 -08:00
Benjamin Kramer 9f20ef0581 Integrate LLVM at llvm/llvm-project@5ce85e6635
Updates LLVM usage to match
[5ce85e66358a](https://github.com/llvm/llvm-project/commit/5ce85e66358a)

PiperOrigin-RevId: 344079124
2020-11-24 09:53:42 -08:00
Benjamin Kramer 80d7342ae3 Integrate LLVM at llvm/llvm-project@fd92c5dbee
Updates LLVM usage to match
[fd92c5dbeeeb](https://github.com/llvm/llvm-project/commit/fd92c5dbeeeb)

PiperOrigin-RevId: 344036276
2020-11-24 05:14:33 -08:00
Smit Hinsu b016b5a219 Fix constant folding of mhlo.convert op with i1 element types
Boolean element values should be fetched as an unsigned integer and not signed integer which would return -1 for true.

Added to a TODO to handle unsigned types correctly as well as we don't seem to be using unsigned types.

PiperOrigin-RevId: 343927564
2020-11-23 14:18:28 -08:00
Benjamin Kramer 2ac41d8cd2 Integrate LLVM at llvm/llvm-project@76bd4444e3
Updates LLVM usage to match
[76bd4444e361](https://github.com/llvm/llvm-project/commit/76bd4444e361)

PiperOrigin-RevId: 343875803
2020-11-23 10:18:01 -08:00
Rahul Joshi ac54c5ccfa [XLA:GPU] Convert Cholesky custom call in XLA HLO to LHLO GPU Dialect.
- Restructured LHLO GPU Cholesky to better match XLA HLO by eliminating the
  untyped buffer and changing is_upper attribute to is_lower.
- Change LhloDialectEmitter to emit LHLO GPU Cholesky operation.

PiperOrigin-RevId: 343873516
2020-11-23 10:06:21 -08:00
Benjamin Kramer aa4d33149a Integrate LLVM at llvm/llvm-project@3c696a212b
Updates LLVM usage to match
[3c696a212ba4](https://github.com/llvm/llvm-project/commit/3c696a212ba4)

PiperOrigin-RevId: 343828493
2020-11-23 04:56:02 -08:00
A. Unique TensorFlower e21731d54a Integrate LLVM at llvm/llvm-project@1b749c0cb5
Updates LLVM usage to match
[1b749c0cb5c0](https://github.com/llvm/llvm-project/commit/1b749c0cb5c0)

PiperOrigin-RevId: 343583560
2020-11-20 16:10:59 -08:00
Tim Shen 5a6edaa588 Roll-forward with fix:
[XLA/GPU] Migrate nested reduce emitter to take LMHLO.

PiperOrigin-RevId: 343582798
2020-11-20 16:06:58 -08:00
A. Unique TensorFlower f849c45b74 Integrate LLVM at llvm/llvm-project@95ce9fbc23
Updates LLVM usage to match
[95ce9fbc235a](https://github.com/llvm/llvm-project/commit/95ce9fbc235a)

PiperOrigin-RevId: 343496733
2020-11-20 08:27:13 -08:00
A. Unique TensorFlower fd758e3efb Integrate LLVM at llvm/llvm-project@0caa82e2ac
Updates LLVM usage to match
[0caa82e2ac53](https://github.com/llvm/llvm-project/commit/0caa82e2ac53)

PiperOrigin-RevId: 343483451
2020-11-20 06:46:41 -08:00
A. Unique TensorFlower 1efc5f9675 Integrate LLVM at llvm/llvm-project@1a7d7f423e
Updates LLVM usage to match
[1a7d7f423e75](https://github.com/llvm/llvm-project/commit/1a7d7f423e75)

PiperOrigin-RevId: 343435160
2020-11-19 22:25:55 -08:00
A. Unique TensorFlower bae7c7bd2b [XLA/GPU] Migrate nested reduce emitter to take LMHLO.
PiperOrigin-RevId: 343388053
2020-11-19 16:00:58 -08:00
Tim Shen 303181924c [XLA/GPU] Migrate nested reduce emitter to take LMHLO.
PiperOrigin-RevId: 343373088
2020-11-19 14:41:55 -08:00
A. Unique TensorFlower 88642f58d3 Integrate LLVM at llvm/llvm-project@620adacf87
Updates LLVM usage to match
[620adacf87a3](https://github.com/llvm/llvm-project/commit/620adacf87a3)

PiperOrigin-RevId: 343307878
2020-11-19 09:37:20 -08:00
A. Unique TensorFlower e742da06dc Integrate LLVM at llvm/llvm-project@1983acce7c
Updates LLVM usage to match
[1983acce7c4d](https://github.com/llvm/llvm-project/commit/1983acce7c4d)

PiperOrigin-RevId: 343260267
2020-11-19 03:45:15 -08:00
A. Unique TensorFlower 7f239c7ba2 Add canonicalizer for Reshape(Broadcast(X)) pattern when it is an identity sequence
PiperOrigin-RevId: 343251257
2020-11-19 02:32:45 -08:00
A. Unique TensorFlower 61537008f4 Integrate LLVM at llvm/llvm-project@6b1341eb5b
Updates LLVM usage to match
[6b1341eb5bb7](https://github.com/llvm/llvm-project/commit/6b1341eb5bb7)

PiperOrigin-RevId: 343243049
2020-11-19 01:18:46 -08:00
A. Unique TensorFlower a4ae116b4d Integrate LLVM at llvm/llvm-project@5824a572ad
Updates LLVM usage to match
[5824a572adb5](https://github.com/llvm/llvm-project/commit/5824a572adb5)

PiperOrigin-RevId: 343223895
2020-11-18 22:17:37 -08:00
A. Unique TensorFlower 1a62ea1b95 Integrate LLVM at llvm/llvm-project@8c1e3cbebf
Updates LLVM usage to match
[8c1e3cbebfe9](https://github.com/llvm/llvm-project/commit/8c1e3cbebfe9)

PiperOrigin-RevId: 343125876
2020-11-18 11:57:00 -08:00