| Age | Commit message (Collapse) | Author |
|---|
|
|
|
Now that C++17 is enabled in LLVM, a lot of the TODOs and patterns to emulate C++17 features can be eliminated.
The steps I have taken were essentially:
```
git grep C++17
git grep c++17
git grep "initializer_list<int>"
```
and address given comments and patterns.
Most of the changes boiled down to just using fold expressions rather than initializer_list.
While doing this I also discovered that Clang by default restricts the depth of fold expressions to 256 elements. I specifically hit this with `TestDialect` in `addOperations`. I opted to not replace it with fold expressions because of that but instead adding a comment documenting the issue.
If any other functions may be called with more than 256 elements in the future we might have to revert other parts as well.
I don't think this is a common occurence besides the `TestDialect` however. If need be, this could potentially be fixed via `mlir-tblgen` in the future.
Differential Revision: https://reviews.llvm.org/D131323
|
|
|
|
|
|
Depends On D131285
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D131291
|
|
Avoid expensive calls to `SymbolTable::lookupNearestSymbolFrom` in verifier
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D131285
|
|
Based on `spv.ISubBorrow` from D127909.
Also resolved some clang-tidy warnings.
Reviewed By: antiagainst, ThomasRaoux
Differential Revision: https://reviews.llvm.org/D131281
|
|
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D131271
|
|
`makeTiledShape(s)`."
This reverts commit 56d94b3b902e21ff79b1ce9a6fb606a3f7c1c4db.
|
|
This patch adds constant folder for Atan2Op which only supports single and double precision floating-point.
Differential Revision: https://reviews.llvm.org/D131050
|
|
We were checking real-part twice, not real/imag-part.
The new test only passes after the bug fix.
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D131190
|
|
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D131080
|
|
This prepares patterns that sometimes are generated by the front-end
and would prohibit fusion of SDDMM flavored kernels.
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D131126
|
|
`SymbolTable::lookupSymbolIn` is an expensive operation and we do not want to do it twice
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D131145
|
|
This change separates computation of the actual parameters of the subset and
the materialization of subview/extract_slice. That way the users can still use
Linalg tiling logic even if they use different operations to materialize the
subsets.
Differential Revision: https://reviews.llvm.org/D131053
|
|
It seems only the default implementation is ever used, so it doesn't seem
necessary to include this method in the interface.
Differential Revision: https://reviews.llvm.org/D130986
|
|
This patch adds constant folder for AtanOp which only supports single and double precision floating-point.
Differential Revision: https://reviews.llvm.org/D130983
|
|
This rewriting was no longer functional after recent migration to one shot
bufferization. However, this revision makes it work again, with a CHECK test
to ensure fusion happens. Note that functionality is tested by several
integration tests.
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D130996
|
|
Although we have plans to support this, and many other, dimension level type(s), currently the tag is not supported. It will be easy to add this back once support is added.
NOTE: based on discussion in https://discourse.llvm.org/t/overcoming-sparsification-limitation-on-level-types/62585
https://github.com/llvm/llvm-project/issues/51658
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D131002
|
|
The generic allocation and deallocation instructions, which are optionally used during the MemRef -> LLVM conversion, should have a name that is specifically bound to their origin, that is the conversion pass itself.
Reviewed By: silvas
Differential Revision: https://reviews.llvm.org/D130588
|
|
This patch adds constant folder for TanhOp which only supports single and double precision floating-point.
Differential Revision: https://reviews.llvm.org/D130960
|
|
This patch adds constant folder for TanOp which only supports single and double precision floating-point.
Differential Revision: https://reviews.llvm.org/D130873
|
|
Adds optional attribute to support tensor cores on F32 datatype by lowering to `mma.sync` with TF32 operands. Since, TF32 is not a native datatype in LLVM we are adding `tf32Enabled` as an attribute to allow the IR to be aware of `MmaSyncOp` datatype. Additionally, this patch adds placeholders for nvgpu-to-nvgpu transformation targeting higher precision tf32x3.
For mma.sync on f32 input using tensor cores there are two possibilites:
(a) tf32 (1 `mma.sync` per warp-level matrix-multiply-accumulate)
(b) tf32x3 (3 `mma.sync` per warp-level matrix-multiply-accumulate)
Typically, tf32 tensor core acceleration comes at a cost of accuracy from missing precision bits. While f32 has 23 precision bits, tf32 has only 10 precision bits. tf32x3 aims to recover the precision bits by splitting each operand into two tf32 values and issue three `mma.sync` tensor core operations.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D130294
|
|
Attempting to apply the range analysis to real code revealed that
trunci wasn't correctly handling the case where truncation would
create wider ranges - for example, if we truncate [255, 257] : i16 to
i8, the result can be 255, 0, or 1, which isn't a contiguous range of
values.
The previous implementation would naively map this to [255, 1], which
would cause issues with unsigned ranges and unification.
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D130501
|
|
This is the follow up on https://reviews.llvm.org/D130730 which goes through upstream code and removes creating constant values in favour of using the constant indices in GEP directly. This leads to less and more readable code and more compact IR as well.
Differential Revision: https://reviews.llvm.org/D130731
|
|
This patch adds constant folder for ExpM1Op which only supports single and double precision floating-point.
Differential Revision: https://reviews.llvm.org/D130567
|
|
This patch removes the `type` field from `Attribute` along with the
`Attribute::getType` accessor.
Going forward, this means that attributes in MLIR will no longer have
types as a first-class concept. This patch lays the groundwork to
incrementally remove or refactor code that relies on generic attributes
being typed. The immediate impact will be on attributes that rely on
`Attribute` containing a type, such as `IntegerAttr`,
`DenseElementsAttr`, and `ml_program::ExternAttr`, which will now need
to define a type parameter on their storage classes. This will save
memory as all other attribute kinds will no longer contain a type.
Moreover, it will not be possible to generically query the type of an
attribute directly. This patch provides an attribute interface
`TypedAttr` that implements only one method, `getType`, which can be
used to generically query the types of attributes that implement the
interface. This interface can be used to retain the concept of a "typed
attribute". The ODS-generated accessor for a `type` parameter
automatically implements this method.
Next steps will be to refactor the assembly formats of certain operations
that rely on `parseAttribute(type)` and `printAttributeWithoutType` to
remove special handling of type elision until `type` can be removed from
the dialect parsing hook entirely; and incrementally remove uses of
`TypedAttr`.
Reviewed By: lattner, rriddle, jpienaar
Differential Revision: https://reviews.llvm.org/D130092
|
|
|
|
ViewLikeInterface.cpp."
This reverts commit e8c2877565149587fd66fbee591b7d44eecd667d.
|
|
Differential Revision: https://reviews.llvm.org/D130706
|
|
|
|
Fix the hardcoded check for `FuncOp` in `getCommonBlock` utility: the
check should have been for an op that starts an affine scope. The
incorrect block returned in turn causes dependence analysis to function
incorrectly.
This change allows affine store-load forwarding to work correctly inside
any ops that start an affine scope.
Reviewed By: ftynse, dcaballe
Differential Revision: https://reviews.llvm.org/D130749
|
|
The implementation and API of GEP Op has gotten a bit convoluted over the time. Issues with it are:
* Misleading naming: `indices` actually only contains the dynamic indices, not all of them. To get the amount of indices you need to query the size of `structIndices`
* Very difficult to iterate over all indices properly: One had to iterate over `structIndices`, check whether it contains the magic constant `kDynamicIndex`, if it does, access the next value in `index` etc.
* Inconvenient to build: One either has create lots of constant ops for every index or have an odd split of passing both a `ValueRange` as well as a `ArrayRef<int32_t>` filled with `kDynamicIndex` at the correct places.
* Implementation doing verification in the build method
and more.
This patch attempts to address all these issues via convenience classes and reworking the way GEP Op works:
* Adds `GEPArg` class which is a sum type of a `int32_t` and `Value` and is used to have a single convenient easy to use `ArrayRef<GEPArg>` in the builders instead of the previous `ValueRange` + `ArrayRef<int32_t>` builders.
* Adds `GEPIndicesAdapter` which is a class used for easy random access and iteration over the indices of a GEP. It is generic and flexible enough to also instead return eg. a corresponding `Attribute` for an operand inside of `fold`.
* Rename `structIndices` to `rawConstantIndices` and `indices` to `dynamicIndices`: `rawConstantIndices` signifies one shouldn't access it directly as it is encoded, and `dynamicIndices` is more accurate and also frees up the `indices` name.
* Add `getIndices` returning a `GEPIndicesAdapter` to easily iterate over the GEP Ops indices.
* Move the verification/asserts out of the build method and into the `verify` method emitting op error messages.
* Add convenient builder methods making use of `GEPArg`.
* Add canonicalizer turning dynamic indices with constant values into constant indices to have a canonical representation.
The only breaking change is for any users building GEPOps that have so far used the old `ValueRange` + `ArrayRef<int32_t>` builder as well as those using the generic syntax.
Another follow up patch then goes through upstream and makes use of the new `ArrayRef<GEPArg>` builder to remove a lot of code building constants for GEP indices.
Differential Revision: https://reviews.llvm.org/D130730
|
|
Adding complex value with 0 for real and imaginary part can be ignored.
NOTE: This type of canonicalization can be written in an easy and tidy format using `complex.number` after constant op supports custom attribute.
Differential Revision: https://reviews.llvm.org/D130748
|
|
A group of functions in the Affine dialect provides a mechanism for
buliding folded-by-construction operations. These functions used to
accept a `RewriterBase` reference because they may need to erase the
operations that were folded and notify the rewriter when called from
rewrite patterns. Adopt a different approach: postpone the builder
notification of the op creation until we are certain that the op will
not be folded away. This removes the need to notify the rewriter about
op deletion following op construction in case of successful folding, and
removes a bunch of one-off `IRRewriter` instances in transform code that
may mess up insertion points.
Reviewed By: springerm, mravishankar
Differential Revision: https://reviews.llvm.org/D130616
|
|
Differential Revision: https://reviews.llvm.org/D130756
|
|
We can canonicalize consecutive complex.conj just by removing all conjugate operations.
Reviewed By: pifon2a
Differential Revision: https://reviews.llvm.org/D130684
|
|
It is more useful to use ComplexType as type of the attribute than to
use the element type as attribute type. This means when using this
attribute in complex::ConstantOp, we just need to check whether
the types match.
Reviewed By: pifon2a
Differential Revision: https://reviews.llvm.org/D130703
|
|
used in payload.
Current implementation of decomposition of Linalg operations wouldnt
work if the `outs` operand values were used within the body of the
operation. Relax this restriction. This potentially sets the stage for
decomposing ops with reduction iterator types (but is not done here
since it requires more study).
Differential Revision: https://reviews.llvm.org/D130527
|
|
the op.
While The tiling interface provides a mechanism for operations to be
tiled into tiled version of the op (or another op at the same level of
abstraction), the `generateScalarImplementation` method added here is
the "exit point" after all transformations have been done. Ops that
implement this method are expected to generate IR that are directly
lowerable to backend dialects like LLVM or SPIR-V dialects.
Differential Revision: https://reviews.llvm.org/D130612
|
|
Differential Revision: https://reviews.llvm.org/D130632
|
|
Handles can be yielded from the ForeachOp.
Differential Revision: https://reviews.llvm.org/D130640
|
|
This commit folds a `tensor.cast` op into a `tensor.collapse_shape` op
when following two conditions meet:
1. the `tensor.collapse_shape` op consumes result of the `tensor.cast` op.
2. `tensor.cast` op casts to a more dynamic version of the source tensor.
This is added as a canonicalization pattern in `tensor.collapse_shape` op.
Signed-Off-By: Gaurav Shukla <gaurav@nod-labs.com>
Reviewed By: mravishankar
Differential Revision: https://reviews.llvm.org/D130650
|
|
* https://discourse.llvm.org/t/rfc-removing-the-quant-dialect/3643/8
* Removes most ops. Leaves casts given final comment (can remove more in a followup).
* There are a few uses in Tosa keeping some of the utilities alive. In a followup, I will probably elect to just move simplified versions of them into Tosa itself vs having this quasi-library dependency.
Differential Revision: https://reviews.llvm.org/D120204
|
|
This commit extends UnifyAliasedResourcePass to handle the case
where aliased resources have different vector sizes. (It still
requires all scalar types to be of the same bitwidth.) This is
effectively reusing the code for handling different-bitwidth
scalar types.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D130671
|
|
This commit fixes spv.CompositeConstruct to assembly to list
operand types to enable vector construction out of smaller vectors.
Validation is also fixed to properly check the cases for vector
construction.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D130669
|
|
Modify an existing test to test the situation.
Reviewed By: Peiming
Differential Revision: https://reviews.llvm.org/D130658
|
|
This patch adds canonicalization conditions for omp.atomic.update thus
eliminating it when it becomes just a write or a no-op due to other
changes during canonicalization.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D126531
|
|
operator.
Sparse compiler failed on the provided test (when the sparse kernel is nested in a scf structrual operator).
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D130609
|
|
Add custom attribute for complex dialect. Although this commit does not have significant impact on the conversion framework, it will lead us to construct complex numbers in a readable and tidy manner.
Related discussion: https://reviews.llvm.org/D127476
Reviewed By: pifon2a, akuegel
Differential Revision: https://reviews.llvm.org/D130149
|
