colin
/
mlir-hlo
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
451 Commits 1 Branch 0 Tags 2.6 MiB
Commit Graph

451 Commits

Author SHA1 Message Date
Rahul Joshi 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
A. Unique TensorFlower 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
A. Unique TensorFlower 6a05893169 [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346725498
 2020-12-10 00:21:36 -08:00
Christian Sigg 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
Alexander Belyaev c36afd275e [HLO] Add a pattern for HLO ConstOp to HLO -> Linalg conversion. 
PiperOrigin-RevId: 346718273
 2020-12-09 23:24:57 -08:00
Tim Shen cfcf741932 [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346716519
 2020-12-09 23:08:13 -08:00
A. Unique TensorFlower 65ebd85563 [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346658288
 2020-12-09 16:05:43 -08:00
Tim Shen 1c10e1fec6 [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346624905
 2020-12-09 13:31:50 -08:00
A. Unique TensorFlower 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
A. Unique TensorFlower 1a5a1b5f41 [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346574093
 2020-12-09 09:54:16 -08:00
Tim Shen 3c33fe4b9e [XLA/GPU] Migrate all unnested elementwise emitters. 
PiperOrigin-RevId: 346559170
 2020-12-09 08:42:51 -08:00
A. Unique TensorFlower 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
Alexander Belyaev e9057cbd93 [mlir] Enable conversion of HLO SignOp to Linalg. 
PiperOrigin-RevId: 346380514
 2020-12-08 12:08:26 -08:00
A. Unique TensorFlower 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
Stephan Herhut 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
Stephan Herhut dd5895d083 Extend unranked hlo transformations to also support and, or and xor. 
PiperOrigin-RevId: 346270393
 2020-12-08 01:00:26 -08:00
A. Unique TensorFlower 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
Benjamin Kramer 5235eceea0 Lower mhlo shifts to linalg 
PiperOrigin-RevId: 346161253
 2020-12-07 13:02:32 -08:00
A. Unique TensorFlower 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
Tres Popp 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
A. Unique TensorFlower 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
Smit Hinsu 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
A. Unique TensorFlower 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
Phoenix Meadowlark 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
Smit Hinsu 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
Rahul Joshi 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
A. Unique TensorFlower 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
A. Unique TensorFlower 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
Smit Hinsu 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
A. Unique TensorFlower 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
Tim Shen 73cb134cca [MLIR] Add XLA HLO -> LMHLO conversion for all elementwise ops. 
PiperOrigin-RevId: 345557248
 2020-12-03 16:12:02 -08:00
A. Unique TensorFlower 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
Tres Popp 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
A. Unique TensorFlower 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
A. Unique TensorFlower 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
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