Commit Graph

6421 Commits

Author SHA1 Message Date
Laszlo Kindrat ac14d5a1db [mlir] Enable perfect forwarding in RewritePatternSet::add
This patch modifies the implementation of `RewritePatternSet::add` to perfectly forward its arguments to pattern constructors. Without this, code like the following compiles but, due to the limited lifetime of the temporary TypeConverter, can produce unexpected behavior:
```
RewritePatternSet patterns(context);
patterns.add<SomeOpConversion, OtherOpConversion>(TypeConverter(), context);

if (failed(applyPartialConversion(getOperation(), target, std::move(patterns))))
  return signalPassFailure();
```

The patch also changes the linalg fusion pattern implementation to correctly fill the test pattern set given the new behavior.

Author: Laszlo Kindrat <laszlokindrat@gmail.com>

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D129601
2022-07-15 19:44:18 +02:00
lorenzo chelini 12929c241a Revert "[mlir] Enable perfect forwarding in RewritePatternSet::add"
Did not preserve author information.

This reverts commit b0afda78f0.
2022-07-15 19:16:09 +02:00
lorenzo chelini b0afda78f0 [mlir] Enable perfect forwarding in RewritePatternSet::add
This patch modifies the implementation of `RewritePatternSet::add` to perfectly forward its arguments to pattern constructors. Without this, code like the following compiles but, due to the limited lifetime of the temporary TypeConverter, can produce unexpected behavior:
```
RewritePatternSet patterns(context);
patterns.add<SomeOpConversion, OtherOpConversion>(TypeConverter(), context);

if (failed(applyPartialConversion(getOperation(), target, std::move(patterns))))
  return signalPassFailure();
```

The patch also changes the linalg fusion pattern implementation to correctly fill the test pattern set given the new behavior.

Author: Laszlo Kindrat <laszlokindrat@gmail.com>

Reviewed By: nicolasvasilache

Differential Revision: https://reviews.llvm.org/D129601
2022-07-15 19:08:23 +02:00
Arjun P 821fe9efa4 [MLIR][Presburger] reintroduce int64_t versions of floorDiv, ceilDiv in mlir::presburger namespace
This is useful because MPInt.h defines identically-named functions that
operate on MPInts, which would otherwie become the only candidates of
overload resolution when calling e.g. ceilDiv from the mlir::presburger
namespace (iff MPInt.h is included). So to access the 64-bit overloads, an
explict call to mlir::ceilDiv would be required. This patch adds `using`
declarations allowing overload resolution to transparently call the right
function.

Reviewed By: Groverkss

Differential Revision: https://reviews.llvm.org/D129820
2022-07-15 17:29:31 +01:00
Arjun P 9390b8d34c [MLIR][Presburger] MPInt: add missing fastpath in ceilDiv
This is not a bug in functionality, just a missed optimization.

Reviewed By: Groverkss

Differential Revision: https://reviews.llvm.org/D129815
2022-07-15 17:26:37 +01:00
Slava Zakharin 451e5e2b28 [mlir][math] Added math::tan operation.
Differential Revision: https://reviews.llvm.org/D129539
2022-07-15 09:17:06 -07:00
Arjun P 86d73c11cf [MLIR][Presburger] SlowMPInt: gcd: assert that operands are non-negative 2022-07-15 15:45:53 +01:00
Arjun P ec5f18e38a [MLIR][Presburger] MPInt: gcd: assert that operands are non-negative 2022-07-15 14:13:41 +01:00
Fangrui Song 3c849d0aef Modernize Optional::{getValueOr,hasValue} 2022-07-15 01:20:39 -07:00
Arjun P d6fbe1394f [MLIR][Presburger] MPInt: use /// for top-level comment, not // (NFC) 2022-07-15 00:30:02 +01:00
Manish Gupta 713d3de5fb [mlir][NVGPU] Verifier for nvgpu.ldmatrix
* Adds verifiers for `nvgpu.ldmatrix` op
* Adds tests to `mlir/test/Dialect/NVGPU/invalid.mlir`

Reviewed By: ThomasRaoux

Differential Revision: https://reviews.llvm.org/D129669
2022-07-14 22:46:38 +00:00
Krzysztof Drewniak bc61cc9a2d [mlir][AMDGPU] Add lds_barrier op
The lds_barrier op allows workgroups to wait at a barrier for
operations to/from their local data store (LDS) to complete without
incurring the performance penalties of a full memory fence.

Reviewed By: nirvedhmeshram

Differential Revision: https://reviews.llvm.org/D129522
2022-07-14 20:45:26 +00:00
Jeff Niu b7f93c2809 [mlir] (NFC) run clang-format on all files 2022-07-14 13:32:13 -07:00
Matthias Springer 74902cc96f [mlir][linalg][NFC] Cleanup: Drop linalg.inplaceable attribute
bufferization.writable is used in most cases instead. All remaining test cases are updated. Some code that is no longer needed is deleted.

Differential Revision: https://reviews.llvm.org/D129739
2022-07-14 15:50:03 +02:00
Nicolas Vasilache 5a0011360c [mlir][Linalg] Retire LinalgPromotion pattern
This revision removes the LinalgPromotion pattern and adds a `transform.structured.promotion` op.
Since the LinalgPromotion transform allows the injection of arbitrary C++ via lambdas, the current
transform op does not handle it.
It is left for future work to decide what the right transform op control is for those cases.

Note the underlying implementation remains unchanged and the mechanism is still controllable by
lambdas from the API.

During this refactoring it was also determined that the `dynamicBuffers` option does not actually
connect to a change of behavior in the algorithm.
This also exhibits that the related test is wrong (and dangerous).
Both the option and the test are therefore removed.

Lastly, a test that connects patterns using the filter-based mechanism is removed: all the independent
pieces are already tested separately.

Context: https://discourse.llvm.org/t/psa-retire-linalg-filter-based-patterns/63785

Differential Revision: https://reviews.llvm.org/D129649
2022-07-14 05:29:27 -07:00
Matthias Springer c66303c287 [mlir][sparse] Switch to One-Shot Bufferize
This change removes the partial bufferization passes from the sparse compilation pipeline and replaces them with One-Shot Bufferize. One-Shot Analysis (and TensorCopyInsertion) is used to resolve all out-of-place bufferizations, dense and sparse. Dense ops are then bufferized with BufferizableOpInterface. Sparse ops are still bufferized in the Sparsification pass.

Details:
* Dense allocations are automatically deallocated, unless they are yielded from a block. (In that case the alloc would leak.) All test cases are modified accordingly. E.g., some funcs now have an "out" tensor argument that is returned from the function. (That way, the allocation happens at the call site.)
* Sparse allocations are *not* automatically deallocated. They must be "released" manually. (No change, this will be addressed in a future change.)
* Sparse tensor copies are not supported yet. (Future change)
* Sparsification no longer has to consider inplacability. If necessary, allocations and/or copies are inserted during TensorCopyInsertion. All tensors are inplaceable by the time Sparsification is running. Instead of marking a tensor as "not inplaceable", it can be marked as "not writable", which will trigger an allocation and/or copy during TensorCopyInsertion.

Differential Revision: https://reviews.llvm.org/D129356
2022-07-14 09:52:48 +02:00
Kazu Hirata c27d815249 [mlir] Use value instead of getValue (NFC) 2022-07-14 00:19:59 -07:00
Manish Gupta f7d42d5149 [mlir][NVGPU] Verifiers for nvgpu.mma.sync Op
- 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
2022-07-13 18:57:07 +00:00
Benoit Jacob 6870a50f43 lowerParallel is also called on unit-size, one-sided reduction dims
See: https://gist.github.com/bjacob/d8be8ec7e70ed0be4b3a5794ced2a7e8

Differential Revision: https://reviews.llvm.org/D129096
2022-07-13 16:21:12 +00:00
Jacques Pienaar 69b6454fff [mlir] Plumb through default attribute populate for extensible dialect. 2022-07-13 09:05:04 -07:00
Kazu Hirata 491d27013d [mlir] Use has_value instead of hasValue (NFC) 2022-07-13 00:57:02 -07:00
Thomas Raoux 5f8cefebd9 [mlir][vector] Fix crash in vector.reduction canonicalization
since vector.reduce support accumulator in all the cases remove the
assert assuming old definition.

Differential Revision: https://reviews.llvm.org/D129602
2022-07-12 23:15:30 +00:00
Jacques Pienaar 0db084d4c7 [mlir] Switch create to use NamedAttrList&&
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
2022-07-12 13:24:09 -07:00
Krzysztof Drewniak d6ef3d20b4 [mlir] Remove VectorToROCDL
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
2022-07-12 15:21:22 +00:00
Alex Zinenko a5c802a429 [mlir] fold more eagerly in structured op splitting
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
2022-07-12 15:06:55 +00:00
Thomas Raoux 051b36ba28 [mlir][vector] Add accumulator operand to MultiDimReduce op
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
2022-07-12 14:28:30 +00:00
Alex Zinenko 81b62f7feb [mlir] Handle linalg.index correctly in TilingInterface
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
2022-07-12 12:36:33 +00:00
Alex Zinenko e15b855e09 [mlir] Use semantically readable functions for transform op effects
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
2022-07-12 12:36:31 +00:00
Alex Zinenko 3963b4d0dc [mlir] Transform op for multitile size generation
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
2022-07-12 12:36:28 +00:00
Alex Zinenko 4e4a4c0576 [mlir] Allow Tile transform op to take dynamic sizes
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
2022-07-12 12:21:54 +00:00
Alex Zinenko 80e17355cd [mlir] assorted fixes in transform dialect documentation
Various typos and formatting fixes that make the generated documentation
hard to follow.
2022-07-12 09:18:51 +00:00
Alex Zinenko 00d1a1a25f [mlir] Add ReplicateOp to the Transform dialect
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
2022-07-12 09:07:59 +00:00
Christopher Bate 609c0e1b9b [mlir] Register linalg external TilingInterface models in InitAllDialects
Differential Revision: https://reviews.llvm.org/D129333
2022-07-11 15:54:37 -06:00
George Petterson 4dc8cf3a86 Fix an issue with grouped conv2d op
Reviewed By: silvas

Differential Revision: https://reviews.llvm.org/D128880
2022-07-11 20:10:58 +00:00
Nirvedh f0cd538985 Revert "Fix an issue with grouped conv2d op"
This reverts commit 45ef20ca71.
2022-07-11 20:03:16 +00:00
George Petterson 45ef20ca71 Fix an issue with grouped conv2d op 2022-07-11 19:59:30 +00:00
Arjun P cdbc5f1e10 [MLIR][Presburger] introduce MPInt to support fast arbitrary precision in 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 d5e31cf38a.

This was previously reverted in 1e10d35ea9 due
to a build failure; relanding now with an attempted fix.

Reviewed By: Groverkss, ftynse

Differential Revision: https://reviews.llvm.org/D128811
2022-07-11 15:46:44 +01:00
Arjun P 1e10d35ea9 Revert "[MLIR][Presburger] introduce MPInt to support fast arbitrary precision in Presburger"
This reverts commit c9035df2fa.
Reverting due to build failure on Windows: https://lab.llvm.org/buildbot/#/builders/172/builds/14767
2022-07-11 14:26:09 +01:00
Arjun P c9035df2fa [MLIR][Presburger] introduce MPInt to support fast arbitrary precision in 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 d5e31cf38a.

Reviewed By: Groverkss, ftynse

Differential Revision: https://reviews.llvm.org/D128811
2022-07-11 11:41:29 +01:00
jacquesguan ad4b7fb3ce [mlir][Math] Support fold Log2Op with constant dense.
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
2022-07-11 10:34:28 +08:00
River Riddle fe4f512be7 [mlir:LSP] Add support for code completing attributes and types
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
2022-07-08 16:24:55 -07:00
River Riddle 2e41ea3247 [mlir:LSP] Add support for keyword code completions
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
2022-07-08 16:24:55 -07:00
Jacques Pienaar 82140ad728 [mlir] Add method to populate default attributes
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
2022-07-08 11:31:13 -07:00
Nicolas Vasilache 69c8319e76 [mlir][Transform] Fix isDefiniteFailure helper
This newly added helper was returning definiteFailure even in the case of silenceableFailure.

Differential Revision: https://reviews.llvm.org/D129347
2022-07-08 00:39:42 -07:00
Mogball c20a581a8d [mlir] Delete ForwardDataFlowAnalysis
With SCCP and integer range analysis ported to the new framework, this old framework is redundant. Delete it.

Depends on D128866

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D128867
2022-07-07 21:08:27 -07:00
Mogball ab701975e7 [mlir] Swap integer range inference to the new framework
Integer range inference has been swapped to the new framework. The integer value range lattices automatically updates the corresponding constant value on update.

Depends on D127173

Reviewed By: krzysz00, rriddle

Differential Revision: https://reviews.llvm.org/D128866
2022-07-07 20:28:13 -07:00
Mogball d80c271c8a [mlir] An implementation of dense data-flow analysis
This patch introduces an implementation of dense data-flow analysis. Dense
data-flow analysis attaches a lattice before and after the execution of every
operation. The lattice state is propagated across operations by a user-defined
transfer function. The state is joined across control-flow and callgraph edges.

Thge patch provides an example pass that uses both a dense and a sparse analysis
together.

Depends on D127139

Reviewed By: rriddle, phisiart

Differential Revision: https://reviews.llvm.org/D127173
2022-07-07 15:12:46 -07:00
Krzysztof Drewniak db590549a9 [mlir][AMDGPU] Use the correct values for OOB_SELECT on gfx10
Differential Revision: https://reviews.llvm.org/D129320
2022-07-07 21:23:38 +00:00
River Riddle ed2fb1736a [mlir:LSP] Add support for MLIR code completions
This commit adds code completion results to the MLIR LSP using
a new code completion context in the MLIR parser. This commit
adds initial completion for dialect, operation, SSA value, and
block names.

Differential Revision: https://reviews.llvm.org/D129183
2022-07-07 13:35:54 -07:00
Robert Suderman b9e642afd1 [mlir][spirv] Add path for math.round to spirv for OCL and GLSL
OpenCL's round function matches `math.round` so we can directly lower to
the op, this includes adding the op definition to the SPIRV OCL ops.
GLSL does not guarantee rounding direction so we include custom rounding
code to guarantee correct rounding direction.

Reviewed By: antiagainst

Differential Revision: https://reviews.llvm.org/D129236
2022-07-07 19:20:20 +00:00