[Mlir-commits] [mlir] [mlir][memref] Define interfaces for ops that access memrefs at an index (PR #177013)
Krzysztof Drewniak
llvmlistbot at llvm.org
Wed Jan 28 10:43:13 PST 2026
https://github.com/krzysz00 updated https://github.com/llvm/llvm-project/pull/177013
>From 514e57373f1288a8bdbc5f779efab16b7b185c7d Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Tue, 20 Jan 2026 17:39:08 +0000
Subject: [PATCH 1/5] [mlir][memref] Define interfaces for ops that access
memrefs at an index
This commit defines interfaces for operations that perform certain
kinds of indexed access on a memref. These interfaces are defined so
that passes like fold-memref-alias-ops and the memref flattener can be
made generic over operations that, informally, have the forms
`op ... %m[%i0, %i1, ...] ...` (an IndexedAccessOpInterface) or the
form `op %src[%s0, %s1, ...], %dst[%d0, %d1, ...] size ...` (an
IndexedMemCopyOpInterface).
These interfaces have been designed such that all the passes under
MemRef/Transforms that currently have a big switch-case on
memref.load, vector.load, nvgpu.ldmatrix, etc. can be migrated to use
them.
(This'll also let us get rid of the awkward fact that we have memref
transforms depending on the GPU and NVGPU dialects)
While the interface doesn't currently contemplate changing element
types (enabling, for example, writing a bf16 => u16 update to be done
in place), future extensions to allow such transformations could be
worth exploring.
This commit only defines the interfaces so that it'll be easier to
review the design - the implementation is in a future PR.
---
.../mlir/Dialect/MemRef/IR/CMakeLists.txt | 1 +
.../MemRef/IR/MemoryAccessOpInterfaces.h | 32 +++
.../MemRef/IR/MemoryAccessOpInterfaces.td | 200 ++++++++++++++++++
mlir/lib/Dialect/MemRef/IR/CMakeLists.txt | 2 +
.../MemRef/IR/MemoryAccessOpInterfaces.cpp | 64 ++++++
5 files changed, 299 insertions(+)
create mode 100644 mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h
create mode 100644 mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
create mode 100644 mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/MemRef/IR/CMakeLists.txt
index b7b12d49f9224..67e27141c18ab 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/MemRef/IR/CMakeLists.txt
@@ -1,2 +1,3 @@
add_mlir_dialect(MemRefOps memref)
+add_mlir_interface(MemoryAccessOpInterfaces)
add_mlir_doc(MemRefOps MemRefOps Dialects/ -gen-op-doc)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h
new file mode 100644
index 0000000000000..0ed9fd99cb78a
--- /dev/null
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h
@@ -0,0 +1,32 @@
+//===- MemoryAccessOpInterfaces.h -------------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_MEMREF_IR_MEMORYACCESSOPINTERFACES_H
+#define MLIR_DIALECT_MEMREF_IR_MEMORYACCESSOPINTERFACES_H
+
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/OpDefinition.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/Support/LLVM.h"
+
+namespace mlir {
+class RewriterBase;
+
+namespace memref::detail {
+LogicalResult verifyIndexedAccessOpInterface(Operation *op);
+LogicalResult verifyIndexedMemCopyOpInterface(Operation *op);
+} // namespace memref::detail
+} // namespace mlir
+
+//===----------------------------------------------------------------------===//
+// Memory Access Op Interfaces
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h.inc"
+
+#endif // MLIR_DIALECT_MEMREF_IR_MEMORYACCESSOPINTERFACES_H
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
new file mode 100644
index 0000000000000..2f2dd30cfe583
--- /dev/null
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
@@ -0,0 +1,200 @@
+//===-- MemoryAccessOpInterfaces.td ------------------------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MEMREF_MEMORY_ACCESS_OP_INTERFACES
+#define MEMREF_MEMORY_ACCESS_OP_INTERFACES
+
+include "mlir/IR/OpBase.td"
+
+def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
+ let description = [{
+ An interface for operations that operate on (by loading from or
+ storing to, atomically modifying, or otherwise) memory located at an
+ index within a memref whose semantics don't depend on the indexing scheme.
+
+ That is, a direct access op is one where, if `%b[%j0, %j1, ..., %jL]`
+ points to the same memory as `%a[%i0, %i1, ... %iK]`, it would be
+ trivial to replace `%a[%i0, ..., %iK]` with %b[%j0, ... %jL]`.
+
+ Operations may impose constraints on allowable reindexings.
+ Returning a non-empty result from `getAccessedShape()` imposes constraints
+ on the dimensions whose strides need to be preserved.
+
+ This interface is intended to enable transformations such as folding in
+ aliasing operations (like memref.subview or memref.collapse_shape) or
+ linearizing memrefs (making them 1-D) to be generic over in-tree and
+ out-of-tree operations.
+ }];
+ let cppNamespace = "::mlir::memref";
+ let methods = [
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the accessed memref. If the operation
+ is still in tensor form, return the null value.
+ }],
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getAccessedMemref",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the indices that are used to access the memref returned by getAccessedMemref().
+
+ The size of this range must be equal to the rank of the memref returned by
+ getAccessedMemref().
+ }],
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the shape of the portion of the memref that is being accessed by
+ this operation, if known, ignoring leading unit dimensions.
+ Reindexing transformations may not modify the *strides* of the trailing
+ N dimensions, where N is the size returned value, and should ensure that
+ at least N indexing dimensions remain after the transformation.
+ }],
+ /*retType=*/"::llvm::SmallVector<int64_t>",
+ /*methodName=*/"getAccessedShape",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
+ return ::llvm::SmallVector<int64_t>{};
+ }]>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Updates the memref being accessed to `newMemref` and the indices to
+ `newIndices`. If `std::nullopt` is returned, the operation was
+ updated in-place (the common case), while if a vector of values
+ is returned, they should be used to replace the operation being
+ updated.
+
+ This implementation of this method shall use the `modifyOpInPlace` method
+ on the provided rewriter when applicable, and may create or clone operations.
+ However, the operation must not replace itself, and should instead return
+ a vector of replacement results to the caller. (This allows for post-processing
+ and prevents stale references.)
+
+ The caller must ensure that the new memref/index pair points to the same
+ location in memory as the existing arguments.
+
+ The element types of the memref may not change.
+ }],
+ /*retType=*/"std::optional<llvm::SmallVector<mlir::Value>>",
+ /*methodName=*/"updateMemrefAndIndices",
+ /*args=*/(ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newMemref,
+ "::mlir::ValueRange":$newIndices)
+ >,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return true if, either by definition or due to some attribute,
+ it's known that all indices are non-negative and less than the size
+ of the dimension they index.
+ }],
+ /*retType=*/"bool",
+ /*methodName=*/"hasInboundsIndices",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
+ return true;
+ }]>
+ ];
+
+ let verify = [{
+ return ::mlir::memref::detail::verifyIndexedAccessOpInterface($_op);
+ }];
+}
+
+def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
+ let description = [{
+ This is an interface for operations that perform a copy of some number
+ of values from `%src[%srcIndices...]` to `%dst[%dstIndices...]`.The copy
+ does not necessarily target some linear sequence of elements (it may, for
+ example, be strided), and tere may be implicit offsets added to the source
+ and/or destination indices (for example, `amdgpu.gather_to_lds` can validly
+ implmenet this interface even though the destination index gets a lane ID *
+ copy size term added to it implicitly).
+
+ The motivating examples for this interface are operations that perform
+ direct loads to workgroup memory on GPUs.
+
+ This allows patterns that reindex memrefs (like folding in subview operations)
+ to treat such operations as a class, just like `IndexedAccessOpInterface`.
+
+ Unlike `IndexedAccessOpInterface`, this interface assumes that the elements
+ being copied are contiguous in memory and that the producers of the
+ operation has ensured this. That is, if the source memref is a
+ `memref<8x3xf32, strided<[9, 1]>>`, it is presumed that a memcopy of
+ 4 floats starting at the indices `[%x, %y]` is intended to read into
+ the space between the length-three rows.
+ }];
+ let cppNamespace = "::mlir::memref";
+ let methods = [
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the source memref for this copy operation.
+ }],
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getSrc",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the indices that are used to access the source memref.
+
+ The size of this range must be equal to the rank of the memref returned by
+ getSrc().
+ }],
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getSrcIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the destination memref for this copy operation.
+ }],
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getDst",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Return the indices that are used to access the destination memref.
+
+ The size of this range must be equal to the rank of the memref returned by
+ getDst().
+ }],
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getDstIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
+ Update the operation with `newSrcMemref` as the new source memref,
+ `newSrcIndices` as the new source indices, `newDst` as the new destination
+ memref, and `newDstIndices` as the new destination indices, leaving
+ other properties unchanged.
+
+ Note that, unlike the method on indexed accessors, this method doesn't
+ currently allow for the possibility of cloning since DMA ops generally
+ do not return results. If needed, such suppport could be added in
+ the future.
+
+ The implementation is responsible for issuing rewriter notifications
+ via modifyOpInPlace().
+
+ The caller must ensure that the new memref/index pairs point to the same
+ location in memory.
+ }],
+ /*retType=*/"void",
+ /*methodName=*/"setMemrefsAndIndices",
+ /*args=*/(ins "::mlir::RewriterBase&":$rewriter,
+ "::mlir::Value":$newSrc, "::mlir::ValueRange":$newSrcIndices,
+ "::mlir::Value":$newDst, "::mlir::ValueRange":$newDstIndices)
+ >
+ ];
+ let verify = [{
+ return ::mlir::memref::detail::verifyIndexedMemCopyOpInterface($_op);
+ }];
+}
+#endif // MEMREF_MEMORY_ACCESS_OP_INTERFACES
diff --git a/mlir/lib/Dialect/MemRef/IR/CMakeLists.txt b/mlir/lib/Dialect/MemRef/IR/CMakeLists.txt
index d358362f1984b..56d98dd431c8e 100644
--- a/mlir/lib/Dialect/MemRef/IR/CMakeLists.txt
+++ b/mlir/lib/Dialect/MemRef/IR/CMakeLists.txt
@@ -1,5 +1,6 @@
add_mlir_dialect_library(MLIRMemRefDialect
MemRefDialect.cpp
+ MemoryAccessOpInterfaces.cpp
MemRefMemorySlot.cpp
MemRefOps.cpp
ValueBoundsOpInterfaceImpl.cpp
@@ -9,6 +10,7 @@ add_mlir_dialect_library(MLIRMemRefDialect
DEPENDS
MLIRMemRefOpsIncGen
+ MLIRMemoryAccessOpInterfacesIncGen
LINK_LIBS PUBLIC
MLIRArithDialect
diff --git a/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
new file mode 100644
index 0000000000000..70eed7daf3962
--- /dev/null
+++ b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
@@ -0,0 +1,64 @@
+//===- MemoryAccessOpInterfaces.cpp ---------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/IR/Operation.h"
+#include "mlir/IR/Value.h"
+
+//===----------------------------------------------------------------------===//
+// IndexedAccessOpInterface and IndexedMemCpyOpInterface
+//===----------------------------------------------------------------------===//
+
+namespace mlir::memref {
+#include "mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp.inc"
+} // namespace mlir::memref
+
+using namespace mlir;
+using namespace mlir::memref;
+
+LogicalResult
+mlir::memref::detail::verifyIndexedAccessOpInterface(Operation *op) {
+ auto iface = dyn_cast<IndexedAccessOpInterface>(op);
+ TypedValue<MemRefType> memref = iface.getAccessedMemref();
+ if (!memref) {
+ // Some operations can carry tensors, this is fine.
+ return success();
+ }
+ if (memref.getType().getRank() !=
+ static_cast<int64_t>(iface.getIndices().size()))
+ return op->emitOpError(
+ "invalid number of indices for accessed memref, expected " +
+ Twine(memref.getType().getRank()) + " but got " +
+ Twine(iface.getIndices().size()));
+ return success();
+}
+
+LogicalResult
+mlir::memref::detail::verifyIndexedMemCopyOpInterface(Operation *op) {
+ auto iface = dyn_cast<IndexedMemCopyOpInterface>(op);
+ TypedValue<MemRefType> src = iface.getSrc();
+ TypedValue<MemRefType> dst = iface.getDst();
+ if (!src || !dst) {
+ // Allow operations to not always have memref arguments.
+ return ::mlir::success();
+ }
+ if (src.getType().getRank() !=
+ static_cast<int64_t>(iface.getSrcIndices().size()))
+ return op->emitOpError(
+ "invalid number of indices for source memref, expected " +
+ Twine(src.getType().getRank()) + ", got " +
+ Twine(iface.getSrcIndices().size()));
+ if (dst.getType().getRank() !=
+ static_cast<int64_t>(iface.getDstIndices().size()))
+ return op->emitOpError(
+ "invalid number of indices for destination memref, expected " +
+ Twine(dst.getType().getRank()) + ", got " +
+ Twine(iface.getDstIndices().size()));
+ return success();
+}
>From a3cc990bc40940c5fea006fb384cc49e2ffca4a9 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Tue, 27 Jan 2026 13:28:13 -0800
Subject: [PATCH 2/5] Feedback and documentation layout fixes
Co-authored-by: Jakub Kuderski <jakub at nod-labs.com>
---
.../MemRef/IR/MemoryAccessOpInterfaces.td | 24 +++++++++----------
.../MemRef/IR/MemoryAccessOpInterfaces.cpp | 8 +++----
2 files changed, 16 insertions(+), 16 deletions(-)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
index 2f2dd30cfe583..957ba9c2c16e5 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
@@ -13,20 +13,20 @@ include "mlir/IR/OpBase.td"
def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
let description = [{
- An interface for operations that operate on (by loading from or
+ An interface for operations that access (by loading from or
storing to, atomically modifying, or otherwise) memory located at an
index within a memref whose semantics don't depend on the indexing scheme.
That is, a direct access op is one where, if `%b[%j0, %j1, ..., %jL]`
points to the same memory as `%a[%i0, %i1, ... %iK]`, it would be
- trivial to replace `%a[%i0, ..., %iK]` with %b[%j0, ... %jL]`.
+ legal to replace `%a[%i0, ..., %iK]` with %b[%j0, ... %jL]`.
Operations may impose constraints on allowable reindexings.
Returning a non-empty result from `getAccessedShape()` imposes constraints
on the dimensions whose strides need to be preserved.
This interface is intended to enable transformations such as folding in
- aliasing operations (like memref.subview or memref.collapse_shape) or
+ aliasing operations (like `memref.subview` or `memref.collapse_shape`) or
linearizing memrefs (making them 1-D) to be generic over in-tree and
out-of-tree operations.
}];
@@ -42,10 +42,10 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
/*args=*/(ins)>,
InterfaceMethod<
/*desc=*/[{
- Return the indices that are used to access the memref returned by getAccessedMemref().
+ Return the indices that are used to access the memref returned by `getAccessedMemref()`.
The size of this range must be equal to the rank of the memref returned by
- getAccessedMemref().
+ `getAccessedMemref()`.
}],
/*retType=*/"::mlir::Operation::operand_range",
/*methodName=*/"getIndices",
@@ -112,11 +112,11 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
let description = [{
This is an interface for operations that perform a copy of some number
- of values from `%src[%srcIndices...]` to `%dst[%dstIndices...]`.The copy
+ of values from `%src[%srcIndices...]` to `%dst[%dstIndices...]`. The copy
does not necessarily target some linear sequence of elements (it may, for
- example, be strided), and tere may be implicit offsets added to the source
+ example, be strided), and there may be implicit offsets added to the source
and/or destination indices (for example, `amdgpu.gather_to_lds` can validly
- implmenet this interface even though the destination index gets a lane ID *
+ implement this interface even though the destination index gets a lane ID *
copy size term added to it implicitly).
The motivating examples for this interface are operations that perform
@@ -127,8 +127,8 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
Unlike `IndexedAccessOpInterface`, this interface assumes that the elements
being copied are contiguous in memory and that the producers of the
- operation has ensured this. That is, if the source memref is a
- `memref<8x3xf32, strided<[9, 1]>>`, it is presumed that a memcopy of
+ operation have ensured this. That is, if the source memref is a
+ `memref<8x3xf32, strided<[9, 1]>>`, it is presumed that a memcpy of
4 floats starting at the indices `[%x, %y]` is intended to read into
the space between the length-three rows.
}];
@@ -146,7 +146,7 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
Return the indices that are used to access the source memref.
The size of this range must be equal to the rank of the memref returned by
- getSrc().
+ `getSrc()`.
}],
/*retType=*/"::mlir::Operation::operand_range",
/*methodName=*/"getSrcIndices",
@@ -163,7 +163,7 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
Return the indices that are used to access the destination memref.
The size of this range must be equal to the rank of the memref returned by
- getDst().
+ `getDst()`.
}],
/*retType=*/"::mlir::Operation::operand_range",
/*methodName=*/"getDstIndices",
diff --git a/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
index 70eed7daf3962..2bd9497e3f54d 100644
--- a/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
+++ b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
@@ -33,9 +33,9 @@ mlir::memref::detail::verifyIndexedAccessOpInterface(Operation *op) {
if (memref.getType().getRank() !=
static_cast<int64_t>(iface.getIndices().size()))
return op->emitOpError(
- "invalid number of indices for accessed memref, expected " +
- Twine(memref.getType().getRank()) + " but got " +
- Twine(iface.getIndices().size()));
+ "invalid number of indices for accessed memref, expected ") <<
+ memref.getType().getRank() << " but got " <<
+ iface.getIndices().size());
return success();
}
@@ -46,7 +46,7 @@ mlir::memref::detail::verifyIndexedMemCopyOpInterface(Operation *op) {
TypedValue<MemRefType> dst = iface.getDst();
if (!src || !dst) {
// Allow operations to not always have memref arguments.
- return ::mlir::success();
+ return success();
}
if (src.getType().getRank() !=
static_cast<int64_t>(iface.getSrcIndices().size()))
>From ca8e98ea45cb1c398e4a45bf8cb3d60162898e71 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Tue, 27 Jan 2026 22:03:16 +0000
Subject: [PATCH 3/5] Revew comments that weren't Github suggestions,
clang-format ran on my .td
---
.../MemRef/IR/MemoryAccessOpInterfaces.td | 133 +++++++++---------
.../MemRef/IR/MemoryAccessOpInterfaces.cpp | 29 ++--
2 files changed, 81 insertions(+), 81 deletions(-)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
index 957ba9c2c16e5..49e0dff68861d 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
@@ -31,42 +31,43 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
out-of-tree operations.
}];
let cppNamespace = "::mlir::memref";
- let methods = [
- InterfaceMethod<
- /*desc=*/[{
- Return the accessed memref. If the operation
- is still in tensor form, return the null value.
+ let methods =
+ [InterfaceMethod<
+ /*desc=*/[{
+ Return the accessed memref. If the operation is still in tensor form, return
+ the null value.
}],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getAccessedMemref",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getAccessedMemref",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the memref returned by `getAccessedMemref()`.
The size of this range must be equal to the rank of the memref returned by
`getAccessedMemref()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the shape of the portion of the memref that is being accessed by
this operation, if known, ignoring leading unit dimensions.
+
Reindexing transformations may not modify the *strides* of the trailing
N dimensions, where N is the size returned value, and should ensure that
at least N indexing dimensions remain after the transformation.
}],
- /*retType=*/"::llvm::SmallVector<int64_t>",
- /*methodName=*/"getAccessedShape",
- /*args=*/(ins),
- /*methodBody=*/[{}],
- /*defaultImplementation=*/[{
+ /*retType=*/"::llvm::SmallVector<int64_t>",
+ /*methodName=*/"getAccessedShape",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
return ::llvm::SmallVector<int64_t>{};
}]>,
- InterfaceMethod<
- /*desc=*/[{
+ InterfaceMethod<
+ /*desc=*/[{
Updates the memref being accessed to `newMemref` and the indices to
`newIndices`. If `std::nullopt` is returned, the operation was
updated in-place (the common case), while if a vector of values
@@ -84,25 +85,24 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
The element types of the memref may not change.
}],
- /*retType=*/"std::optional<llvm::SmallVector<mlir::Value>>",
- /*methodName=*/"updateMemrefAndIndices",
- /*args=*/(ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newMemref,
- "::mlir::ValueRange":$newIndices)
- >,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"std::optional<llvm::SmallVector<mlir::Value>>",
+ /*methodName=*/"updateMemrefAndIndices",
+ /*args=*/
+ (ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newMemref,
+ "::mlir::ValueRange":$newIndices)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return true if, either by definition or due to some attribute,
it's known that all indices are non-negative and less than the size
of the dimension they index.
}],
- /*retType=*/"bool",
- /*methodName=*/"hasInboundsIndices",
- /*args=*/(ins),
- /*methodBody=*/[{}],
- /*defaultImplementation=*/[{
+ /*retType=*/"bool",
+ /*methodName=*/"hasInboundsIndices",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
return true;
- }]>
- ];
+ }]>];
let verify = [{
return ::mlir::memref::detail::verifyIndexedAccessOpInterface($_op);
@@ -133,43 +133,43 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
the space between the length-three rows.
}];
let cppNamespace = "::mlir::memref";
- let methods = [
- InterfaceMethod<
- /*desc=*/[{
+ let methods =
+ [InterfaceMethod<
+ /*desc=*/[{
Return the source memref for this copy operation.
- }],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getSrc",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ }],
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getSrc",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the source memref.
The size of this range must be equal to the rank of the memref returned by
`getSrc()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getSrcIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getSrcIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the destination memref for this copy operation.
}],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getDst",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getDst",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the destination memref.
The size of this range must be equal to the rank of the memref returned by
`getDst()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getDstIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getDstIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Update the operation with `newSrcMemref` as the new source memref,
`newSrcIndices` as the new source indices, `newDst` as the new destination
memref, and `newDstIndices` as the new destination indices, leaving
@@ -186,13 +186,12 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
The caller must ensure that the new memref/index pairs point to the same
location in memory.
}],
- /*retType=*/"void",
- /*methodName=*/"setMemrefsAndIndices",
- /*args=*/(ins "::mlir::RewriterBase&":$rewriter,
- "::mlir::Value":$newSrc, "::mlir::ValueRange":$newSrcIndices,
- "::mlir::Value":$newDst, "::mlir::ValueRange":$newDstIndices)
- >
- ];
+ /*retType=*/"void",
+ /*methodName=*/"setMemrefsAndIndices",
+ /*args=*/
+ (ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newSrc,
+ "::mlir::ValueRange":$newSrcIndices, "::mlir::Value":$newDst,
+ "::mlir::ValueRange":$newDstIndices)>];
let verify = [{
return ::mlir::memref::detail::verifyIndexedMemCopyOpInterface($_op);
}];
diff --git a/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
index 2bd9497e3f54d..c71df9a2015f7 100644
--- a/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
+++ b/mlir/lib/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp
@@ -17,14 +17,12 @@
namespace mlir::memref {
#include "mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.cpp.inc"
-} // namespace mlir::memref
-
-using namespace mlir;
-using namespace mlir::memref;
-LogicalResult
-mlir::memref::detail::verifyIndexedAccessOpInterface(Operation *op) {
+LogicalResult detail::verifyIndexedAccessOpInterface(Operation *op) {
auto iface = dyn_cast<IndexedAccessOpInterface>(op);
+ if (!iface)
+ return failure();
+
TypedValue<MemRefType> memref = iface.getAccessedMemref();
if (!memref) {
// Some operations can carry tensors, this is fine.
@@ -33,15 +31,17 @@ mlir::memref::detail::verifyIndexedAccessOpInterface(Operation *op) {
if (memref.getType().getRank() !=
static_cast<int64_t>(iface.getIndices().size()))
return op->emitOpError(
- "invalid number of indices for accessed memref, expected ") <<
- memref.getType().getRank() << " but got " <<
- iface.getIndices().size());
+ "invalid number of indices for accessed memref, expected ")
+ << memref.getType().getRank() << " but got "
+ << iface.getIndices().size();
return success();
}
-LogicalResult
-mlir::memref::detail::verifyIndexedMemCopyOpInterface(Operation *op) {
+LogicalResult detail::verifyIndexedMemCopyOpInterface(Operation *op) {
auto iface = dyn_cast<IndexedMemCopyOpInterface>(op);
+ if (!iface)
+ return failure();
+
TypedValue<MemRefType> src = iface.getSrc();
TypedValue<MemRefType> dst = iface.getDst();
if (!src || !dst) {
@@ -57,8 +57,9 @@ mlir::memref::detail::verifyIndexedMemCopyOpInterface(Operation *op) {
if (dst.getType().getRank() !=
static_cast<int64_t>(iface.getDstIndices().size()))
return op->emitOpError(
- "invalid number of indices for destination memref, expected " +
- Twine(dst.getType().getRank()) + ", got " +
- Twine(iface.getDstIndices().size()));
+ "invalid number of indices for destination memref, expected ")
+ << dst.getType().getRank() << ", got "
+ << iface.getDstIndices().size();
return success();
}
+} // namespace mlir::memref
>From 4ae320b2d3bb77bfbb63cb43a9acca9f79908bed Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Wed, 28 Jan 2026 18:27:04 +0000
Subject: [PATCH 4/5] Undo clang-format damage
---
.../MemRef/IR/MemoryAccessOpInterfaces.td | 123 +++++++++---------
1 file changed, 61 insertions(+), 62 deletions(-)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
index 49e0dff68861d..5dca94b5ddd4e 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
@@ -32,26 +32,26 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
}];
let cppNamespace = "::mlir::memref";
let methods =
- [InterfaceMethod<
- /*desc=*/[{
+ [InterfaceMethod<
+ /*desc=*/[{
Return the accessed memref. If the operation is still in tensor form, return
- the null value.
+ the null value.
}],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getAccessedMemref",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getAccessedMemref",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the memref returned by `getAccessedMemref()`.
The size of this range must be equal to the rank of the memref returned by
`getAccessedMemref()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the shape of the portion of the memref that is being accessed by
this operation, if known, ignoring leading unit dimensions.
@@ -59,15 +59,15 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
N dimensions, where N is the size returned value, and should ensure that
at least N indexing dimensions remain after the transformation.
}],
- /*retType=*/"::llvm::SmallVector<int64_t>",
- /*methodName=*/"getAccessedShape",
- /*args=*/(ins),
- /*methodBody=*/[{}],
- /*defaultImplementation=*/[{
+ /*retType=*/"::llvm::SmallVector<int64_t>",
+ /*methodName=*/"getAccessedShape",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
return ::llvm::SmallVector<int64_t>{};
}]>,
- InterfaceMethod<
- /*desc=*/[{
+ InterfaceMethod<
+ /*desc=*/[{
Updates the memref being accessed to `newMemref` and the indices to
`newIndices`. If `std::nullopt` is returned, the operation was
updated in-place (the common case), while if a vector of values
@@ -85,24 +85,24 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
The element types of the memref may not change.
}],
- /*retType=*/"std::optional<llvm::SmallVector<mlir::Value>>",
- /*methodName=*/"updateMemrefAndIndices",
- /*args=*/
- (ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newMemref,
- "::mlir::ValueRange":$newIndices)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"std::optional<llvm::SmallVector<mlir::Value>>",
+ /*methodName=*/"updateMemrefAndIndices",
+ /*args=*/(ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newMemref,
+ "::mlir::ValueRange":$newIndices)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return true if, either by definition or due to some attribute,
it's known that all indices are non-negative and less than the size
of the dimension they index.
}],
- /*retType=*/"bool",
- /*methodName=*/"hasInboundsIndices",
- /*args=*/(ins),
- /*methodBody=*/[{}],
- /*defaultImplementation=*/[{
+ /*retType=*/"bool",
+ /*methodName=*/"hasInboundsIndices",
+ /*args=*/(ins),
+ /*methodBody=*/[{}],
+ /*defaultImplementation=*/[{
return true;
- }]>];
+ }]>
+ ];
let verify = [{
return ::mlir::memref::detail::verifyIndexedAccessOpInterface($_op);
@@ -134,42 +134,42 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
}];
let cppNamespace = "::mlir::memref";
let methods =
- [InterfaceMethod<
- /*desc=*/[{
+ [InterfaceMethod<
+ /*desc=*/[{
Return the source memref for this copy operation.
- }],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getSrc",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ }],
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getSrc",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the source memref.
The size of this range must be equal to the rank of the memref returned by
`getSrc()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getSrcIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getSrcIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the destination memref for this copy operation.
}],
- /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
- /*methodName=*/"getDst",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::TypedValue<::mlir::MemRefType>",
+ /*methodName=*/"getDst",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Return the indices that are used to access the destination memref.
The size of this range must be equal to the rank of the memref returned by
`getDst()`.
}],
- /*retType=*/"::mlir::Operation::operand_range",
- /*methodName=*/"getDstIndices",
- /*args=*/(ins)>,
- InterfaceMethod<
- /*desc=*/[{
+ /*retType=*/"::mlir::Operation::operand_range",
+ /*methodName=*/"getDstIndices",
+ /*args=*/(ins)>,
+ InterfaceMethod<
+ /*desc=*/[{
Update the operation with `newSrcMemref` as the new source memref,
`newSrcIndices` as the new source indices, `newDst` as the new destination
memref, and `newDstIndices` as the new destination indices, leaving
@@ -186,12 +186,11 @@ def IndexedMemCopyOpInterface : OpInterface<"IndexedMemCopyOpInterface"> {
The caller must ensure that the new memref/index pairs point to the same
location in memory.
}],
- /*retType=*/"void",
- /*methodName=*/"setMemrefsAndIndices",
- /*args=*/
- (ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newSrc,
- "::mlir::ValueRange":$newSrcIndices, "::mlir::Value":$newDst,
- "::mlir::ValueRange":$newDstIndices)>];
+ /*retType=*/"void",
+ /*methodName=*/"setMemrefsAndIndices",
+ /*args=*/(ins "::mlir::RewriterBase&":$rewriter, "::mlir::Value":$newSrc,
+ "::mlir::ValueRange":$newSrcIndices, "::mlir::Value":$newDst,
+ "::mlir::ValueRange":$newDstIndices)>];
let verify = [{
return ::mlir::memref::detail::verifyIndexedMemCopyOpInterface($_op);
}];
>From 7851b6513f655ce1f1b5f2ba0c02240e4fb34e45 Mon Sep 17 00:00:00 2001
From: Krzysztof Drewniak <Krzysztof.Drewniak at amd.com>
Date: Wed, 28 Jan 2026 18:42:57 +0000
Subject: [PATCH 5/5] Clarify docs
---
mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
index 5dca94b5ddd4e..7fc69b4fabca6 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemoryAccessOpInterfaces.td
@@ -75,7 +75,7 @@ def IndexedAccessOpInterface : OpInterface<"IndexedAccessOpInterface"> {
updated.
This implementation of this method shall use the `modifyOpInPlace` method
- on the provided rewriter when applicable, and may create or clone operations.
+ on the provided rewriter when applicable. It may also create or clone operations.
However, the operation must not replace itself, and should instead return
a vector of replacement results to the caller. (This allows for post-processing
and prevents stale references.)
More information about the Mlir-commits
mailing list