Commit Graph

209 Commits

Author SHA1 Message Date
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 64d9a465ee Integrate LLVM at llvm/llvm-project@a70f2eb3e3
Updates LLVM usage to match
[a70f2eb3e39a](https://github.com/llvm/llvm-project/commit/a70f2eb3e39a)

PiperOrigin-RevId: 330468316
2020-09-08 02:03:22 -07:00
A. Unique TensorFlower e391762513 Integrate LLVM at llvm/llvm-project@7cfc8f0c7c
Updates LLVM usage to match
[7cfc8f0c7c24](https://github.com/llvm/llvm-project/commit/7cfc8f0c7c24)

PiperOrigin-RevId: 330443088
2020-09-07 21:14:16 -07:00
A. Unique TensorFlower 8e2b874fe8 Integrate LLVM at llvm/llvm-project@7d53fecb67
Updates LLVM usage to match
[7d53fecb6792](https://github.com/llvm/llvm-project/commit/7d53fecb6792)

PiperOrigin-RevId: 330418199
2020-09-07 15:14:42 -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
Eugene Burmako dde1ed56cc Add support for legalizing mhlo.slice to lmhlo.slice
PiperOrigin-RevId: 330153599
2020-09-04 21:43:01 -07:00
Eugene Burmako 2a9d7ac084 Add support for legalizing lmhlo.transpose to linalg.generic
PiperOrigin-RevId: 330130704
2020-09-04 15:37:42 -07:00
Eugene Burmako f5d12604ed Add support for legalizing mhlo.transpose to lmhlo.transpose
PiperOrigin-RevId: 330127758
2020-09-04 14:59:08 -07:00
Mehdi Amini b6b30698ad Integrate LLVM at llvm/llvm-project@2dd9a4d855
Updates LLVM usage to match
[2dd9a4d855f6](https://github.com/llvm/llvm-project/commit/2dd9a4d855f6)

PiperOrigin-RevId: 330048272
2020-09-03 21:28:38 -07:00
Prakalp Srivastava b7248424ae Fix return type of replica-id to unsigned int 32 tensor.
PiperOrigin-RevId: 330012181
2020-09-03 16:08:38 -07:00
A. Unique TensorFlower 12b221f459 Integrate LLVM at llvm/llvm-project@dc8d7d23d8
Updates LLVM usage to match
[dc8d7d23d8d2](https://github.com/llvm/llvm-project/commit/dc8d7d23d8d2)

PiperOrigin-RevId: 329992306
2020-09-03 14:21:30 -07:00
Mehdi Amini 56689ce908 Temporarily mark hlo-legalize-to-lhlo.mlir as failing, pending fixes
Upstream changed the buffer assignment API and this needs adjustment on our side.

PiperOrigin-RevId: 329959756
2020-09-03 11:36:20 -07:00
A. Unique TensorFlower 801262688c Integrate LLVM at llvm/llvm-project@1426ac0482
Updates LLVM usage to match
[1426ac048295](https://github.com/llvm/llvm-project/commit/1426ac048295)

PiperOrigin-RevId: 329839640
2020-09-02 20:03:41 -07:00
A. Unique TensorFlower aedcea6acf Integrate LLVM at llvm/llvm-project@3445ec9ba7
Updates LLVM usage to match
[3445ec9ba718](https://github.com/llvm/llvm-project/commit/3445ec9ba718)

PiperOrigin-RevId: 329691486
2020-09-02 05:08:34 -07:00
A. Unique TensorFlower 65b0613491 Integrate LLVM at llvm/llvm-project@202766947e
Updates LLVM usage to match
[202766947edb](https://github.com/llvm/llvm-project/commit/202766947edb)

PiperOrigin-RevId: 329673065
2020-09-02 02:27:52 -07:00
Robert Suderman 7c93352a40 Scalar / Trivial folding for mhlo.select
This covers the case where the predicate is a splat or the on_true/on_false
values are the same.

PiperOrigin-RevId: 329622785
2020-09-01 18:34:12 -07:00
A. Unique TensorFlower 158dbba4e5 Integrate LLVM at llvm/llvm-project@5ffd940ac0
Updates LLVM usage to match
[5ffd940ac02a](https://github.com/llvm/llvm-project/commit/5ffd940ac02a)

PiperOrigin-RevId: 329606435
2020-09-01 16:40:41 -07:00
A. Unique TensorFlower a622bf479b Fix mhlo::SliceOp::fold to not crash on unknown shapes
PiperOrigin-RevId: 329504383
2020-09-01 07:37:36 -07:00
A. Unique TensorFlower dcab119cec Integrate LLVM at llvm/llvm-project@ffd0b31c7c
Updates LLVM usage to match
[ffd0b31c7cba](https://github.com/llvm/llvm-project/commit/ffd0b31c7cba)

PiperOrigin-RevId: 329466048
2020-09-01 02:06:53 -07:00
Mehdi Amini 0e6e2f06e7 Integrate LLVM at llvm/llvm-project@1d3d9b9cd8
Updates LLVM usage to match
[1d3d9b9cd808](https://github.com/llvm/llvm-project/commit/1d3d9b9cd808)

PiperOrigin-RevId: 329432069
2020-08-31 20:34:04 -07:00
A. Unique TensorFlower 3e870929b5 Integrate LLVM at llvm/llvm-project@646f19bb9d
Updates LLVM usage to match
[646f19bb9dc8](https://github.com/llvm/llvm-project/commit/646f19bb9dc8)

PiperOrigin-RevId: 329406157
2020-08-31 16:55:04 -07:00
A. Unique TensorFlower a4857fad84 Integrate LLVM at llvm/llvm-project@bc3d4d9ed7
Updates LLVM usage to match
[bc3d4d9ed783](https://github.com/llvm/llvm-project/commit/bc3d4d9ed783)

PiperOrigin-RevId: 329359501
2020-08-31 12:54:51 -07:00
Benjamin Kramer 049f034116 Lower mhlo.floor to lmhlo to linalg
PiperOrigin-RevId: 329304882
2020-08-31 08:16:36 -07:00
A. Unique TensorFlower acf7413d40 Integrate LLVM at llvm/llvm-project@1d01fc100b
Updates LLVM usage to match
[1d01fc100bb5](https://github.com/llvm/llvm-project/commit/1d01fc100bb5)

PiperOrigin-RevId: 329266181
2020-08-31 02:47:45 -07:00
A. Unique TensorFlower f619da8a5c Integrate LLVM at llvm/llvm-project@11cf6346fd
Updates LLVM usage to match
[11cf6346fd49](https://github.com/llvm/llvm-project/commit/11cf6346fd49)

PiperOrigin-RevId: 329251819
2020-08-31 00:26:30 -07:00
A. Unique TensorFlower 29d7f01fbb Integrate LLVM at llvm/llvm-project@90166c2563
Updates LLVM usage to match
[90166c256310](https://github.com/llvm/llvm-project/commit/90166c256310)

PiperOrigin-RevId: 329065383
2020-08-28 23:07:36 -07:00
Jacques Pienaar 344c500fca [mhlo] Add legalize to SCF pass
Start of pass to legalize MHLO control flow to SCF for further optimization in common form. The current version just matches a very simple instance (which also happens to occur a few times). Exposes some further canonicalization opportunities that aren't yet addressed.

PiperOrigin-RevId: 329017723
2020-08-28 15:11:58 -07:00
Benjamin Kramer 7176fb1839 Integrate LLVM at llvm/llvm-project@85dacca29f
Updates LLVM usage to match
[85dacca29f82](https://github.com/llvm/llvm-project/commit/85dacca29f82)

PiperOrigin-RevId: 328919376
2020-08-28 04:58:12 -07:00
Benjamin Kramer 848ca244d2 Integrate LLVM at llvm/llvm-project@b5924a8e27
Updates LLVM usage to match
[b5924a8e2753](https://github.com/llvm/llvm-project/commit/b5924a8e2753)

PiperOrigin-RevId: 328760781
2020-08-27 10:12:46 -07:00
Benjamin Kramer 011ecf41d3 Integrate LLVM at llvm/llvm-project@1d8af682ef
Updates LLVM usage to match
[1d8af682ef1d](https://github.com/llvm/llvm-project/commit/1d8af682ef1d)

PiperOrigin-RevId: 328733270
2020-08-27 07:28:58 -07:00
A. Unique TensorFlower ca33164bb6 Integrate LLVM at llvm/llvm-project@95848ea101
Updates LLVM usage to match
[95848ea10127](https://github.com/llvm/llvm-project/commit/95848ea10127)

PiperOrigin-RevId: 328605350
2020-08-26 14:16:41 -07:00
Benjamin Kramer e8db30a663 Integrate LLVM at llvm/llvm-project@642cb7865f
Updates LLVM usage to match
[642cb7865f35](https://github.com/llvm/llvm-project/commit/642cb7865f35)

PiperOrigin-RevId: 328515095
2020-08-26 05:48:57 -07:00
Mehdi Amini 36ddbeb6b2 Remove the dependency on global dialect registry from mlir-hlo
PiperOrigin-RevId: 328457105
2020-08-25 20:30:42 -07:00
A. Unique TensorFlower ebe7857fb1 Integrate LLVM at llvm/llvm-project@9500a72091
Updates LLVM usage to match
[9500a7209163](https://github.com/llvm/llvm-project/commit/9500a7209163)

PiperOrigin-RevId: 328436293
2020-08-25 17:24:17 -07:00
A. Unique TensorFlower 7bc22bfa11 Integrate LLVM at llvm/llvm-project@ef76686916
Updates LLVM usage to match
[ef76686916d4](https://github.com/llvm/llvm-project/commit/ef76686916d4)

PiperOrigin-RevId: 328404846
2020-08-25 14:33:30 -07:00
Benjamin Kramer dd5f6fb550 Integrate LLVM at llvm/llvm-project@f8454d60b8
Updates LLVM usage to match
[f8454d60b829](https://github.com/llvm/llvm-project/commit/f8454d60b829)

PiperOrigin-RevId: 328352040
2020-08-25 10:07:12 -07:00
Benjamin Kramer 3ce5337881 Integrate LLVM at llvm/llvm-project@c6fb72de4f
Updates LLVM usage to match
[c6fb72de4f55](https://github.com/llvm/llvm-project/commit/c6fb72de4f55)

PiperOrigin-RevId: 328304428
2020-08-25 04:05:03 -07:00
A. Unique TensorFlower e4383a1da8 Integrate LLVM at llvm/llvm-project@77e5a195f8
Updates LLVM usage to match
[77e5a195f818](https://github.com/llvm/llvm-project/commit/77e5a195f818)

PiperOrigin-RevId: 328236932
2020-08-24 17:12:41 -07:00
Uday Bondhugula 94296bb7ec PR #42509: [MLIR] Add folder for mhlo get_dimension_size
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/42509

Add folder for mhlo GetDimensionSizeOp (`mhlo.get_dimension_size`).
`get_dimension_size` folds to a constant when the corresponding tensor
dimension size is statically known / constant.
Copybara import of the project:

--
5994915525ec2e932125aa1f133ce2260ba100af by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Add folder for mhlo get_dimension_size

Add folder for mhlo GetDimensionSizeOp. get_dimension_size folds to a
constant when the corresponding tensor dimension size is statically
known / constant.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/42509 from polymage-labs:get_dimension_size_fold 5994915525ec2e932125aa1f133ce2260ba100af
PiperOrigin-RevId: 328222517
2020-08-24 15:36:30 -07:00
Mehdi Amini 73b5a44f33 Update third_party/tensorflow/compiler/mlir/tensorflow/utils/... to not depend on the global Dialect Registry (NFC)
PiperOrigin-RevId: 328171679
2020-08-24 10:59:22 -07:00
Benjamin Kramer 6434448813 Integrate LLVM at llvm/llvm-project@bad7d6b373
Updates LLVM usage to match
[bad7d6b3735d](https://github.com/llvm/llvm-project/commit/bad7d6b3735d)

PiperOrigin-RevId: 328142991
2020-08-24 08:20:25 -07:00
Mehdi Amini 12b4bc75b3 Integrate LLVM at llvm/llvm-project@b999400a4f
Updates LLVM usage to match
[b999400a4fb6](https://github.com/llvm/llvm-project/commit/b999400a4fb6)

PiperOrigin-RevId: 328096769
2020-08-24 01:20:00 -07:00
Mehdi Amini a229bc977d Integrate LLVM at llvm/llvm-project@f6decfa36d
Updates LLVM usage to match
[f6decfa36d89](https://github.com/llvm/llvm-project/commit/f6decfa36d89)

PiperOrigin-RevId: 328073633
2020-08-23 20:40:59 -07:00
Mehdi Amini ab84a780d0 Integrate LLVM at llvm/llvm-project@f164534ca8
Updates LLVM usage to match
[f164534ca8e0](https://github.com/llvm/llvm-project/commit/f164534ca8e0)

PiperOrigin-RevId: 328046788
2020-08-23 13:39:12 -07:00
Uday Bondhugula 282dba6d38 PR #42508: [MLIR] Erase dead lmhlo.constant ops
Imported from GitHub PR https://github.com/tensorflow/tensorflow/pull/42508

An lmhlo.constant op on an memref that is locally allocated and with
no users other than dealloc's can be deleted. Add a canonicalization
pattern for this.
Copybara import of the project:

--
8758c409a15f567e7cb8e1077faa020f5705c85a by Uday Bondhugula <uday@polymagelabs.com>:

[MLIR] Erase dead lmhlo.constant ops

An lmhlo.constant op on an memref that is locally allocated and with
no other users (other than dealloc's) can be deleted. Add a
canonicalization patter for this.

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/tensorflow/pull/42508 from polymage-labs:lhlo_constant_erase 8758c409a15f567e7cb8e1077faa020f5705c85a
PiperOrigin-RevId: 328042416
2020-08-23 12:28:54 -07:00
Hanhan Wang bfd629ecb0 Enhance lowering reshape op to Linalg.
Handle non-expansion and non-collapsion cases by rewriting it to two reshape
ops.

PiperOrigin-RevId: 327926863
2020-08-21 23:27:34 -07:00
A. Unique TensorFlower d2c9d03f31 Integrate LLVM at llvm/llvm-project@02bf5632a9
Updates LLVM usage to match
[02bf5632a94d](https://github.com/llvm/llvm-project/commit/02bf5632a94d)

PiperOrigin-RevId: 327894378
2020-08-21 16:34:38 -07:00
A. Unique TensorFlower 0a8096d91d Integrate LLVM at llvm/llvm-project@5e3fd471ac
Updates LLVM usage to match
[5e3fd471acb7](https://github.com/llvm/llvm-project/commit/5e3fd471acb7)

PiperOrigin-RevId: 327878947
2020-08-21 14:51:58 -07:00
A. Unique TensorFlower 72eb42a908 Integrate LLVM at llvm/llvm-project@50aae46331
Updates LLVM usage to match
[50aae463315d](https://github.com/llvm/llvm-project/commit/50aae463315d)

PiperOrigin-RevId: 327849475
2020-08-21 12:02:23 -07:00
A. Unique TensorFlower a4092162cd Integrate LLVM at llvm/llvm-project@c1dd5df425
Updates LLVM usage to match
[c1dd5df4255c](https://github.com/llvm/llvm-project/commit/c1dd5df4255c)

PiperOrigin-RevId: 327818600
2020-08-21 09:04:28 -07:00