Commit Graph

960 Commits

Author SHA1 Message Date
Stephan Herhut b0e0ca830c Integrate LLVM at llvm/llvm-project@c8a914db5c
Updates LLVM usage to match
[c8a914db5c60](https://github.com/llvm/llvm-project/commit/c8a914db5c60)

PiperOrigin-RevId: 351363798
2021-01-12 07:09:38 -08:00
A. Unique TensorFlower d52092a5cc Integrate LLVM at llvm/llvm-project@aefeb5f136
Updates LLVM usage to match
[aefeb5f136e7](https://github.com/llvm/llvm-project/commit/aefeb5f136e7)

PiperOrigin-RevId: 351291118
2021-01-11 20:24:58 -08:00
Stephan Herhut 5359ea620a Integrate LLVM at llvm/llvm-project@89878e8c96
Updates LLVM usage to match
[89878e8c966a](https://github.com/llvm/llvm-project/commit/89878e8c966a)

PiperOrigin-RevId: 351192927
2021-01-11 11:08:35 -08:00
Hanhan Wang 8f58f844e5 Upstream mhlo.dot lowering to Linalg to MHLO repo.
We prototyped the lowering from mhlo.dot to linalg.matmul in IREE. Since Linalg
now supports matmul in tensors world, we can move the lowering logic to tensors
world, and upstream to legalize_to_linalg.cc. The patch lowers the mhlo.dot to
the linalg.matmul/matvec/dot in tensors world.

PiperOrigin-RevId: 351184911
2021-01-11 10:35:24 -08:00
Alexander Belyaev 180f917446 [KERNEL_GEN] Add a pattern for hlo.dyn_broadcast->linalg to enable is_inf kernel.
PiperOrigin-RevId: 351179620
2021-01-11 10:13:31 -08:00
Alexander Belyaev ecf1bf5132 [KERNEL_GEN] Add a canonicalization pattern to drop a redundant dynamic reshape.
PiperOrigin-RevId: 351141868
2021-01-11 06:38:03 -08:00
Geoffrey Martin-Noble 47848764a5 Change hardcoded external/ include to match repository name
This is the preferred name for the Bazel repository when included from
another project.

Note: These hardcoded "external/" includes are a horrible hack that we
can hopefully get rid of soon with more robust Bazel tablegen rules.
PiperOrigin-RevId: 350810768
2021-01-08 12:07:53 -08:00
Alexander Belyaev ac38478350 Integrate LLVM at llvm/llvm-project@bcbdeafa9c
Updates LLVM usage to match
[bcbdeafa9cb3](https://github.com/llvm/llvm-project/commit/bcbdeafa9cb3)

PiperOrigin-RevId: 350763293
2021-01-08 07:57:20 -08:00
Alexander Belyaev 6c42f54298 [KERNEL_GEN] Restrict broadcast -> reshape canonicalization to identity dims.
This is needed to avoid the case, when the broadcast_in_dims also performs permutation.

PiperOrigin-RevId: 350650342
2021-01-07 15:30:28 -08:00
Alexander Belyaev a8c0f2b944 Integrate LLVM at llvm/llvm-project@7f7b0dc4e1
Updates LLVM usage to match
[7f7b0dc4e15f](https://github.com/llvm/llvm-project/commit/7f7b0dc4e15f)

PiperOrigin-RevId: 350576840
2021-01-07 09:24:01 -08:00
Alexander Belyaev 49a540bf72 Integrate LLVM at llvm/llvm-project@d38a0258a5
Updates LLVM usage to match
[d38a0258a5f4](https://github.com/llvm/llvm-project/commit/d38a0258a5f4)

PiperOrigin-RevId: 350140950
2021-01-05 08:03:11 -08:00
Alexander Belyaev 095dc28e5c [KERNEL_GEN] Add canonicalizaton pattern to drop a redundant broadcast op.
PiperOrigin-RevId: 350105790
2021-01-05 03:01:00 -08:00
A. Unique TensorFlower 2727ed4cf2 Integrate LLVM at llvm/llvm-project@8bee4d4e8f
Updates LLVM usage to match
[8bee4d4e8f54](https://github.com/llvm/llvm-project/commit/8bee4d4e8f54)

PiperOrigin-RevId: 349678834
2020-12-31 07:35:45 -08:00
A. Unique TensorFlower f8fbfdd061 Integrate LLVM at llvm/llvm-project@51a292d994
Updates LLVM usage to match
[51a292d99453](https://github.com/llvm/llvm-project/commit/51a292d99453)

PiperOrigin-RevId: 349610524
2020-12-30 15:45:59 -08:00
A. Unique TensorFlower 181f19881b Integrate LLVM at llvm/llvm-project@2016f2c8a7
Updates LLVM usage to match
[2016f2c8a76d](https://github.com/llvm/llvm-project/commit/2016f2c8a76d)

PiperOrigin-RevId: 349565754
2020-12-30 09:11:29 -08:00
A. Unique TensorFlower 2e215f75a5 Integrate LLVM at llvm/llvm-project@8001dcbd50
Updates LLVM usage to match
[8001dcbd50ce](https://github.com/llvm/llvm-project/commit/8001dcbd50ce)

PiperOrigin-RevId: 349514172
2020-12-29 22:42:21 -08:00
A. Unique TensorFlower a9e2b45bd5 Integrate LLVM at llvm/llvm-project@930c74f12d
Updates LLVM usage to match
[930c74f12d79](https://github.com/llvm/llvm-project/commit/930c74f12d79)

PiperOrigin-RevId: 349426692
2020-12-29 08:32:16 -08:00
Adrian Kuegel 50fc56a208 Fix header include.
It should not have the third_party prefix, this doesn't work in open source.

PiperOrigin-RevId: 348905548
2020-12-24 01:40:31 -08:00
A. Unique TensorFlower b0bf2ef45b Integrate LLVM at llvm/llvm-project@c3acda0798
Updates LLVM usage to match
[c3acda0798f9](https://github.com/llvm/llvm-project/commit/c3acda0798f9)

PiperOrigin-RevId: 348896724
2020-12-23 23:53:54 -08:00
Rahul Joshi e3754d7b5c [MLIR:HLO] Change LHLO Infeed to support multiple outputs.
PiperOrigin-RevId: 348851612
2020-12-23 14:57:08 -08:00
Rahul Joshi bd8e768e23 [MLIR:HLO] Rename LMHLO Infeed and Outfeed to InfeedOp and OutfeedOp.
- To be consistent with other operation names.

PiperOrigin-RevId: 348824379
2020-12-23 11:38:05 -08:00
Rahul Joshi 8252eafa99 [NFC] Factor out repeated code out of InferFusibilityOpInterface.
PiperOrigin-RevId: 348671671
2020-12-22 12:04:29 -08:00
Rahul Joshi bc367971ec [MLIR:LHLO_GPU] Add additional constraints for batchnorm
- Constrain batchnorm inputs and outputs to be fp memrefs.

PiperOrigin-RevId: 348665747
2020-12-22 11:30:40 -08:00
Stephan Herhut ccdd07f8e4 Prepare to remove tensor_load and tensor_store special handling from hlo to lhlo legalization.
This updates the tests to no longer rely on tensor_store. Once all users of this behavior have adopted, the tensor_store support will be removed.

PiperOrigin-RevId: 348624899
2020-12-22 06:29:12 -08:00
A. Unique TensorFlower c4accdcc41 Integrate LLVM at llvm/llvm-project@1b97cdf885
Updates LLVM usage to match
[1b97cdf885d6](https://github.com/llvm/llvm-project/commit/1b97cdf885d6)

PiperOrigin-RevId: 348587513
2020-12-21 23:49:18 -08:00
Tres Popp a42213b870 Define lowering of [l]mhlo.pow.
For floating point operations, this uses std.pow.
For integer operations, this lowers to a loop.
This adds a dependency on scf.

PiperOrigin-RevId: 348537232
2020-12-21 15:27:40 -08:00
A. Unique TensorFlower 99a0ee378c Integrate LLVM at llvm/llvm-project@511cfe9441
Updates LLVM usage to match
[511cfe944195](https://github.com/llvm/llvm-project/commit/511cfe944195)

PiperOrigin-RevId: 348482094
2020-12-21 09:50:18 -08:00
A. Unique TensorFlower 9fd87f5939 Integrate LLVM at llvm/llvm-project@f4c8b80318
Updates LLVM usage to match
[f4c8b8031800](https://github.com/llvm/llvm-project/commit/f4c8b8031800)

PiperOrigin-RevId: 348453089
2020-12-21 05:34:47 -08:00
A. Unique TensorFlower ee74c519a6 Integrate LLVM at llvm/llvm-project@17b3ff511c
Updates LLVM usage to match
[17b3ff511c0a](https://github.com/llvm/llvm-project/commit/17b3ff511c0a)

PiperOrigin-RevId: 348245062
2020-12-18 21:59:48 -08:00
A. Unique TensorFlower 1e36f6d38d Integrate LLVM at llvm/llvm-project@9ca67d7f44
Updates LLVM usage to match
[9ca67d7f4467](https://github.com/llvm/llvm-project/commit/9ca67d7f4467)

PiperOrigin-RevId: 348211935
2020-12-18 13:09:17 -08:00
Smit Hinsu 9466cffaf3 Restrict CHLO Acos and Sinh op lowering to non complex types
These are failing for complex types. Complex types require special handling. We have a fallback lowering for these ops so we can disable complex element types for now.

PiperOrigin-RevId: 348205002
2020-12-18 11:32:10 -08:00
Smit Hinsu 8d051723c0 Use InferTypeOpInterface for HLO AbsOp and fix result shape inference
Shape inference in case of ops with complex element types need to use the element type of complex as the result element type and not the full operand type.

Before:
"mhlo.abs"(%arg0) : (tensor<4xcomplex<f32>>) -> tensor<4xtensor<4xcomplex<f32>>>
After:
"mhlo.abs"(%arg0) : (tensor<4xcomplex<f32>>) -> tensor<4xf32>
PiperOrigin-RevId: 348123967
2020-12-17 17:37:07 -08:00
A. Unique TensorFlower 43ede42ce1 Integrate LLVM at llvm/llvm-project@e113317958
Updates LLVM usage to match
[e1133179587d](https://github.com/llvm/llvm-project/commit/e1133179587d)

PiperOrigin-RevId: 348115823
2020-12-17 16:35:29 -08:00
Smit Hinsu 737d15ded5 Handle operands with zero elements in HLO PadOp folder
PiperOrigin-RevId: 348034821
2020-12-17 09:27:36 -08:00
Rahul Joshi 8134bff98d [XLA:GPU] Add layout attributes to LHLO_GPU Convolution operations.
- MLIR MemRefs do not preserve layout information correctly when unit dimensions
  are involved. Operations like convolution that use cuDNN however need the correct
  layout to be preserved so that we do not end up creating an incompatible combination
  of input/filter/output layout that is not supported by cuDNN.
- Add these layouts to convolution attributes in the form of I32ArrayAttr for representing
  the layout in "minor_to_major" form similar to XLA.

PiperOrigin-RevId: 348034757
2020-12-17 09:26:28 -08:00
Christian Sigg dc7e63f74c Use mlir::OpState::operator->() to get to methods of mlir::Operation.
This is a preparation step to remove those methods from OpState.

PiperOrigin-RevId: 348010582
2020-12-17 06:29:33 -08:00
Adrian Kuegel 1f244c3e2c Fix SignOp lowering for floating point values.
It didn't return 0 for 0.0 and -0.0.
Currently we emit -0.0 for -0.0 which is correct according to the HLO dialect.
For the TF_SignOp we should emit 0.0 in that case, we will leave that as a TODO.
Enable the tests which work now, and add another one for Int64.
Also improve the registration code, we should not register the Int32 kernel.

PiperOrigin-RevId: 347981124
2020-12-17 01:45:54 -08:00
Christian Sigg 099c130daf Fix MLIR include paths.
PiperOrigin-RevId: 347976151
2020-12-17 00:56:04 -08:00
A. Unique TensorFlower 5da9190dd9 Integrate LLVM at llvm/llvm-project@0cf7e4b252
Updates LLVM usage to match
[0cf7e4b252fe](https://github.com/llvm/llvm-project/commit/0cf7e4b252fe)

PiperOrigin-RevId: 347948887
2020-12-16 20:30:17 -08:00
Phoenix Meadowlark 5a080ad795 Add Bazel build configuration to MLIR HLO standalone repo.
Tested that this works in OSS in tree with `bazel build //...` and as a submodule in IREE.

PiperOrigin-RevId: 347919106
2020-12-16 16:23:31 -08:00
Alexander Belyaev 65222893ae [KERNEL_GEN] Convert LHLO AddOp, SubOp (ComplexType) to complex ops.
PiperOrigin-RevId: 347805898
2020-12-16 05:45:06 -08:00
Adrian Kuegel 61244b136c Try to avoid a segfault if we don't support a lowering.
It can happen that a lowering for a certain type is not implemented yet.
We should not segfault in such a case, but instead return a failure().

PiperOrigin-RevId: 347801106
2020-12-16 04:58:17 -08:00
Alexander Belyaev e6e8920921 [KERNEL_GEN] Switch the pipeline to Linalg-on-Tensors.
PiperOrigin-RevId: 347781190
2020-12-16 01:51:15 -08:00
River Riddle 5ab94a00a7 [mlir][NFC] Replace usages or mlir/IR/StandardTypes.h with mlir/IR/BuiltinTypes.h
StandardTypes.h was moved to BuiltinTypes.h and is being removed.

PiperOrigin-RevId: 347751511
2020-12-15 21:28:18 -08:00
A. Unique TensorFlower ccbf39842f Integrate LLVM at llvm/llvm-project@7aeb3804c4
Updates LLVM usage to match
[7aeb3804c46c](https://github.com/llvm/llvm-project/commit/7aeb3804c46c)

PiperOrigin-RevId: 347727361
2020-12-15 17:40:25 -08:00
Tres Popp 6d7812bbc5 Correct non ConversionRewriter transformation in LegalizeToLinalg
PiperOrigin-RevId: 347622657
2020-12-15 08:56:58 -08:00
A. Unique TensorFlower f0c2695d31 Fix SignOp lowering for floating point values.
It didn't return 0 for 0.0 and -0.0.
Currently we emit -0.0 for -0.0 which is correct according to the HLO dialect.
For the TF_SignOp we should emit 0.0 in that case, we will leave that as a TODO.
Enable the tests which work now, and add another one for Int64.
Also improve the registration code, we should not register the Int32 kernel.

PiperOrigin-RevId: 347602378
2020-12-15 06:49:48 -08:00
Alexander Belyaev ddda2699fb [KERNEL_GEN] Switch the pipeline to Linalg-on-Tensors.
PiperOrigin-RevId: 347600145
2020-12-15 06:32:25 -08:00
Adrian Kuegel 79fa36bcbc Fix SignOp lowering for floating point values.
It didn't return 0 for 0.0 and -0.0.
Currently we emit -0.0 for -0.0 which is correct according to the HLO dialect.
For the TF_SignOp we should emit 0.0 in that case, we will leave that as a TODO.
Enable the tests which work now, and add another one for Int64.
Also improve the registration code, we should not register the Int32 kernel.

PiperOrigin-RevId: 347590340
2020-12-15 05:12:48 -08:00
River Riddle 9540e51617 [mlir][NFC] Replace usages or mlir/IR/StandardTypes.h with mlir/IR/BuiltinTypes.h
StandardTypes.h was moved to BuiltinTypes.h and is being removed.

PiperOrigin-RevId: 347559927
2020-12-15 00:59:29 -08:00