summaryrefslogtreecommitdiff
path: root/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
AgeCommit message (Collapse)Author
2025-10-04[mlir][vector] Simplify rewrite pattern inheriting constructors. NFC. (#161966)Jakub Kuderski
Use the `Base` type alias from https://github.com/llvm/llvm-project/pull/158433.
2025-09-28[MLIR] Apply clang-tidy fixes for bugprone-argument-comment in ↵Mehdi Amini
VectorToGPU.cpp (NFC)
2025-09-17[mlir][gpu][vector] Lower Vector dialect to GPU for element-wise ops only ↵Hsiangkai Wang
(#159091) Current convertVectorToMMAOps starts from vector.contract and finds its dependencies as the targets to convert. In GPU dialect, we have gpu.subgroup_mma_elementwise operation. We should be able to lower element-wise operations to GPU MMA operations without vector.contract. This patch adds this case to the pattern.
2025-09-15[mlir][vector] Use `source` as the source argument name (#158258)Andrzej Warzyński
This patch updates the following ops to use `source` (instead of `vector`) as the name for their source argument: * `vector.extract` * `vector.scalable.extract` * `vector.extract_strided_slice` This change ensures naming consistency with the "builders" for these Ops that already use the name `source` rather than `vector`. It also addresses part of: * https://github.com/llvm/llvm-project/issues/131602 Specifically, it ensures that we use `source` and `dest` for read and write operations, respectively (as opposed to `vector` and `dest`).
2025-07-30[MLIR] Migrate some conversion passes and dialects to LDBG() macro (NFC) ↵Mehdi Amini
(#151349)
2025-07-25[mlir][NFC] update `mlir` create APIs (34/n) (#150660)Maksim Levental
See https://github.com/llvm/llvm-project/pull/147168 for more info.
2025-07-22[mlir][NFC] update `Conversion` create APIs (7/n) (#149889)Maksim Levental
See https://github.com/llvm/llvm-project/pull/147168 for more info.
2025-07-21[mlir][vector] Support direct broadcast conversion (LLVM & SPIRV) (#148027)James Newling
Add conversion for broadcast from scalar for LLVM and SPIRV. Also some miscellaneous replacements of vector.splat with vector.broadcast in VectorToGPU and ArithToAMDGPU. Part of deprecation of vector.splat RFC: https://discourse.llvm.org/t/rfc-mlir-vector-deprecate-then-remove-vector-splat/87143/4
2025-07-04[mlir] Remove unused includes (NFC) (#147101)Kazu Hirata
These are identified by misc-include-cleaner. I've filtered out those that break builds. Also, I'm staying away from llvm-config.h, config.h, and Compiler.h, which likely cause platform- or compiler-specific build failures.
2025-05-22[mlir] Fix unused-variable warningsKazu Hirata
This patch fixes warnings of the form: mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:320:19: error: unused variable 'result' [-Werror,-Wunused-variable]
2025-05-22[MLIR] Change getBackwardSlice to return a logicalresult rather than crash ↵William Moses
(#140961) The current implementation of getBackwardSlice will crash if an operation in the dependency chain is defined by an operation with multiple regions or blocks. Crashing is bad (and forbids many analyses from using getBackwardSlice, as well as causing existing users of getBackwardSlice to fail for IR with this property). This PR instead causes the analysis to return a failure, rather than crash in the cases it cannot compute the full slice --------- Co-authored-by: Oleksandr "Alex" Zinenko <git@ozinenko.com>
2025-05-12[mlir][vector] Standardize `base` Naming Across Vector Ops (NFC) (#137859)Andrzej Warzyński
[mlir][vector] Standardize base Naming Across Vector Ops (NFC) This change standardizes the naming convention for the argument representing the value to read from or write to in Vector ops that interface with Tensors or MemRefs. Specifically, it ensures that all such ops use the name `base` (i.e., the base address or location to which offsets are applied). Updated operations: * `vector.transfer_read`, * `vector.transfer_write`. For reference, these ops already use `base`: * `vector.load`, `vector.store`, `vector.scatter`, `vector.gather`, `vector.expandload`, `vector.compressstore`, `vector.maskedstore`, `vector.maskedload`. This is a non-functional change (NFC) and does not alter the semantics of these operations. However, it does require users of the XFer ops to switch from `op.getSource()` to `op.getBase()`. To ease the transition, this PR temporarily adds a `getSource()` interface method for compatibility. This is intended for downstream use only and should not be relied on upstream. The method will be removed prior to the LLVM 21 release. Implements #131602
2025-05-06[mlir] Remove unused local variables (NFC) (#138642)Kazu Hirata
2025-03-20[mlir] Use *Set::insert_range (NFC) (#132326)Kazu Hirata
DenseSet, SmallPtrSet, SmallSet, SetVector, and StringSet recently gained C++23-style insert_range. This patch replaces: Dest.insert(Src.begin(), Src.end()); with: Dest.insert_range(Src); This patch does not touch custom begin like succ_begin for now.
2025-01-21[mlir][IR][NFC] Move free-standing functions to `MemRefType` (#123465)Matthias Springer
Turn free-standing `MemRefType`-related helper functions in `BuiltinTypes.h` into member functions.
2024-12-20[mlir] Enable decoupling two kinds of greedy behavior. (#104649)Jacques Pienaar
The greedy rewriter is used in many different flows and it has a lot of convenience (work list management, debugging actions, tracing, etc). But it combines two kinds of greedy behavior 1) how ops are matched, 2) folding wherever it can. These are independent forms of greedy and leads to inefficiency. E.g., cases where one need to create different phases in lowering and is required to applying patterns in specific order split across different passes. Using the driver one ends up needlessly retrying folding/having multiple rounds of folding attempts, where one final run would have sufficed. Of course folks can locally avoid this behavior by just building their own, but this is also a common requested feature that folks keep on working around locally in suboptimal ways. For downstream users, there should be no behavioral change. Updating from the deprecated should just be a find and replace (e.g., `find ./ -type f -exec sed -i 's|applyPatternsAndFoldGreedily|applyPatternsGreedily|g' {} \;` variety) as the API arguments hasn't changed between the two.
2024-10-24[mlir][vector] Fix a crash in `VectorToGPU` (#113454)Longsheng Mou
This PR fixes a crash in `VectorToGPU` when the operand of `extOp` is a function argument, which cannot be retrieved using `getDefiningOp`. Fixes #107967.
2024-08-04[mlir] Construct SmallVector with ArrayRef (NFC) (#101896)Kazu Hirata
2024-07-02mlir/LogicalResult: move into llvm (#97309)Ramkumar Ramachandra
This patch is part of a project to move the Presburger library into LLVM.
2024-06-17[mlir][GPU] Improve handling of GPU bounds (#95166)Krzysztof Drewniak
This change reworks how range information for GPU dispatch IDs (block IDs, thread IDs, and so on) is handled. 1. `known_block_size` and `known_grid_size` become inherent attributes of GPU functions. This makes them less clunky to work with. As a consequence, the `gpu.func` lowering patterns now only look at the inherent attributes when setting target-specific attributes on the `llvm.func` that they lower to. 2. At the same time, `gpu.known_block_size` and `gpu.known_grid_size` are made official dialect-level discardable attributes which can be placed on arbitrary functions. This allows for progressive lowerings (without this, a lowering for `gpu.thread_id` couldn't know about the bounds if it had already been moved from a `gpu.func` to an `llvm.func`) and allows for range information to be provided even when `gpu.*_{id,dim}` are being used outside of a `gpu.func` context. 3. All of these index operations have gained an optional `upper_bound` attribute, allowing for an alternate mode of operation where the bounds are specified locally and not inherited from the operation's context. These also allow handling of cases where the precise launch sizes aren't known, but can be bounded more precisely than the maximum of what any platform's API allows. (I'd like to thank @benvanik for pointing out that this could be useful.) When inferring bounds (either for range inference or for setting `range` during lowering) these sources of information are consulted in order of specificity (`upper_bound` > inherent attribute > discardable attribute, except that dimension sizes check for `known_*_bounds` to see if they can be constant-folded before checking their `upper_bound`). This patch also updates the documentation about the bounds and inference behavior to clarify what these attributes do when set and the consequences of setting them up incorrectly. --------- Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
2024-05-22[MLIR][Analysis] Consolidate topological sort utilities (#92563)Christian Ulmann
This PR attempts to consolidate the different topological sort utilities into one place. It adds them to the analysis folder because the `SliceAnalysis` uses some of these. There are now two different sorting strategies: 1. Sort only according to SSA use-def chains 2. Sort while taking regions into account. This requires a much more elaborate traversal and cannot be applied on graph regions that easily. This additionally reimplements the region aware topological sorting because the previous implementation had an exponential space complexity. I'm open to suggestions on how to combine this further or how to fuse the test passes.
2024-05-13[mlir][gpu] Support extf before contract when converting to MMA ops (#91988)Lei Zhang
This commit allows `inferFragType` to see through all arith.ext op and other elementwise users before reaching contract op for figuring out the fragment type.
2024-05-13[NFC] Make NVGPU casing consistent (#91903)tyb0807
2024-04-19Switch member calls to `isa/dyn_cast/cast/...` to free function calls. (#89356)Christian Sigg
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-01[mlir][NFC] Simplify type checks with isa predicates (#87183)Jakub Kuderski
For more context on isa predicates, see: https://github.com/llvm/llvm-project/pull/83753.
2024-02-08[MLIR] Fix crash in AffineMap::replace for zero result maps (#80930)Uday Bondhugula
Fix obvious bug in AffineMap::replace for the case of zero result maps. Extend/complete inferExprsFromList to work with empty expression lists.
2024-01-21[mlir][IR] Add `notifyBlockRemoved` callback to listener (#78306)Matthias Springer
There is already a "block inserted" notification (in `OpBuilder::Listener`), so there should also be a "block removed" notification. The purpose of this change is to make the listener API more mature. There is currently a gap between what kind of IR changes can be made and what IR changes can be listened to. At the moment, the only way to inform listeners about "block removal" is to send a manual `notifyOperationModified` for the parent op (e.g., by wrapping the `eraseBlock(b)` method call in `updateRootInPlace(b->getParentOp())`). This tells the listener that *something* has changed, but it is somewhat of an API abuse.
2024-01-15Apply clang-tidy fixes for readability-simplify-boolean-expr in ↵Mehdi Amini
VectorToGPU.cpp (NFC)
2023-12-07Apply clang-tidy fixes for performance-unnecessary-value-param in ↵Mehdi Amini
VectorToGPU.cpp (NFC)
2023-12-07Apply clang-tidy fixes for llvm-qualified-auto in VectorToGPU.cpp (NFC)Mehdi Amini
2023-12-07[mlir] Extend CombineTransferReadOpTranspose pattern to handle extf ops. ↵harsh-nod
(#74754) This patch modifies the CombineTransferReadOpTranspose pattern to handle extf ops. Also adds a test which shows the transpose getting folded into the transfer_read.
2023-11-20[mlir][vector] Modernize `vector.transpose` op (#72594)Matthias Springer
* Declare arguments/results with `let` statements. * Rename `transp` to `permutation`. * Change type of `transp` from `I64ArrayAttr` to `DenseI64ArrayAttr` (provides direct access to `ArrayRef<int64_t>` instead of `ArrayAttr`).
2023-11-14[mlir][affine][nfc] cleanup deprecated T.cast style functions (#71269)long.chen
detail see the docment: https://mlir.llvm.org/deprecation/ Not all changes are made manually, most of them are made through a clang tool I wrote https://github.com/lipracer/cpp-refactor.
2023-09-19[mlir][Interfaces] `LoopLikeOpInterface`: Support ops with multiple regions ↵Matthias Springer
(#66754) This commit implements `LoopLikeOpInterface` on `scf.while`. This enables LICM (and potentially other transforms) on `scf.while`. `LoopLikeOpInterface::getLoopBody()` is renamed to `getLoopRegions` and can now return multiple regions. Also fix a bug in the default implementation of `LoopLikeOpInterface::isDefinedOutsideOfLoop()`, which returned "false" for some values that are defined outside of the loop (in a nested op, in such a way that the value does not dominate the loop). This interface is currently only used for LICM and there is no way to trigger this bug, so no test is added.
2023-09-18[mlir][SCF] scf.for: Consistent API around `initArgs` (#66512)Matthias Springer
* Always use the auto-generated `getInitArgs` function. Remove the hand-written `getInitOperands` duplicate. * Remove `hasIterOperands` and `getNumIterOperands`. The names were inconsistent because the "arg" is called `initArgs` in TableGen. Use `getInitArgs().size()` instead. * Fix verification around ops with no results.
2023-09-14[mlir][VectorToGPU] Update memref stride preconditions on `nvgpu.mma.sync` pathChristopher Bate
This change removes the requirement that the row stride be statically known when converting `vector.transfer_read` and `vector.transfer_write` to distributed SIMT operations in the `nvgpu` lowering path. It also adds a check to verify that the last dimension of the source memref is statically known to have stride 1 since this is assumed in the conversion logic. No other change should be required since the generated `vector.load` operations are never created across dimensions other than the last. The routines for checking preconditions on `vector.transfer_read/write` are moved to under nvgpu utilities. The change is NFC with respect to the GPU dialect lowering path. Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D155753
2023-09-11[mlir][arith] Rename operations: `maxf` → `maximumf`, `minf` → ↵Daniil Dudkin
`minimumf` (#65800) This patch is part of a larger initiative aimed at fixing floating-point `max` and `min` operations in MLIR: https://discourse.llvm.org/t/rfc-fix-floating-point-max-and-min-operations-in-mlir/72671. This commit addresses Task 1.2 of the mentioned RFC. By renaming these operations, we align their names with LLVM intrinsics that have corresponding semantics.
2023-08-01[mlir][gpu] Support arith.extf in subgroup MMA elementwise opsLei Zhang
This commit adds support for arith.extf in the supported list of elementwise ops for subgroup MMA ops, and enables lowering to SPIR-V. Reviewed By: mravishankar Differential Revision: https://reviews.llvm.org/D156847
2023-07-31[mlir][vector] Use DenseI64ArrayAttr for ExtractOp/InsertOp positionsMatthias Springer
`DenseI64ArrayAttr` provides a better API than `I64ArrayAttr`. E.g., accessors returning `ArrayRef<int64_t>` (instead of `ArrayAttr`) are generated. Differential Revision: https://reviews.llvm.org/D156684
2023-06-22[mlir][affine] More efficient `makeComposedFolded...` helpersMatthias Springer
The old code used to materialize constants as ops, immediately folded them into the resulting affine map and then deleted the constant ops again. Instead, directly fold the attributes into the affine map. Furthermore, all helpers accept `OpFoldResult` instead of `Value` now. This makes the code at call sites more efficient, because it is no longer necessary to materialize a `Value`, just to be able to use these helper functions. Note: The API has changed (accepts OpFoldResult instead of Value), otherwise this change is NFC. Differential Revision: https://reviews.llvm.org/D153324
2023-06-08[mlir][SliceAnalysis] Add an options object to forward and backward slice.Mahesh Ravishankar
Add an options object to allow control of the slice computation (for both forward and backward slice). This makes the ABI stable, and also allows avoiding an assert that makes the slice analysis unusable for operations with multiple blocks. Reviewed By: hanchung, nicolasvasilache Differential Revision: https://reviews.llvm.org/D151520
2023-05-12[mlir] Move casting calls from methods to function callsTres Popp
The MLIR classes Type/Attribute/Operation/Op/Value support cast/dyn_cast/isa/dyn_cast_or_null functionality through llvm's doCast functionality in addition to defining methods with the same name. This change begins the migration of uses of the method to the corresponding function call as has been decided as more consistent. Note that there still exist classes that only define methods directly, such as AffineExpr, and this does not include work currently to support a functional cast/isa call. Caveats include: - This clang-tidy script probably has more problems. - This only touches C++ code, so nothing that is being generated. Context: - https://mlir.llvm.org/deprecation/ at "Use the free function variants for dyn_cast/cast/isa/…" - Original discussion at https://discourse.llvm.org/t/preferred-casting-style-going-forward/68443 Implementation: This first patch was created with the following steps. The intention is to only do automated changes at first, so I waste less time if it's reverted, and so the first mass change is more clear as an example to other teams that will need to follow similar steps. Steps are described per line, as comments are removed by git: 0. Retrieve the change from the following to build clang-tidy with an additional check: https://github.com/llvm/llvm-project/compare/main...tpopp:llvm-project:tidy-cast-check 1. Build clang-tidy 2. Run clang-tidy over your entire codebase while disabling all checks and enabling the one relevant one. Run on all header files also. 3. Delete .inc files that were also modified, so the next build rebuilds them to a pure state. 4. Some changes have been deleted for the following reasons: - Some files had a variable also named cast - Some files had not included a header file that defines the cast functions - Some files are definitions of the classes that have the casting methods, so the code still refers to the method instead of the function without adding a prefix or removing the method declaration at the same time. ``` ninja -C $BUILD_DIR clang-tidy run-clang-tidy -clang-tidy-binary=$BUILD_DIR/bin/clang-tidy -checks='-*,misc-cast-functions'\ -header-filter=mlir/ mlir/* -fix rm -rf $BUILD_DIR/tools/mlir/**/*.inc git restore mlir/lib/IR mlir/lib/Dialect/DLTI/DLTI.cpp\ mlir/lib/Dialect/Complex/IR/ComplexDialect.cpp\ mlir/lib/**/IR/\ mlir/lib/Dialect/SparseTensor/Transforms/SparseVectorization.cpp\ mlir/lib/Dialect/Vector/Transforms/LowerVectorMultiReduction.cpp\ mlir/test/lib/Dialect/Test/TestTypes.cpp\ mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp\ mlir/test/lib/Dialect/Test/TestAttributes.cpp\ mlir/unittests/TableGen/EnumsGenTest.cpp\ mlir/test/python/lib/PythonTestCAPI.cpp\ mlir/include/mlir/IR/ ``` Differential Revision: https://reviews.llvm.org/D150123
2023-04-20[mlir][Affine][NFC] Wrap dialect in "affine" namespaceMatthias Springer
This cleanup aligns the affine dialect with all the other dialects. Differential Revision: https://reviews.llvm.org/D148687
2023-04-11[Updated commit] Fix Transpose Check in MMA.SYNC Path.Manish Gupta
Pushed a stale commit for the same review in my previous commit. I am updating the main-line with the latest commit including review commits. Apologies for the redundant commit. Differential Revision: https://reviews.llvm.org/D147749
2023-04-11Fix Transpose Check in MMA.SYNC PathManish Gupta
Differential Revision: https://reviews.llvm.org/D147749
2023-03-29[mlir][Vector] Remove lhs and rhs masks from vector.contractDiego Caballero
This patch removes the historical lhs and rhs masks in vector.contract, now that vector.mask supports vector.contract and the lhs and rhs masks are barely supported by all the vector.contract lowerings and transformations. Reviewed By: nicolasvasilache Differential Revision: https://reviews.llvm.org/D144430
2023-03-09[mlir][vector][nvgpu] Move MMA contraction preparation to VectorUtilsJakub Kuderski
This pattern is not specific to nvgpu; I intend to use in SPIR-V codegen. `VectorTransforms` seems like a more generally useful place. In addition: - Fix a bug in the second condition (the dimensions were swapped for RHS). - Add tests. - Add support for externally provided filter functions, similar to other vector transforms. - Prefer to transpose before zero/sign-extending inputs. Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D145638
2023-02-16[mlir][vector] NFC: Improve vector type accessor methodsLei Zhang
Plain `getVectorType()` can be quite confusing and error-prone given that, well, vector ops always work on vector types, and it can commonly involve both source and result vectors. So this commit makes various such accessor methods to be explicit w.r.t. source or result vectors. Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D144159
2023-02-15[mlir][vectorToGPU] Fix type used when folding transpose into read opThomas Raoux
Pick the right result type when folding transpose op into a read Differential Revision: https://reviews.llvm.org/D144113
2023-02-15[mlir] Fix two build warnings in VectorToGPU (NFC)Jie Fu
In file included from /data/llvm-project/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:13: /data/llvm-project/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h:15:1: error: class 'LogicalResult' was previously declared as a struct; this is valid, but may result in linker errors under the Microsoft C++ ABI [-Werror,-Wmismatched-tags] class LogicalResult; ^ /data/llvm-project/mlir/include/mlir/Support/LogicalResult.h:26:22: note: previous use is here struct [[nodiscard]] LogicalResult { ^ /data/llvm-project/mlir/include/mlir/Conversion/VectorToGPU/VectorToGPU.h:15:1: note: did you mean struct here? class LogicalResult; ^~~~~ struct /data/llvm-project/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp:724:5: error: ignoring return value of function declared with 'nodiscard' attribute [-Werror,-Wunused-result] rewriter.notifyMatchFailure( ^~~~~~~~~~~~~~~~~~~~~~~~~~~ 2 errors generated.