Also generate the kernels for all types of casts between signed int and float types.
This requires some adaptations to our build macros so that we can also specify the
output type of a kernel.
PiperOrigin-RevId: 354067727
Allow for relative tolerance in unary kernel tests. In case of the cosh kernels,
this allows to accept an observed difference of 5.6e-8 between the kernel and
the `std::cosh` reference (32829984.568665262 vs. 32829984.568665318) in one of
the test cases.
PiperOrigin-RevId: 351983698
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
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
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
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
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
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
- 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
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
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
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
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
- 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
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
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
Use constant to generate the correct assertion message. This avoids
confusion when lowering the max rank specialization for debugging.
PiperOrigin-RevId: 344769021
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
M_PI and other math constants (used in chlo_legalize_hlo_patterns.td)
are not part of the C++ standard and must be enabled on MSVC
(similar to _GNU_SOURCE adding glibc symbols to posix headers).
PiperOrigin-RevId: 342432987
- Extend MHLO CustomCall to have multiple tensors as results.
- Extend LHLO CustomCall to have multiple memrefs for output operands.
- Fix HLO->LHLO and XLA HLO->LHLO mapping for CustomCall to setup the
operand_segment_sizes attribute correctly.
PiperOrigin-RevId: 342067762
This specific pattern can be replaced with the shape
passed to dynamic_reshape. This is implemented as a
canonicalization on mhlo.dynamic_reshape to fit in
the infrastructure of canonicalization.
PiperOrigin-RevId: 342009365
- Extract code to create result memref's into a ConvertResults function.
- Also fix a bug when using reifyReturnTypes: use correct index for result_shape instead
of always using the first element.
PiperOrigin-RevId: 341852227
The conversion had a bug in computation of strides and sizes args for std.memref_reinterpret_cast. The previous version also relied on linalg::ReshapeOp to do broadcasting when the rank of the output was higher than the rank of the input. Now the broadcasting is entirely done via descriptor modification and linalg::ReshapeOp was replaced with CopyOp.
PiperOrigin-RevId: 341379871
This is to match with HLO semantics and general dimension semantics in MLIR.
Also,
* Define minimal verifier for these ops.
* Add folder for SetDimensionSize op on static shaped dimension.
* Fix assumption of ranked shape in GetDimensionSize op.
PiperOrigin-RevId: 341150923
Lowerings that depended on operations between real and complex types may
not infer the correct intermediate type. Removing these operations as
they are not technically legally generated operations. Updated tests
to validate this.
PiperOrigin-RevId: 341128903
Previously this started at rank 2 after checking for scalars and equal shapes. This resulted in cases such as <1xf32> + <2xf32> being treated as impossible.
PiperOrigin-RevId: 341043965
A. Unique TensorFlower
Adrian Kuegel
Rahul Joshi
Tim Shen
Hanhan Wang
Alexander Belyaev
Stephan Herhut
Jacques Pienaar
Tres Popp
Smit Hinsu
Christian Sigg
River Riddle
Benjamin Kramer
Phoenix Meadowlark
Lucy Fox
Ben Vanik
Robert Suderman