[mlir] Flip LinAlg dialect to _Both

This one required more changes than ideal due to overlapping generated name
with different return types. Changed getIndexingMaps to getIndexingMapsArray to
move i

[mlir] Flip LinAlg dialect to _Both

This one required more changes than ideal due to overlapping generated name
with different return types. Changed getIndexingMaps to getIndexingMapsArray to
move it out of the way/highlight that it returns (more expensively) a
SmallVector and uses the prefixed name for the Attribute.

Differential Revision: https://reviews.llvm.org/D129919

show more ...

[mlir][VectorToGPU] Fix support for i4, col-major operand support

For the conversion to nvgpu `mma.sync` and `ldmatrix` pathways, the code
was missing support for the `i4` data type. While fixing th

[mlir][VectorToGPU] Fix support for i4, col-major operand support

For the conversion to nvgpu `mma.sync` and `ldmatrix` pathways, the code
was missing support for the `i4` data type. While fixing this, another
bug was discoverd that caused the number of ldmatrix tiles calculated for
certain operand types and configurations to be incorrect. This change
fixes both issues and adds additional tests.

Differential Revision: https://reviews.llvm.org/D128074

show more ...

Don't use Optional::hasValue (NFC)

[mlir] move SCF headers to SCF/{IR,Transforms} respectively

This aligns the SCF dialect file layout with the majority of the dialects.

Reviewed By: jpienaar

Differential Revision: https://reviews.

[mlir] move SCF headers to SCF/{IR,Transforms} respectively

This aligns the SCF dialect file layout with the majority of the dialects.

Reviewed By: jpienaar

Differential Revision: https://reviews.llvm.org/D128049

show more ...

[mlir][nvgpu] shared memory access optimization pass

This change adds a transformation and pass to the NvGPU dialect that
attempts to optimize reads/writes from a memref representing GPU shared
mem

[mlir][nvgpu] shared memory access optimization pass

This change adds a transformation and pass to the NvGPU dialect that
attempts to optimize reads/writes from a memref representing GPU shared
memory in order to avoid bank conflicts. Given a value representing a
shared memory memref, it traverses all reads/writes within the parent op
and, subject to suitable conditions, rewrites all last dimension index
values such that element locations in the final (col) dimension are
given by
`newColIdx = col % vecSize + perm[row](col/vecSize,row)`
where `perm` is a permutation function indexed by `row` and `vecSize`
is the vector access size in elements (currently assumes 128bit
vectorized accesses, but this can be made a parameter). This specific
transformation can help optimize typical distributed & vectorized accesses
common to loading matrix multiplication operands to/from shared memory.

Differential Revision: https://reviews.llvm.org/D127457

show more ...

[mlir][gpu] Move GPU headers into IR/ and Transforms/

Depends on D127350

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D127352

[mlir][VectorToGPU] Fix bug generating incorrect ldmatrix ops

ldmatrix transpose can only be used with types that are 16bits wide.

Differential Revision: https://reviews.llvm.org/D126846

[MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass

This changes adds the option to lower to NvGpu dialect ops during the
VectorToGPU convsersion pass. Because this transformation reuses
exi

[MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass

This changes adds the option to lower to NvGpu dialect ops during the
VectorToGPU convsersion pass. Because this transformation reuses
existing VectorToGPU logic, a seperate VectorToNvGpu conversion pass is
not created. The option `use-nvgpu` is added to the VectorToGPU pass.
When this is true, the pass will attempt to convert slices rooted at
`vector.contract` operations into `nvgpu.mma.sync` ops, and
`vector.transfer_read` ops are converted to either `nvgpu.ldmatrix` or
one or more `vector.load` operations. The specific data loaded will
depend on the thread id within a subgroup (warp). These index
calculations depend on data type and shape of the MMA op
according to the downstream PTX specification. The code for supporting
these details is separated into `NvGpuSupport.cpp|h`.

Differential Revision: https://reviews.llvm.org/D122940

show more ...

[mlir] Flip Vector dialect accessors used to prefixed form.

This has been on _Both for a couple of weeks. Flip usages in core with
intention to flip flag to _Prefixed in follow up. Needed to add a c

[mlir] Flip Vector dialect accessors used to prefixed form.

This has been on _Both for a couple of weeks. Flip usages in core with
intention to flip flag to _Prefixed in follow up. Needed to add a couple
of helper methods in AffineOps and Linalg to facilitate a pure flag flip
in follow up as some of these classes are used in templates and so
sensitive to Vector dialect changes.

Differential Revision: https://reviews.llvm.org/D122151

show more ...

[mlir][gpu] Relax restriction on mma load/store op

Those ops can support more complex layout as long as the most inner
dimension is contiguous.

Differential Revision: https://reviews.llvm.org/D1224

[mlir][gpu] Relax restriction on mma load/store op

Those ops can support more complex layout as long as the most inner
dimension is contiguous.

Differential Revision: https://reviews.llvm.org/D122452

show more ...

[mlir] Update FuncOp conversion passes to Pass/InterfacePass<FunctionOpInterface>

These passes generally don't rely on any special aspects of FuncOp, and moving allows
for these passes to be used in

[mlir] Update FuncOp conversion passes to Pass/InterfacePass<FunctionOpInterface>

These passes generally don't rely on any special aspects of FuncOp, and moving allows
for these passes to be used in many more situations. The passes that obviously weren't
relying on invariants guaranteed by a "function" were updated to be generic pass, the
rest were updated to be FunctionOpinterface InterfacePasses.

The test updates are NFC switching from implicit nesting (-pass -pass2) form to
the -pass-pipeline form (generic passes do not implicitly nest as op-specific passes do).

Differential Revision: https://reviews.llvm.org/D121190

show more ...

[mlir][vector][NFC] Split into IR, Transforms and Utils

This reduces the dependencies of the MLIRVector target and makes the dialect consistent with other dialects.

Differential Revision: https://r

[mlir][vector][NFC] Split into IR, Transforms and Utils

This reduces the dependencies of the MLIRVector target and makes the dialect consistent with other dialects.

Differential Revision: https://reviews.llvm.org/D118533

show more ...

[VectorToGPU] Fix horizontal stride calculation for N-D memref

Fix a bug in how we calculate the stride of mma load/store ops for N-D
memrefs

Differential Revision: https://reviews.llvm.org/D118378

[mlir] Make locations required when adding/creating block arguments

BlockArguments gained the ability to have locations attached a while ago, but they
have always been optional. This goes against th

[mlir] Make locations required when adding/creating block arguments

BlockArguments gained the ability to have locations attached a while ago, but they
have always been optional. This goes against the core tenant of MLIR where location
information is a requirement, so this commit updates the API to require locations.

Fixes #53279

Differential Revision: https://reviews.llvm.org/D117633

show more ...

[mlir][Pass] Deprecate FunctionPass in favor of OperationPass<FuncOp>

The only benefit of FunctionPass is that it filters out function
declarations. This isn't enough to justify carrying it around,

[mlir][Pass] Deprecate FunctionPass in favor of OperationPass<FuncOp>

The only benefit of FunctionPass is that it filters out function
declarations. This isn't enough to justify carrying it around, as we can
simplify filter out declarations when necessary within the pass. We can
also explore with better scheduling primitives to filter out declarations
at the pipeline level in the future.

The definition of FunctionPass is left intact for now to allow time for downstream
users to migrate.

Differential Revision: https://reviews.llvm.org/D117182

show more ...

Apply clang-tidy fixes for performance-for-range-copy to MLIR (NFC)

Apply clang-tidy fixes for readability-simplify-boolean-expr to MLIR (NFC)

Reviewed By: rriddle, Mogball

Differential Revision: https://reviews.llvm.org/D116253

[mlir] Switching accessors to prefixed form (NFC)

Makes eventual prefixing flag flip smaller change.

[mlir][Vector] Thread 0-d vectors through vector.transfer ops

This revision adds 0-d vector support to vector.transfer ops.
In the process, numerous cleanups are applied, in particular around normal

[mlir][Vector] Thread 0-d vectors through vector.transfer ops

This revision adds 0-d vector support to vector.transfer ops.
In the process, numerous cleanups are applied, in particular around normalizing
and reducing the number of builders.

Reviewed By: ThomasRaoux, springerm

Differential Revision: https://reviews.llvm.org/D114803

show more ...

[mlir] Move min/max ops from Std to Arith.

Differential Revision: https://reviews.llvm.org/D113881

[mlir][VectorToGPU] Support more cases in conversion to MMA ops

Support load with broadcast, elementwise divf op and remove the
hardcoded restriction on the vector size. Picking the right size shoul

[mlir][VectorToGPU] Support more cases in conversion to MMA ops

Support load with broadcast, elementwise divf op and remove the
hardcoded restriction on the vector size. Picking the right size should
be enfored by user and will fail conversion to llvm/spirv if it is not
supported.

Differential Revision: https://reviews.llvm.org/D113618

show more ...

[mlir] Remove the non-templated DenseElementsAttr::getSplatValue

This predates the templated variant, and has been simply forwarding
to getSplatValue<Attribute> for some time. Removing this makes th

[mlir] Remove the non-templated DenseElementsAttr::getSplatValue

This predates the templated variant, and has been simply forwarding
to getSplatValue<Attribute> for some time. Removing this makes the
API a bit more uniform, and also helps prevent users from thinking
it is "cheap".

show more ...

[mlir][VectorToGPU] Add support for elementwise mma to vector to GPU

Differential Revision: https://reviews.llvm.org/D112960

[mlir] Switch arith, llvm, std & shape dialects to accessors prefixed both form.

Following
https://llvm.discourse.group/t/psa-ods-generated-accessors-will-change-to-have-a-get-prefix-update-you-apis

[mlir] Switch arith, llvm, std & shape dialects to accessors prefixed both form.

Following
https://llvm.discourse.group/t/psa-ods-generated-accessors-will-change-to-have-a-get-prefix-update-you-apis/4476,
this follows flipping these dialects to _Both prefixed form. This
changes the accessors to have a prefix. This was possibly mostly without
breaking breaking changes if the existing convenience methods were used.

(https://github.com/jpienaar/llvm-project/blob/main/clang-tools-extra/clang-tidy/misc/AddGetterCheck.cpp
was used to migrate the callers post flipping, using the output from
Operator.cpp)

Differential Revision: https://reviews.llvm.org/D112383

show more ...

[MLIR] Replace std ops with arith dialect ops

Precursor: https://reviews.llvm.org/D110200

Removed redundant ops from the standard dialect that were moved to the
`arith` or `math` dialects.

Renamed

[MLIR] Replace std ops with arith dialect ops

Precursor: https://reviews.llvm.org/D110200

Removed redundant ops from the standard dialect that were moved to the
`arith` or `math` dialects.

Renamed all instances of operations in the codebase and in tests.

Reviewed By: rriddle, jpienaar

Differential Revision: https://reviews.llvm.org/D110797

show more ...