| Age | Commit message (Collapse) | Author |
|---|
|
|
|
- Adds verification for `nvgpu.mma.sync` op
- Adds tests to `mlir/test/Dialect/NVGPU/invalid.mlir`
- `nvgpu.mma.sync` verifier caught a bug and triggered a failure in m16n8k4_tf32_f32 variant in `mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir`
- The output shape of vector holding thread-level accumulators was inconsistent and fixed in this change
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D129400
|
|
See: https://gist.github.com/bjacob/d8be8ec7e70ed0be4b3a5794ced2a7e8
Differential Revision: https://reviews.llvm.org/D129096
|
|
|
|
This pass tests patterns that are already tested elsewhere by applying them in a semi-targeted
fashion using anchor function and op names.
From now on, targeted tests should use the transform dialect interpreter.
Differential Revision: https://reviews.llvm.org/D129627
|
|
|
|
Differential Revision: https://reviews.llvm.org/D129552
|
|
Replace some redundant indices checks with the correct checks
Reviewed By: aartbik
Differential Revision: https://reviews.llvm.org/D129101
|
|
since vector.reduce support accumulator in all the cases remove the
assert assuming old definition.
Differential Revision: https://reviews.llvm.org/D129602
|
|
The lowering of fmax/fmin reduce was ignoring the optional accumulator.
Differential Revision: https://reviews.llvm.org/D129597
|
|
Avoids needing the two parallel functions as NamedAttrList already takes care
of caching DictionaryAttr and implicitly can convert from either.
Differential Revision: https://reviews.llvm.org/D129527
|
|
Between issues such as
https://github.com/llvm/llvm-project/issues/56323, the fact that this
lowering (unlike the code in amdgpu-to-rocdl) does not correctly set
up bounds checks (and thus will cause page faults on reads that might
need to be padded instead), and that fixing these problems would,
essentially, involve replicating amdgpu-to-rocdl, remove
--vector-to-rocdl for being broken. In addition, the lowering does not
support many aspects of transfer_{read,write}, like supervectors, and
may not work correctly in their presence.
We (the MLIR-based convolution generator at AMD) do not use this
conversion pass, nor are we aware of any other clients.
Migration strategies:
- Use VectorToLLVM
- If buffer ops are particularly needed in your application, use
amdgpu.raw_buffer_{load,store}
A VectorToAMDGPU pass may be introduced in the future.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D129308
|
|
Existing implementation of structured op splitting creates several
affine.apply and affine.min operations in its subshape computation.
As these shapes are further used in data slice extraction, this may lead
to slice shapes being dynamic even when the original shapes and the
splitting point are static. This is particularly visible when splitting
is combined with further subsetting transformations such as tiling. Use
composition and folding more aggressively in splitting to avoid this.
In particular, introduce a `createComposedAffineMin` function that the
affine map used in "min" with the maps used by any `affine.apply` that
may be feeding the operands to the "min". This enables production of
more static shapes. Also introduce a `createComposedFoldedAffineApply`
function that combines the existing `createComposedAffineApply` with
in-place folding to propagate constants produced by zero-input affine
maps. Using these when splitting allows the subsequent canonicalizer
pass to recover static shapes for structured ops.
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D129379
|
|
This allows vectorizing linalg reductions without changing the operation
order. Therefore this produce a valid vectorization even if operations
are not associative.
Differential Revision: https://reviews.llvm.org/D129535
|
|
The existing implementation of the TilingInterface for Linalg ops was not
modifying the `linalg.index` ops contained within other Linalg ops (they need
to be summed up with the values of respective tile loop induction variables),
which led to the interface-based tiling being incorrect for any Linalg op with
index semantics.
In the process, fix the function performing the index offsetting to use the
pattern rewriter API instead of RAUW as it is being called from patterns and
may mess up the internal state of the rewriter. Also rename the function to
clearly catch all uses.
Depends On D129365
Reviewed By: mravishankar
Differential Revision: https://reviews.llvm.org/D129366
|
|
A recent commit introduced helper functions with semantically meaningful names
to populate the lists of memory effects in transform ops, use them whenever
possible.
Depends On D129287
Reviewed By: springerm
Differential Revision: https://reviews.llvm.org/D129365
|
|
Introduce a structured transform op that emits IR computing the multi-tile
sizes with requested parameters (target size and divisor) for the given
structured op. The sizes may fold to arithmetic constant operations when the
shape is constant. These operations may then be used to call the existing
tiling transformation with a single non-zero dynamic size (i.e. perform
strip-mining) for each of the dimensions separately, thus achieving multi-size
tiling with optional loop interchange. A separate test exercises the entire
script.
Depends On D129217
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D129287
|
|
Extend the definition of the Tile structured transform op to enable it
accepting handles to operations that produce tile sizes at runtime. This is
useful by itself and prepares for more advanced tiling strategies. Note that
the changes are relevant only to the transform dialect, the tiling
transformation itself already supports dynamic sizes.
Depends On D129216
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D129217
|
|
Various typos and formatting fixes that make the generated documentation
hard to follow.
|
|
This handle manipulation operation allows one to define a new handle that is
associated with a the same payload IR operations N times, where N can be driven
by the size of payload IR operation list associated with another handle. This
can be seen as a sort of broadcast that can be used to ensure the lists
associated with two handles have equal numbers of payload IR ops as expected by
many pairwise transform operations.
Introduce an additional "expensive" check that guards against consuming a
handle that is assocaited with the same payload IR operation more than once as
this is likely to lead to double-free or other undesired effects.
Depends On D129110
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D129216
|
|
This patch uses constFoldUnaryOpConditional to replace current folder in order to support constant dense.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D129459
|
|
This is still failing as endianness of binary blob external
resources is still not handled correctly.
|
|
https://reviews.llvm.org/D126446 added support for encoding
binary blobs in MLIR assembly. To enable cross-architecture
compatibility, these need to be encoded in little-endian format.
This patch is a first step in that direction by reading and
writing the alignment tag that those blobs are prefixed by
in little-endian format. This fixes assertion failures in
several test cases on big-endian platforms.
The actual content of the blob is not yet handled here.
Differential Revision: https://reviews.llvm.org/D129483
|
|
dot is currently miscompiled for booleans (uses add instead of or).
Reviewed By: bkramer
Differential Revision: https://reviews.llvm.org/D129292
|
|
This can be enabled by using a `-re` suffix when defining the expected line,
e.g. `expected-error-re`. This support is similar to what clang provides in its "expected"
diagnostic framework(e.g. the `-re` is also the same). The regex definitions themselves are
similar to FileCheck in that regex blocks are specified within `{{` `}}` blocks.
Differential Revision: https://reviews.llvm.org/D129343
|
|
Differential Revision: https://reviews.llvm.org/D129333
|
|
A previous revision implemented expand/collapse reshaping between
dense and sparse tensors for sparse2dense and dense2sparse since those
could use the "cheap" view reshape on the already materialized
dense tensor (at either the input or output side), and do some
reshuffling from or to sparse. The dense2dense case, as always,
is handled with a "cheap" view change.
This revision implements the sparse2sparse cases. Lacking any "view"
support on sparse tensors this operation necessarily has to perform
data reshuffling on both ends.
Tracker for improving this:
https://github.com/llvm/llvm-project/issues/56477
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D129416
|
|
Reviewed By: silvas
Differential Revision: https://reviews.llvm.org/D128880
|
|
This reverts commit 45ef20ca71aaba9ad50c4641fe7fcbb786724af8.
|
|
|
|
Lower complex.log to corresponding function call with libm.
Reviewed By: bixia
Differential Revision: https://reviews.llvm.org/D129417
|
|
This patch adds OMPIRBuilder support for the simdlen clause for the
simd directive. It uses the simdlen support in OpenMPIRBuilder when
it is enabled in Clang. Simdlen is lowered by OpenMPIRBuilder by
generating the loop.vectorize.width metadata.
Reviewed By: jdoerfert, Meinersbur
Differential Revision: https://reviews.llvm.org/D129149
|
|
Tosa to Linalg conversion crashes when input tensor is a float type other than fp32.
Because tosa.clamp and tosa.reluN have fp32 min/max attribute which is converted as arith.constant with the attribute type.
This commit fixes the crash by correctly setting the float constant type from the input tensor.
Reviewed By: eric-k256
Differential Revision: https://reviews.llvm.org/D128630
|
|
Distribute splat constant out of WarpExecuteOnLane0Op region.
Differential Revision: https://reviews.llvm.org/D129467
|
|
Prevent creating multiple output for the same Value when distributing
operations out of WarpExecuteOnLane0Op. This avoid creating combinatory
explosion of outputs.
Differential Revision: https://reviews.llvm.org/D129465
|
|
Presburger
This uses an int64_t-based fastpath for the common case and falls back to
SlowMPInt to handle the rare cases where larger numbers occur.
It uses `__builtin_*` for performance through the support in LLVM MathExtras.
Using this in the Presburger library results in a minor performance
*improvement* over any commit hash before sequence of patches
starting at d5e31cf38adfc2c240fb9717989792537cc9e819.
This was previously reverted in 1e10d35ea9c02e9b5694836fd3dcc0b9baf28b48 due
to a build failure; relanding now with an attempted fix.
Reviewed By: Groverkss, ftynse
Differential Revision: https://reviews.llvm.org/D128811
|
|
precision in Presburger"
This reverts commit c9035df2fad4da9ea75b9211b3a9b0a230925000.
Reverting due to build failure on Windows: https://lab.llvm.org/buildbot/#/builders/172/builds/14767
|
|
Presburger
This uses an int64_t-based fastpath for the common case and falls back to
SlowMPInt to handle the rare cases where larger numbers occur.
It uses `__builtin_*` for performance through the support in LLVM MathExtras.
Using this in the Presburger library results in a minor performance
*improvement* over any commit hash before sequence of patches
starting at d5e31cf38adfc2c240fb9717989792537cc9e819.
Reviewed By: Groverkss, ftynse
Differential Revision: https://reviews.llvm.org/D128811
|
|
Another mechanical sweep to keep diff small for flip to _Prefixed.
|
|
This patch is similar to D129108, it adds a conditional unary constant folder which allow to exit when the constants not meet the fold condition. And use it for Log2Op to make it able to fold the constant dense.
Differential Revision: https://reviews.llvm.org/D129251
|
|
In D128230, we accidentally moved the install for Python sources outside of the loop, having one install() per group of files. While it would be nice if we could do this, it means that we flatten the relative directory tree and every source ends up in the root. The right way to do this is to use FILE_SETS, which preserve the relative directory tree, but they are not available until CMake 3.23.
Differential Revision: https://reviews.llvm.org/D129434
|
|
Support distributing reductions with vector size multiple of the warp
size.
Differential Revision: https://reviews.llvm.org/D129387
|
|
This is a partial revert of D128615.
to_memref(to_tensor(x)) always be folded to x. But to_tensor(to_memref(x)) cannot be folded in the general case because writes to the intermediary memref may go unnoticed.
Differential Revision: https://reviews.llvm.org/D129354
|
|
These can result in accidentally accepting a completion when it isn't intended.
|
|
|
|
This required changing a bit of how attributes/types are parsed. A new
`KeywordSwitch` class was added to AsmParser that provides a StringSwitch
like API for parsing keywords with a set of potential matches. It intends to
both provide a cleaner API, and enable injection for code completion. This
required changing the API of `generated(Attr|Type)Parser` to handle the
parsing of the keyword, instead of having the user do it. Most upstream
dialects use the autogenerated handling and didn't require a direct update.
Differential Revision: https://reviews.llvm.org/D129267
|
|
This commit adds code completion results to the MLIR LSP when
parsing keywords. Keyword support is currently limited to the
case where the expected keyword is provided, but a followup will
work on expanding the set of keyword cases we handle (e.g. to
allow capturing attribute/type mnemonics).
Differential Revision: https://reviews.llvm.org/D129184
|
|
Previously default attributes were only usable by way of the ODS generated
accessors, but this was undesirable as
1. The ODS getters could construct Attribute each get request;
2. For non-C++ uses this would require either duplicating some of tee default
attribute generating or generating additional bindings to generate methods;
3. Accessing op.getAttr("foo") and op.getFoo() would return different results;
Generate method to populate default attributes that can be used to address
these.
This merely adds this facility but does not employ by default on any path.
Differential Revision: https://reviews.llvm.org/D128962
|
|
Add support for colorizing angle brackets "<>" in TableGen files.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D128229
|
|
Restricted to just 1 dynamic input dim as that worked all the way through to
codegen.
Differential Revision: https://reviews.llvm.org/D129334
|
