Commit Graph

19586 Commits

Author SHA1 Message Date
Christian Sigg
a5757c5b65 Switch member calls to isa/dyn_cast/cast/... to free function calls. (#89356)
This change cleans up call sites. Next step is to mark the member
functions deprecated.

See https://mlir.llvm.org/deprecation and
https://discourse.llvm.org/t/preferred-casting-style-going-forward.
2024-04-19 15:58:27 +02:00
Christian Ulmann
090d03d1c7 [MLIR][LLVM] Add flag to skip import of DICompositeType's elems (#89355)
This commit introduces a flag to allow skipping the potentially
recursive import of DICompositeType elements. This patch is essentially a
bandaid for the still broken recursive debug type import.

Some of our downstream inputs are produced by excessive usage of
template meta programming, and thus contain tens of thousands of types
that all participate in such recursions. Unfortunately, the series of
patches that introduces type support is not easily revertible due to
being around for a while now and Modular depending on it.

We can consider to revert this change once the type importer has show to
be very performant, but for now we are talking second vs hours to import
specific files.
2024-04-19 14:03:59 +02:00
Christian Ulmann
df411fbac6 [MLIR][DataLayout] Add support for scalable vectors (#89349)
This commit extends the data layout to support scalable vectors. For
scalable vectors, the `TypeSize`'s scalable field is set accordingly,
and the alignment information remains the same as for normal vectors.
This behavior is in sync with what LLVM's data layout queries are
producing.

Before this change, scalable vectors incorrectly returned the same size
as "normal" vectors.
2024-04-19 13:46:10 +02:00
Tina Jung
95ffa8a2ac [mlir][emitc] Restrict types in EmitC (#88391)
Restrict the types which are valid for EmitC operations. Use what is
currently supported by the emitter as restriction. Define a utility
functions for valid types, such that they can be used to restrict the
operations in table gen as well as being available for reuse in dialect
conversions.
2024-04-19 10:52:31 +02:00
Valentin Clement (バレンタイン クレメン)
13beabe5e9 Revert "[mlir][llvm] Port overflowFlags to a native operation property" (#89344)
Reverts llvm/llvm-project#89312 as it breaks all flang buildbots
2024-04-18 22:15:42 -07:00
Maksim Levental
8c212a3f41 [mlir][python] allow passing in PYTHONPATH to lit tests (#89296) 2024-04-18 18:06:11 -05:00
Will Dietz
8fb7ddf75f [Tosa] TosaToLinalg: link in Index dialect (#89318)
This is needed now that this depends on it:
ee0284ec10

Co-authored-by: Will Dietz <w@wdtz.org>
2024-04-19 00:46:01 +02:00
Jeff Niu
0c41eea474 [mlir][llvm] Port overflowFlags to a native operation property (#89312)
This PR changes the LLVM dialect's IntegerOverflowFlags to be stored on
operations as native properties.
2024-04-18 15:35:39 -07:00
Jeff Niu
2132ebf0fe [mlir] Add def to Properties.td to help with enum properties (#89311)
This is useful for defining operation properties that are enums.
2024-04-18 15:35:31 -07:00
Maksim Levental
6e6da74c8b [mlir][python] add binding to #gpu.object (#88992) 2024-04-18 16:31:55 -05:00
Durgadoss R
b3383d7769 [MLIR][NVVM] Fix a typo in barrier.arrive Op (#89273) 2024-04-18 20:26:23 +02:00
Charitha Saumya
c577f91d26 [mlir][vector] Add support for linearizing Extract, ExtractStridedSlice, Shuffle VectorOps in VectorLinearize (#88204)
This PR adds support for converting `vector.extract_strided_slice` and
`vector.extract` operations to equivalent `vector.shuffle` operations
that operates on linearized (1-D) vectors. `vector.shuffle` operations
operating on n-D (n > 1) are also converted to equivalent shuffle
operations working on linearized vectors.
2024-04-18 21:13:49 +03:00
Dmitriy Smirnov
ee0284ec10 [TOSA] FFT2D/RFFT2D accuracy increased (#88510)
This PR increases accurasy of FFT2D/RFFT2D calculation by removing
periodic part of sin/cos
2024-04-18 17:53:00 +01:00
Bharathi Ramana Joshi
24da7fa029 [MLIR][Presburger] Use Identifiers outside Presburger library (#77316)
The pull request includes the following changes.
1. Refactors the interface to `PresburgerSpace::identifiers` to `setId` and a
const `getId`, instead of previous `getId` which returned a mutable
reference. `resetIds` does not need to be called to use identifiers, `setId`
calls `resetIds` if identifiers are not enabled.
2. Deprecates `FlatAffineRelation` by refactoring all usages of
`FlatAffineRelation` to `IntegerRelation`. To achieve this,
`FlatAffineRelation::compose` is refactored into
`IntegerRelation::mergeAndCompose`.
3. Deletes unneeded overrides of virtual functions `hasConsistentState`,
`clearAndCopyFrom` and `fourierMotzkinEliminate` from
`FlatLinearValueConstraints` as these were only used through
`FlatAffineRelation` and we now use `IntegerRelation`'s member functions
instead.
4. Fixes an existing bug in FlatLinearValueConstraints' constructor
which caused
identifiers set by superclass FlatLinearConstraints' constructor to be
erased.
5. Fixes `IntegerRelation::convertVarKind` not preserving identifiers.
2024-04-18 21:56:53 +05:30
Corentin Ferry
44c876f9e0 [mlir][emitc] Add EmitC lowering for arith.cmpi (#88700) 2024-04-18 16:42:13 +02:00
tomnatan30
bc5536469d [mlir][python] Fix PyOperationBase::walk not catching exception in python callback (#89225)
If the python callback throws an error, the c++ code will throw a
py::error_already_set that needs to be caught and handled in the c++
code .

This change is inspired by the similar solution in
PySymbolTable::walkSymbolTables.
2024-04-18 16:09:31 +02:00
Matthias Gehre
c515c78024 [mlir][Bufferization] castOrReallocMemRefValue: Use BufferizationOptions (#89175)
This allows to configure both the op used for allocation and copy of
memrefs.
It also changes the default behavior because the default allocation in
`BufferizationOptions` creates `memref.alloc` with `alignment = 64`
where we used to create `memref.alloca` without any alignment before.
Fixes
```
// TODO: Use alloc/memcpy callback from BufferizationOptions if called via
// BufferizableOpInterface impl of ToMemrefOp.
```
2024-04-18 15:47:08 +02:00
Sergio Afonso
97c60d61d0 [MLIR][OpenMP] Check that omp.loop_nest represents at least one loop (#89082) 2024-04-18 12:58:58 +01:00
Benjamin Kramer
3f0d52b2db [mlir][llvm] Clean out f8 types from compatible types list
These were added in faf697e49b so things
can flow through non-opaque LLVM ptrs. Those ptrs are gone so there is
no reason for this to be around anymore. LLVM doesn't support f8 types,
they get converted to i8 when lowering to LLVM dialect.

Removing the f8 types makes LLVM::isCompatibleType and
LLVM::isCompatibleFloatingPointType consistent again.
2024-04-18 13:50:19 +02:00
Matthias Springer
2a4e61b342 [mlir][NFC] Move and improve ownership-based buffer dellocation docs (#89196)
Move the documentation of the ownership-based buffer deallocation pass
to a separate file. Also improve the documentation a bit and insert a
figure that explains the `bufferization.dealloc` op (copied from the
tutorial at the LLVM Dev Summit 2023).
2024-04-18 13:26:15 +02:00
Christian Ulmann
ac39fa740b [MLIR][Mem2Reg][LLVM] Enhance partial load support (#89094)
This commit improves LLVM dialect's Mem2Reg interfaces to support
promotions of partial loads from larger memory slots. To support this,
the Mem2Reg interface methods are extended with additional data layout
parameters. The data layout is required to determine type sizes to
produce correct conversion sequences.

Note: There will be additional followups that introduce a similar
functionality for stores, and there are plans to support accesses into
the middle of memory slots.
2024-04-18 13:09:16 +02:00
Guray Ozen
c82f45f9de [mlir][nvgpu] Simplify TMA IR generation (#87153)
This PR add `TmaDescriptorBuilder` 
- class simplifies TMA generation.
- Makes the code ready to support various Tma configurations 
   - removes strings and use the enums from `mlir.nvgpu.ENUMs`.
- Example "swizzle = swizzle_128b, l2promo=none, oob=zero,
interleave=none" to enums in `mlir.nvgpu` dialects.
- Enums have string equivalent that are used during the IR writing and
generation (see `TmaDescriptorBuilder::tensormap_descriptor_ty`).
- Improves readability and abstracts out TMA descriptor builders in
reusable component.

---------

Co-authored-by: Manish Gupta <manigupta@google.com>
2024-04-18 09:58:24 +02:00
Durgadoss R
6e30d97e89 [MLIR][NVVM] [NFC] Update Docs for shfl.sync Op (#89044)
The first argument to the nvvm_shfl_sync_* family
of intrinsics is the thread_mask (aka member_mask). 
This patch renames the corresponding operand in the Op
to reflect the same i.e. `dst` -> `thread_mask`.

While we are there, add summary and description
for this Op.

Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
2024-04-18 09:40:49 +02:00
Dmitrii Agibov
b5aff11aa1 [mlir][tosa] Add folding for TOSA ArgMax operator (#88871) 2024-04-18 07:33:51 +01:00
Jakub Kuderski
3d72c44fed [mlir][gpu] Improve gpu.shuffle documentation. NFC. (#89168)
* Make the wording around lanes / threads / work items more consistent.
* Add examples for all shufle modes.
* Also clean up `gpu.subgroup_reduce`.
2024-04-18 00:49:38 -04:00
Abdul Raheem
693a458287 [MLIR] Update doc comment in ViewLikeInterface.td (NFC) (#89074)
Signed-off: Abdul Raheem Beigh <abdulraheembeigh@gmail.com>
2024-04-17 19:24:04 +02:00
fabrizio-indirli
4edeaffbf2 [mlir][tosa] Fix tosa.Resize-to-linalg lowering (#88514) 2024-04-17 17:43:22 +01:00
Guray Ozen
4f88c23111 [mlir][py] Add NVGPU's TensorMapDescriptorType in py bindings (#88855)
This PR adds NVGPU dialects' TensorMapDescriptorType in the py bindings.

This is a follow-up issue from [this
PR](https://github.com/llvm/llvm-project/pull/87153#discussion_r1546193095)
2024-04-17 15:59:18 +02:00
Oleksandr "Alex" Zinenko
73140daebb [mlir] expose transform dialect symbol merge to python (#87690)
This functionality is available in C++, make it available in Python
directly to operate on transform modules.
2024-04-17 15:01:59 +02:00
Oleksandr "Alex" Zinenko
37b26bf48b [mlir] transform.apply_patterns support more config options (#88484)
Greedy rewrite driver has options to control the number of rewrites
applies. Expose those via the corresponding transform op.
2024-04-17 14:24:51 +02:00
MbjYjbpivj
631c5e818e [mlir] fix intNEQValue summary (#89029)
Fix the summary of intNEQValue.
2024-04-17 14:15:30 +02:00
Sergio Afonso
16b0be6132 [MLIR][OpenMP] NFC: Remove LoopControl parsing/printing code (#88909)
This patch removes the LoopControl parsing/printing functions that are
no longer used after transitioning `omp.simdloop` and `omp.taskloop`
into loop wrapper operations.
2024-04-17 11:30:11 +01:00
Sergio Afonso
3eb0ba34b0 [MLIR][Flang][OpenMP] Make omp.simdloop into a loop wrapper (#87365)
This patch updates the definition of `omp.simdloop` to enforce the
restrictions of a wrapper operation. It has been renamed to `omp.simd`,
to better reflect the naming used in the spec. All uses of "simdloop" in
function names have been updated accordingly.

Some changes to Flang lowering and OpenMP to LLVM IR translation are
introduced to prevent the introduction of compilation/test failures. The
eventual long term solution might be different.
2024-04-17 11:28:30 +01:00
Matthias Springer
9f3334e993 [mlir][SparseTensor] Add missing dependent dialect to pass (#88870)
This commit fixes the following error when stopping the sparse compiler
pipeline after bufferization (e.g., with `test-analysis-only`):

```
LLVM ERROR: Building op `vector.print` but it isn't known in this MLIRContext: the dialect may not be loaded or this operation hasn't been added by the dialect. See also https://mlir.llvm.org/getting_started/Faq/#registered-loaded-dependent-whats-up-with-dialects-management
```
2024-04-17 09:20:55 +02:00
Hideto Ueno
47148832d4 [mlir][python] Add walk method to PyOperationBase (#87962)
This commit adds `walk` method to PyOperationBase that uses a python
object as a callback, e.g. `op.walk(callback)`. Currently callback must
return a walk result explicitly.

We(SiFive) have implemented walk method with python in our internal
python tool for a while. However the overhead of python is expensive and
it didn't scale well for large MLIR files. Just replacing walk with this
version reduced the entire execution time of the tool by 30~40% and
there are a few configs that the tool takes several hours to finish so
this commit significantly improves tool performance.
2024-04-17 15:09:47 +09:00
Kai Sasaki
8c9d814b66 [mlir][complex] Fastmath flag for complex angle (#88658)
See
https://discourse.llvm.org/t/rfc-fastmath-flags-support-in-complex-dialect/71981
2024-04-17 09:19:15 +09:00
Prashant Kumar
ce5381e22a [mlir][vector] Determine vector sizes from the result shape in the ca… (#88249)
…se of tensor pack

When the vector sizes are not passed as inputs to the vector transform
operation, the vector sizes are queried from the static result shape in
the case of tensor.pack op.
2024-04-17 05:36:40 +05:30
Peiming Liu
8aa061ffc7 [mlir][sparse][NFC] switching to using let argments/results in td files (#88994)
followed the same style used in "TensorOps.td".
2024-04-16 15:18:47 -07:00
Jeremy Kun
ed7038ef33 specify dialect in polynomial docs (#88933)
I figured out how to test this with `make mlir-doc doxygen-mlir`
2024-04-16 12:24:19 -07:00
Kazu Hirata
c9731a3dcc [mlir] Fix a warning about an extraneous semicolon
This patch fixes:

  mlir/lib/Dialect/XeGPU/IR/XeGPUOps.cpp:58:2: error: extra ';'
  outside of a function is incompatible with C++98
  [-Werror,-Wc++98-compat-extra-semi]
2024-04-16 12:10:05 -07:00
Chao Chen
b01879ec1f [MLIR][XeGPU] Add XeGPU scattered ops (#86594)
- Extended TensorDescAttr with scattered attribute
- Add scattered ops: CreateDescOp, PrefetchOp, LoadGatherOp,
StoreScatterOp, UpdateOffsetOp
- Add a block op: UpdateNdOffsetOp

---------

Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
Co-authored-by: Adam Siemieniuk <adam.siemieniuk@intel.com>
2024-04-16 13:44:14 -05:00
Peiming Liu
481bd5d416 [mlir][sparse] introduce sparse_tensor.extract_iteration_space operation. (#88554)
A `sparse_tensor.extract_space %tensor at %iterator` extracts a *sparse*
iteration space defined `%tensor`, the operation to traverse the
iteration space will be introduced in following PRs.
2024-04-16 11:32:30 -07:00
Peiming Liu
b9556532c7 Revert "[mlir][sparse] introduce sparse_tensor.iterate operation" (#88953)
Reverts llvm/llvm-project#88807 (merged by mistake)
2024-04-16 11:31:33 -07:00
Peiming Liu
8debcf03c5 [mlir][sparse] introduce sparse_tensor.iterate operation (#88807)
A `sparse_tensor.iterate` iterates over a sparse iteration space
extracted from `sparse_tensor.extract_iteration_space` operation
introduced in https://github.com/llvm/llvm-project/pull/88554.

*DO NOT MERGE* before https://github.com/llvm/llvm-project/pull/88554
2024-04-16 11:31:09 -07:00
Jeremy Kun
5a34ff12b8 fix Polynomial.td doc filename (#88900)
Not sure how best to test this, but I think it fixes the error
https://github.com/llvm/mlir-www/actions/runs/8699908058/job/23859264085#step:7:1111

Co-authored-by: Jeremy Kun <j2kun@users.noreply.github.com>
Co-authored-by: Jacques Pienaar <jpienaar@google.com>
2024-04-16 07:35:36 -07:00
Spenser Bauman
1c076b43c2 [mlir][tosa] Implement dynamic shape support for tosa.max_pool2d lowering (#87538)
The existing lowering for tosa.max_pool2d only supports dynamic
dimensions when the dynamic dimension is the batch dimension. This
change updates the lowering to support arbitrary dynamic dimensions on
the inputs and outputs of the tosa.max_pool2d operation.

This change also fixes a bug in the implementation of implicit
broadcasting in the tosa-to-linalg pass, which was introducing uses of
constant ops that violated dominance requirements.
2024-04-16 08:14:22 -04:00
Frederik Harwath
ac1f2de7b5 [MLIR][docs] Mention declarePromisedInterface in Interfaces doc (#88689)
Co-authored-by: Frederik Harwath <fharwath@amd.com>
Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
2024-04-16 14:09:25 +02:00
Andrzej Warzyński
711df7b0ae [vector][mlir] Restrict vector.shuffle to fixed-width vectors (#88733)
At the moment there is no support for vector.shuffle for scalable
vectors - various hooks/helpers related to `vector.shuffle` simply
ignore the scalable flags (e.g. ` ShuffleOp::inferReturnTypes`).

This is unlikely to change any time soon (vector shuffles are known to
be tricky for scalable vectors), hence this patch restricts
`vector.shuffle` to fixed width vectors.
2024-04-16 12:57:59 +01:00
Benjamin Maxwell
dadcaf8227 [mlir][ArmSME] Support decomposing constant splats into ArmSME tiles (#88762)
This adds a simple rewrite/legalization to decompose constant splats
larger than a single ArmSME tile into multiple SME virtual tile sized
splats. E.g. a constant splat to `vector<[8]x[8]xi32>` would decompose
into four `vector<[4]x[4]xi32>` splats.
2024-04-16 12:54:01 +01:00
Jon Chesterfield
ac6b4c618a Reapply "[Verifier] Reject va_start in non-variadic function (#88809)"
This reverts commit f4960da602.
Includes a fix for the MLIR test case.
2024-04-16 12:26:08 +01:00