9930c20c31
[mlir][hlo] Fix lowering of NE comparison. It should return true if either side is NaN
...
PiperOrigin-RevId: 346988987
2020-12-11 06:46:14 -08:00
ab6ee11813
Fix folding of HLO SliceOp with zero elements
...
This was causing division by zero in this case.
PiperOrigin-RevId: 346920942
2020-12-10 20:22:48 -08:00
f232da1f9d
[MLIR:HLO] Add window_reversal attribute to convolution attributes.
...
- Add this attribute to match the corresponding XLA HLO attribute on convolution
operations.
- A true value indicates a reversal of the corresponding kernel spatial dimension.
- Since XLA builder does not support this attribute, use a custom HLO converted to map
from mlir::mhlo::ConvOp to XLA.
PiperOrigin-RevId: 346891737
2020-12-10 16:39:19 -08:00
3442ac270d
Integrate LLVM at llvm/llvm-project@997a719d5a
...
Updates LLVM usage to match
[997a719d5a70](https://github.com/llvm/llvm-project/commit/997a719d5a70 )
PiperOrigin-RevId: 346834025
2020-12-10 11:58:16 -08:00
6a05893169
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346725498
2020-12-10 00:21:36 -08:00
51f535454d
Use OpState::operator->() to get to member functions in Operation so we can remove the corresponding methods from OpState.
...
PiperOrigin-RevId: 346721668
2020-12-09 23:54:07 -08:00
c36afd275e
[HLO] Add a pattern for HLO ConstOp to HLO -> Linalg conversion.
...
PiperOrigin-RevId: 346718273
2020-12-09 23:24:57 -08:00
cfcf741932
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346716519
2020-12-09 23:08:13 -08:00
65ebd85563
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346658288
2020-12-09 16:05:43 -08:00
1c10e1fec6
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346624905
2020-12-09 13:31:50 -08:00
afea1e1897
Integrate LLVM at llvm/llvm-project@6883042528
...
Updates LLVM usage to match
[6883042528d0](https://github.com/llvm/llvm-project/commit/6883042528d0 )
PiperOrigin-RevId: 346606080
2020-12-09 12:06:47 -08:00
1a5a1b5f41
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346574093
2020-12-09 09:54:16 -08:00
3c33fe4b9e
[XLA/GPU] Migrate all unnested elementwise emitters.
...
PiperOrigin-RevId: 346559170
2020-12-09 08:42:51 -08:00
6cb9cc53ec
Integrate LLVM at llvm/llvm-project@d553243fe4
...
Updates LLVM usage to match
[d553243fe4b5](https://github.com/llvm/llvm-project/commit/d553243fe4b5 )
PiperOrigin-RevId: 346393264
2020-12-08 13:09:13 -08:00
e9057cbd93
[mlir] Enable conversion of HLO SignOp to Linalg.
...
PiperOrigin-RevId: 346380514
2020-12-08 12:08:26 -08:00
4e6367e9d6
Integrate LLVM at llvm/llvm-project@a1344779ab
...
Updates LLVM usage to match
[a1344779ab01](https://github.com/llvm/llvm-project/commit/a1344779ab01 )
PiperOrigin-RevId: 346312316
2020-12-08 06:46:57 -08:00
c3790af758
Add plumbing for or and xor to hlo to lhlo and linalg lowerings.
...
PiperOrigin-RevId: 346311314
2020-12-08 06:39:02 -08:00
dd5895d083
Extend unranked hlo transformations to also support and, or and xor.
...
PiperOrigin-RevId: 346270393
2020-12-08 01:00:26 -08:00
812221db97
Integrate LLVM at llvm/llvm-project@2ac4d0f45a
...
Updates LLVM usage to match
[2ac4d0f45a2a](https://github.com/llvm/llvm-project/commit/2ac4d0f45a2a )
PiperOrigin-RevId: 346203743
2020-12-07 16:20:57 -08:00
5235eceea0
Lower mhlo shifts to linalg
...
PiperOrigin-RevId: 346161253
2020-12-07 13:02:32 -08:00
eaa21130e8
Integrate LLVM at llvm/llvm-project@ecaff13fc0
...
Updates LLVM usage to match
[ecaff13fc0bc](https://github.com/llvm/llvm-project/commit/ecaff13fc0bc )
PiperOrigin-RevId: 346083459
2020-12-07 07:00:58 -08:00
d327fc5737
[kernel_gen] Lower max rank specialization from 6 to 5
...
We don't care much about rank 6 broadcasting operations and this lowers compile times significantly.
PiperOrigin-RevId: 346046601
2020-12-07 02:18:38 -08:00
4c0c9ab119
Integrate LLVM at llvm/llvm-project@db226cdf4c
...
Updates LLVM usage to match
[db226cdf4cf9](https://github.com/llvm/llvm-project/commit/db226cdf4cf9 )
PiperOrigin-RevId: 346045832
2020-12-07 02:13:36 -08:00
bc7b6374c8
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: 345902167
2020-12-05 18:55:41 -08:00
55268f9ee8
Integrate LLVM at llvm/llvm-project@7f6f9f4cf9
...
Updates LLVM usage to match
[7f6f9f4cf966](https://github.com/llvm/llvm-project/commit/7f6f9f4cf966 )
PiperOrigin-RevId: 345745888
2020-12-04 13:33:48 -08:00
c33bdcbd03
Remove fold of `mhlo.compare(%arg0, %arg0)` for floating types.
...
Two tensors having the same SSA-value isn't sufficient for equality for floating types, as `NaN != NaN`. As written this causes `tf.IsNan` to [miscompile](https://github.com/google/iree/issues/4061 ).
PiperOrigin-RevId: 345730640
2020-12-04 12:15:02 -08:00
9bd1995f90
Legalize XlaReplicaId to HLO replica-id op
...
Also, define shape inference function for HLO replica-id op.
PiperOrigin-RevId: 345714342
2020-12-04 11:04:40 -08:00
e48881af81
[MLIR:LHLO_GPU] Add fused convolution operation without any side inputs.
...
- Add a variant of the fused convolution that does not need a side input and side input scale.
- Rename the existing one to `ConvForwardFusedSideInputOp`.
- Update tests to exercise all variants of the convolution ops in the GPU dialect.
- Eliminate unused `LHLO_ExtentBuffer` and changed LHLO_Buffer to allow any integer element
type to match what XLA can generate sometimes for scratch buffers.
PiperOrigin-RevId: 345701569
2020-12-04 10:09:27 -08:00
3691e39f62
Integrate LLVM at llvm/llvm-project@f5d52916ce
...
Updates LLVM usage to match
[f5d52916ce34](https://github.com/llvm/llvm-project/commit/f5d52916ce34 )
PiperOrigin-RevId: 345672828
2020-12-04 07:24:17 -08:00
e87d53742b
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: 345618958
2020-12-04 00:04:10 -08:00
9456af5880
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: 345605910
2020-12-03 22:09:56 -08:00
a7506fab4f
Integrate LLVM at llvm/llvm-project@55db6ec1cc
...
Updates LLVM usage to match
[55db6ec1cc20](https://github.com/llvm/llvm-project/commit/55db6ec1cc20 )
PiperOrigin-RevId: 345565090
2020-12-03 16:55:25 -08:00
73cb134cca
[MLIR] Add XLA HLO -> LMHLO conversion for all elementwise ops.
...
PiperOrigin-RevId: 345557248
2020-12-03 16:12:02 -08:00
4b6a38da54
Integrate LLVM at llvm/llvm-project@1860331932
...
Updates LLVM usage to match
[18603319321a](https://github.com/llvm/llvm-project/commit/18603319321a )
PiperOrigin-RevId: 345513348
2020-12-03 12:45:16 -08:00
7c3f049c8e
[kernel_gen] Lower max rank specialization from 6 to 5
...
We don't care much about rank 6 broadcasting operations and this lowers compile times significantly.
PiperOrigin-RevId: 345466476
2020-12-03 09:19:25 -08:00
d2e3797d7d
Integrate LLVM at llvm/llvm-project@6627a3c287
...
Updates LLVM usage to match
[6627a3c2873f](https://github.com/llvm/llvm-project/commit/6627a3c2873f )
PiperOrigin-RevId: 345428622
2020-12-03 05:10:04 -08:00
1c6c04f76b
Integrate LLVM at llvm/llvm-project@baa005c96c
...
Updates LLVM usage to match
[baa005c96ce6](https://github.com/llvm/llvm-project/commit/baa005c96ce6 )
PiperOrigin-RevId: 345327906
2020-12-02 15:41:44 -08:00
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
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
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
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
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
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
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
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
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
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
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
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
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
Benjamin Kramer
Smit Hinsu
Rahul Joshi
A. Unique TensorFlower
Christian Sigg
Alexander Belyaev
Tim Shen
Stephan Herhut
Tres Popp
Phoenix Meadowlark
River Riddle