Merge google -> main (#7984)
* 6ea9e5f Remove note about IREE_PLATFORM_GOOGLE.
* 121a187 Fix MetalSPIRV smoke test to match mlir changes and enable the test across all..
* 74841ca Integrate LLVM at llvm/llvm-project@128c6ed
* daf0700 Integrate LLVM at llvm/llvm-project@9ebeac8
* 20b1f52 Integrate LLVM at llvm/llvm-project@5da6d26
* f9b7a1d [mlir] Switching accessors to prefixed form (NFC)
* 4959d11 Integrate LLVM at llvm/llvm-project@09f43c1
diff --git a/SUBMODULE_VERSIONS.txt b/SUBMODULE_VERSIONS.txt
index 4ca6920..dc1dcfe 100644
--- a/SUBMODULE_VERSIONS.txt
+++ b/SUBMODULE_VERSIONS.txt
@@ -4,7 +4,7 @@
aa533abfd4232b01f9e57041d70114d5a77e6de0 third_party/googletest
88b845dee001723c4a0db1fe5477de735b6d3bb0 third_party/liburing
f8f760f7387d2cc56a2fc7b1be313a3bf3f7f58c third_party/libyaml
-505d57486e57eb61e29bed6517de5152d208fede third_party/llvm-project
+128c6ed73b8f906a13ae908008c6f415415964bb third_party/llvm-project
78236bbc93fbf56ce048f1781a7610dfc86ec4c5 third_party/mlir-hlo
3f701faace7addc75d16dea8a6cd769fa5b3f260 third_party/musl
59aa99860c60bd171b9565e9920f125fdb749267 third_party/pybind11
diff --git a/iree/base/target_platform.h b/iree/base/target_platform.h
index 8052967..a15f80c 100644
--- a/iree/base/target_platform.h
+++ b/iree/base/target_platform.h
@@ -51,9 +51,6 @@
// IREE_PLATFORM_LINUX
// IREE_PLATFORM_MACOS
// IREE_PLATFORM_WINDOWS
-//
-// The special define IREE_PLATFORM_GOOGLE will be specified if the build
-// is being performed within the internal Google repository.
//==============================================================================
// IREE_ARCH_*
diff --git a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
index 3b94ddc..9b6f240 100644
--- a/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
+++ b/iree/compiler/Codegen/Common/BufferizationAnalysis.cpp
@@ -110,13 +110,6 @@
/// equivalence class as the source.
static LogicalResult analyseInterfaceLoadTensorOp(
IREE::Flow::DispatchTensorLoadOp loadOp, BufferizationPlan &plan) {
- if (!(loadOp.getMixedOffsets().empty() && loadOp.getMixedSizes().empty() &&
- loadOp.getMixedStrides().empty()) &&
- !canUsersHandleSubviews(loadOp)) {
- plan.insert(loadOp.source());
- plan.insert(loadOp.result());
- return success();
- }
plan.unionSets(loadOp.result(), loadOp.source());
return success();
}
@@ -150,22 +143,6 @@
static bool canSetStoreValueAndTargetAsEquivalent(
IREE::Flow::DispatchTensorStoreOp storeOp, BufferizationPlan &plan) {
Value value = storeOp.value();
- if (!(storeOp.getMixedOffsets().empty() && storeOp.getMixedSizes().empty() &&
- storeOp.getMixedStrides().empty())) {
- SmallVector<Value> mappedTensors = plan.getTensorsMappedToSameSet(value);
- for (auto v : mappedTensors) {
- // TODO(ravishankarm): At this point it is not clear why the following
- // restriction exists. It might have something to do with subviews and
- // reshapes not working well together, but there is no comment about why
- // this was added with the change that added this.
- Operation *op = v.getDefiningOp();
- if (op && isa<tensor::CollapseShapeOp, tensor::ExpandShapeOp>(
- v.getDefiningOp())) {
- return false;
- }
- }
- }
-
Value target = storeOp.target();
auto targetInterfaceOp =
getEquivalentOpOfType<IREE::HAL::InterfaceBindingSubspanOp>(target, plan);
@@ -357,19 +334,19 @@
static LogicalResult analyseScfForOp(scf::ForOp forOp,
BufferizationPlan &plan) {
- if (forOp.results().empty()) return success();
+ if (forOp.getResults().empty()) return success();
if (!llvm::all_of(forOp->getResultTypes(), [](Type resultType) {
return resultType.isa<RankedTensorType>();
})) {
return success();
}
- auto yeildOp = cast<scf::YieldOp>(forOp.getBody()->getTerminator());
+ auto yieldOp = cast<scf::YieldOp>(forOp.getBody()->getTerminator());
auto regionArgs = forOp.getRegionIterArgs();
- auto initArgs = forOp.initArgs();
- for (int i = 0; i < yeildOp.results().size(); ++i) {
- Value yieldTensor = yeildOp.results()[i];
- Value resultTensor = forOp.results()[i];
+ auto initArgs = forOp.getInitArgs();
+ for (int i = 0; i < yieldOp.getResults().size(); ++i) {
+ Value yieldTensor = yieldOp.getResults()[i];
+ Value resultTensor = forOp.getResults()[i];
Value initArg = initArgs[i];
Value arg = regionArgs[i];
// Always tie the yield, the result tensor, and the region arg
diff --git a/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp b/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
index 213421b..7b54aca 100644
--- a/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
+++ b/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
@@ -115,9 +115,9 @@
initArgs[it.index()] = rewriter.clone(*op, mapping)->getResult(0);
}
if (iteratorFolded.empty()) return failure();
- auto newLoop =
- rewriter.create<scf::ForOp>(forOp.getLoc(), forOp.lowerBound(),
- forOp.upperBound(), forOp.step(), initArgs);
+ auto newLoop = rewriter.create<scf::ForOp>(
+ forOp.getLoc(), forOp.getLowerBound(), forOp.getUpperBound(),
+ forOp.getStep(), initArgs);
transferBody(forOp.getBody(), newLoop.getBody(), returnValues, rewriter);
// Replace the operation by the new one.
@@ -170,8 +170,8 @@
// Create a new loop with the casted init values. This also creates
// induction variables with proper type.
auto newLoop = rewriter.create<scf::ForOp>(
- forOp.getLoc(), forOp.lowerBound(), forOp.upperBound(), forOp.step(),
- ivInitValues);
+ forOp.getLoc(), forOp.getLowerBound(), forOp.getUpperBound(),
+ forOp.getStep(), ivInitValues);
// Move all operations to the new for op. This also replaces block
// arguments. to the new block arguments.
diff --git a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
index 9b0540e..39e3820 100644
--- a/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
+++ b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
@@ -144,6 +144,68 @@
layout, memorySpace);
}
+/// Checks if the offsets, sizes and strides with src, form a no-op
+/// subview. This is true if
+/// 1) The offsets are 0
+/// 2) The strides are 1
+/// 3) The sizes are same as that of the src.
+/// For (3) when the shape is dynamic if the `src` is defined using an operation
+/// that implements the `ShapeAwareOpInterface` (like
+/// `hal.interface.binding.subspan`) then we can use that to check dynamic
+/// equality.
+static bool generatesNoOpSubView(Value src, ArrayRef<OpFoldResult> offsets,
+ ArrayRef<OpFoldResult> sizes,
+ ArrayRef<OpFoldResult> strides) {
+ auto interfaceOp =
+ dyn_cast_or_null<IREE::Util::ShapeAwareOpInterface>(src.getDefiningOp());
+ if (!interfaceOp) {
+ return false;
+ }
+ /// Check offsets are 0.
+ if (llvm::any_of(offsets, [](OpFoldResult ofr) {
+ Optional<int64_t> intValue = getConstantIntValue(ofr);
+ return !intValue || intValue.getValue() != 0;
+ })) {
+ return false;
+ }
+ /// Check strides are 0.
+ if (llvm::any_of(strides, [](OpFoldResult ofr) {
+ Optional<int64_t> intValue = getConstantIntValue(ofr);
+ return !intValue || intValue.getValue() != 1;
+ })) {
+ return false;
+ }
+ /// Check sizes are same as the source.
+ auto dynamicDims = interfaceOp.getResultDynamicDims(0);
+ unsigned dynamicDimsPos = 0;
+ ArrayRef<int64_t> srcShape = src.getType().cast<MemRefType>().getShape();
+ for (auto size : enumerate(sizes)) {
+ if (Optional<int64_t> intValue = getConstantIntValue(size.value())) {
+ if (intValue != srcShape[size.index()]) {
+ return false;
+ }
+ continue;
+ }
+ if (size.value().get<Value>() == dynamicDims[dynamicDimsPos]) {
+ dynamicDimsPos++;
+ continue;
+ }
+ auto loadConstOp1 =
+ size.value()
+ .get<Value>()
+ .getDefiningOp<IREE::HAL::InterfaceConstantLoadOp>();
+ auto loadConstOp2 =
+ dynamicDims[dynamicDimsPos]
+ .getDefiningOp<IREE::HAL::InterfaceConstantLoadOp>();
+ if (!loadConstOp1 || !loadConstOp2 ||
+ loadConstOp1.index() != loadConstOp2.index()) {
+ return false;
+ }
+ dynamicDimsPos++;
+ }
+ return true;
+}
+
/// Creates a subview operation given the `src`, `offsets`, `sizes` and
/// `strides`. Handles the corner case where the `offsets`, `sizes` and
/// `strides` are empty in which case just forward the `src` value. If the
@@ -152,7 +214,9 @@
Value src, ArrayRef<OpFoldResult> offsets,
ArrayRef<OpFoldResult> sizes,
ArrayRef<OpFoldResult> strides) {
- if (offsets.empty() && sizes.empty() && strides.empty()) return src;
+ if (generatesNoOpSubView(src, offsets, sizes, strides)) {
+ return src;
+ }
MemRefType resultType;
MemRefType srcType = src.getType().cast<MemRefType>();
if (srcType.getRank() != resultRank) {
@@ -269,20 +333,25 @@
return subview;
}
-/// Gets the reverse of a `tensor.expand_shape`/`tensor.collapse_shape` op to
-/// get a memref type that can be used for in-place computation of the result
-/// of a dispatch region.
-template <typename TensorReshapeOpTy>
-static Value getReverseOfReshapeOp(OpBuilder &b, TensorReshapeOpTy reshapeOp,
+/// Gets the reverse of a `tensor.collapse_shape` op to get a memref type that
+/// can be used for in-place computation of the result of a dispatch region.
+static Value getReverseOfReshapeOp(OpBuilder &b,
+ tensor::CollapseShapeOp reshapeOp,
Value resultBuffer) {
auto memrefType = getMemrefTypeForTensor(
reshapeOp.getSrcType(), {},
resultBuffer.getType().cast<MemRefType>().getMemorySpace());
- using ReverseReshapeOpTy = typename std::conditional<
- std::is_same<TensorReshapeOpTy, tensor::CollapseShapeOp>::value,
- memref::ExpandShapeOp, memref::CollapseShapeOp>::type;
- return b.create<ReverseReshapeOpTy>(reshapeOp.getLoc(), memrefType,
- resultBuffer, reshapeOp.reassociation());
+ return b.create<memref::ExpandShapeOp>(
+ reshapeOp.getLoc(), memrefType, resultBuffer, reshapeOp.reassociation());
+}
+
+/// Gets the reverse of a `tensor.expand_shape` op to get a memref type that can
+/// be used for in-place computation of the result of a dispatch region.
+static Value getReverseOfReshapeOp(OpBuilder &b,
+ tensor::ExpandShapeOp reshapeOp,
+ Value resultBuffer) {
+ return b.create<memref::CollapseShapeOp>(reshapeOp.getLoc(), resultBuffer,
+ reshapeOp.getReassociationIndices());
}
/// Gets the reverse of a `tensor.cast` op to get a memref type that
@@ -413,12 +482,26 @@
loadOp.getMixedSizes(), loadOp.getMixedStrides());
}
-/// Converts a `tensor.collapse/expand_shape` operation to a
-/// `linalg.collapse/expand_shape` operation with the result aliasing the buffer
-/// for the operand.
-template <typename TensorReshapeOpTy>
+/// Converts a `tensor.collapse_shape` operation to a `memref.collapse_shape`
+/// operation with the result aliasing the buffer for the operand.
static Value getAliasingBufferForReshapeResult(OpBuilder &b,
- TensorReshapeOpTy op,
+ tensor::CollapseShapeOp op,
+ BlockAndValueMapping &bvm) {
+ Location loc = op.getLoc();
+ Value srcTensor = op.src();
+ Value inputBuffer = bvm.lookup(srcTensor);
+
+ // Create the reshape op.
+ Value bufferReshape = b.create<memref::CollapseShapeOp>(
+ loc, inputBuffer, op.getReassociationIndices());
+ return bufferReshape;
+}
+
+/// Converts a `tensor.expand_shape` operation to a
+/// `memref.expand_shape` operation with the result aliasing the buffer
+/// for the operand.
+static Value getAliasingBufferForReshapeResult(OpBuilder &b,
+ tensor::ExpandShapeOp op,
BlockAndValueMapping &bvm) {
Location loc = op.getLoc();
Value srcTensor = op.src();
@@ -429,11 +512,8 @@
MemRefType inputBufferType = inputBuffer.getType().cast<MemRefType>();
auto reshapeResultType = getMemrefTypeForTensor(
resultTensorType, {}, inputBufferType.getMemorySpace());
- using ReshapeOpTy = typename std::conditional<
- std::is_same<TensorReshapeOpTy, tensor::CollapseShapeOp>::value,
- memref::CollapseShapeOp, memref::ExpandShapeOp>::type;
- Value bufferReshape = b.create<ReshapeOpTy>(loc, reshapeResultType,
- inputBuffer, op.reassociation());
+ Value bufferReshape = b.create<memref::ExpandShapeOp>(
+ loc, reshapeResultType, inputBuffer, op.reassociation());
return bufferReshape;
}
@@ -452,9 +532,9 @@
/// Returns output buffers that aliases inputs.
static SmallVector<Value> getAliasingBuffersForResult(
scf::ForOp scfFor, BlockAndValueMapping &bvm) {
- SmallVector<Value> aliasedBuffers(scfFor.results().size(), nullptr);
- for (int i = 0; i < scfFor.results().size(); ++i) {
- Value inputTensor = scfFor.initArgs()[i];
+ SmallVector<Value> aliasedBuffers(scfFor.getResults().size(), nullptr);
+ for (int i = 0; i < scfFor.getResults().size(); ++i) {
+ Value inputTensor = scfFor.getInitArgs()[i];
if (!inputTensor.getType().isa<RankedTensorType>()) continue;
Value inputBuffer = bvm.lookup(inputTensor);
aliasedBuffers[i] = inputBuffer;
diff --git a/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp b/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp
index 462a047..7bfc774 100644
--- a/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp
+++ b/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp
@@ -70,7 +70,7 @@
/// Return true if the given tiled loop is distributed to workgroups.
static bool isWorkgroupLoop(const LoopTilingAndDistributionInfo &info) {
auto forOp = cast<scf::ForOp>(info.loop);
- Operation *lbOp = forOp.lowerBound().getDefiningOp();
+ Operation *lbOp = forOp.getLowerBound().getDefiningOp();
if (isa<IREE::HAL::InterfaceWorkgroupIDOp>(lbOp)) return true;
auto applyOp = dyn_cast<AffineApplyOp>(lbOp);
return applyOp && llvm::any_of(applyOp.getMapOperands(), [](Value operand) {
diff --git a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
index 3524a6c..3787471 100644
--- a/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
+++ b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
@@ -8,7 +8,7 @@
%1 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:3x3x1x96xf32>
%2 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x3x96xf32>
// CHECK: %[[LOAD:.+]] = flow.dispatch.tensor.load %[[ARG]], {{.*}} : !flow.dispatch.tensor<readonly:3x3x96xf32> -> tensor<3x3x96xf32>
- %3 = flow.dispatch.tensor.load %1, offsets=[], sizes =[], strides=[] : !flow.dispatch.tensor<readonly:3x3x1x96xf32> -> tensor<3x3x1x96xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets=[0, 0, 0, 0], sizes =[3, 3, 1, 96], strides=[1, 1, 1, 1] : !flow.dispatch.tensor<readonly:3x3x1x96xf32> -> tensor<3x3x1x96xf32>
%4 = tensor.collapse_shape %3 [[0, 1, 2, 3]] : tensor<3x3x1x96xf32> into tensor<864xf32>
%5 = tensor.expand_shape %4 [[0, 1, 2]] : tensor<864xf32> into tensor<3x3x96xf32>
// CHECK: flow.dispatch.tensor.store %[[LOAD]], {{.*}}
@@ -46,7 +46,7 @@
%dim2 = hal.interface.constant.load[2] : index
%1 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?x96xf32>{%dim0, %dim1}
%2 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x12x8xf32>{%dim2}
- %3 = flow.dispatch.tensor.load %1, offsets=[], sizes =[], strides=[] : !flow.dispatch.tensor<readonly:?x?x96xf32> -> tensor<?x?x96xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets=[0, 0, 0], sizes =[%dim0, %dim1, 96], strides=[1, 1, 1] : !flow.dispatch.tensor<readonly:?x?x96xf32> -> tensor<?x?x96xf32>
// CHECK: tensor.collapse_shape
// CHECK: tensor.expand_shape
%4 = tensor.collapse_shape %3 [[0, 1], [2]] : tensor<?x?x96xf32> into tensor<?x96xf32>
diff --git a/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir b/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
index ca0818c..91070fd 100644
--- a/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
+++ b/iree/compiler/Codegen/Common/test/convert_to_destination_passing_style.mlir
@@ -153,9 +153,9 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
- flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %3, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
// CHECK: func @reshape_simple()
@@ -175,7 +175,7 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
%4 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%5 = linalg.generic {
@@ -186,7 +186,7 @@
%6 = arith.addi %arg0, %arg0 : i32
linalg.yield %6 : i32
} -> tensor<3x4xi32>
- flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %5, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
// CHECK: func @reshape_fused_source()
@@ -211,7 +211,7 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %3 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%4 = tensor.expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
%5 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%6 = linalg.generic {
@@ -222,8 +222,8 @@
%7 = arith.addi %arg0, %arg0 : i32
linalg.yield %7 : i32
} -> tensor<3x4xi32>
- flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
- flow.dispatch.tensor.store %4, %2, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %6, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %4, %2, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
// CHECK: func @reshape_fused_source_and_copyout()
@@ -249,7 +249,7 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:3x4xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:12xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
%3 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%4 = linalg.generic {
indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -260,7 +260,7 @@
linalg.yield %5 : i32
} -> tensor<3x4xi32>
%5 = tensor.collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
- flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
+ flow.dispatch.tensor.store %5, %1, offsets = [0], sizes = [12], strides = [1] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
return
}
// CHECK: func @reshape_fused_target()
diff --git a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
index 296e65a..4df4597 100644
--- a/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
+++ b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
@@ -589,9 +589,9 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
- flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %3, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
@@ -611,7 +611,7 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%3 = tensor.expand_shape %2 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
%4 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%5 = linalg.generic {
@@ -622,7 +622,7 @@
%6 = arith.addi %arg0, %arg0 : i32
linalg.yield %6 : i32
} -> tensor<3x4xi32>
- flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %5, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
@@ -645,7 +645,7 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:12xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
- %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
+ %3 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [12], strides = [1] : !flow.dispatch.tensor<readonly:12xi32> -> tensor<12xi32>
%4 = tensor.expand_shape %3 [[0, 1]] : tensor<12xi32> into tensor<3x4xi32>
%5 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%6 = linalg.generic {
@@ -656,8 +656,8 @@
%7 = arith.addi %arg0, %arg0 : i32
linalg.yield %7 : i32
} -> tensor<3x4xi32>
- flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
- flow.dispatch.tensor.store %4, %2, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %6, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ flow.dispatch.tensor.store %4, %2, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
@@ -681,7 +681,7 @@
%c12 = arith.constant 12 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:3x4xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:12xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
%3 = linalg.init_tensor [3, 4] : tensor<3x4xi32>
%4 = linalg.generic {
indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>],
@@ -692,7 +692,7 @@
linalg.yield %5 : i32
} -> tensor<3x4xi32>
%5 = tensor.collapse_shape %4 [[0, 1]] : tensor<3x4xi32> into tensor<12xi32>
- flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
+ flow.dispatch.tensor.store %5, %1, offsets = [0], sizes = [12], strides = [1] : tensor<12xi32> -> !flow.dispatch.tensor<writeonly:12xi32>
return
}
@@ -715,7 +715,7 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:1x1x2xf32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:2x3xf32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:1x3xf32>
- %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x1x2xf32> -> tensor<1x1x2xf32>
+ %3 = flow.dispatch.tensor.load %0, offsets = [0, 0, 0], sizes = [1, 1, 2], strides = [1, 1, 1] : !flow.dispatch.tensor<readonly:1x1x2xf32> -> tensor<1x1x2xf32>
%4 = tensor.collapse_shape %3 [[0, 1], [2]] : tensor<1x1x2xf32> into tensor<1x2xf32>
%workgroup_size_x = hal.interface.workgroup.size[0] : index
%workgroup_size_y = hal.interface.workgroup.size[1] : index
@@ -767,9 +767,9 @@
%5 = hal.interface.constant.load[3] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xi32>{%2, %3}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xi32>{%4, %5}
- %6 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
+ %6 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%2, %3], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
%7 = tensor.extract_slice %6[%2, %3] [%4, %5] [1, 1] : tensor<?x?xi32> to tensor<?x?xi32>
- flow.dispatch.tensor.store %7, %1, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
+ flow.dispatch.tensor.store %7, %1, offsets = [0, 0], sizes = [%4, %5], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
@@ -790,9 +790,9 @@
%8 = hal.interface.constant.load[4] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?x?xi32>{%8, %8, %8}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xi32>{%4, %5}
- %6 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?x?xi32> -> tensor<?x?x?xi32>
+ %6 = flow.dispatch.tensor.load %0, offsets = [0, 0, 0], sizes = [%8, %8, %8], strides = [1, 1, 1] : !flow.dispatch.tensor<readonly:?x?x?xi32> -> tensor<?x?x?xi32>
%7 = tensor.extract_slice %6[%2, %2, %3] [%4, 1, %5] [1, 1, 1] : tensor<?x?x?xi32> to tensor<?x?xi32>
- flow.dispatch.tensor.store %7, %1, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
+ flow.dispatch.tensor.store %7, %1, offsets = [0, 0], sizes = [%4, %5], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
@@ -816,10 +816,10 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?x?xi32>{%12, %12, %12}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?x?xi32>{%6, %7, %8}
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xi32>{%6, %8}
- %9 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?x?xi32> -> tensor<?x?x?xi32>
+ %9 = flow.dispatch.tensor.load %0, offsets = [0, 0, 0], sizes = [%12, %12, %12], strides = [1, 1, 1] : !flow.dispatch.tensor<readonly:?x?x?xi32> -> tensor<?x?x?xi32>
%10 = tensor.extract_slice %9[%3, %4, %5] [%6, %7, %8] [1, 1, 1] : tensor<?x?x?xi32> to tensor<?x?x?xi32>
%11 = tensor.extract_slice %9[%3, %4, %5] [%6, 1, %8] [1, 1, 1] : tensor<?x?x?xi32> to tensor<?x?xi32>
- flow.dispatch.tensor.store %10, %1, offsets = [], sizes = [], strides = [] : tensor<?x?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?x?xi32>
+ flow.dispatch.tensor.store %10, %1, offsets = [0, 0, 0], sizes = [%6, %7, %8], strides = [1, 1, 1] : tensor<?x?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?x?xi32>
flow.dispatch.tensor.store %11, %2, offsets = [%3, %5], sizes = [%6, %8], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
@@ -844,8 +844,8 @@
%2 = hal.interface.constant.load[0] : index
%3 = hal.interface.constant.load[1] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:?x?xi32>{%2, %3}
- %6 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:?x?xi32> -> tensor<?x?xi32>
- flow.dispatch.tensor.store %6, %0, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<readwrite:?x?xi32>
+ %6 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%2, %3], strides = [1, 1] : !flow.dispatch.tensor<readwrite:?x?xi32> -> tensor<?x?xi32>
+ flow.dispatch.tensor.store %6, %0, offsets = [0, 0], sizes = [%2, %3], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<readwrite:?x?xi32>
return
}
@@ -863,17 +863,18 @@
%dim3 = hal.interface.constant.load[3] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xi32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xi32>{%dim2, %dim3}
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
%3 = tensor.extract_slice %2[1, 0] [1, 4] [1, 1] : tensor<?x?xi32> to tensor<1x4xi32>
- flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<1x4xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
+ flow.dispatch.tensor.store %3, %1, offsets = [0, 0], sizes = [1, 4], strides = [1, 1] : tensor<1x4xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
// CHECK-LABEL: func @slice_whole_stride_dispatch_0()
// CHECK-DAG: %[[INPUT:.+]] = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer)
// CHECK-DAG: %[[OUTPUT:.+]] = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer)
-// CHECK: %[[SUBVIEW:.+]] = memref.subview %[[INPUT]]
-// CHECK: linalg.copy(%[[SUBVIEW]], %[[OUTPUT]])
+// CHECK-DAG: %[[SUBVIEW_INPUT:.+]] = memref.subview %[[INPUT]]
+// CHECK-DAG: %[[SUBVIEW_OUTPUT:.+]] = memref.subview %[[OUTPUT]]
+// CHECK: linalg.copy(%[[SUBVIEW_INPUT]], %[[SUBVIEW_OUTPUT]])
// -----
@@ -889,12 +890,12 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xi32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xi32>{%dim2, %dim3}
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xi32>{%dim4, %dim5}
- %3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
- %4 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
+ %3 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
+ %4 = flow.dispatch.tensor.load %1, offsets = [0, 0], sizes = [%dim2, %dim3], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xi32> -> tensor<?x?xi32>
%5 = tensor.dim %3, %c0 : tensor<?x?xi32>
%6 = tensor.dim %3, %c1 : tensor<?x?xi32>
%7 = tensor.insert_slice %3 into %4[3, 4] [%5, %6] [1, 1] : tensor<?x?xi32> into tensor<?x?xi32>
- flow.dispatch.tensor.store %7, %2, offsets = [], sizes = [], strides = [] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
+ flow.dispatch.tensor.store %7, %2, offsets = [0, 0], sizes = [%dim4, %dim5], strides = [1, 1] : tensor<?x?xi32> -> !flow.dispatch.tensor<writeonly:?x?xi32>
return
}
@@ -918,7 +919,7 @@
%3 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:i32> -> tensor<i32>
%4 = tensor.extract %3[] : tensor<i32>
%5 = linalg.fill(%4, %2) : i32, tensor<3x9xi32> -> tensor<3x9xi32>
- flow.dispatch.tensor.store %5, %1, offsets = [], sizes = [], strides = [] : tensor<3x9xi32> -> !flow.dispatch.tensor<writeonly:3x9xi32>
+ flow.dispatch.tensor.store %5, %1, offsets = [0, 0], sizes = [3, 9], strides = [1, 1] : tensor<3x9xi32> -> !flow.dispatch.tensor<writeonly:3x9xi32>
return
}
@@ -934,8 +935,8 @@
%c0 = arith.constant 0 : index
%1 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<writeonly:3x4xi32>
%2 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:3x4xi32>
- %3 = flow.dispatch.tensor.load %2, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
- flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
+ %3 = flow.dispatch.tensor.load %2, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : !flow.dispatch.tensor<readonly:3x4xi32> -> tensor<3x4xi32>
+ flow.dispatch.tensor.store %3, %1, offsets = [0, 0], sizes = [3, 4], strides = [1, 1] : tensor<3x4xi32> -> !flow.dispatch.tensor<writeonly:3x4xi32>
return
}
@@ -950,7 +951,7 @@
%c0 = arith.constant 0 : index
%cst = arith.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32>
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<writeonly:2x2x3xi32>
- flow.dispatch.tensor.store %cst, %0, offsets = [], sizes = [], strides = [] : tensor<2x2x3xi32> -> !flow.dispatch.tensor<writeonly:2x2x3xi32>
+ flow.dispatch.tensor.store %cst, %0, offsets = [0, 0, 0], sizes = [2, 2, 3], strides = [1, 1, 1] : tensor<2x2x3xi32> -> !flow.dispatch.tensor<writeonly:2x2x3xi32>
return
}
@@ -970,7 +971,7 @@
%c1 = arith.constant 1 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:1x5x3x1xf32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:5x5xf32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x5x3x1xf32> -> tensor<1x5x3x1xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0, 0, 0], sizes = [1, 5, 3, 1], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:1x5x3x1xf32> -> tensor<1x5x3x1xf32>
%3 = tensor.collapse_shape %2 [[0, 1], [2, 3]] : tensor<1x5x3x1xf32> into tensor<5x3xf32>
%workgroup_size_x = hal.interface.workgroup.size[0] : index
%workgroup_size_y = hal.interface.workgroup.size[1] : index
@@ -1026,8 +1027,8 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:?xi32>{%dim2}
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xf32>{%dim3, %dim4}
- %4 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = []: !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
- %5 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?xi32> -> tensor<?xi32>
+ %4 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1]: !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %5 = flow.dispatch.tensor.load %1, offsets = [0], sizes = [%dim2], strides = [1] : !flow.dispatch.tensor<readonly:?xi32> -> tensor<?xi32>
%d0 = tensor.dim %5, %c0 : tensor<?xi32>
%d1 = tensor.dim %4, %c1 : tensor<?x?xf32>
%3 = linalg.init_tensor [%d0, %d1] : tensor<?x?xf32>
@@ -1038,7 +1039,7 @@
%9 = tensor.extract %4[%8, %iv1] : tensor<?x?xf32>
linalg.yield %9 : f32
} -> tensor<?x?xf32>
- flow.dispatch.tensor.store %7, %2, offsets = [], sizes = [], strides = [] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32>
+ flow.dispatch.tensor.store %7, %2, offsets = [0, 0], sizes = [%dim3, %dim4], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32>
return
}
@@ -1062,7 +1063,7 @@
%3 = linalg.init_tensor [2, 3] : tensor<2x3xf32>
%4 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:f32> -> tensor<f32>
%5 = tensor.extract %4[] : tensor<f32>
- %6 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x4x6x1xf32> -> tensor<1x4x6x1xf32>
+ %6 = flow.dispatch.tensor.load %1, offsets = [0, 0, 0, 0], sizes = [1, 4, 6, 1], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:1x4x6x1xf32> -> tensor<1x4x6x1xf32>
%7 = linalg.init_tensor [1, 2, 2, 1] : tensor<1x2x2x1xf32>
%8 = linalg.fill(%5, %7) : f32, tensor<1x2x2x1xf32> -> tensor<1x2x2x1xf32>
%9 = linalg.pooling_nhwc_sum {
@@ -1070,7 +1071,7 @@
strides = dense<[2, 3]> : vector<2xi64>
} ins(%6, %3 : tensor<1x4x6x1xf32>, tensor<2x3xf32>)
outs(%8 : tensor<1x2x2x1xf32>) -> tensor<1x2x2x1xf32>
- flow.dispatch.tensor.store %9, %2, offsets = [], sizes = [], strides = [] : tensor<1x2x2x1xf32> -> !flow.dispatch.tensor<writeonly:1x2x2x1xf32>
+ flow.dispatch.tensor.store %9, %2, offsets = [0, 0, 0, 0], sizes = [1, 2, 2, 1], strides = [1, 1, 1, 1] : tensor<1x2x2x1xf32> -> !flow.dispatch.tensor<writeonly:1x2x2x1xf32>
return
}
@@ -1100,9 +1101,9 @@
%pc5 = hal.interface.constant.load[5] : index
%0 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xf32>{%pc0, %pc1}
%1 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%pc2, %pc3}
- %2 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %2 = flow.dispatch.tensor.load %1, offsets = [0, 0], sizes = [%pc2, %pc3], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
%3 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%pc4, %pc5}
- %4 = flow.dispatch.tensor.load %3, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %4 = flow.dispatch.tensor.load %3, offsets = [0, 0], sizes = [%pc4, %pc5], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
%workgroup_size_x = hal.interface.workgroup.size[0] : index
%workgroup_size_y = hal.interface.workgroup.size[1] : index
%workgroup_id_x = hal.interface.workgroup.id[0] : index
@@ -1168,7 +1169,7 @@
%dim2 = hal.interface.constant.load[2] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?xf32>{%dim2}
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
%3 = tensor.collapse_shape %2 [[0, 1]]
: tensor<?x?xf32> into tensor<?xf32>
%4 = tensor.dim %3, %c0 : tensor<?xf32>
@@ -1181,7 +1182,7 @@
%7 = arith.addf %arg0, %arg0 : f32
linalg.yield %7 : f32
} -> tensor<?xf32>
- flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = []: tensor<?xf32> -> !flow.dispatch.tensor<writeonly:?xf32>
+ flow.dispatch.tensor.store %6, %1, offsets = [0], sizes = [%dim2], strides = [1]: tensor<?xf32> -> !flow.dispatch.tensor<writeonly:?xf32>
return
}
@@ -1203,9 +1204,9 @@
%offset_subspan = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<readonly:32xf32>
%output_subspan = hal.interface.binding.subspan set(0) binding(3) type(storage_buffer) : !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
- %input = flow.dispatch.tensor.load %input_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> -> tensor<1x225x225x16xf32>
- %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32>
- %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32>
+ %input = flow.dispatch.tensor.load %input_subspan, offsets = [0, 0, 0, 0], sizes = [1, 255, 255, 16], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> -> tensor<1x225x225x16xf32>
+ %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [0, 0, 0, 0], sizes = [3, 3, 16, 32], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32>
+ %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32>
%cst = arith.constant 0.0 : f32
%0 = linalg.init_tensor [1, 112, 112, 32] : tensor<1x112x112x32xf32>
@@ -1227,7 +1228,7 @@
%sub = arith.subf %a, %b : f32
linalg.yield %sub : f32
} -> tensor<1x112x112x32xf32>
- flow.dispatch.tensor.store %3, %output_subspan, offsets = [], sizes = [], strides = [] : tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
+ flow.dispatch.tensor.store %3, %output_subspan, offsets = [0, 0, 0, 0], sizes = [1, 112, 112, 32], strides = [1, 1, 1, 1] : tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
return
}
@@ -1252,9 +1253,9 @@
%offset_subspan = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<readonly:32xf32>
%output_subspan = hal.interface.binding.subspan set(0) binding(3) type(storage_buffer) : !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
- %input = flow.dispatch.tensor.load %input_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> -> tensor<1x225x225x16xf32>
- %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32>
- %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32>
+ %input = flow.dispatch.tensor.load %input_subspan, offsets = [0, 0, 0, 0], sizes = [1, 225, 225, 16], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:1x225x225x16xf32> -> tensor<1x225x225x16xf32>
+ %filter = flow.dispatch.tensor.load %filter_subspan, offsets = [0, 0, 0, 0], sizes = [3, 3, 16, 32], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:3x3x16x32xf32> -> tensor<3x3x16x32xf32>
+ %offset = flow.dispatch.tensor.load %offset_subspan, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readonly:32xf32> -> tensor<32xf32>
%cst0 = arith.constant 0.0 : f32
%cst1 = arith.constant 1.0 : f32
@@ -1279,7 +1280,7 @@
%add = arith.addf %sub, %c : f32
linalg.yield %add : f32
} -> tensor<1x112x112x32xf32>
- flow.dispatch.tensor.store %4, %output_subspan, offsets = [], sizes = [], strides = []: tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
+ flow.dispatch.tensor.store %4, %output_subspan, offsets = [0, 0, 0, 0], sizes = [1, 112, 112, 32], strides = [1, 1, 1, 1]: tensor<1x112x112x32xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
return
}
@@ -1303,7 +1304,7 @@
%cst5 = arith.constant dense<[1, 2, 3, 4, 5]> : tensor<5xi32>
%input = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:5xf32>
%output = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:i32>
- %1 = flow.dispatch.tensor.load %input, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:5xf32> -> tensor<5xf32>
+ %1 = flow.dispatch.tensor.load %input, offsets=[0], sizes=[5], strides=[1] : !flow.dispatch.tensor<readonly:5xf32> -> tensor<5xf32>
%2 = linalg.generic {
indexing_maps = [affine_map<(d0) -> (-d0 + 4)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>],
iterator_types = ["reduction"]}
@@ -1396,12 +1397,12 @@
%dim4 = hal.interface.constant.load[4] : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:?x?x?xf32>{%dim2, %dim3, %dim4}
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
- %3 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:?x?x?xf32> -> tensor<?x?x?xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets = [0, 0, 0], sizes = [%dim2, %dim3, %dim4], strides = [1, 1, 1] : !flow.dispatch.tensor<readwrite:?x?x?xf32> -> tensor<?x?x?xf32>
%4 = tensor.dim %3, %c1 : tensor<?x?x?xf32>
%5 = tensor.dim %3, %c2 : tensor<?x?x?xf32>
%6 = tensor.insert_slice %2 into %3[0, 0, 0] [1, %4, %5] [1, 1, 1] : tensor<?x?xf32> into tensor<?x?x?xf32>
- flow.dispatch.tensor.store %6, %1, offsets = [], sizes = [], strides = [] : tensor<?x?x?xf32> -> !flow.dispatch.tensor<readwrite:?x?x?xf32>
+ flow.dispatch.tensor.store %6, %1, offsets = [0, 0, 0], sizes = [%dim2, %dim3, %dim4], strides = [1, 1, 1] : tensor<?x?x?xf32> -> !flow.dispatch.tensor<readwrite:?x?x?xf32>
return
}
@@ -1732,85 +1733,6 @@
// -----
-func @im2col() {
- %c0 = arith.constant 0 : index
- %cst = arith.constant 0.000000e+00 : f32
- %c112 = arith.constant 112 : index
- %c32 = arith.constant 32 : index
- %c16 = arith.constant 16 : index
- %c4 = arith.constant 4 : index
- %0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:1x225x225x8xf32>
- %1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:3x3x8x32xf32>
- %2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
- %workgroup_id_x = hal.interface.workgroup.id[0] : index
- %workgroup_count_x = hal.interface.workgroup.count[0] : index
- %workgroup_id_y = hal.interface.workgroup.id[1] : index
- %workgroup_count_y = hal.interface.workgroup.count[1] : index
- %workgroup_id_z = hal.interface.workgroup.id[2] : index
- %workgroup_count_z = hal.interface.workgroup.count[2] : index
- %3 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_id_z]
- %4 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_count_z]
- scf.for %arg0 = %3 to %c112 step %4 {
- %5 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_id_y]
- %6 = affine.apply affine_map<()[s0] -> (s0 * 16)>()[%workgroup_count_y]
- scf.for %arg1 = %5 to %c112 step %6 {
- %7 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%workgroup_id_x]
- %8 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%workgroup_count_x]
- scf.for %arg2 = %7 to %c32 step %8 {
- %9 = affine.apply affine_map<(d0) -> (d0 * 2)>(%arg0)
- %10 = affine.min affine_map<(d0) -> (33, d0 * -2 + 225)>(%arg0)
- %11 = affine.apply affine_map<(d0) -> (d0 * 2)>(%arg1)
- %12 = affine.min affine_map<(d0) -> (33, d0 * -2 + 225)>(%arg1)
- %13 = flow.dispatch.tensor.load %0, offsets = [0, %9, %11, 0], sizes = [1, %10, %12, 8], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:1x225x225x8xf32> -> tensor<1x?x?x8xf32>
- %14 = flow.dispatch.tensor.load %1, offsets = [0, 0, 0, %arg2], sizes = [3, 3, 8, 4], strides = [1, 1, 1, 1] : !flow.dispatch.tensor<readonly:3x3x8x32xf32> -> tensor<3x3x8x4xf32>
- %15 = linalg.init_tensor [1, 16, 16, 4] : tensor<1x16x16x4xf32>
- %16 = linalg.fill(%cst, %15) {__internal_linalg_transform__ = "workgroup"} : f32, tensor<1x16x16x4xf32> -> tensor<1x16x16x4xf32>
- %17 = linalg.init_tensor [1, 16, 16, 3, 3, 8] : tensor<1x16x16x3x3x8xf32>
- %18 = linalg.generic {indexing_maps = [affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1 * 2 + d3, d2 * 2 + d4, d5)>, affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3, d4, d5)>], iterator_types = ["parallel", "parallel", "parallel", "parallel", "parallel", "parallel"]} ins(%13 : tensor<1x?x?x8xf32>) outs(%17 : tensor<1x16x16x3x3x8xf32>) {
- ^bb0(%arg3: f32, %arg4: f32): // no predecessors
- linalg.yield %arg3 : f32
- } -> tensor<1x16x16x3x3x8xf32>
- %19 = tensor.collapse_shape %18 [[0, 1, 2], [3, 4, 5]] : tensor<1x16x16x3x3x8xf32> into tensor<256x72xf32>
- %20 = tensor.collapse_shape %14 [[0, 1, 2], [3]] : tensor<3x3x8x4xf32> into tensor<72x4xf32>
- %21 = tensor.collapse_shape %16 [[0, 1, 2], [3]] : tensor<1x16x16x4xf32> into tensor<256x4xf32>
- %22 = linalg.matmul ins(%19, %20 : tensor<256x72xf32>, tensor<72x4xf32>) outs(%21 : tensor<256x4xf32>) -> tensor<256x4xf32>
- %23 = tensor.expand_shape %22 [[0, 1, 2], [3]] : tensor<256x4xf32> into tensor<1x16x16x4xf32>
- %24 = tensor.cast %23 : tensor<1x16x16x4xf32> to tensor<1x?x?x?xf32>
- flow.dispatch.tensor.store %24, %2, offsets = [0, %arg0, %arg1, %arg2], sizes = [1, %c16, %c16, %c4], strides = [1, 1, 1, 1] : tensor<1x?x?x?xf32> -> !flow.dispatch.tensor<writeonly:1x112x112x32xf32>
- }
- }
- }
- return
-}
-
-// CHECK-LABEL: func @im2col
-// CHECK-DAG: %[[ARG0:.+]] = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer)
-// CHECK-DAG: %[[ARG1:.+]] = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer)
-// CHECK-DAG: %[[RET0:.+]] = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer)
-// CHECK-DAG: %[[ALLOC_ARG0:.+]] = memref.alloc() : memref<1x16x16x3x3x8xf32>
-// CHECK-DAG: %[[ALLOC_ARG1:.+]] = memref.alloc() : memref<3x3x8x4xf32>
-// CHECK-DAG: %[[ALLOC_RET0:.+]] = memref.alloc() : memref<1x16x16x4xf32>
-// CHECK: scf.for
-// CHECK: scf.for
-// CHECK: scf.for
-// CHECK-DAG: %[[ARG0_SV:.+]] = memref.subview %[[ARG0]]
-// CHECK-DAG: %[[ARG1_SV:.+]] = memref.subview %[[ARG1]]
-// CHECK-DAG: linalg.copy(%[[ARG1_SV]], %[[ALLOC_ARG1]])
-// CHECK-DAG: linalg.fill(%{{.*}}, %[[ALLOC_RET0]]
-// CHECK: linalg.generic
-// CHECK-SAME: ins(%[[ARG0_SV]]
-// CHECK-SAME: outs(%[[ALLOC_ARG0]]
-// CHECK-DAG: %[[ALLOC_ARG0_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_ARG0]]
-// CHECK-DAG: %[[ALLOC_ARG1_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_ARG1]]
-// CHECK-DAG: %[[ALLOC_RET0_RESHAPE:.+]] = memref.collapse_shape %[[ALLOC_RET0]]
-// CHECK: linalg.matmul
-// CHECK-SAME: ins(%[[ALLOC_ARG0_RESHAPE]], %[[ALLOC_ARG1_RESHAPE]]
-// CHECK-SAME: outs(%[[ALLOC_RET0_RESHAPE]]
-// CHECK: %[[RET0_SV:.+]] = memref.subview %[[RET0]]
-// CHECK: linalg.copy(%[[ALLOC_RET0]], %[[RET0_SV]])
-
-// -----
-
func @multi_result_reduce() {
%c0 = arith.constant 0 : index
%c0_i32 = arith.constant 0 : i32
@@ -2161,7 +2083,7 @@
%c1 = arith.constant 1 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:4xi32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:4xi32> -> tensor<4xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [4], strides = [1] : !flow.dispatch.tensor<readonly:4xi32> -> tensor<4xi32>
%3 = scf.for %arg0 = %c0 to %c4 step %c1 iter_args(%arg1 = %2) -> (tensor<4xi32>) {
%4 = scf.for %arg2 = %c0 to %c3 step %c1 iter_args(%arg3 = %arg1) -> (tensor<4xi32>) {
%5 = arith.addi %arg2, %c1 : index
@@ -2179,7 +2101,7 @@
}
scf.yield %4 : tensor<4xi32>
}
- flow.dispatch.tensor.store %3, %1, offsets = [], sizes = [], strides = [] : tensor<4xi32> -> !flow.dispatch.tensor<writeonly:4xi32>
+ flow.dispatch.tensor.store %3, %1, offsets = [0], sizes = [4], strides = [1] : tensor<4xi32> -> !flow.dispatch.tensor<writeonly:4xi32>
return
}
@@ -2204,7 +2126,7 @@
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:4xi32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:4xi32> -> tensor<4xi32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [4], strides = [1] : !flow.dispatch.tensor<readwrite:4xi32> -> tensor<4xi32>
%3 = scf.for %arg0 = %c0 to %c4 step %c1 iter_args(%arg1 = %2) -> (tensor<4xi32>) {
%4 = scf.for %arg2 = %c0 to %c3 step %c1 iter_args(%arg3 = %arg1) -> (tensor<4xi32>) {
%5 = arith.addi %arg2, %c1 : index
@@ -2222,7 +2144,7 @@
}
scf.yield %4 : tensor<4xi32>
}
- flow.dispatch.tensor.store %3, %0, offsets = [], sizes = [], strides = [] : tensor<4xi32> -> !flow.dispatch.tensor<readwrite:4xi32>
+ flow.dispatch.tensor.store %3, %0, offsets = [0], sizes = [4], strides = [1] : tensor<4xi32> -> !flow.dispatch.tensor<readwrite:4xi32>
return
}
@@ -2242,13 +2164,13 @@
func @iree_linalg_ext_sort_1d() {
%c0 = arith.constant 0 : index
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:128xi32>
- %1 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:128xi32> -> tensor<128xi32>
+ %1 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [128], strides = [1] : !flow.dispatch.tensor<readwrite:128xi32> -> tensor<128xi32>
%2 = iree_linalg_ext.sort dimension(0) outs(%1 : tensor<128xi32>) {
^bb0(%arg0: i32, %arg1: i32): // no predecessors
%3 = arith.cmpi sgt, %arg0, %arg1 : i32
iree_linalg_ext.yield %3 : i1
} -> tensor<128xi32>
- flow.dispatch.tensor.store %2, %0, offsets = [], sizes = [], strides = [] : tensor<128xi32> -> !flow.dispatch.tensor<readwrite:128xi32>
+ flow.dispatch.tensor.store %2, %0, offsets = [0], sizes = [128], strides = [1] : tensor<128xi32> -> !flow.dispatch.tensor<readwrite:128xi32>
return
}
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
index 1c3f2c9..067cf93 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -92,8 +92,8 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:?x?xf32>{%dim0, %dim1}
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:?xf32>{%dim1}
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:?x?xf32>{%dim0, %dim1}
- %3 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
- %4 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:?xf32> -> tensor<?xf32>
+ %3 = flow.dispatch.tensor.load %0, offsets=[0, 0], sizes=[%dim0, %dim1], strides=[1, 1] : !flow.dispatch.tensor<readonly:?x?xf32> -> tensor<?x?xf32>
+ %4 = flow.dispatch.tensor.load %1, offsets=[0], sizes=[%dim1], strides=[1] : !flow.dispatch.tensor<readonly:?xf32> -> tensor<?xf32>
%5 = linalg.init_tensor [%dim0, %dim1] : tensor<?x?xf32>
%6 = linalg.generic {
indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
@@ -105,7 +105,7 @@
%7 = arith.addf %arg0, %arg1 : f32
linalg.yield %7 : f32
} -> tensor<?x?xf32>
- flow.dispatch.tensor.store %6, %2, offsets = [], sizes = [], strides = [] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32>
+ flow.dispatch.tensor.store %6, %2, offsets = [0, 0], sizes = [%dim0, %dim1], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:?x?xf32>
return
}
}
@@ -452,11 +452,11 @@
%cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32>
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
- %3 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
%4:2 = iree_linalg_ext.fft {__internal_linalg_transform__ = "workgroup"} ins(%c2, %cst, %cst_0 : index, tensor<2xf32>, tensor<2xf32>) outs(%2, %3 : tensor<32xf32>, tensor<32xf32>) : tensor<32xf32>, tensor<32xf32>
- flow.dispatch.tensor.store %4#0, %0, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
- flow.dispatch.tensor.store %4#1, %1, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#0, %0, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#1, %1, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
return
}
}
diff --git a/iree/compiler/Codegen/LLVMCPU/test/tile_fuse_and_vectorize.mlir b/iree/compiler/Codegen/LLVMCPU/test/tile_fuse_and_vectorize.mlir
index 7e3399f..b2e297e 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/tile_fuse_and_vectorize.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/tile_fuse_and_vectorize.mlir
@@ -92,7 +92,7 @@
%3 = hal.interface.binding.subspan set(0) binding(3) type(storage_buffer) : !flow.dispatch.tensor<readonly:384x512xf32>
%4 = hal.interface.binding.subspan set(0) binding(4) type(storage_buffer) offset(%c1835008) : !flow.dispatch.tensor<readonly:2x512xf32>
%5 = hal.interface.binding.subspan set(0) binding(5) type(storage_buffer) : !flow.dispatch.tensor<writeonly:384x512xf32>
- %6 = flow.dispatch.tensor.load %4, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:2x512xf32> -> tensor<2x512xf32>
+ %6 = flow.dispatch.tensor.load %4, offsets = [0, 0], sizes = [2, 512], strides = [1, 1] : !flow.dispatch.tensor<readonly:2x512xf32> -> tensor<2x512xf32>
%workgroup_id_x = hal.interface.workgroup.id[0] : index
%workgroup_count_x = hal.interface.workgroup.count[0] : index
%workgroup_id_y = hal.interface.workgroup.id[1] : index
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
index 4e9213e..bd0b42b 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -69,16 +69,18 @@
// CHECK-DAG: %[[C4:.+]] = arith.constant 4 : index
// CHECK-DAG: %[[C256:.+]] = arith.constant 256 : index
// CHECK-DAG: %[[C1024:.+]] = arith.constant 1024 : index
+// CHECK-DAG: %[[BUFFER0:.+]] = memref.get_global @__shared_memory___0 : memref<4x256xf32, 3>
+// CHECK-DAG: %[[BUFFER1:.+]] = memref.get_global @__shared_memory__ : memref<2x4xf32, 3>
// CHECK: scf.for %[[K:.+]] = %[[C0]] to %[[C1024]] step %[[C4]] {
// CHECK: gpu.barrier
-// CHECK: linalg.copy(%{{.*}}, %{{.*}}) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<2x4xf32, #{{.*}}>, memref<2x4xf32, #{{.*}}, 3>
+// CHECK: linalg.copy(%{{.*}}, %{{.*}}) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<2x4xf32, #{{.*}}>, memref<2x4xf32, 3>
// CHECK-NOT: gpu.barrier
-// CHECK: linalg.copy(%{{.*}}, %{{.*}}) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<4x256xf32, #{{.*}}>, memref<4x256xf32, #{{.*}}, 3>
+// CHECK: linalg.copy(%{{.*}}, %{{.*}}) {__internal_linalg_transform__ = "copy_to_workgroup_memory"} : memref<4x256xf32, #{{.*}}>, memref<4x256xf32, 3>
// CHECK: gpu.barrier
// CHECK: scf.for %[[IND0:.+]] = %{{.*}} to %[[C2]] step %[[C2]] {
// CHECK: scf.for %[[IND1:.+]] = %{{.*}} to %[[C256]] step %[[C256]] {
-// CHECK-DAG: %[[A:.+]] = memref.subview %17[%[[IND0]], 0] [2, 4] [1, 1] : memref<2x4xf32, #{{.*}}, 3> to memref<2x4xf32, #{{.*}}, 3>
-// CHECK-DAG: %[[B:.+]] = memref.subview %18[0, %[[IND1]]] [4, 4] [1, 1] : memref<4x256xf32, #{{.*}}, 3> to memref<4x4xf32, #{{.*}}, 3>
+// CHECK-DAG: %[[A:.+]] = memref.subview %[[BUFFER1]][%[[IND0]], 0] [2, 4] [1, 1] : memref<2x4xf32, 3> to memref<2x4xf32, #{{.*}}, 3>
+// CHECK-DAG: %[[B:.+]] = memref.subview %[[BUFFER0]][0, %[[IND1]]] [4, 4] [1, 1] : memref<4x256xf32, 3> to memref<4x4xf32, #{{.*}}, 3>
// CHECK-DAG: %[[C:.+]] = memref.subview %11[%[[IND0]], %[[IND1]]] [2, 4] [1, 1] : memref<2x256xf32, #{{.*}}> to memref<2x4xf32, #{{.*}}>
// CHECK: linalg.matmul {__internal_linalg_transform__ = "vectorize", {{.*}}} ins(%[[A]], %[[B]] : memref<2x4xf32, #{{.*}}, 3>, memref<4x4xf32, #{{.*}}, 3>) outs(%[[C]] : memref<2x4xf32, #{{.*}}>)
// CHECK: }
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
index 60319c5..9186308 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -17,14 +17,14 @@
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:16384xf32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:16384xf32>
%3 = linalg.init_tensor [16384] : tensor<16384xf32>
- %4 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16384xf32> -> tensor<16384xf32>
- %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16384xf32> -> tensor<16384xf32>
+ %4 = flow.dispatch.tensor.load %0, offsets=[0], sizes=[16384], strides=[1] : !flow.dispatch.tensor<readonly:16384xf32> -> tensor<16384xf32>
+ %5 = flow.dispatch.tensor.load %1, offsets=[0], sizes=[16384], strides=[1] : !flow.dispatch.tensor<readonly:16384xf32> -> tensor<16384xf32>
%6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16384xf32>, tensor<16384xf32>) outs(%3 : tensor<16384xf32>) {
^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors
%7 = arith.addf %arg0, %arg1 : f32
linalg.yield %7 : f32
} -> tensor<16384xf32>
- flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16384xf32> -> !flow.dispatch.tensor<writeonly:16384xf32>
+ flow.dispatch.tensor.store %6, %2, offsets=[0], sizes=[16384], strides=[1] : tensor<16384xf32> -> !flow.dispatch.tensor<writeonly:16384xf32>
return
}
}
@@ -290,11 +290,11 @@
%cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32>
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
- %3 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
%4:2 = iree_linalg_ext.fft {__internal_linalg_transform__ = "workgroup"} ins(%c2, %cst, %cst_0 : index, tensor<2xf32>, tensor<2xf32>) outs(%2, %3 : tensor<32xf32>, tensor<32xf32>) : tensor<32xf32>, tensor<32xf32>
- flow.dispatch.tensor.store %4#0, %0, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
- flow.dispatch.tensor.store %4#1, %1, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#0, %0, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#1, %1, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
return
}
}
diff --git a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
index a03bcc1..cd5e3b7 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
@@ -20,14 +20,14 @@
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:16xf32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:16xf32>
%3 = linalg.init_tensor [16] : tensor<16xf32>
- %4 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %4 = flow.dispatch.tensor.load %0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %5 = flow.dispatch.tensor.load %1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) {
^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors
%7 = arith.addf %arg0, %arg1 : f32
linalg.yield %7 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %6, %2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
@@ -277,14 +277,14 @@
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readonly:16xf32>
%2 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<writeonly:16xf32>
%3 = linalg.init_tensor [16] : tensor<16xf32>
- %4 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %4 = flow.dispatch.tensor.load %0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%5 = arith.constant dense<[1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0]> : tensor<16xf32>
%6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) {
^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors
%7 = arith.addf %arg0, %arg1 : f32
linalg.yield %7 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %6, %2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
index c732a5c..bb1cbcb 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
@@ -20,14 +20,14 @@
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readonly:16xf32>
%2 = hal.interface.binding.subspan set(0) binding(2) type(storage_buffer) : !flow.dispatch.tensor<writeonly:16xf32>
%3 = linalg.init_tensor [16] : tensor<16xf32>
- %4 = flow.dispatch.tensor.load %0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %5 = flow.dispatch.tensor.load %1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %4 = flow.dispatch.tensor.load %0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %5 = flow.dispatch.tensor.load %1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%6 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%4, %5 : tensor<16xf32>, tensor<16xf32>) outs(%3 : tensor<16xf32>) {
^bb0(%arg0: f32, %arg1: f32, %arg2: f32): // no predecessors
%7 = arith.addf %arg0, %arg1 : f32
linalg.yield %7 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %6, %2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %6, %2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp b/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp
index fe82ffe..16f291b 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVCopyToWorkgroupMemory.cpp
@@ -88,8 +88,8 @@
forBounds.reserve(numLoops);
permutation.reserve(numLoops);
Location loc = pLoopOp.getLoc();
- auto lbs = pLoopOp.lowerBound(), ubs = pLoopOp.upperBound(),
- steps = pLoopOp.step();
+ auto lbs = pLoopOp.getLowerBound(), ubs = pLoopOp.getUpperBound(),
+ steps = pLoopOp.getStep();
for (unsigned i : llvm::seq<unsigned>(0, procInfo.size())) {
Value mappedLb = rewriter.create<arith::AddIOp>(
loc, lbs[i],
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVDistribute.cpp b/iree/compiler/Codegen/SPIRV/SPIRVDistribute.cpp
index 16a0248..846b7a7 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVDistribute.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVDistribute.cpp
@@ -51,12 +51,13 @@
auto mulMap = AffineMap::get(0, 2, {sym0 * sym1}, context);
auto newLb = rewriter.create<AffineApplyOp>(
- loc, mulAddMap, ValueRange{idOp, forOp.step(), forOp.lowerBound()});
+ loc, mulAddMap,
+ ValueRange{idOp, forOp.getStep(), forOp.getLowerBound()});
auto newStep = rewriter.create<AffineApplyOp>(
- loc, mulMap, ValueRange{countOp, forOp.step()});
+ loc, mulMap, ValueRange{countOp, forOp.getStep()});
- forOp.lowerBoundMutable().assign(newLb);
- forOp.stepMutable().assign(newStep);
+ forOp.getLowerBoundMutable().assign(newLb);
+ forOp.getStepMutable().assign(newStep);
// Remove the attribute to avoid endless recursion.
forOp->removeAttr(getSPIRVDistributeAttrName());
return success();
diff --git a/iree/compiler/Codegen/SPIRV/Utils.cpp b/iree/compiler/Codegen/SPIRV/Utils.cpp
index f3b3bf0..9eda38f 100644
--- a/iree/compiler/Codegen/SPIRV/Utils.cpp
+++ b/iree/compiler/Codegen/SPIRV/Utils.cpp
@@ -188,9 +188,9 @@
// iterations of the inner loops.
SmallVector<Value, 2> iterationStride;
iterationStride.resize(pLoopOp.getNumLoops());
- auto lbs = pLoopOp.lowerBound();
- auto ubs = pLoopOp.upperBound();
- auto steps = pLoopOp.step();
+ auto lbs = pLoopOp.getLowerBound();
+ auto ubs = pLoopOp.getUpperBound();
+ auto steps = pLoopOp.getStep();
for (int i = numLoops - 1; i >= 0; --i) {
Value lb = lbs[i], ub = ubs[i], step = steps[i];
Value iterCount = rewriter.create<arith::DivSIOp>(
@@ -324,8 +324,8 @@
assert(numLoops == procInfo.size() &&
"expected as many ids as number of loops");
- auto lbs = pLoopOp.lowerBound();
- auto step = pLoopOp.step();
+ auto lbs = pLoopOp.getLowerBound();
+ auto step = pLoopOp.getStep();
SmallVector<Value, 2> ivReplacements;
for (unsigned i : llvm::seq<unsigned>(0, numLoops)) {
Value iterValue = rewriter.create<arith::AddIOp>(
@@ -338,7 +338,7 @@
if (generateGuard) {
TypeConverter::SignatureConversion signatureConverter(numLoops);
Value cond = nullptr;
- auto ubs = pLoopOp.upperBound();
+ auto ubs = pLoopOp.getUpperBound();
for (unsigned i : llvm::seq<unsigned>(0, numLoops)) {
Value cmp = rewriter.create<arith::CmpIOp>(loc, arith::CmpIPredicate::slt,
ivReplacements[i], ubs[i]);
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
index 287e6e6..bbfbb2a 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_default_linalg_ext_ops.mlir
@@ -18,13 +18,13 @@
builtin.func @static_1d_sort() {
%c0 = arith.constant 0 : index
%0 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:1000xi32>
- %1 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:1000xi32> -> tensor<1000xi32>
+ %1 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [1000], strides = [1] : !flow.dispatch.tensor<readwrite:1000xi32> -> tensor<1000xi32>
%2 = iree_linalg_ext.sort dimension(0) {__internal_linalg_transform__ = "workgroup"} outs(%1 : tensor<1000xi32>) {
^bb0(%arg0: i32, %arg1: i32): // no predecessors
%3 = arith.cmpi slt, %arg0, %arg1 : i32
iree_linalg_ext.yield %3 : i1
} -> tensor<1000xi32>
- flow.dispatch.tensor.store %2, %0, offsets = [], sizes = [], strides = [] : tensor<1000xi32> -> !flow.dispatch.tensor<readwrite:1000xi32>
+ flow.dispatch.tensor.store %2, %0, offsets = [0], sizes = [1000], strides = [1] : tensor<1000xi32> -> !flow.dispatch.tensor<readwrite:1000xi32>
return
}
}
@@ -141,11 +141,11 @@
%cst_0 = arith.constant dense<[-0.000000e+00, -1.000000e+00]> : tensor<2xf32>
%0 = hal.interface.binding.subspan set(0) binding(0) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
%1 = hal.interface.binding.subspan set(0) binding(1) type(storage_buffer) : !flow.dispatch.tensor<readwrite:32xf32>
- %2 = flow.dispatch.tensor.load %0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
- %3 = flow.dispatch.tensor.load %1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %2 = flow.dispatch.tensor.load %0, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
+ %3 = flow.dispatch.tensor.load %1, offsets = [0], sizes = [32], strides = [1] : !flow.dispatch.tensor<readwrite:32xf32> -> tensor<32xf32>
%4:2 = iree_linalg_ext.fft {__internal_linalg_transform__ = "workgroup"} ins(%c2, %cst, %cst_0 : index, tensor<2xf32>, tensor<2xf32>) outs(%2, %3 : tensor<32xf32>, tensor<32xf32>) : tensor<32xf32>, tensor<32xf32>
- flow.dispatch.tensor.store %4#0, %0, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
- flow.dispatch.tensor.store %4#1, %1, offsets = [], sizes = [], strides = [] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#0, %0, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
+ flow.dispatch.tensor.store %4#1, %1, offsets = [0], sizes = [32], strides = [1] : tensor<32xf32> -> !flow.dispatch.tensor<readwrite:32xf32>
return
}
}
diff --git a/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
index cfc5e9f..cbdd4ff 100644
--- a/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
+++ b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
@@ -38,23 +38,25 @@
SmallVectorImpl<Value> &dims,
SmallVectorImpl<Value> &symbols) {
MLIRContext *ctx = forOp.getContext();
- auto lbConstant = forOp.lowerBound().getDefiningOp<arith::ConstantIndexOp>();
+ auto lbConstant =
+ forOp.getLowerBound().getDefiningOp<arith::ConstantIndexOp>();
AffineExpr lb = lbConstant ? getAffineConstantExpr(lbConstant.value(), ctx)
: getAffineDimExpr(dims.size(), ctx);
- auto stepConstant = forOp.step().getDefiningOp<arith::ConstantIndexOp>();
+ auto stepConstant = forOp.getStep().getDefiningOp<arith::ConstantIndexOp>();
AffineExpr step = stepConstant
? getAffineConstantExpr(stepConstant.value(), ctx)
: getAffineSymbolExpr(symbols.size(), ctx);
- if (!lbConstant) dims.push_back(forOp.lowerBound());
- if (!stepConstant) symbols.push_back(forOp.step());
+ if (!lbConstant) dims.push_back(forOp.getLowerBound());
+ if (!stepConstant) symbols.push_back(forOp.getStep());
exprs.push_back(lb + step * getAffineDimExpr(dims.size(), ctx));
- auto ubConstant = forOp.upperBound().getDefiningOp<arith::ConstantIndexOp>();
+ auto ubConstant =
+ forOp.getUpperBound().getDefiningOp<arith::ConstantIndexOp>();
AffineExpr ub = ubConstant ? getAffineConstantExpr(ubConstant.value(), ctx)
: getAffineDimExpr(dims.size(), ctx);
- if (!ubConstant) dims.push_back(forOp.upperBound());
+ if (!ubConstant) dims.push_back(forOp.getUpperBound());
exprs.push_back(ub);
dims.push_back(forOp.getInductionVar());
diff --git a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
index 8796495..e223e80 100644
--- a/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
+++ b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
@@ -46,9 +46,9 @@
auto forOp = dyn_cast_or_null<scf::ForOp>(containingOp);
if (forOp && forOp.getInductionVar() == dim) {
iv = dim;
- ub = forOp.upperBound();
- lb = forOp.lowerBound();
- step = forOp.step();
+ ub = forOp.getUpperBound();
+ lb = forOp.getLowerBound();
+ step = forOp.getStep();
break;
}
auto parallelOp = dyn_cast_or_null<scf::ParallelOp>(containingOp);
@@ -56,9 +56,9 @@
for (auto inductionVar : llvm::enumerate(parallelOp.getInductionVars())) {
if (inductionVar.value() == dim) {
iv = dim;
- ub = parallelOp.upperBound()[inductionVar.index()];
- lb = parallelOp.lowerBound()[inductionVar.index()];
- step = parallelOp.step()[inductionVar.index()];
+ ub = parallelOp.getUpperBound()[inductionVar.index()];
+ lb = parallelOp.getLowerBound()[inductionVar.index()];
+ step = parallelOp.getStep()[inductionVar.index()];
break;
}
}
diff --git a/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp b/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
index 50d7aeb..36c2536 100644
--- a/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
+++ b/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
@@ -119,9 +119,9 @@
SmallVector<Value, 4> dims;
SmallVector<Value, 4> symbols;
AffineExpr lb = getAffineDimExpr(dims.size(), ctx);
- dims.push_back(op.lowerBound());
+ dims.push_back(op.getLowerBound());
AffineExpr ub = getAffineDimExpr(dims.size(), ctx);
- dims.push_back(op.upperBound());
+ dims.push_back(op.getUpperBound());
AffineExpr iterZero = ub - lb;
auto map = AffineMap::get(dims.size(), 0, iterZero);
AffineMap simplifiedMap = substituteMin(map, dims, symbols, getMinMax);
@@ -141,11 +141,11 @@
SmallVector<Value, 4> dims;
SmallVector<Value, 4> symbols;
AffineExpr lb = getAffineDimExpr(dims.size(), ctx);
- dims.push_back(op.lowerBound());
+ dims.push_back(op.getLowerBound());
AffineExpr ub = getAffineDimExpr(dims.size(), ctx);
- dims.push_back(op.upperBound());
+ dims.push_back(op.getUpperBound());
AffineExpr step = getAffineDimExpr(dims.size(), ctx);
- dims.push_back(op.step());
+ dims.push_back(op.getStep());
AffineExpr iterOne = lb + step - ub;
auto map = AffineMap::get(dims.size(), 0, iterOne);
@@ -177,7 +177,7 @@
// so the loop always have 1 iteration. Inline its body and remove the loop.
SmallVector<Value, 4> blockArgs;
blockArgs.reserve(op.getNumIterOperands() + 1);
- blockArgs.push_back(op.lowerBound());
+ blockArgs.push_back(op.getLowerBound());
llvm::append_range(blockArgs, op.getIterOperands());
replaceOpWithRegion(rewriter, op, op.getLoopBody(), blockArgs);
return success();
diff --git a/iree/compiler/Codegen/Utils/Utils.cpp b/iree/compiler/Codegen/Utils/Utils.cpp
index e2171ab..357db58 100644
--- a/iree/compiler/Codegen/Utils/Utils.cpp
+++ b/iree/compiler/Codegen/Utils/Utils.cpp
@@ -489,23 +489,23 @@
scf::ForOp forOp) {
LoopTilingAndDistributionInfo loopInfo;
loopInfo.loop = forOp;
- loopInfo.untiledUpperBound = getAsOpFoldResult(forOp.upperBound());
+ loopInfo.untiledUpperBound = getAsOpFoldResult(forOp.getUpperBound());
- auto lbApplyOp = forOp.lowerBound().getDefiningOp<AffineApplyOp>();
- auto stepApplyOp = forOp.step().getDefiningOp<AffineApplyOp>();
+ auto lbApplyOp = forOp.getLowerBound().getDefiningOp<AffineApplyOp>();
+ auto stepApplyOp = forOp.getStep().getDefiningOp<AffineApplyOp>();
if (!lbApplyOp || !stepApplyOp) {
// Try to see if this s a specical case where we have:
// scf.for %iv = %id to %ub step %count
Optional<unsigned> idDim;
if (auto ifx = dyn_cast_or_null<ProcessorIDInterface>(
- forOp.lowerBound().getDefiningOp())) {
+ forOp.getLowerBound().getDefiningOp())) {
idDim = ifx.getDimIndex();
}
Optional<unsigned> countDim;
if (auto ifx = dyn_cast_or_null<ProcessorCountInterface>(
- forOp.step().getDefiningOp())) {
+ forOp.getStep().getDefiningOp())) {
countDim = ifx.getDimIndex();
}
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
index 20db870..6b5f604 100644
--- a/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
+++ b/iree/compiler/Dialect/Flow/IR/FlowOps.cpp
@@ -76,6 +76,33 @@
}
}
+/// Implements default offset, sizes and strides, for
+/// `flow.dispatch.tensor.load/store` ops. When no offsets, sizes and strides
+/// are specified, the offsets are all zeros, sizes are same as the dispatch
+/// tensor and strides are all 1.
+static void getDefaultOffsetSizeAndStrides(
+ OpBuilder &builder, IREE::Flow::DispatchTensorType dispatchTensorType,
+ ValueRange dynamicDims, SmallVectorImpl<OpFoldResult> &offsets,
+ SmallVectorImpl<OpFoldResult> &sizes,
+ SmallVectorImpl<OpFoldResult> &strides) {
+ auto zeroAttr = builder.getI64IntegerAttr(0);
+ auto oneAttr = builder.getI64IntegerAttr(1);
+ int64_t dispatchTensorRank = dispatchTensorType.getRank();
+ offsets.assign(dispatchTensorRank, zeroAttr);
+ strides.assign(dispatchTensorRank, oneAttr);
+ sizes.resize(dispatchTensorRank);
+ unsigned pos = 0;
+ for (auto dim : llvm::enumerate(dispatchTensorType.getShape())) {
+ if (ShapedType::isDynamic(dim.value())) {
+ assert(pos < dynamicDims.size() && "missing dynamic dims specifications");
+ sizes[dim.index()] = dynamicDims[pos++];
+ continue;
+ }
+ sizes[dim.index()] = builder.getI64IntegerAttr(dim.value());
+ }
+ return;
+}
+
RankedTensorType DispatchTensorLoadOp::inferRankReducedResultType(
unsigned resultRank, IREE::Flow::DispatchTensorType sourceType,
ArrayRef<OpFoldResult> mixedSizes) {
@@ -123,10 +150,12 @@
RankedTensorType returnType, Value source,
ValueRange sourceDynamicDims,
ArrayRef<NamedAttribute> attributes) {
- build(builder, state, returnType, source, sourceDynamicDims,
- ArrayRef<Value>(), ArrayRef<Value>(), ArrayRef<Value>(),
- builder.getI64ArrayAttr({}), builder.getI64ArrayAttr({}),
- builder.getI64ArrayAttr({}));
+ SmallVector<OpFoldResult> offsets, strides, sizes;
+ getDefaultOffsetSizeAndStrides(
+ builder, source.getType().cast<IREE::Flow::DispatchTensorType>(),
+ sourceDynamicDims, offsets, sizes, strides);
+ build(builder, state, returnType, source, sourceDynamicDims, offsets, sizes,
+ strides, attributes);
}
void DispatchTensorLoadOp::build(OpBuilder &builder, OperationState &state,
@@ -154,6 +183,7 @@
strides, builder.getI64ArrayAttr(staticOffsets),
builder.getI64ArrayAttr(staticSizes),
builder.getI64ArrayAttr(staticStrides));
+ state.addAttributes(attributes);
}
void DispatchTensorLoadOp::build(OpBuilder &builder, OperationState &state,
@@ -177,8 +207,8 @@
shape = llvm::to_vector<6>(llvm::map_range(
getMixedSizes(), [&](OpFoldResult valueOrAttr) -> Value {
if (auto attr = valueOrAttr.dyn_cast<Attribute>()) {
- return b.create<arith::ConstantOp>(getLoc(),
- attr.cast<IntegerAttr>());
+ return b.create<arith::ConstantIndexOp>(
+ getLoc(), attr.cast<IntegerAttr>().getInt());
} else {
return valueOrAttr.dyn_cast<Value>();
}
@@ -206,10 +236,35 @@
Value value, Value target,
ValueRange targetDynamicDims,
ArrayRef<NamedAttribute> attributes) {
+ SmallVector<OpFoldResult> offsets, sizes, strides;
+ getDefaultOffsetSizeAndStrides(
+ builder, target.getType().cast<IREE::Flow::DispatchTensorType>(),
+ targetDynamicDims, offsets, sizes, strides);
+ build(builder, state, value, target, targetDynamicDims, offsets, sizes,
+ strides, attributes);
+}
+
+void DispatchTensorStoreOp::build(OpBuilder &builder, OperationState &state,
+ Value value, Value target,
+ ValueRange targetDynamicDims,
+ ArrayRef<OpFoldResult> mixedOffsets,
+ ArrayRef<OpFoldResult> mixedSizes,
+ ArrayRef<OpFoldResult> mixedStrides,
+ ArrayRef<NamedAttribute> attributes) {
+ SmallVector<Value> offsets, sizes, strides;
+ SmallVector<int64_t> staticOffsets, staticSizes, staticStrides;
+ processMixedOperands(mixedOffsets, offsets, staticOffsets,
+ ShapedType::kDynamicStrideOrOffset);
+ processMixedOperands(mixedSizes, sizes, staticSizes,
+ ShapedType::kDynamicSize);
+ processMixedOperands(mixedStrides, strides, staticStrides,
+ ShapedType::kDynamicStrideOrOffset);
+
build(builder, state, ArrayRef<Type>(), value, target, targetDynamicDims,
- ArrayRef<Value>(), ArrayRef<Value>(), ArrayRef<Value>(),
- builder.getI64ArrayAttr({}), builder.getI64ArrayAttr({}),
- builder.getI64ArrayAttr({}));
+ offsets, sizes, strides, builder.getI64ArrayAttr(staticOffsets),
+ builder.getI64ArrayAttr(staticSizes),
+ builder.getI64ArrayAttr(staticStrides));
+ state.addAttributes(attributes);
}
//===----------------------------------------------------------------------===//
diff --git a/iree/compiler/Dialect/Flow/IR/FlowOps.td b/iree/compiler/Dialect/Flow/IR/FlowOps.td
index af19305..687a390 100644
--- a/iree/compiler/Dialect/Flow/IR/FlowOps.td
+++ b/iree/compiler/Dialect/Flow/IR/FlowOps.td
@@ -393,6 +393,16 @@
"ValueRange":$targetDynamicDims,
CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs
)>,
+ // Builder for tensor.store with mixed static and dynamic offset, sizes and strides.
+ OpBuilder<(ins
+ "Value":$value,
+ "Value":$target,
+ "ValueRange":$targetDynamicDims,
+ "ArrayRef<OpFoldResult>":$mixedOffsets,
+ "ArrayRef<OpFoldResult>":$mixedSizes,
+ "ArrayRef<OpFoldResult>":$mixedStrides,
+ CArg<"ArrayRef<NamedAttribute>", "{}">:$attrs
+ )>
];
let extraClassDeclaration = [{
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir
index 171cf5d..8af91e6 100644
--- a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir
+++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups.mlir
@@ -44,7 +44,7 @@
// Load tensors (optional offsets/sizes/strides):
// CHECK: %[[ARG0_VALUE:.+]] = flow.dispatch.tensor.load %[[INNER_ARG0]], {{.*}} : !flow.dispatch.tensor<readonly:?x4xf32>{%[[INNER_ARG0_DIM0]]} -> tensor<?x4xf32>
- %arg0_value = flow.dispatch.tensor.load %arg0_capture, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:?x4xf32>{%arg0_dim0} -> tensor<?x4xf32>
+ %arg0_value = flow.dispatch.tensor.load %arg0_capture, offsets=[0, 0], sizes=[%arg0_dim0, 4], strides=[1, 1] : !flow.dispatch.tensor<readonly:?x4xf32>{%arg0_dim0} -> tensor<?x4xf32>
// Operate on tensors:
@@ -54,7 +54,7 @@
// Store tensors (optional offsets/sizes/strides):
// CHECK: flow.dispatch.tensor.store %[[RET0_VALUE]], %[[INNER_RET0]], {{.*}} : tensor<4x?xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%[[INNER_RET0_DIM1]]}
- flow.dispatch.tensor.store %ret0_value, %ret0, offsets=[], sizes=[], strides=[] : tensor<4x?xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%ret0_dim1}
+ flow.dispatch.tensor.store %ret0_value, %ret0, offsets=[0, 0], sizes=[4, %ret0_dim1], strides=[1, 1] : tensor<4x?xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%ret0_dim1}
// CHECK-NEXT: flow.return
flow.return
@@ -85,9 +85,9 @@
// CHECK-SAME: %[[INNER_ARG1:.+]]: index) {
(%arg0_capture: !flow.dispatch.tensor<readwrite:?x4xf32>, %arg1_capture: index) {
// CHECK: %[[VALUE:.+]] = flow.dispatch.tensor.load %[[INNER_ARG0]], {{.*}} : !flow.dispatch.tensor<readwrite:?x4xf32> -> tensor<?x4xf32>
- %t = flow.dispatch.tensor.load %arg0_capture, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readwrite:?x4xf32> -> tensor<?x4xf32>
+ %t = flow.dispatch.tensor.load %arg0_capture, offsets=[0, 0], sizes=[%arg1_capture, 4], strides=[1, 1] : !flow.dispatch.tensor<readwrite:?x4xf32> -> tensor<?x4xf32>
// CHECK: flow.dispatch.tensor.store %[[VALUE]], %[[INNER_ARG0]], {{.*}}: tensor<?x4xf32> -> !flow.dispatch.tensor<readwrite:?x4xf32>
- flow.dispatch.tensor.store %t, %arg0_capture, offsets=[], sizes=[], strides=[] : tensor<?x4xf32> -> !flow.dispatch.tensor<readwrite:?x4xf32>
+ flow.dispatch.tensor.store %t, %arg0_capture, offsets=[0, 0], sizes=[%arg1_capture, 4], strides=[1, 1] : tensor<?x4xf32> -> !flow.dispatch.tensor<readwrite:?x4xf32>
// CHECK-NEXT: flow.return
flow.return
}
diff --git a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir
index 2c4db65..95ed2ac 100644
--- a/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir
+++ b/iree/compiler/Dialect/Flow/IR/test/dispatch_workgroups_folding.mlir
@@ -89,9 +89,9 @@
%arg1_capture: !flow.dispatch.tensor<readwrite:4x8xf32>
) {
"test.sink"(%arg0_capture) : (!flow.dispatch.tensor<readonly:1x4xf32>) -> ()
- %load = flow.dispatch.tensor.load %arg1_capture, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readwrite:4x8xf32> -> tensor<4x8xf32>
+ %load = flow.dispatch.tensor.load %arg1_capture, offsets=[0, 0], sizes=[4, 8], strides=[1, 1] : !flow.dispatch.tensor<readwrite:4x8xf32> -> tensor<4x8xf32>
%0 = "test.do_work"(%load) : (tensor<4x8xf32>) -> (tensor<4x8xf32>)
- flow.dispatch.tensor.store %0, %arg1_capture, offsets=[], sizes=[], strides=[] : tensor<4x8xf32> -> !flow.dispatch.tensor<readwrite:4x8xf32>
+ flow.dispatch.tensor.store %0, %arg1_capture, offsets=[0, 0], sizes=[4, 8], strides=[1, 1] : tensor<4x8xf32> -> !flow.dispatch.tensor<readwrite:4x8xf32>
flow.return
}
return %0 : tensor<4x8xf32>
@@ -110,8 +110,8 @@
(%arg0: !flow.dispatch.tensor<readonly:9xi32>, %arg1: !flow.dispatch.tensor<readonly:9xi32>, %arg2: !flow.dispatch.tensor<writeonly:i32>, %arg3: !flow.dispatch.tensor<writeonly:i32>) {
%c0_i32 = arith.constant 0 : i32
%c-2147483648_i32 = arith.constant -2147483648 : i32
- %0 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
- %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
+ %0 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[9], strides = [1] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
+ %1 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[9], strides = [1] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
%2 = linalg.init_tensor [] : tensor<i32>
%3 = linalg.fill(%c-2147483648_i32, %2) : i32, tensor<i32> -> tensor<i32>
%4 = linalg.fill(%c0_i32, %2) : i32, tensor<i32> -> tensor<i32>
@@ -135,8 +135,8 @@
(%arg0: !flow.dispatch.tensor<readonly:9xi32>, %arg1: !flow.dispatch.tensor<readonly:9xi32>, %arg2: !flow.dispatch.tensor<writeonly:i32>, %arg3: !flow.dispatch.tensor<readwrite:i32>) {
%c0_i32 = arith.constant 0 : i32
%c-2147483648_i32 = arith.constant -2147483648 : i32
- %0 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
- %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
+ %0 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[9], strides = [1] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
+ %1 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[9], strides = [1] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
%2 = linalg.init_tensor [] : tensor<i32>
%3 = linalg.fill(%c-2147483648_i32, %2) : i32, tensor<i32> -> tensor<i32>
%4 = linalg.fill(%c0_i32, %2) : i32, tensor<i32> -> tensor<i32>
@@ -159,7 +159,7 @@
%c-2147483648_i32 = arith.constant -2147483648 : i32
%0 = flow.dispatch.tensor.load %arg3, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readwrite:i32> -> tensor<i32>
%val = tensor.extract %0[] : tensor<i32>
- %1 = flow.dispatch.tensor.load %arg1, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
+ %1 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[9], strides = [1] : !flow.dispatch.tensor<readonly:9xi32> -> tensor<9xi32>
%2 = linalg.init_tensor [] : tensor<i32>
%3 = linalg.fill(%c-2147483648_i32, %2) : i32, tensor<i32> -> tensor<i32>
%4 = linalg.fill(%val, %2) : i32, tensor<i32> -> tensor<i32>
diff --git a/iree/compiler/Dialect/Flow/Transforms/DestructiveUpdateUtils.cpp b/iree/compiler/Dialect/Flow/Transforms/DestructiveUpdateUtils.cpp
index 25303c3..16b3027 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DestructiveUpdateUtils.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DestructiveUpdateUtils.cpp
@@ -166,7 +166,7 @@
operand->getOperandNumber() - innerForOp.getNumControlOperands();
Value innerForOpResultTensor = innerForOp.getResult(innerIterArgIdx);
Value yieldValue =
- scfForOp.region().front().getTerminator()->getOperand(idx);
+ scfForOp.getRegion().front().getTerminator()->getOperand(idx);
// Check that the return position of dk and the yield position of dk
// agree (in the loop structure below). This avoids ping-pong effects
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
index 81729ca..e2b95b1 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
@@ -250,9 +250,7 @@
rewriter.create<IREE::Flow::DispatchTensorStoreOp>(
loc, std::get<0>(it), std::get<1>(it),
resultDynamicDims.slice(dynamicDimIdx,
- resultType.getNumDynamicDims()),
- llvm::None, llvm::None, llvm::None, rewriter.getArrayAttr({}),
- rewriter.getArrayAttr({}), rewriter.getArrayAttr({}));
+ resultType.getNumDynamicDims()));
dynamicDimIdx += resultType.getNumDynamicDims();
}
rewriter.create<IREE::Flow::ReturnOp>(loc);
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
index 8354744..3812ed0 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/dispatch_linalg_on_tensors.mlir
@@ -1009,9 +1009,11 @@
// CHECK: %[[RESULT:.+]]:2 = flow.dispatch.workgroups[%[[C1]], %[[C1]], %[[C1]]]
// CHECK-SAME: (%[[ARG0]], %[[ARG0_D0]], %[[ARG1]], %[[ARG1_D0]])
// CHECK-NEXT: (%[[ARG0_CAPTURE:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readwrite:?xi32>
+// CHECK-SAME: %[[ARG0_D0_CAPTURE:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1_CAPTURE:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readwrite:?xf32>
-// CHECK-DAG: %[[OUT1_TILE:.+]] = flow.dispatch.tensor.load %[[ARG0_CAPTURE]], offsets = [], sizes = []
-// CHECK-DAG: %[[OUT2_TILE:.+]] = flow.dispatch.tensor.load %[[ARG1_CAPTURE]], offsets = [], sizes = []
+// CHECK-SAME: %[[ARG1_D0_CAPTURE:[a-zA-Z0-9_]+]]: index
+// CHECK-DAG: %[[OUT1_TILE:.+]] = flow.dispatch.tensor.load %[[ARG0_CAPTURE]], offsets = [0], sizes = [%[[ARG0_D0_CAPTURE]]]
+// CHECK-DAG: %[[OUT2_TILE:.+]] = flow.dispatch.tensor.load %[[ARG1_CAPTURE]], offsets = [0], sizes = [%[[ARG1_D0_CAPTURE]]]
// CHECK: %[[RESULT_TILE:.+]]:2 = iree_linalg_ext.sort dimension(0)
// CHECK-SAME: outs(%[[OUT1_TILE]], %[[OUT2_TILE]] : tensor<?xi32>, tensor<?xf32>)
// CHECK-DAG: flow.dispatch.tensor.store %[[RESULT_TILE]]#0, %[[ARG0_CAPTURE]]
@@ -1046,7 +1048,7 @@
// CHECK-SAME: %[[ARG5:[a-zA-Z0-9_]+]]: !flow.dispatch.tensor<readwrite:8xi32>
// CHECK: scf.for %[[IV:.+]] = %{{.+}} to %{{.+}} step %{{.+}} {
// CHECK: %[[SCATTER_TILE:.+]] = iree_linalg_ext.scatter
-// CHECK: flow.dispatch.tensor.store %[[SCATTER_TILE]], %[[ARG5]], offsets = [], sizes = [], strides = []
+// CHECK: flow.dispatch.tensor.store %[[SCATTER_TILE]], %[[ARG5]], offsets = [0], sizes = [8], strides = [1]
// CHECK-NEXT: }
// CHECK: return %[[RESULT]]
diff --git a/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir b/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir
index 7be983f..8d3781f 100644
--- a/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir
+++ b/iree/compiler/Dialect/Flow/Transforms/test/outline_dispatch_regions.mlir
@@ -25,9 +25,9 @@
%0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<8x4xf32>) -> tensor<4x8xf32> = (
%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>
) {
- %arg_value = flow.dispatch.tensor.load %arg, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:8x4xf32> -> tensor<8x4xf32>
+ %arg_value = flow.dispatch.tensor.load %arg, offsets=[0, 0], sizes=[8, 4], strides=[1, 1] : !flow.dispatch.tensor<readonly:8x4xf32> -> tensor<8x4xf32>
%ret_value = "test.sink"(%arg_value) : (tensor<8x4xf32>) -> (tensor<4x8xf32>)
- flow.dispatch.tensor.store %ret_value, %ret, offsets=[], sizes=[], strides=[] : tensor<4x8xf32> -> !flow.dispatch.tensor<writeonly:4x8xf32>
+ flow.dispatch.tensor.store %ret_value, %ret, offsets=[0, 0], sizes=[4, 8], strides=[1, 1] : tensor<4x8xf32> -> !flow.dispatch.tensor<writeonly:4x8xf32>
flow.return
}
// CHECK-NEXT: return %[[RET]]
@@ -59,9 +59,9 @@
%0 = flow.dispatch.workgroups[%x, %y](%arg0) : (tensor<8x4xf32>) -> (tensor<4x8xf32>) = (
%arg: !flow.dispatch.tensor<readonly:8x4xf32>, %ret: !flow.dispatch.tensor<writeonly:4x8xf32>
) {
- %arg_value = flow.dispatch.tensor.load %arg, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:8x4xf32> -> tensor<8x4xf32>
+ %arg_value = flow.dispatch.tensor.load %arg, offsets=[0, 0], sizes=[8, 4], strides=[1, 1] : !flow.dispatch.tensor<readonly:8x4xf32> -> tensor<8x4xf32>
%ret_value = "test.sink1"(%arg_value) : (tensor<8x4xf32>) -> (tensor<4x8xf32>)
- flow.dispatch.tensor.store %ret_value, %ret, offsets=[], sizes=[], strides=[] : tensor<4x8xf32> -> !flow.dispatch.tensor<writeonly:4x8xf32>
+ flow.dispatch.tensor.store %ret_value, %ret, offsets=[0, 0], sizes=[4, 8], strides=[1, 1] : tensor<4x8xf32> -> !flow.dispatch.tensor<writeonly:4x8xf32>
flow.return
}
// CHECK: %[[RET1:.+]] = flow.dispatch @dispatchFnMuli_dispatch_1::@dispatchFnMuli_dispatch_1[
@@ -70,9 +70,9 @@
%1 = flow.dispatch.workgroups[%y, %x](%0) : (tensor<4x8xf32>) -> (tensor<8x4xf32>) = (
%arg: !flow.dispatch.tensor<readonly:4x8xf32>, %ret: !flow.dispatch.tensor<writeonly:8x4xf32>
) {
- %arg_value = flow.dispatch.tensor.load %arg, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:4x8xf32> -> tensor<8x4xf32>
+ %arg_value = flow.dispatch.tensor.load %arg, offsets=[0, 0], sizes=[4, 8], strides=[1, 1] : !flow.dispatch.tensor<readonly:4x8xf32> -> tensor<8x4xf32>
%ret_value = "test.sink2"(%arg_value) : (tensor<8x4xf32>) -> (tensor<8x4xf32>)
- flow.dispatch.tensor.store %ret_value, %ret, offsets=[], sizes=[], strides=[] : tensor<8x4xf32> -> !flow.dispatch.tensor<writeonly:8x4xf32>
+ flow.dispatch.tensor.store %ret_value, %ret, offsets=[0, 0], sizes=[8, 4], strides=[1, 1] : tensor<8x4xf32> -> !flow.dispatch.tensor<writeonly:8x4xf32>
flow.return
}
// CHECK-NEXT: return %[[RET1]]
@@ -150,9 +150,9 @@
%dim1_capture: index, %dim3_capture: index,
%ret: !flow.dispatch.tensor<writeonly:?x?x1024xf32>
) {
- %arg_tile = flow.dispatch.tensor.load %arg, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:7x?x24x?xf32>{%dim1_capture, %dim3_capture} -> tensor<7x?x24x?xf32>
+ %arg_tile = flow.dispatch.tensor.load %arg, offsets=[0, 0, 0, 0], sizes=[7, %dim1_capture, 24, %dim3_capture], strides=[1, 1, 1, 1] : !flow.dispatch.tensor<readonly:7x?x24x?xf32>{%dim1_capture, %dim3_capture} -> tensor<7x?x24x?xf32>
%ret_tile = "test.tile_math"(%arg_tile) : (tensor<7x?x24x?xf32>) -> (tensor<?x?x1024xf32>)
- flow.dispatch.tensor.store %ret_tile, %ret, offsets=[], sizes=[], strides=[] : tensor<?x?x1024xf32> -> !flow.dispatch.tensor<writeonly:?x?x1024xf32>{%dim3_capture, %dim1_capture}
+ flow.dispatch.tensor.store %ret_tile, %ret, offsets=[0, 0, 0], sizes=[%dim3_capture, %dim1_capture, 1024], strides=[1, 1, 1] : tensor<?x?x1024xf32> -> !flow.dispatch.tensor<writeonly:?x?x1024xf32>{%dim3_capture, %dim1_capture}
flow.return
}
// CHECK-NEXT: return %[[RET0]]
diff --git a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp
index 748a5a0..3923e96 100644
--- a/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp
+++ b/iree/compiler/Dialect/HAL/Conversion/StandardToHAL/ConvertStructuralOps.cpp
@@ -146,17 +146,17 @@
ifOp.getResultTypes(),
[&](Type type) { return getTypeConverter()->convertType(type); }));
auto newOp = rewriter.create<scf::IfOp>(ifOp.getLoc(), resultTypes,
- adaptor.condition(),
+ adaptor.getCondition(),
ifOp.elseBlock() != nullptr);
- rewriter.inlineRegionBefore(ifOp.thenRegion(), newOp.thenRegion(),
- newOp.thenRegion().end());
- rewriter.eraseBlock(&newOp.thenRegion().front());
+ rewriter.inlineRegionBefore(ifOp.getThenRegion(), newOp.getThenRegion(),
+ newOp.getThenRegion().end());
+ rewriter.eraseBlock(&newOp.getThenRegion().front());
if (ifOp.elseBlock()) {
- rewriter.inlineRegionBefore(ifOp.elseRegion(), newOp.elseRegion(),
- newOp.elseRegion().end());
- rewriter.eraseBlock(&newOp.elseRegion().front());
+ rewriter.inlineRegionBefore(ifOp.getElseRegion(), newOp.getElseRegion(),
+ newOp.getElseRegion().end());
+ rewriter.eraseBlock(&newOp.getElseRegion().front());
}
- rewriter.replaceOp(ifOp, newOp.results());
+ rewriter.replaceOp(ifOp, newOp.getResults());
return success();
}
};
@@ -166,7 +166,7 @@
LogicalResult matchAndRewrite(
scf::YieldOp yieldOp, OpAdaptor adaptor,
ConversionPatternRewriter &rewriter) const override {
- rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, adaptor.results());
+ rewriter.replaceOpWithNewOp<scf::YieldOp>(yieldOp, adaptor.getResults());
return success();
}
};
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
index 8ea58fd..fef60d5 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/test/smoketest.mlir
@@ -22,14 +22,14 @@
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg2 = stream.binding.subspan %arg2_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:16xf32>
%0 = linalg.init_tensor [16] : tensor<16xf32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %2 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) {
^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors
%4 = arith.addf %arg3, %arg4 : f32
linalg.yield %4 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %3, %arg2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
index 4ee1d57..158d0d4 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/test/smoketest.mlir
@@ -22,14 +22,14 @@
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg2 = stream.binding.subspan %arg2_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:16xf32>
%0 = linalg.init_tensor [16] : tensor<16xf32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %2 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) {
^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors
%4 = arith.addf %arg3, %arg4 : f32
linalg.yield %4 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %3, %arg2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/BUILD b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/BUILD
new file mode 100644
index 0000000..569c32f
--- /dev/null
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/BUILD
@@ -0,0 +1,26 @@
+# Copyright 2021 The IREE Authors
+#
+# Licensed 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
+
+load("//iree:lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ ["smoketest.mlir"],
+ include = ["*.mlir"],
+ ),
+ data = [
+ "//iree/tools:IreeFileCheck",
+ "//iree/tools:iree-opt",
+ ],
+)
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/CMakeLists.txt
index d5cc7b4..bc008da 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/CMakeLists.txt
@@ -1,18 +1,23 @@
-# Copyright 2020 The IREE Authors
-#
-# Licensed 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
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
iree_add_all_subdirs()
-file(GLOB _GLOB_X_MLIR LIST_DIRECTORIES false RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} CONFIGURE_DEPENDS *.mlir)
iree_lit_test_suite(
NAME
lit
SRCS
- "${_GLOB_X_MLIR}"
+ "smoketest.mlir"
DATA
iree::tools::IreeFileCheck
iree::tools::iree-opt
)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
index fa5345e..cf67808 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/test/smoketest.mlir
@@ -20,7 +20,7 @@
%arg0 = stream.binding.subspan %arg0_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:f32>
%0 = linalg.init_tensor [] : tensor<f32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor<f32>) {
^bb0(%arg2: f32, %arg3: f32):
%4 = arith.addf %arg2, %arg3 : f32
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
index e20aaac..7442f7b 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/test/smoketest.mlir
@@ -19,14 +19,14 @@
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg2 = stream.binding.subspan %arg2_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:16xf32>
%0 = linalg.init_tensor [16] : tensor<16xf32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %2 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) {
^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors
%4 = arith.addf %arg3, %arg4 : f32
linalg.yield %4 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %3, %arg2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
index fa0f28a..53eb702 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/test/smoketest.mlir
@@ -19,14 +19,14 @@
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg2 = stream.binding.subspan %arg2_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:16xf32>
%0 = linalg.init_tensor [16] : tensor<16xf32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
- %2 = flow.dispatch.tensor.load %arg1, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %2 = flow.dispatch.tensor.load %arg1, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>, affine_map<(d0) -> (d0)>], iterator_types = ["parallel"]} ins(%1, %2 : tensor<16xf32>, tensor<16xf32>) outs(%0 : tensor<16xf32>) {
^bb0(%arg3: f32, %arg4: f32, %arg5: f32): // no predecessors
%4 = arith.addf %arg3, %arg4 : f32
linalg.yield %4 : f32
} -> tensor<16xf32>
- flow.dispatch.tensor.store %3, %arg2, offsets=[], sizes=[], strides=[] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
+ flow.dispatch.tensor.store %3, %arg2, offsets=[0], sizes=[16], strides=[1] : tensor<16xf32> -> !flow.dispatch.tensor<writeonly:16xf32>
return
}
}
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
index c0c089b..b255a4e 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/test/smoketest.mlir
@@ -20,7 +20,7 @@
%arg0 = stream.binding.subspan %arg0_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:f32>
%0 = linalg.init_tensor [] : tensor<f32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor<f32>) {
^bb0(%arg2: f32, %arg3: f32):
%4 = arith.addf %arg2, %arg3 : f32
diff --git a/iree/compiler/Dialect/Stream/Transforms/PackConstants.cpp b/iree/compiler/Dialect/Stream/Transforms/PackConstants.cpp
index 8abbce5..cc789e4 100644
--- a/iree/compiler/Dialect/Stream/Transforms/PackConstants.cpp
+++ b/iree/compiler/Dialect/Stream/Transforms/PackConstants.cpp
@@ -393,8 +393,8 @@
ifResults.push_back(stagingResult.timepoint);
elseBuilder.create<scf::YieldOp>(loc, ifResults);
});
- auto ifTimepoint = ifOp.results().back();
- auto ifResources = ifOp.results().slice(0, ifOp.results().size() - 1);
+ auto ifTimepoint = ifOp.getResults().back();
+ auto ifResources = ifOp.getResults().slice(0, ifOp.getResults().size() - 1);
// Use the result of either the direct mapping or the staging upload.
UploadResult uploadResult;
diff --git a/iree/compiler/Dialect/Stream/Transforms/test/convert_to_stream.mlir b/iree/compiler/Dialect/Stream/Transforms/test/convert_to_stream.mlir
index 47bba67..3b64c1e 100644
--- a/iree/compiler/Dialect/Stream/Transforms/test/convert_to_stream.mlir
+++ b/iree/compiler/Dialect/Stream/Transforms/test/convert_to_stream.mlir
@@ -11,11 +11,11 @@
// CHECK: %[[ARG0_TENSOR:.+]] = stream.binding.subspan %arg0[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:?x4xf32>{%[[ARG0_DIM0]]}
// CHECK: %[[ARG1_TENSOR:.+]] = stream.binding.subspan %arg1[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:4x?xf32>{%[[ARG1_DIM1]]}
- // CHECK: %[[TILE:.+]] = flow.dispatch.tensor.load %[[ARG0_TENSOR]], offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x4xf32>{%[[ARG0_DIM0]]} -> tensor<?x4xf32>
- %0 = flow.dispatch.tensor.load %arg0, offsets = [], sizes = [], strides = [] : !flow.dispatch.tensor<readonly:?x4xf32>{%arg0_dim0} -> tensor<?x4xf32>
+ // CHECK: %[[TILE:.+]] = flow.dispatch.tensor.load %[[ARG0_TENSOR]], offsets = [0, 0], sizes = [%[[ARG0_DIM0]], 4], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x4xf32>{%[[ARG0_DIM0]]} -> tensor<?x4xf32>
+ %0 = flow.dispatch.tensor.load %arg0, offsets = [0, 0], sizes = [%arg0_dim0, 4], strides = [1, 1] : !flow.dispatch.tensor<readonly:?x4xf32>{%arg0_dim0} -> tensor<?x4xf32>
- // CHECK: flow.dispatch.tensor.store %[[TILE]], %[[ARG1_TENSOR]], offsets = [], sizes = [], strides = [] : tensor<?x4xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%[[ARG1_DIM1]]}
- flow.dispatch.tensor.store %0, %arg1, offsets = [], sizes = [], strides = [] : tensor<?x4xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%arg1_dim1}
+ // CHECK: flow.dispatch.tensor.store %[[TILE]], %[[ARG1_TENSOR]], offsets = [0, 0], sizes = [4, %[[ARG1_DIM1]]], strides = [1, 1] : tensor<?x4xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%[[ARG1_DIM1]]}
+ flow.dispatch.tensor.store %0, %arg1, offsets = [0, 0], sizes = [4, %arg1_dim1], strides = [1, 1] : tensor<?x4xf32> -> !flow.dispatch.tensor<writeonly:4x?xf32>{%arg1_dim1}
return
}
diff --git a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
index 4a6604f..20ce81f 100644
--- a/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
+++ b/iree/compiler/Dialect/Vulkan/Utils/test/target_env_conversion.mlir
@@ -22,7 +22,7 @@
%arg0 = stream.binding.subspan %arg0_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<readonly:16xf32>
%arg1 = stream.binding.subspan %arg1_binding[%c0] : !stream.binding -> !flow.dispatch.tensor<writeonly:f32>
%0 = linalg.init_tensor [] : tensor<f32>
- %1 = flow.dispatch.tensor.load %arg0, offsets=[], sizes=[], strides=[] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
+ %1 = flow.dispatch.tensor.load %arg0, offsets=[0], sizes=[16], strides=[1] : !flow.dispatch.tensor<readonly:16xf32> -> tensor<16xf32>
%3 = linalg.generic {indexing_maps = [affine_map<(d0) -> (d0)>, affine_map<(d0) -> ()>], iterator_types = ["reduction"]} ins(%1 : tensor<16xf32>) outs(%0 : tensor<f32>) {
^bb0(%arg2: f32, %arg3: f32):
%4 = arith.addf %arg2, %arg3 : f32
diff --git a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/LinalgExtOps.cpp b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/LinalgExtOps.cpp
index 16e98ba..abe37c0 100644
--- a/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/LinalgExtOps.cpp
+++ b/llvm-external-projects/iree-dialects/lib/Dialect/LinalgExt/IR/LinalgExtOps.cpp
@@ -471,7 +471,7 @@
});
auto &srcBlock = region().front();
- Region ®ion = scfFor.region();
+ Region ®ion = scfFor.getRegion();
BlockAndValueMapping bvm;
{
OpBuilder::InsertionGuard guard(b);
diff --git a/third_party/llvm-project b/third_party/llvm-project
index 505d574..128c6ed 160000
--- a/third_party/llvm-project
+++ b/third_party/llvm-project
@@ -1 +1 @@
-Subproject commit 505d57486e57eb61e29bed6517de5152d208fede
+Subproject commit 128c6ed73b8f906a13ae908008c6f415415964bb
