[Mlir-commits] [mlir] [mlir][memref] Add static_sizes attribute to memref.view (PR #183795)
Amir Bishara
llvmlistbot at llvm.org
Fri Feb 27 10:42:26 PST 2026
https://github.com/amirBish updated https://github.com/llvm/llvm-project/pull/183795
>From 21b2a943800cf223a4d53f6fa2cba471bb8d2bcd Mon Sep 17 00:00:00 2001
From: Amir Bishara <amir.bishara at mobileye.com>
Date: Fri, 27 Feb 2026 18:56:38 +0200
Subject: [PATCH] [mlir][memref] Add static_sizes attribute to memref.view
Add a `DenseI64ArrayAttr:$static_sizes` argument to `memref.view`,
aligning it with `memref.subview` and `memref.reinterpret_cast` which
already represent sizes as a mix of static attribute values and dynamic
SSA operands.
Previously, `memref.view` only accepted dynamic SSA values for sizes,
and static size information was only encoded implicitly in the result
memref type. This made the op inconsistent with the rest of the dialect
and limited the ability to express mixed static/dynamic sizes directly.
Also, Not having such an attribute made it hard to make a verification
in the ViewOp's verifier without inspecting the producer's ops, therfore
having such encoded attribute will make it possible.
The assembly format now uses `custom<DynamicIndexList>` to print/parse
the sizes list, where static sizes appear as integer literals and
dynamic sizes as SSA values:
// Before:
memref.view %buf[%off][] : memref<2048xi8> to memref<64x4xf32>
memref.view %buf[%off][%s0] : memref<2048xi8> to memref<4x?xf32>
// After:
memref.view %buf[%off][64, 4] : memref<2048xi8> to memref<64x4xf32>
memref.view %buf[%off][4, %s0] : memref<2048xi8> to memref<4x?xf32>
Three convenience builders are added accepting `ArrayRef<OpFoldResult>`,
`ArrayRef<int64_t>`, and `ValueRange` for the sizes argument.
---
.../mlir/Dialect/MemRef/IR/MemRefOps.td | 38 +++--
.../Conversion/MemRefToLLVM/MemRefToLLVM.cpp | 5 +-
mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp | 132 +++++++++++-------
mlir/test/Analysis/test-alias-analysis.mlir | 2 +-
.../GPUCommon/lower-memory-space-attrs.mlir | 4 +-
.../MemRefToLLVM/memref-to-llvm.mlir | 8 +-
.../XeGPUToXeVM/loadstore_matrix.mlir | 2 +-
.../Dialect/Func/func-transform-invalid.mlir | 44 +++---
mlir/test/Dialect/Func/func-transform.mlir | 80 +++++------
.../Dialect/GPU/dynamic-shared-memory.mlir | 8 +-
mlir/test/Dialect/Linalg/canonicalize.mlir | 2 +-
.../Linalg/forward-vector-transfers.mlir | 6 +-
mlir/test/Dialect/Linalg/loops.mlir | 24 ++--
mlir/test/Dialect/Linalg/promote.mlir | 12 +-
.../Dialect/Linalg/promotion_options.mlir | 4 +-
mlir/test/Dialect/Linalg/roundtrip.mlir | 4 +-
.../Dialect/Linalg/transform-promotion.mlir | 12 +-
mlir/test/Dialect/MemRef/canonicalize.mlir | 14 +-
mlir/test/Dialect/MemRef/invalid.mlir | 2 +-
mlir/test/Dialect/MemRef/ops.mlir | 4 +-
mlir/test/Dialect/NVGPU/canonicalization.mlir | 4 +-
.../Dialect/OpenACC/acc-implicit-data.mlir | 2 +-
mlir/test/Transforms/canonicalize.mlir | 10 +-
23 files changed, 241 insertions(+), 182 deletions(-)
diff --git a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
index 70180c101407a..e0e5a28bc9504 100644
--- a/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
+++ b/mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td
@@ -2422,8 +2422,11 @@ def MemRef_ViewOp : MemRef_Op<"view", [
* A single dynamic byte-shift operand must be specified which represents a
a shift of the base 1-D memref pointer from which to create the resulting
contiguous memref view with identity layout.
- * A dynamic size operand that must be specified for each dynamic dimension
- in the resulting view memref type.
+ * A size operand for each dimension in the resulting view memref type.
+ Sizes can be either static (specified as integer literals in the
+ `static_sizes` attribute) or dynamic (specified as SSA values in
+ `sizes`). The `static_sizes` attribute uses `ShapedType::kDynamic` as
+ sentinel for dynamic entries.
The "view" operation gives a structured indexing form to a flat 1-D buffer.
Unlike "subview" it can perform a type change. The type change behavior
@@ -2443,24 +2446,40 @@ def MemRef_ViewOp : MemRef_Op<"view", [
%0 = memref.alloc() : memref<2048xi8>
// ViewOp with dynamic offset and static sizes.
- %1 = memref.view %0[%offset_1024][] : memref<2048xi8> to memref<64x4xf32>
+ %1 = memref.view %0[%offset_1024][64, 4] : memref<2048xi8> to memref<64x4xf32>
- // ViewOp with dynamic offset and two dynamic size.
- %2 = memref.view %0[%offset_1024][%size0, %size1] :
+ // ViewOp with dynamic offset and mixed static/dynamic sizes.
+ %2 = memref.view %0[%offset_1024][%size0, 4, %size1] :
memref<2048xi8> to memref<?x4x?xf32>
```
}];
let arguments = (ins MemRefRankOf<[I8], [1]>:$source,
Index:$byte_shift,
- Variadic<Index>:$sizes);
+ Variadic<Index>:$sizes,
+ DenseI64ArrayAttr:$static_sizes);
let results = (outs AnyMemRef);
+ let builders = [
+ // Build a ViewOp with mixed static and dynamic entries.
+ OpBuilder<(ins "MemRefType":$resultType, "Value":$source,
+ "Value":$byte_shift, "ArrayRef<OpFoldResult>":$sizes,
+ CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs)>,
+ // Build a ViewOp with static entries.
+ OpBuilder<(ins "MemRefType":$resultType, "Value":$source,
+ "Value":$byte_shift, "ArrayRef<int64_t>":$sizes,
+ CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs)>,
+ // Build a ViewOp with dynamic entries.
+ OpBuilder<(ins "MemRefType":$resultType, "Value":$source,
+ "Value":$byte_shift, "ValueRange":$sizes,
+ CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs)>
+ ];
+
let extraClassDeclaration = [{
/// The result of a view is always a memref.
MemRefType getType() { return ::llvm::cast<MemRefType>(getResult().getType()); }
- // Return both static and dynamic sizes as a list of `OpFoldResult`.
+ /// Return both static and dynamic sizes as a list of `OpFoldResult`.
SmallVector<OpFoldResult> getMixedSizes();
/// Returns the dynamic sizes for this view operation. This is redundant
@@ -2476,8 +2495,9 @@ def MemRef_ViewOp : MemRef_Op<"view", [
}];
let assemblyFormat = [{
- $source `[` $byte_shift `]` `` `[` $sizes `]` attr-dict
- `:` type($source) `to` type(results)
+ $source `[` $byte_shift `]` ``
+ custom<DynamicIndexList>($sizes, $static_sizes)
+ attr-dict `:` type($source) `to` type(results)
}];
let hasCanonicalizer = 1;
diff --git a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
index 91a0c4b55fa84..ceff719d81995 100644
--- a/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
+++ b/mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp
@@ -1863,10 +1863,11 @@ struct ViewOpLowering : public ConvertOpToLLVMPattern<memref::ViewOp> {
// Fields 4 and 5: Update sizes and strides.
Value stride = nullptr, nextSize = nullptr;
+ ArrayRef<int64_t> staticSizes = viewOp.getStaticSizes();
for (int i = viewMemRefType.getRank() - 1; i >= 0; --i) {
// Update size.
- Value size = getSize(rewriter, loc, viewMemRefType.getShape(),
- adaptor.getSizes(), i, indexType);
+ Value size =
+ getSize(rewriter, loc, staticSizes, adaptor.getSizes(), i, indexType);
targetMemRef.setSize(rewriter, loc, i, size);
// Update stride.
stride =
diff --git a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
index 844e6183cff06..465886507fc90 100644
--- a/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
+++ b/mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp
@@ -3741,6 +3741,33 @@ void ViewOp::getAsmResultNames(function_ref<void(Value, StringRef)> setNameFn) {
setNameFn(getResult(), "view");
}
+void ViewOp::build(OpBuilder &b, OperationState &result, MemRefType resultType,
+ Value source, Value byte_shift, ArrayRef<OpFoldResult> sizes,
+ ArrayRef<NamedAttribute> attrs) {
+ SmallVector<int64_t> staticSizes;
+ SmallVector<Value> dynamicSizes;
+ dispatchIndexOpFoldResults(sizes, dynamicSizes, staticSizes);
+ result.addAttributes(attrs);
+ build(b, result, resultType, source, byte_shift, dynamicSizes,
+ b.getDenseI64ArrayAttr(staticSizes));
+}
+
+void ViewOp::build(OpBuilder &b, OperationState &result, MemRefType resultType,
+ Value source, Value byte_shift, ArrayRef<int64_t> sizes,
+ ArrayRef<NamedAttribute> attrs) {
+ SmallVector<OpFoldResult> sizeValues = llvm::map_to_vector<4>(
+ sizes, [&](int64_t v) -> OpFoldResult { return b.getI64IntegerAttr(v); });
+ build(b, result, resultType, source, byte_shift, sizeValues, attrs);
+}
+
+void ViewOp::build(OpBuilder &b, OperationState &result, MemRefType resultType,
+ Value source, Value byte_shift, ValueRange sizes,
+ ArrayRef<NamedAttribute> attrs) {
+ SmallVector<OpFoldResult> sizeValues =
+ llvm::map_to_vector<4>(sizes, [](Value v) -> OpFoldResult { return v; });
+ build(b, result, resultType, source, byte_shift, sizeValues, attrs);
+}
+
LogicalResult ViewOp::verify() {
auto baseType = llvm::cast<MemRefType>(getOperand(0).getType());
auto viewType = getType();
@@ -3759,9 +3786,31 @@ LogicalResult ViewOp::verify() {
"type ")
<< baseType << " and view memref type " << viewType;
- // Verify that we have the correct number of sizes for the result type.
- if (failed(verifyDynamicDimensionCount(getOperation(), viewType, getSizes())))
- return failure();
+ // Verify that static_sizes length matches the result type rank.
+ unsigned resultRank = viewType.getRank();
+ if (getStaticSizes().size() != resultRank)
+ return emitError("incorrect number of size values, has ")
+ << getStaticSizes().size() << ", expected " << resultRank;
+
+ // Verify that the number of dynamic sizes matches the number of dynamic
+ // entries in static_sizes.
+ unsigned numDynamicSizes = llvm::count_if(
+ getStaticSizes(), [](int64_t v) { return ShapedType::isDynamic(v); });
+ if (getSizes().size() != numDynamicSizes)
+ return emitError("incorrect number of dynamic sizes, has ")
+ << getSizes().size() << ", expected " << numDynamicSizes;
+
+ // Verify consistency between static_sizes and result type shape.
+ unsigned idx = 0;
+ for (auto [resultSize, expectedSize] :
+ llvm::zip_equal(viewType.getShape(), getStaticSizes())) {
+ if (!ShapedType::isDynamic(expectedSize) &&
+ !ShapedType::isDynamic(resultSize) && resultSize != expectedSize)
+ return emitError("expected result type with size = ")
+ << expectedSize << " instead of " << resultSize
+ << " in dim = " << idx;
+ ++idx;
+ }
return success();
}
@@ -3780,82 +3829,71 @@ OpFoldResult ViewOp::fold(FoldAdaptor adaptor) {
}
SmallVector<OpFoldResult> ViewOp::getMixedSizes() {
- SmallVector<OpFoldResult> result;
- unsigned ctr = 0;
Builder b(getContext());
- for (int64_t dim : getType().getShape()) {
- if (ShapedType::isDynamic(dim)) {
- result.push_back(getSizes()[ctr++]);
- } else {
- result.push_back(b.getIndexAttr(dim));
- }
- }
- return result;
+ return getMixedValues(getStaticSizes(), getSizes(), b);
}
namespace {
-/// Given a memref type and a range of values that defines its dynamic
-/// dimension sizes, turn all dynamic sizes that have a constant value into
-/// static dimension sizes.
+/// Given a memref type and its associated static_sizes attribute and the
+/// dynamic size operands, turn all dynamic sizes that have a constant value
+/// into static sizes. Returns a new MemRefType with the folded shape and
+/// populates `foldedSizes` with the surviving mixed static/dynamic sizes.
static MemRefType
-foldDynamicToStaticDimSizes(MemRefType type, ValueRange dynamicSizes,
- SmallVectorImpl<Value> &foldedDynamicSizes) {
- SmallVector<int64_t> staticShape(type.getShape());
+foldDynamicToStaticDimSizes(MemRefType type, ArrayRef<int64_t> staticSizes,
+ ValueRange dynamicSizes,
+ SmallVectorImpl<OpFoldResult> &foldedSizes) {
+ SmallVector<int64_t> newShape(staticSizes);
assert(type.getNumDynamicDims() == dynamicSizes.size() &&
"incorrect number of dynamic sizes");
+ Builder builder(type.getContext());
- // Compute new static and dynamic sizes.
- unsigned ctr = 0;
- for (auto [dim, dimSize] : llvm::enumerate(type.getShape())) {
- if (ShapedType::isStatic(dimSize))
+ unsigned dynamicIdx = 0;
+ for (auto [dim, dimSize] : llvm::enumerate(staticSizes)) {
+ if (ShapedType::isStatic(dimSize)) {
+ foldedSizes.push_back(builder.getI64IntegerAttr(dimSize));
+ continue;
+ }
+ assert(ShapedType::isDynamic(dimSize) && "expected dynamic size");
+ Value dynVal = dynamicSizes[dynamicIdx++];
+ std::optional<int64_t> cst = getConstantIntValue(dynVal);
+ if (!cst.has_value() || cst.value() < 0) {
+ foldedSizes.push_back(dynVal);
continue;
-
- Value dynamicSize = dynamicSizes[ctr++];
- if (auto cst = getConstantIntValue(dynamicSize)) {
- // Dynamic size must be non-negative.
- if (cst.value() < 0) {
- foldedDynamicSizes.push_back(dynamicSize);
- continue;
- }
- staticShape[dim] = cst.value();
- } else {
- foldedDynamicSizes.push_back(dynamicSize);
}
+ newShape[dim] = cst.value();
+ foldedSizes.push_back(builder.getI64IntegerAttr(cst.value()));
}
- return MemRefType::Builder(type).setShape(staticShape);
+ return MemRefType::Builder(type).setShape(newShape);
}
-/// Change the result type of a `memref.view` by making originally dynamic
-/// dimensions static when their sizes come from `constant` ops.
+/// Fold dynamic sizes into static_sizes when they are constants.
/// Example:
/// ```
/// %c5 = arith.constant 5: index
-/// %0 = memref.view %src[%offset][%c5] : memref<?xi8> to memref<?x4xf32>
+/// %0 = memref.view %src[%offset][%c5, 4] : memref<?xi8> to memref<?x4xf32>
/// ```
/// to
/// ```
-/// %0 = memref.view %src[%offset][] : memref<?xi8> to memref<5x4xf32>
+/// %0 = memref.view %src[%offset][5, 4] : memref<?xi8> to memref<5x4xf32>
/// ```
struct ViewOpShapeFolder : public OpRewritePattern<ViewOp> {
using Base::Base;
LogicalResult matchAndRewrite(ViewOp viewOp,
PatternRewriter &rewriter) const override {
- SmallVector<Value> foldedDynamicSizes;
+ SmallVector<OpFoldResult> foldedSizes;
MemRefType resultType = viewOp.getType();
MemRefType foldedMemRefType = foldDynamicToStaticDimSizes(
- resultType, viewOp.getSizes(), foldedDynamicSizes);
+ resultType, viewOp.getStaticSizes(), viewOp.getSizes(), foldedSizes);
// Stop here if no dynamic size was promoted to static.
if (foldedMemRefType == resultType)
return failure();
- // Create new ViewOp.
- auto newViewOp = ViewOp::create(rewriter, viewOp.getLoc(), foldedMemRefType,
- viewOp.getSource(), viewOp.getByteShift(),
- foldedDynamicSizes);
- // Insert a cast so we have the same type as the old memref type.
+ auto newViewOp =
+ ViewOp::create(rewriter, viewOp.getLoc(), foldedMemRefType,
+ viewOp.getSource(), viewOp.getByteShift(), foldedSizes);
rewriter.replaceOpWithNewOp<CastOp>(viewOp, resultType, newViewOp);
return success();
}
@@ -3873,7 +3911,7 @@ struct ViewOpMemrefCastFolder : public OpRewritePattern<ViewOp> {
rewriter.replaceOpWithNewOp<ViewOp>(
viewOp, viewOp.getType(), memrefCastOp.getSource(),
- viewOp.getByteShift(), viewOp.getSizes());
+ viewOp.getByteShift(), viewOp.getMixedSizes());
return success();
}
};
diff --git a/mlir/test/Analysis/test-alias-analysis.mlir b/mlir/test/Analysis/test-alias-analysis.mlir
index d71adee05c7a3..cba9abc20c652 100644
--- a/mlir/test/Analysis/test-alias-analysis.mlir
+++ b/mlir/test/Analysis/test-alias-analysis.mlir
@@ -228,7 +228,7 @@ func.func @view_like(%arg: memref<2xf32>, %size: index) attributes {test.ptr = "
%c0 = arith.constant 0 : index
%2 = memref.alloca (%size) {test.ptr = "alloca_1"} : memref<?xi8>
- %3 = memref.view %2[%c0][] {test.ptr = "view"} : memref<?xi8> to memref<8x64xf32>
+ %3 = memref.view %2[%c0][8, 64] {test.ptr = "view"} : memref<?xi8> to memref<8x64xf32>
return
}
diff --git a/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir b/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
index 771f3185904bb..9d2c31c48ffa8 100644
--- a/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
+++ b/mlir/test/Conversion/GPUCommon/lower-memory-space-attrs.mlir
@@ -53,7 +53,7 @@ gpu.module @kernel {
gpu.func @dynamic_shmem_with_vector(%arg1: memref<1xf32>) {
%0 = arith.constant 0 : index
%1 = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %2 = memref.view %1[%0][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<1xf32, #gpu.address_space<workgroup>>
+ %2 = memref.view %1[%0][1] : memref<?xi8, #gpu.address_space<workgroup>> to memref<1xf32, #gpu.address_space<workgroup>>
%3 = vector.load %2[%0] : memref<1xf32, #gpu.address_space<workgroup>>, vector<1xf32>
vector.store %3, %arg1[%0] : memref<1xf32>, vector<1xf32>
gpu.return
@@ -73,7 +73,7 @@ gpu.module @kernel {
gpu.func @dynamic_shmem(%arg0: f32) {
%0 = arith.constant 0 : index
%1 = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %2 = memref.view %1[%0][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<4xf32, #gpu.address_space<workgroup>>
+ %2 = memref.view %1[%0][4] : memref<?xi8, #gpu.address_space<workgroup>> to memref<4xf32, #gpu.address_space<workgroup>>
memref.store %arg0, %2[%0] : memref<4xf32, #gpu.address_space<workgroup>>
gpu.return
}
diff --git a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
index 0cbe064572911..dedbd3ee15576 100644
--- a/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
+++ b/mlir/test/Conversion/MemRefToLLVM/memref-to-llvm.mlir
@@ -51,7 +51,7 @@ func.func @view(%arg0 : index, %arg1 : index, %arg2 : index) {
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[3, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
// CHECK: llvm.mul %{{.*}}, %[[ARG1]]
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[4, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
- %3 = memref.view %0[%arg2][%arg1] : memref<2048xi8> to memref<4x?xf32>
+ %3 = memref.view %0[%arg2][4, %arg1] : memref<2048xi8> to memref<4x?xf32>
// Test static sizes.
// CHECK: llvm.mlir.poison : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
@@ -68,7 +68,7 @@ func.func @view(%arg0 : index, %arg1 : index, %arg2 : index) {
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[3, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
// CHECK: llvm.mlir.constant(4 : index) : i64
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[4, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
- %5 = memref.view %0[%arg2][] : memref<2048xi8> to memref<64x4xf32>
+ %5 = memref.view %0[%arg2][64, 4] : memref<2048xi8> to memref<64x4xf32>
// Test view memory space.
// CHECK: llvm.mlir.constant(2048 : index) : i64
@@ -89,7 +89,7 @@ func.func @view(%arg0 : index, %arg1 : index, %arg2 : index) {
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[3, 0] : !llvm.struct<(ptr<4>, ptr<4>, i64, array<2 x i64>, array<2 x i64>)>
// CHECK: llvm.mlir.constant(4 : index) : i64
// CHECK: llvm.insertvalue %{{.*}}, %{{.*}}[4, 0] : !llvm.struct<(ptr<4>, ptr<4>, i64, array<2 x i64>, array<2 x i64>)>
- %7 = memref.view %6[%arg2][] : memref<2048xi8, 4> to memref<64x4xf32, 4>
+ %7 = memref.view %6[%arg2][64, 4] : memref<2048xi8, 4> to memref<64x4xf32, 4>
return
}
@@ -128,7 +128,7 @@ func.func @view_empty_memref(%offset: index, %mem: memref<0xi8>) {
// CHECK-INTERFACE: llvm.insertvalue %{{.*}}, %{{.*}}[3, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
// CHECK-INTERFACE: llvm.mlir.constant(4 : index) : i64
// CHECK-INTERFACE: = llvm.insertvalue %{{.*}}, %{{.*}}[4, 0] : !llvm.struct<(ptr, ptr, i64, array<2 x i64>, array<2 x i64>)>
- %0 = memref.view %mem[%offset][] : memref<0xi8> to memref<0x4xf32>
+ %0 = memref.view %mem[%offset][0, 4] : memref<0xi8> to memref<0x4xf32>
return
}
diff --git a/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir b/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir
index fa683175693be..ec1e67f653951 100644
--- a/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir
+++ b/mlir/test/Conversion/XeGPUToXeVM/loadstore_matrix.mlir
@@ -37,7 +37,7 @@ gpu.module @test_kernel [#xevm.target<chip = "pvc">] {
//CHECK-LABEL: load_store_matrix_plain_2d_input
gpu.func @load_store_matrix_plain_2d_input(%arg0: memref<1024xi8, 3>) -> f32 {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][]: memref<1024xi8, 3> to memref<64x32xf32, 3>
+ %view = memref.view %arg0[%c0][64, 32] : memref<1024xi8, 3> to memref<64x32xf32, 3>
%subview = memref.subview %view[32, 0] [32, 32] [1, 1] : memref<64x32xf32, 3> to memref<32x32xf32, strided<[32, 1], offset: 1024>, 3>
diff --git a/mlir/test/Dialect/Func/func-transform-invalid.mlir b/mlir/test/Dialect/Func/func-transform-invalid.mlir
index 29bd58ab52742..b58cb7a3255f4 100644
--- a/mlir/test/Dialect/Func/func-transform-invalid.mlir
+++ b/mlir/test/Dialect/Func/func-transform-invalid.mlir
@@ -3,9 +3,9 @@
module {
func.func private @func_with_reverse_order_no_result_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
}
@@ -25,9 +25,9 @@ module attributes {transform.with_named_sequence} {
module {
func.func private @func_with_reverse_order_no_result_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
}
@@ -47,9 +47,9 @@ module attributes {transform.with_named_sequence} {
module {
func.func private @func_with_reverse_order_no_result_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
}
@@ -69,9 +69,9 @@ module attributes {transform.with_named_sequence} {
module {
func.func private @func_with_reverse_order_no_result_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
}
@@ -90,9 +90,9 @@ module attributes {transform.with_named_sequence} {
func.func private @func_with_no_duplicate_args(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
@@ -113,9 +113,9 @@ module attributes {transform.with_named_sequence} {
func.func private @func_not_found(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
@@ -131,8 +131,8 @@ module attributes {transform.with_named_sequence} {
func.func private @func_with_multiple_calls(%arg0: memref<1xi8, 1>, %arg1: memref<1xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
return
}
@@ -158,8 +158,8 @@ module attributes {transform.with_named_sequence} {
func.func private @func_with_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<1xi8, 1>) {
%c0 = arith.constant 0 : index
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg1[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
return
}
diff --git a/mlir/test/Dialect/Func/func-transform.mlir b/mlir/test/Dialect/Func/func-transform.mlir
index 8a71511e3ed5b..2a69d285ca800 100644
--- a/mlir/test/Dialect/Func/func-transform.mlir
+++ b/mlir/test/Dialect/Func/func-transform.mlir
@@ -126,12 +126,12 @@ module {
func.func private @func_with_reverse_order_no_result_no_calls(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
// CHECK: %[[C0:.*]] = arith.constant 0 : index
%c0 = arith.constant 0 : index
- // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<3xi8, 1> to memref<3xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}3] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
}
@@ -152,12 +152,12 @@ module {
func.func private @func_with_reverse_order_no_result(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) {
// CHECK: %[[C0:.*]] = arith.constant 0 : index
%c0 = arith.constant 0 : index
- // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<3xi8, 1> to memref<3xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}3] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
return
}
@@ -186,12 +186,12 @@ module {
func.func private @func_with_reverse_order(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<3xi8, 1>) -> (memref<1xi8, 1>, memref<2xi8, 1>) {
// CHECK: %[[C0:.*]] = arith.constant 0 : index
%c0 = arith.constant 0 : index
- // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[RET_1:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<3xi8, 1> to memref<3xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[RET_1:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}3] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
// CHECK: return %[[RET_1]], %[[RET_0]] : memref<2xi8, 1>, memref<1xi8, 1>
return %view, %view0 : memref<1xi8, 1>, memref<2xi8, 1>
}
@@ -222,12 +222,12 @@ module {
func.func private @func_with_reverse_order_with_attr(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>{transform.readonly}, %arg2: memref<3xi8, 1>) -> (memref<1xi8, 1>, memref<2xi8, 1>) {
// CHECK: %[[C0:.*]] = arith.constant 0 : index
%c0 = arith.constant 0 : index
- // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[RET_1:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<3xi8, 1> to memref<3xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
+ // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[RET_1:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[VAL_6:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}3] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view1 = memref.view %arg2[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
// CHECK: return %[[RET_1]], %[[RET_0]] : memref<2xi8, 1>, memref<1xi8, 1>
return %view, %view0 : memref<1xi8, 1>, memref<2xi8, 1>
}
@@ -256,12 +256,12 @@ module attributes {transform.with_named_sequence} {
// CHECK: func.func private @func_with_duplicate_args(%[[ARG0:.*]]: memref<1xi8, 1>, %[[ARG1:.*]]: memref<2xi8, 1>) {
func.func private @func_with_duplicate_args(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<1xi8, 1>) {
%c0 = arith.constant 0 : index
- // CHECK: %[[VAL_3:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0:.*]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view0 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view1 = memref.view %arg2[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[VAL_3:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0:.*]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[VAL_4:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view0 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[VAL_5:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view1 = memref.view %arg2[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
return
}
@@ -284,16 +284,16 @@ module attributes {transform.with_named_sequence} {
// CHECK: func.func private @func_with_complex_duplicate_args(%[[ARG0:.*]]: memref<1xi8, 1>, %[[ARG1:.*]]: memref<2xi8, 1>, %[[ARG2:.*]]: memref<3xi8, 1>) -> (memref<1xi8, 1>, memref<2xi8, 1>, memref<1xi8, 1>, memref<3xi8, 1>, memref<2xi8, 1>) {
func.func private @func_with_complex_duplicate_args(%arg0: memref<1xi8, 1>, %arg1: memref<2xi8, 1>, %arg2: memref<1xi8, 1>, %arg3: memref<3xi8, 1>, %arg4: memref<2xi8, 1>) -> (memref<1xi8, 1>, memref<2xi8, 1>, memref<1xi8, 1>, memref<3xi8, 1>, memref<2xi8, 1>) {
%c0 = arith.constant 0 : index
- // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0:.*]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view0 = memref.view %arg0[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[RET_1:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view1 = memref.view %arg1[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
- // CHECK: %[[RET_2:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]][] : memref<1xi8, 1> to memref<1xi8, 1>
- %view2 = memref.view %arg2[%c0][] : memref<1xi8, 1> to memref<1xi8, 1>
- // CHECK: %[[RET_3:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]][] : memref<3xi8, 1> to memref<3xi8, 1>
- %view3 = memref.view %arg3[%c0][] : memref<3xi8, 1> to memref<3xi8, 1>
- // CHECK: %[[RET_4:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]][] : memref<2xi8, 1> to memref<2xi8, 1>
- %view4 = memref.view %arg4[%c0][] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[RET_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0:.*]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view0 = memref.view %arg0[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[RET_1:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view1 = memref.view %arg1[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
+ // CHECK: %[[RET_2:.*]] = memref.view %[[ARG0]]{{\[}}%[[C0]]]{{\[}}1] : memref<1xi8, 1> to memref<1xi8, 1>
+ %view2 = memref.view %arg2[%c0][1] : memref<1xi8, 1> to memref<1xi8, 1>
+ // CHECK: %[[RET_3:.*]] = memref.view %[[ARG2]]{{\[}}%[[C0]]]{{\[}}3] : memref<3xi8, 1> to memref<3xi8, 1>
+ %view3 = memref.view %arg3[%c0][3] : memref<3xi8, 1> to memref<3xi8, 1>
+ // CHECK: %[[RET_4:.*]] = memref.view %[[ARG1]]{{\[}}%[[C0]]]{{\[}}2] : memref<2xi8, 1> to memref<2xi8, 1>
+ %view4 = memref.view %arg4[%c0][2] : memref<2xi8, 1> to memref<2xi8, 1>
// CHECK: return %[[RET_0]], %[[RET_1]], %[[RET_2]], %[[RET_3]], %[[RET_4]] : memref<1xi8, 1>, memref<2xi8, 1>, memref<1xi8, 1>, memref<3xi8, 1>, memref<2xi8, 1>
return %view0, %view1, %view2, %view3, %view4 : memref<1xi8, 1>, memref<2xi8, 1>, memref<1xi8, 1>, memref<3xi8, 1>, memref<2xi8, 1>
}
diff --git a/mlir/test/Dialect/GPU/dynamic-shared-memory.mlir b/mlir/test/Dialect/GPU/dynamic-shared-memory.mlir
index 810f7d2c3ebda..b6ba4796cfedc 100644
--- a/mlir/test/Dialect/GPU/dynamic-shared-memory.mlir
+++ b/mlir/test/Dialect/GPU/dynamic-shared-memory.mlir
@@ -14,10 +14,10 @@ gpu.module @modules {
%shmem = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
%shmem2 = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %0 = memref.view %shmem[%c8192][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
+ %0 = memref.view %shmem[%c8192][32, 64] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
"test.use.shared.memory"(%0) : (memref<32x64xf32, #gpu.address_space<workgroup>>) -> ()
- %1 = memref.view %shmem[%c16384][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
+ %1 = memref.view %shmem[%c16384][32, 64] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
"test.use.shared.memory"(%1) : (memref<32x64xf32, #gpu.address_space<workgroup>>) -> ()
// CHECK-DAG: %[[S0:.+]] = llvm.mlir.constant(32 : index) : i64
@@ -52,7 +52,7 @@ gpu.module @modules {
gpu.func @gpu_device_function() {
%c8192 = arith.constant 8192 : index
%shmem = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %0 = memref.view %shmem[%c8192][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
+ %0 = memref.view %shmem[%c8192][32, 64] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
"test.use.shared.memory"(%0) : (memref<32x64xf32, #gpu.address_space<workgroup>>) -> ()
// CHECK-DAG: %[[S0:.+]] = llvm.mlir.constant(32 : index) : i64
// CHECK-DAG: %[[S1:.+]] = llvm.mlir.constant(64 : index) : i64
@@ -78,7 +78,7 @@ gpu.module @modules {
func.func @func_device_function() {
%c8192 = arith.constant 8192 : index
%shmem = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %0 = memref.view %shmem[%c8192][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
+ %0 = memref.view %shmem[%c8192][32, 64] : memref<?xi8, #gpu.address_space<workgroup>> to memref<32x64xf32, #gpu.address_space<workgroup>>
"test.use.shared.memory"(%0) : (memref<32x64xf32, #gpu.address_space<workgroup>>) -> ()
// CHECK-DAG: %[[S0:.+]] = llvm.mlir.constant(32 : index) : i64
// CHECK-DAG: %[[S1:.+]] = llvm.mlir.constant(64 : index) : i64
diff --git a/mlir/test/Dialect/Linalg/canonicalize.mlir b/mlir/test/Dialect/Linalg/canonicalize.mlir
index 77c1c3da17166..a4e5f2c5e6315 100644
--- a/mlir/test/Dialect/Linalg/canonicalize.mlir
+++ b/mlir/test/Dialect/Linalg/canonicalize.mlir
@@ -7,7 +7,7 @@ func.func @memref_cast(%a: index, %b: index) -> memref<?x?xf32> {
%c8 = arith.constant 8 : index
%c16 = arith.constant 16 : index
%1 = memref.alloc (%b) : memref<?xi8>
- %2 = memref.view %1[%c0][] : memref<?xi8> to memref<16x16xf32>
+ %2 = memref.view %1[%c0][16, 16] : memref<?xi8> to memref<16x16xf32>
%3 = memref.cast %2 : memref<16x16xf32> to memref<?x?xf32>
// CHECK: linalg.matmul ins({{.*}}memref<16x16xf32>, memref<16x16xf32>) outs({{.*}}memref<16x16xf32>)
diff --git a/mlir/test/Dialect/Linalg/forward-vector-transfers.mlir b/mlir/test/Dialect/Linalg/forward-vector-transfers.mlir
index 3530770580782..4f6c101b1ef9b 100644
--- a/mlir/test/Dialect/Linalg/forward-vector-transfers.mlir
+++ b/mlir/test/Dialect/Linalg/forward-vector-transfers.mlir
@@ -48,7 +48,7 @@ func.func @testViewRead(%in: memref<? x f32>) -> vector<32 x f32> {
%c0 = arith.constant 0: index
%f0 = arith.constant 0.0: f32
%alloc = memref.alloc() : memref<128 x i8>
- %view = memref.view %alloc[%c0][] : memref<128 x i8> to memref<32 x f32>
+ %view = memref.view %alloc[%c0][32] : memref<128 x i8> to memref<32 x f32>
%subview = memref.subview %view[0][16][1] : memref<32 x f32> to memref<16 x f32>
memref.copy %in, %subview : memref<? x f32> to memref<16 x f32>
%0 = vector.transfer_read %view[%c0], %f0 {in_bounds = [true]} : memref<32 x f32>, vector<32 x f32>
@@ -67,7 +67,7 @@ func.func @testViewFillRead(%in: memref<? x f32>) -> vector<32 x f32> {
%c0 = arith.constant 0: index
%f0 = arith.constant 0.0: f32
%alloc = memref.alloc() : memref<128 x i8>
- %view = memref.view %alloc[%c0][] : memref<128 x i8> to memref<32 x f32>
+ %view = memref.view %alloc[%c0][32] : memref<128 x i8> to memref<32 x f32>
%subview = memref.subview %view[0][16][1] : memref<32 x f32> to memref<16 x f32>
linalg.fill ins(%f0 : f32) outs(%view : memref<32 x f32>)
memref.copy %in, %subview : memref<? x f32> to memref<16 x f32>
@@ -105,7 +105,7 @@ func.func @testViewWrite(%vec: vector<32 x f32>, %out: memref<? x f32>) {
%c0 = arith.constant 0: index
%f0 = arith.constant 0.0: f32
%alloc = memref.alloc() : memref<128 x i8>
- %view = memref.view %alloc[%c0][] : memref<128 x i8> to memref<32 x f32>
+ %view = memref.view %alloc[%c0][32] : memref<128 x i8> to memref<32 x f32>
%subview = memref.subview %view[0][16][1] : memref<32 x f32> to memref<16 x f32>
vector.transfer_write %vec, %view[%c0] {in_bounds = [true]} : vector<32 x f32>, memref<32 x f32>
memref.copy %subview, %out : memref<16 x f32> to memref<? x f32>
diff --git a/mlir/test/Dialect/Linalg/loops.mlir b/mlir/test/Dialect/Linalg/loops.mlir
index efe8010cffc91..86c9a8dc89696 100644
--- a/mlir/test/Dialect/Linalg/loops.mlir
+++ b/mlir/test/Dialect/Linalg/loops.mlir
@@ -22,9 +22,9 @@ func.func @matmul(%arg0: memref<?xi8>, %M: index, %N: index, %K: index) {
// CHECK-SAME: [[M:arg[0-9]+]]: index
// CHECK-SAME: [[N:arg[0-9]+]]: index
// CHECK-SAME: [[K:arg[0-9]+]]: index
-// CHECK: %[[A:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECK: %[[B:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECK: %[[C:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECK: %[[A:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECK: %[[B:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECK: %[[C:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
// CHECK: scf.for {{.*}} to %[[M]]
// CHECK: scf.for {{.*}} to %[[N]]
// CHECK: scf.for {{.*}} to %[[K]]
@@ -39,9 +39,9 @@ func.func @matmul(%arg0: memref<?xi8>, %M: index, %N: index, %K: index) {
// CHECKPARALLEL-SAME: [[M:arg[0-9]+]]: index
// CHECKPARALLEL-SAME: [[N:arg[0-9]+]]: index
// CHECKPARALLEL-SAME: [[K:arg[0-9]+]]: index
-// CHECKPARALLEL: %[[A:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECKPARALLEL: %[[B:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECKPARALLEL: %[[C:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECKPARALLEL: %[[A:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECKPARALLEL: %[[B:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECKPARALLEL: %[[C:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
// CHECKPARALLEL: scf.parallel {{.*}} to (%[[M]], %[[N]]) step (%{{.*}}, %{{.*}} {
// CHECKPARALLEL: scf.for {{.*}} to %[[K]]
// CHECKPARALLEL-DAG: %[[a:.*]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<?x?xf32>
@@ -66,9 +66,9 @@ func.func @matvec(%arg0: memref<?xi8>, %M: index, %N: index) {
// CHECK-LABEL: func @matvec(%{{.*}}: memref<?xi8>,
// CHECK-SAME: [[M:arg[0-9]+]]: index
// CHECK-SAME: [[K:arg[0-9]+]]: index
-// CHECK: %[[A:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECK: %[[B:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?xf32>
-// CHECK: %[[C:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?xf32>
+// CHECK: %[[A:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECK: %[[B:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}] : memref<?xi8> to memref<?xf32>
+// CHECK: %[[C:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}] : memref<?xi8> to memref<?xf32>
// CHECK: scf.for {{.*}} to %[[M]]
// CHECK: scf.for {{.*}} to %[[K]]
// CHECK-DAG: %[[a:.*]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<?x?xf32>
@@ -81,9 +81,9 @@ func.func @matvec(%arg0: memref<?xi8>, %M: index, %N: index) {
// CHECKPARALLEL-LABEL: func @matvec(%{{.*}}: memref<?xi8>,
// CHECKPARALLEL-SAME: [[M:arg[0-9]+]]: index
// CHECKPARALLEL-SAME: [[K:arg[0-9]+]]: index
-// CHECKPARALLEL: %[[A:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?x?xf32>
-// CHECKPARALLEL: %[[B:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?xf32>
-// CHECKPARALLEL: %[[C:.*]] = memref.view %{{.*}}[{{.*}}] : memref<?xi8> to memref<?xf32>
+// CHECKPARALLEL: %[[A:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}, {{.*}}] : memref<?xi8> to memref<?x?xf32>
+// CHECKPARALLEL: %[[B:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}] : memref<?xi8> to memref<?xf32>
+// CHECKPARALLEL: %[[C:.*]] = memref.view %{{.*}}[{{.*}}][{{.*}}] : memref<?xi8> to memref<?xf32>
// CHECKPARALLEL: scf.parallel (%{{.*}}) = (%{{.*}}) to (%[[M]]) step (%{{.*}}) {
// CHECKPARALLEL: scf.for {{.*}} to %[[K]]
// CHECKPARALLEL-DAG: %[[a:.*]] = memref.load %[[A]][%{{.*}}, %{{.*}}] : memref<?x?xf32>
diff --git a/mlir/test/Dialect/Linalg/promote.mlir b/mlir/test/Dialect/Linalg/promote.mlir
index bab606c3a8169..707d4fb68929a 100644
--- a/mlir/test/Dialect/Linalg/promote.mlir
+++ b/mlir/test/Dialect/Linalg/promote.mlir
@@ -41,15 +41,15 @@ func.func @matmul_f32(%A: memref<?xi8>, %M: index, %N: index, %K: index) {
// CHECK: %[[vC:.*]] = memref.subview {{.*}} : memref<?x?xf32>
///
// CHECK: %[[tmpA:.*]] = memref.alloca() : memref<32xi8>
-// CHECK: %[[fullA:.*]] = memref.view %[[tmpA]][{{.*}}][{{.*}}] : memref<32xi8> to memref<?x?xf32>
+// CHECK: %[[fullA:.*]] = memref.view %[[tmpA]][{{.*}}][{{.*}}, {{.*}}] : memref<32xi8> to memref<?x?xf32>
// CHECK: %[[partialA:.*]] = memref.subview %[[fullA]]{{.*}} : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
///
// CHECK: %[[tmpB:.*]] = memref.alloca() : memref<48xi8>
-// CHECK: %[[fullB:.*]] = memref.view %[[tmpB]][{{.*}}][{{.*}}] : memref<48xi8> to memref<?x?xf32>
+// CHECK: %[[fullB:.*]] = memref.view %[[tmpB]][{{.*}}][{{.*}}, {{.*}}] : memref<48xi8> to memref<?x?xf32>
// CHECK: %[[partialB:.*]] = memref.subview %[[fullB]]{{.*}} : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
///
// CHECK: %[[tmpC:.*]] = memref.alloca() : memref<24xi8>
-// CHECK: %[[fullC:.*]] = memref.view %[[tmpC]][{{.*}}][{{.*}}] : memref<24xi8> to memref<?x?xf32>
+// CHECK: %[[fullC:.*]] = memref.view %[[tmpC]][{{.*}}][{{.*}}, {{.*}}] : memref<24xi8> to memref<?x?xf32>
// CHECK: %[[partialC:.*]] = memref.subview %[[fullC]]{{.*}} : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK: linalg.copy ins(%[[vA]] : memref<?x?xf32, strided<[?, 1], offset: ?>>) outs(%[[partialA]] : memref<?x?xf32, strided<[?, 1]>>)
@@ -111,15 +111,15 @@ func.func @matmul_f64(%A: memref<?xi8>, %M: index, %N: index, %K: index) {
// CHECK: %[[vC_f64:.*]] = memref.subview {{.*}} : memref<?x?xf64>
///
// CHECK: %[[tmpA_f64:.*]] = memref.alloc() : memref<64xi8>
-// CHECK: %[[fullA_f64:.*]] = memref.view %[[tmpA_f64]][{{.*}}][{{.*}}] : memref<64xi8> to memref<?x?xf64>
+// CHECK: %[[fullA_f64:.*]] = memref.view %[[tmpA_f64]][{{.*}}][{{.*}}, {{.*}}] : memref<64xi8> to memref<?x?xf64>
// CHECK: %[[partialA_f64:.*]] = memref.subview %[[fullA_f64]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf64> to memref<?x?xf64, strided<[?, 1]>>
///
// CHECK: %[[tmpB_f64:.*]] = memref.alloc() : memref<96xi8>
-// CHECK: %[[fullB_f64:.*]] = memref.view %[[tmpB_f64]][{{.*}}][{{.*}}] : memref<96xi8> to memref<?x?xf64>
+// CHECK: %[[fullB_f64:.*]] = memref.view %[[tmpB_f64]][{{.*}}][{{.*}}, {{.*}}] : memref<96xi8> to memref<?x?xf64>
// CHECK: %[[partialB_f64:.*]] = memref.subview %[[fullB_f64]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf64> to memref<?x?xf64, strided<[?, 1]>>
///
// CHECK: %[[tmpC_f64:.*]] = memref.alloc() : memref<48xi8>
-// CHECK: %[[fullC_f64:.*]] = memref.view %[[tmpC_f64]][{{.*}}][{{.*}}] : memref<48xi8> to memref<?x?xf64>
+// CHECK: %[[fullC_f64:.*]] = memref.view %[[tmpC_f64]][{{.*}}][{{.*}}, {{.*}}] : memref<48xi8> to memref<?x?xf64>
// CHECK: %[[partialC_f64:.*]] = memref.subview %[[fullC_f64]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf64> to memref<?x?xf64, strided<[?, 1]>>
// CHECK: linalg.copy ins(%[[vA_f64]] : memref<?x?xf64, strided<[?, 1], offset: ?>>) outs(%[[partialA_f64]] : memref<?x?xf64, strided<[?, 1]>>)
diff --git a/mlir/test/Dialect/Linalg/promotion_options.mlir b/mlir/test/Dialect/Linalg/promotion_options.mlir
index dbc073c2665f9..59a72797c021e 100644
--- a/mlir/test/Dialect/Linalg/promotion_options.mlir
+++ b/mlir/test/Dialect/Linalg/promotion_options.mlir
@@ -20,11 +20,11 @@ func.func @gemm(%a : memref<?x?xf32>, %b : memref<?x?xf32>, %c : memref<?x?xf32>
// CHECK: %[[svC:.+]] = memref.subview %[[ARG2]]
// CHECK: %[[tmpA:.*]] = memref.alloc() : memref<1024xi8>
-// CHECK: %[[VA:.*]] = memref.view %[[tmpA]][%[[C0]]][] : memref<1024xi8> to memref<16x16xf32>
+// CHECK: %[[VA:.*]] = memref.view %[[tmpA]][%[[C0]]][16, 16] : memref<1024xi8> to memref<16x16xf32>
// CHECK: %[[svAA:.+]] = memref.subview %[[VA]]
// CHECK: %[[tmpC:.*]] = memref.alloc() : memref<1024xi8>
-// CHECK: %[[VC:.*]] = memref.view %[[tmpC]][%[[C0]]][] : memref<1024xi8> to memref<16x16xf32>
+// CHECK: %[[VC:.*]] = memref.view %[[tmpC]][%[[C0]]][16, 16] : memref<1024xi8> to memref<16x16xf32>
// CHECK: %[[svCC:.+]] = memref.subview %[[VC]]
// CHECK: linalg.copy ins(%[[svA]] : memref<?x?xf32, strided<[?, 1], offset: ?>>) outs(%[[svAA]] : memref<?x?xf32, strided<[16, 1]>>)
diff --git a/mlir/test/Dialect/Linalg/roundtrip.mlir b/mlir/test/Dialect/Linalg/roundtrip.mlir
index bfb92c3289a49..b980273bcb9cd 100644
--- a/mlir/test/Dialect/Linalg/roundtrip.mlir
+++ b/mlir/test/Dialect/Linalg/roundtrip.mlir
@@ -18,9 +18,9 @@ func.func @views(%arg0: index) {
// CHECK-LABEL: func @views
// CHECK: arith.muli %{{.*}}, %{{.*}} : index
// CHECK-NEXT: memref.alloc(%{{.*}}) : memref<?xi8>
-// CHECK-NEXT: memref.view %{{.*}}[%{{.*}}][%{{.*}}] :
+// CHECK-NEXT: memref.view %{{.*}}[%{{.*}}][%{{.*}}, %{{.*}}] :
// CHECK-SAME: memref<?xi8> to memref<?x?xf32>
-// CHECK-NEXT: memref.view %{{.*}}[%{{.*}}][%{{.*}}] :
+// CHECK-NEXT: memref.view %{{.*}}[%{{.*}}][%{{.*}}, %{{.*}}] :
// CHECK-SAME: memref<?xi8> to memref<?x?xvector<4x4xf32>>
// CHECK-NEXT: memref.dealloc %{{.*}} : memref<?xi8>
diff --git a/mlir/test/Dialect/Linalg/transform-promotion.mlir b/mlir/test/Dialect/Linalg/transform-promotion.mlir
index 7c4cd623c742d..68babeb53b342 100644
--- a/mlir/test/Dialect/Linalg/transform-promotion.mlir
+++ b/mlir/test/Dialect/Linalg/transform-promotion.mlir
@@ -40,15 +40,15 @@ func.func @promote_subview_matmul(%arg0: memref<?x?xf32, strided<[?, 1], offset:
// CHECK: %[[s1:.*]] = memref.subview {{.*}}: memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
// CHECK: %[[s2:.*]] = memref.subview {{.*}}: memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
// CHECK: %[[a0:.*]] = memref.alloc() : memref<32000000xi8>
-// CHECK: %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<32000000xi8> to memref<?x?xf32>
+// CHECK: %[[v0:.*]] = memref.view %[[a0]][{{.*}}][{{.*}}, {{.*}}] : memref<32000000xi8> to memref<?x?xf32>
// CHECK: %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1]
// CHECK-SAME: memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK: %[[a1:.*]] = memref.alloc() : memref<48000000xi8>
-// CHECK: %[[v1:.*]] = memref.view %[[a1]]{{.*}} : memref<48000000xi8> to memref<?x?xf32>
+// CHECK: %[[v1:.*]] = memref.view %[[a1]][{{.*}}][{{.*}}, {{.*}}] : memref<48000000xi8> to memref<?x?xf32>
// CHECK: %[[l1:.*]] = memref.subview %[[v1]][0, 0] [%{{.*}}, %{{.*}}] [1, 1]
// CHECK-SAME: memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK: %[[a2:.*]] = memref.alloc() : memref<24000000xi8>
-// CHECK: %[[v2:.*]] = memref.view %[[a2]]{{.*}} : memref<24000000xi8> to memref<?x?xf32>
+// CHECK: %[[v2:.*]] = memref.view %[[a2]][{{.*}}][{{.*}}, {{.*}}] : memref<24000000xi8> to memref<?x?xf32>
// CHECK: %[[l2:.*]] = memref.subview %[[v2]][0, 0] [%{{.*}}, %{{.*}}] [1, 1]
// CHECK-SAME: memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK: linalg.copy ins(%[[s0]] : memref<?x?xf32, strided{{.*}}>) outs(%[[l0]] : memref<?x?xf32, strided{{.*}}>)
@@ -109,7 +109,7 @@ func.func @promote_first_subview_matmul(%arg0: memref<?x?xf32, strided<[?, 1], o
// CHECK: %[[s1:.*]] = memref.subview {{.*}}: memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
// CHECK: %[[s2:.*]] = memref.subview {{.*}}: memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
// CHECK: %[[a0:.*]] = memref.alloc() : memref<32000000xi8>
-// CHECK: %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<32000000xi8> to memref<?x?xf32>
+// CHECK: %[[v0:.*]] = memref.view %[[a0]][{{.*}}][{{.*}}, {{.*}}] : memref<32000000xi8> to memref<?x?xf32>
// CHECK: %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK-NOT: memref.alloc
// CHECK-NOT: memref.view
@@ -146,7 +146,7 @@ func.func @aligned_promote_fill(%arg0: memref<?x?xf32, strided<[?, 1], offset: ?
// CHECK: %[[cf:.*]] = arith.constant 1.{{.*}} : f32
// CHECK: %[[s0:.*]] = memref.subview {{.*}}: memref<?x?xf32, strided{{.*}}> to memref<?x?xf32, strided{{.*}}>
// CHECK: %[[a0:.*]] = memref.alloc() {alignment = 32 : i64} : memref<32000000xi8>
-// CHECK: %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<32000000xi8> to memref<?x?xf32>
+// CHECK: %[[v0:.*]] = memref.view %[[a0]][{{.*}}][{{.*}}, {{.*}}] : memref<32000000xi8> to memref<?x?xf32>
// CHECK: %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xf32> to memref<?x?xf32, strided<[?, 1]>>
// CHECK: linalg.fill ins({{.*}} : f32) outs(%[[v0]] : memref<?x?xf32>)
// CHECK: linalg.copy ins(%[[s0]] : memref<?x?xf32, strided{{.*}}>) outs(%[[l0]] : memref<?x?xf32, strided{{.*}}>)
@@ -179,7 +179,7 @@ func.func @aligned_promote_fill_complex(%arg0: memref<?x?xcomplex<f32>, strided<
// CHECK: %[[cc:.*]] = complex.create {{.*}} : complex<f32>
// CHECK: %[[s0:.*]] = memref.subview {{.*}}: memref<?x?xcomplex<f32>, strided{{.*}}> to memref<?x?xcomplex<f32>, strided{{.*}}>
// CHECK: %[[a0:.*]] = memref.alloc() {alignment = 32 : i64} : memref<64000000xi8>
-// CHECK: %[[v0:.*]] = memref.view %[[a0]]{{.*}} : memref<64000000xi8> to memref<?x?xcomplex<f32>>
+// CHECK: %[[v0:.*]] = memref.view %[[a0]][{{.*}}][{{.*}}, {{.*}}] : memref<64000000xi8> to memref<?x?xcomplex<f32>>
// CHECK: %[[l0:.*]] = memref.subview %[[v0]][0, 0] [%{{.*}}, %{{.*}}] [1, 1] : memref<?x?xcomplex<f32>> to memref<?x?xcomplex<f32>, strided<[?, 1]>>
// CHECK: linalg.fill ins({{.*}} : complex<f32>) outs(%[[v0]] : memref<?x?xcomplex<f32>>)
// CHECK: linalg.copy ins(%[[s0]] : memref<?x?xcomplex<f32>, strided{{.*}}>) outs(%[[l0]] : memref<?x?xcomplex<f32>, strided{{.*}}>)
diff --git a/mlir/test/Dialect/MemRef/canonicalize.mlir b/mlir/test/Dialect/MemRef/canonicalize.mlir
index 3cfea1e8cd961..74faf3bf91a9d 100644
--- a/mlir/test/Dialect/MemRef/canonicalize.mlir
+++ b/mlir/test/Dialect/MemRef/canonicalize.mlir
@@ -1582,7 +1582,7 @@ func.func @fold_view_same_source_result_types(%0: memref<128xi8>) -> memref<128x
%c0 = arith.constant 0 : index
// CHECK-NOT: memref.view
// CHECK: return %[[ARG]]
- %res = memref.view %0[%c0][] : memref<128xi8> to memref<128xi8>
+ %res = memref.view %0[%c0][128] : memref<128xi8> to memref<128xi8>
return %res : memref<128xi8>
}
@@ -1593,9 +1593,9 @@ func.func @fold_view_same_source_result_types(%0: memref<128xi8>) -> memref<128x
func.func @non_fold_view_non_zero_offset(%0: memref<128xi8>) -> memref<128xi8> {
%c1 = arith.constant 1 : index
// CHECK: %[[C1:.*]] = arith.constant 1 : index
- // CHECK: %[[RES:.*]] = memref.view %[[ARG]][%[[C1]]][] : memref<128xi8> to memref<128xi8>
+ // CHECK: %[[RES:.*]] = memref.view %[[ARG]][%[[C1]]][128] : memref<128xi8> to memref<128xi8>
// CHECK: return %[[RES]]
- %res = memref.view %0[%c1][] : memref<128xi8> to memref<128xi8>
+ %res = memref.view %0[%c1][128] : memref<128xi8> to memref<128xi8>
return %res : memref<128xi8>
}
@@ -1616,13 +1616,13 @@ func.func @non_fold_view_same_source_dynamic_size(%0: memref<?xi8>, %arg0 : inde
// CHECK-LABEL: func.func @replace_view_static_dims(
// CHECK-SAME: %[[ARG0:.*]]: memref<?xi8>, %[[ARG1:.*]]: index) -> memref<?x4xi32> {
-// CHECK: %[[VIEW_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[ARG1]]][] : memref<?xi8> to memref<5x4xi32>
+// CHECK: %[[VIEW_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[ARG1]]][5, 4] : memref<?xi8> to memref<5x4xi32>
// CHECK: %[[CAST_0:.*]] = memref.cast %[[VIEW_0]] : memref<5x4xi32> to memref<?x4xi32>
// CHECK: return %[[CAST_0]] : memref<?x4xi32>
// CHECK: }
func.func @replace_view_static_dims(%src: memref<?xi8>, %offset : index) -> memref<?x4xi32> {
%c5 = arith.constant 5: index
- %res = memref.view %src[%offset][%c5] : memref<?xi8> to memref<?x4xi32>
+ %res = memref.view %src[%offset][%c5, 4] : memref<?xi8> to memref<?x4xi32>
return %res : memref<?x4xi32>
}
@@ -1631,11 +1631,11 @@ func.func @replace_view_static_dims(%src: memref<?xi8>, %offset : index) -> memr
// CHECK-LABEL: func.func @non_replace_view_negative_static_dims(
// CHECK-SAME: %[[ARG0:.*]]: memref<?xi8>, %[[ARG1:.*]]: index) -> memref<?x4xi32> {
// CHECK: %[[CONSTANT_0:.*]] = arith.constant -1 : index
-// CHECK: %[[VIEW_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[ARG1]]]{{\[}}%[[CONSTANT_0]]] : memref<?xi8> to memref<?x4xi32>
+// CHECK: %[[VIEW_0:.*]] = memref.view %[[ARG0]]{{\[}}%[[ARG1]]]{{\[}}%[[CONSTANT_0]], 4] : memref<?xi8> to memref<?x4xi32>
// CHECK: return %[[VIEW_0]] : memref<?x4xi32>
// CHECK: }
func.func @non_replace_view_negative_static_dims(%src: memref<?xi8>, %offset : index) -> memref<?x4xi32> {
%c-1 = arith.constant -1: index
- %res = memref.view %src[%offset][%c-1] : memref<?xi8> to memref<?x4xi32>
+ %res = memref.view %src[%offset][%c-1, 4] : memref<?xi8> to memref<?x4xi32>
return %res : memref<?x4xi32>
}
diff --git a/mlir/test/Dialect/MemRef/invalid.mlir b/mlir/test/Dialect/MemRef/invalid.mlir
index af068d8ca8e95..d8fc9cf79a6bc 100644
--- a/mlir/test/Dialect/MemRef/invalid.mlir
+++ b/mlir/test/Dialect/MemRef/invalid.mlir
@@ -632,7 +632,7 @@ func.func @invalid_view(%arg0 : index, %arg1 : index, %arg2 : index) {
func.func @invalid_view(%arg0 : index, %arg1 : index, %arg2 : index) {
%0 = memref.alloc() : memref<2048xi8>
- // expected-error at +1 {{incorrect number of dynamic sizes, has 1, expected 2}}
+ // expected-error at +1 {{incorrect number of size values, has 1, expected 2}}
%1 = memref.view %0[%arg2][%arg0]
: memref<2048xi8> to memref<?x?xf32>
return
diff --git a/mlir/test/Dialect/MemRef/ops.mlir b/mlir/test/Dialect/MemRef/ops.mlir
index 14ac6a03d6ae0..3d9cb4885f949 100644
--- a/mlir/test/Dialect/MemRef/ops.mlir
+++ b/mlir/test/Dialect/MemRef/ops.mlir
@@ -291,12 +291,12 @@ func.func @memref_view(%arg0 : index, %arg1 : index, %arg2 : index) {
// Test one dynamic size and dynamic offset.
// CHECK: memref.view {{.*}} : memref<2048xi8> to memref<4x?xf32>
- %3 = memref.view %0[%arg2][%arg1] : memref<2048xi8> to memref<4x?xf32>
+ %3 = memref.view %0[%arg2][4, %arg1] : memref<2048xi8> to memref<4x?xf32>
// Test static sizes and static offset.
// CHECK: memref.view {{.*}} : memref<2048xi8> to memref<64x4xf32>
%c0 = arith.constant 0: index
- %5 = memref.view %0[%c0][] : memref<2048xi8> to memref<64x4xf32>
+ %5 = memref.view %0[%c0][64, 4] : memref<2048xi8> to memref<64x4xf32>
return
}
diff --git a/mlir/test/Dialect/NVGPU/canonicalization.mlir b/mlir/test/Dialect/NVGPU/canonicalization.mlir
index a7fbfd8067395..fc8c72e3dea90 100644
--- a/mlir/test/Dialect/NVGPU/canonicalization.mlir
+++ b/mlir/test/Dialect/NVGPU/canonicalization.mlir
@@ -15,13 +15,13 @@ gpu.module @main_kernel {
// CHECK: %[[S0:.+]] = gpu.thread_id x
// CHECK: %[[S1:.+]] = arith.cmpi eq, %[[S0]], %[[c0]] : index
// CHECK: %[[S2:.+]] = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- // CHECK: %[[S3:.+]] = memref.view %[[S2]][%[[c0]]][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x32xf32, #gpu.address_space<workgroup>>
+ // CHECK: %[[S3:.+]] = memref.view %[[S2]][%[[c0]]][128, 32] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x32xf32, #gpu.address_space<workgroup>>
// CHECK: nvgpu.tma.async.store %[[S3]] to %[[arg0]][%[[c0]], %[[c0]]], predicate = %[[S1]] : memref<128x32xf32, #gpu.address_space<workgroup>> -> <tensor = memref<128x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>
%c0 = arith.constant 0 : index
%0 = gpu.thread_id x
%1 = arith.cmpi eq, %0, %c0 : index
%2 = gpu.dynamic_shared_memory : memref<?xi8, #gpu.address_space<workgroup>>
- %view = memref.view %2[%c0][] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x32xf32, #gpu.address_space<workgroup>>
+ %view = memref.view %2[%c0][128, 32] : memref<?xi8, #gpu.address_space<workgroup>> to memref<128x32xf32, #gpu.address_space<workgroup>>
nvgpu.tma.async.store %view to %arg0[%c0, %c0], predicate = %1 : memref<128x32xf32, #gpu.address_space<workgroup>> -> <tensor = memref<128x32xf32, 3>, swizzle = none, l2promo = none, oob = zero, interleave = none>
nvvm.cp.async.bulk.commit.group
nvvm.cp.async.bulk.wait_group 0
diff --git a/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir b/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
index df0dbbfee8b1d..1e30178520398 100644
--- a/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-implicit-data.mlir
@@ -208,7 +208,7 @@ func.func @test_memref_view(%size: index) {
%c0 = arith.constant 0 : index
%buffer = memref.alloca(%size) : memref<?xi8>
%copyin = acc.copyin varPtr(%buffer : memref<?xi8>) -> memref<?xi8> {name = "buffer"}
- %view = memref.view %buffer[%c0][] : memref<?xi8> to memref<8x64xf32>
+ %view = memref.view %buffer[%c0][8, 64] : memref<?xi8> to memref<8x64xf32>
acc.kernels dataOperands(%copyin : memref<?xi8>) {
%c0_0 = arith.constant 0 : index
%c0_1 = arith.constant 0 : index
diff --git a/mlir/test/Transforms/canonicalize.mlir b/mlir/test/Transforms/canonicalize.mlir
index 8e02c06a0a293..b8ada88d2d59d 100644
--- a/mlir/test/Transforms/canonicalize.mlir
+++ b/mlir/test/Transforms/canonicalize.mlir
@@ -709,22 +709,22 @@ func.func @view(%arg0 : index) -> (f32, f32, f32, f32) {
%c15 = arith.constant 15 : index
// Test: fold constant sizes.
- // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][] : memref<2048xi8> to memref<7x11xf32>
+ // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][7, 11] : memref<2048xi8> to memref<7x11xf32>
%1 = memref.view %0[%c15][%c7, %c11] : memref<2048xi8> to memref<?x?xf32>
%r0 = memref.load %1[%c0, %c0] : memref<?x?xf32>
// Test: fold one constant size.
- // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][%arg0, %arg0] : memref<2048xi8> to memref<?x?x7xf32>
+ // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][%arg0, %arg0, 7] : memref<2048xi8> to memref<?x?x7xf32>
%2 = memref.view %0[%c15][%arg0, %arg0, %c7] : memref<2048xi8> to memref<?x?x?xf32>
%r1 = memref.load %2[%c0, %c0, %c0] : memref<?x?x?xf32>
// Test: preserve an existing static size.
- // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][] : memref<2048xi8> to memref<7x4xf32>
- %3 = memref.view %0[%c15][%c7] : memref<2048xi8> to memref<?x4xf32>
+ // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][7, 4] : memref<2048xi8> to memref<7x4xf32>
+ %3 = memref.view %0[%c15][%c7, 4] : memref<2048xi8> to memref<?x4xf32>
%r2 = memref.load %3[%c0, %c0] : memref<?x4xf32>
// Test: folding static alloc and memref.cast into a view.
- // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][] : memref<2048xi8> to memref<15x7xf32>
+ // CHECK: memref.view %[[ALLOC_MEM]][%[[C15]]][15, 7] : memref<2048xi8> to memref<15x7xf32>
%4 = memref.cast %0 : memref<2048xi8> to memref<?xi8>
%5 = memref.view %4[%c15][%c15, %c7] : memref<?xi8> to memref<?x?xf32>
%r3 = memref.load %5[%c0, %c0] : memref<?x?xf32>
More information about the Mlir-commits
mailing list