Remove use of iree.executable.workload attribute from SPIR-V translation.
With on-going work on dynamic shapes, use of attributes to specify the
workload is not feasible. Instead treat the shape of the outputs as
the launch size. Also move helper functions to get launch size,
workgroup size, etc. into translation folder, deprecating
IREECodegenUtils.h/.cpp
PiperOrigin-RevId: 295999426
diff --git a/iree/compiler/Translation/CMakeLists.txt b/iree/compiler/Translation/CMakeLists.txt
index 9dc0dc7..e08c48c 100644
--- a/iree/compiler/Translation/CMakeLists.txt
+++ b/iree/compiler/Translation/CMakeLists.txt
@@ -12,6 +12,7 @@
# See the License for the specific language governing permissions and
# limitations under the License.
+add_subdirectory(CodegenUtils)
add_subdirectory(Interpreter)
add_subdirectory(SPIRV)
add_subdirectory(XLAToLinalg)
diff --git a/iree/compiler/Translation/CodegenUtils/BUILD b/iree/compiler/Translation/CodegenUtils/BUILD
new file mode 100644
index 0000000..5c1dc39
--- /dev/null
+++ b/iree/compiler/Translation/CodegenUtils/BUILD
@@ -0,0 +1,38 @@
+# Copyright 2020 Google LLC
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# https://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+# Utilities for working with IREE MLIR types.
+
+package(
+ default_visibility = ["//visibility:public"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+cc_library(
+ name = "CodegenUtils",
+ srcs = [
+ "CodegenUtils.cpp",
+ ],
+ hdrs = [
+ "CodegenUtils.h",
+ ],
+ deps = [
+ "//iree/compiler/Dialect/Flow/Utils",
+ "//iree/compiler/Dialect/IREE/IR",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:StandardOps",
+ "@llvm-project//mlir:Support",
+ ],
+ alwayslink = 1,
+)
diff --git a/iree/compiler/Translation/CodegenUtils/CMakeLists.txt b/iree/compiler/Translation/CodegenUtils/CMakeLists.txt
new file mode 100644
index 0000000..125089a
--- /dev/null
+++ b/iree/compiler/Translation/CodegenUtils/CMakeLists.txt
@@ -0,0 +1,30 @@
+# Copyright 2020 Google LLC
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# https://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+iree_cc_library(
+ NAME
+ CodegenUtils
+ HDRS
+ "CodegenUtils.h"
+ SRCS
+ "CodegenUtils.cpp"
+ DEPS
+ MLIRIR
+ MLIRStandardOps
+ MLIRSupport
+ iree::compiler::Dialect::Flow::Utils
+ iree::compiler::Dialect::IREE::IR
+ ALWAYSLINK
+ PUBLIC
+)
diff --git a/iree/compiler/Translation/CodegenUtils/CodegenUtils.cpp b/iree/compiler/Translation/CodegenUtils/CodegenUtils.cpp
new file mode 100644
index 0000000..2b92b93
--- /dev/null
+++ b/iree/compiler/Translation/CodegenUtils/CodegenUtils.cpp
@@ -0,0 +1,130 @@
+// Copyright 2020 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "iree/compiler/Translation/CodegenUtils/CodegenUtils.h"
+
+#include "iree/compiler/Dialect/Flow/Utils/WorkloadUtils.h"
+#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
+#include "mlir/Dialect/StandardOps/Ops.h"
+#include "mlir/IR/StandardTypes.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+ArrayRef<int64_t> dropTrailingOnes(ArrayRef<int64_t> vector) {
+ if (vector.empty()) return vector;
+ auto numTrailingOnes = 0;
+ for (unsigned i = vector.size() - 1; i > 0; --i) {
+ if (vector[i] != 1) {
+ break;
+ }
+ numTrailingOnes++;
+ }
+ return vector.drop_back(numTrailingOnes);
+}
+
+bool isDispatchFunction(FuncOp funcOp) {
+ return funcOp.getAttr("iree.executable.export") != nullptr;
+}
+
+/// Helper function to check shapes are equal.
+static bool areShapesEqual(ArrayRef<int64_t> lhs, ArrayRef<int64_t> rhs) {
+ return lhs == rhs;
+}
+
+/// Get the shape to use for a type. For now this is returning shapes as static
+/// value.
+// TODO(ravishankarm) : Modify this to return the Values to use for the extent.
+static LogicalResult getExtentFromStoreOpSrc(Operation *storeOp,
+ SmallVectorImpl<int64_t> &extent) {
+ extent.clear();
+ extent.resize(3, 1);
+ Value srcVal = storeOp->getOperand(0);
+ if (srcVal.getType().isIntOrFloat()) return success();
+ auto workload = dyn_cast_or_null<ConstantOp>(
+ IREE::Flow::calculateWorkload(srcVal.getDefiningOp(), srcVal)
+ .getDefiningOp());
+ if (!workload) return failure();
+ auto extentVal = workload.getValue().cast<DenseIntElementsAttr>();
+ for (auto val : enumerate(extentVal.getValues<APInt>()))
+ extent[val.index()] = val.value().getSExtValue();
+ workload.erase();
+ return success();
+}
+
+LogicalResult getLaunchSize(FuncOp funcOp,
+ SmallVectorImpl<int64_t> &launchSize) {
+ auto &body = funcOp.getBody();
+ if (!mlir::has_single_element(body)) {
+ return funcOp.emitError(
+ "unhandled multiple blocks within dispatch function");
+ }
+ SmallVector<Operation *, 1> storeOperations;
+ auto storeOps = body.front().getOps<IREE::StoreOutputOp>();
+ storeOperations.assign(storeOps.begin(), storeOps.end());
+ auto storeReduceOps = body.front().getOps<IREE::StoreReduceOp>();
+ storeOperations.append(storeReduceOps.begin(), storeReduceOps.end());
+ if (storeOperations.begin() == storeOperations.end()) {
+ return funcOp.emitError(
+ "expected dispatch function to have at least one iree.store_output "
+ "instruction");
+ }
+
+ Operation *firstStoreOp = *storeOperations.begin();
+ if (failed(getExtentFromStoreOpSrc(firstStoreOp, launchSize))) {
+ return firstStoreOp->emitError("unhandled type of the output tensor");
+ }
+ for (auto it = std::next(storeOperations.begin()), ie = storeOperations.end();
+ it != ie; ++it) {
+ SmallVector<int64_t, 3> checkShape;
+ Operation *storeOp = *it;
+ if (failed(getExtentFromStoreOpSrc(storeOp, checkShape))) {
+ return storeOp->emitError("unhandled type of the output tensor");
+ }
+ if (!areShapesEqual(launchSize, checkShape)) {
+ return storeOp->emitError("mismatch in shapes of the output tensors");
+ }
+ }
+ return success();
+}
+
+/// Gets the workgroup size.
+template <typename intType>
+LogicalResult getWorkGroupSize(FuncOp funcOp,
+ SmallVectorImpl<intType> &workGroupSize) {
+ if (!funcOp.getAttr("iree.executable.export")) {
+ return funcOp.emitError(
+ "expected operation to be in dispatch function to get launch size");
+ }
+ auto workGroupSizeAttr =
+ funcOp.getAttrOfType<DenseElementsAttr>("iree.executable.workgroup_size");
+ if (!workGroupSizeAttr) {
+ return funcOp.emitError(
+ "unable to find workload size, missing attribute "
+ "iree.executable.workload in dispatch function");
+ }
+ workGroupSize.clear();
+ for (auto value : workGroupSizeAttr.getValues<APInt>()) {
+ workGroupSize.push_back(value.getSExtValue());
+ }
+ return success();
+}
+
+template LogicalResult getWorkGroupSize<int32_t>(
+ FuncOp funcOp, SmallVectorImpl<int32_t> &workGroupSize);
+template LogicalResult getWorkGroupSize<int64_t>(
+ FuncOp funcOp, SmallVectorImpl<int64_t> &workGroupSize);
+
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Translation/CodegenUtils/CodegenUtils.h b/iree/compiler/Translation/CodegenUtils/CodegenUtils.h
new file mode 100644
index 0000000..bc5be2a
--- /dev/null
+++ b/iree/compiler/Translation/CodegenUtils/CodegenUtils.h
@@ -0,0 +1,43 @@
+// Copyright 2020 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef IREE_COMPILER_TRANSLATION_CODEGENUTILS_CODEGENUTILS_H
+#define IREE_COMPILER_TRANSLATION_CODEGENUTILS_CODEGENUTILS_H
+
+#include "mlir/IR/Function.h"
+#include "mlir/Support/LogicalResult.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+/// Drop trailing ones.
+ArrayRef<int64_t> dropTrailingOnes(ArrayRef<int64_t> vector);
+
+/// Checks that a given function is a dispatch function.
+bool isDispatchFunction(FuncOp funcOp);
+
+/// The launch size is the size of the outputs of the kernel. For now all
+/// outputs have to be the same shape and static shaped.
+LogicalResult getLaunchSize(FuncOp funcOp,
+ SmallVectorImpl<int64_t> &launchSize);
+
+/// Gets the workgroup size. Has to be a static constant.
+template <typename intType>
+LogicalResult getWorkGroupSize(FuncOp funcOp,
+ SmallVectorImpl<intType> &workGroupSize);
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_TRANSLATION_CODEGENUTILS_CODEGENUTILS_H
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/BUILD b/iree/compiler/Translation/SPIRV/IndexComputation/BUILD
index 25a9d61..da35230 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/BUILD
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/BUILD
@@ -60,8 +60,8 @@
],
deps = [
"//iree/compiler/Dialect/IREE/IR",
+ "//iree/compiler/Translation/CodegenUtils",
"//iree/compiler/Translation/SPIRV/IndexComputation:IndexComputationAttrGen",
- "//iree/compiler/Utils",
"@llvm-project//llvm:support",
"@llvm-project//mlir:IR",
"@llvm-project//mlir:Pass",
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/CMakeLists.txt b/iree/compiler/Translation/SPIRV/IndexComputation/CMakeLists.txt
index a94c898..760a1d2 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/CMakeLists.txt
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/CMakeLists.txt
@@ -48,8 +48,8 @@
MLIRSPIRV
MLIRStandardOps
iree::compiler::Dialect::IREE::IR
+ iree::compiler::Translation::CodegenUtils
iree::compiler::Translation::SPIRV::IndexComputation::IndexComputationAttrGen
- iree::compiler::Utils
tensorflow::mlir_xla
ALWAYSLINK
PUBLIC
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.cpp b/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.cpp
index 9199416..3d7b72b 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.cpp
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.cpp
@@ -49,7 +49,7 @@
}
SmallVector<int64_t, 3> launchSize;
- if (failed(getLegacyLaunchSize(funcOp, launchSize))) {
+ if (failed(getLaunchSize(funcOp, launchSize))) {
return failure();
}
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.h b/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.h
index 3a1ef6d..0919ea6 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.h
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/IREEIndexComputation.h
@@ -20,8 +20,9 @@
#ifndef IREE_COMPILER_TRANSLATION_SPIRV_INDEXCOMPUTATION_IREEINDEXCOMP_H
#define IREE_COMPILER_TRANSLATION_SPIRV_INDEXCOMPUTATION_IREEINDEXCOMP_H
+#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
+#include "iree/compiler/Translation/CodegenUtils/CodegenUtils.h"
#include "iree/compiler/Translation/SPIRV/IndexComputation/XLAIndexPropagation.h"
-#include "iree/compiler/Utils/IREECodegenUtils.h"
#include "mlir/IR/Function.h"
namespace mlir {
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/IndexComputation.h b/iree/compiler/Translation/SPIRV/IndexComputation/IndexComputation.h
index 236d84b..dbbcf66 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/IndexComputation.h
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/IndexComputation.h
@@ -21,8 +21,8 @@
#ifndef IREE_COMPILER_TRANSLATION_SPIRV_INDEXCOMPUTATION_INDEXCOMPUTATION_H
#define IREE_COMPILER_TRANSLATION_SPIRV_INDEXCOMPUTATION_INDEXCOMPUTATION_H
+#include "iree/compiler/Translation/CodegenUtils/CodegenUtils.h"
#include "iree/compiler/Translation/SPIRV/IndexComputation/IndexComputationAttribute.h"
-#include "iree/compiler/Utils/IREECodegenUtils.h"
#include "llvm/ADT/ArrayRef.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
@@ -298,7 +298,7 @@
// Set the attribute for the number of launch dims.
auto funcOp = region.getParentOfType<FuncOp>();
SmallVector<int64_t, 3> launchSize;
- if (failed(getLegacyLaunchSize(funcOp, launchSize))) {
+ if (failed(getLaunchSize(funcOp, launchSize))) {
return emitError(region.getLoc(),
"expected region of index propagation to be in dispatch "
"function to get launch size");
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast.mlir
index 3017f54..5d26e92 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast.mlir
@@ -9,7 +9,7 @@
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
func @broadcast_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
// CHECK: xla_hlo.broadcast
// CHECK-SAME: iree.index_computation_info
@@ -34,7 +34,7 @@
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
func @broadcast_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32>
// CHECK: xla_hlo.broadcast
// CHECK-SAME: iree.index_computation_info
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast_in_dim.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast_in_dim.mlir
index 0d21047..a6b01c2 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast_in_dim.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/broadcast_in_dim.mlir
@@ -9,7 +9,7 @@
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
func @broadcast_in_dim_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
// CHECK: xla_hlo.broadcast_in_dim
// CHECK-SAME: iree.index_computation_info
@@ -34,7 +34,7 @@
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
func @broadcast_in_dim_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32>
// CHECK: xla_hlo.broadcast_in_dim
// CHECK-SAME: iree.index_computation_info
@@ -52,20 +52,20 @@
module {
func @const_float_splat(%arg0: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: constant
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1)>]
%0 = constant dense<1.0> : tensor<12xf32>
// CHECK: xla_hlo.broadcast_in_dim
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[0]> : tensor<1xi64>} : (tensor<12xf32>) -> tensor<12x42xf32>
iree.store_output(%1 : tensor<12x42xf32>, %arg0 : memref<12x42xf32>)
iree.return
@@ -76,20 +76,20 @@
module {
func @const_int_splat(%arg0: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: constant
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1)>]
%0 = constant dense<42> : tensor<12xi32>
// CHECK: xla_hlo.broadcast_in_dim
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[0]> : tensor<1xi64>} : (tensor<12xi32>) -> tensor<12x42xi32>
iree.store_output(%1 : tensor<12x42xi32>, %arg0 : memref<12x42xi32>)
iree.return
@@ -100,7 +100,7 @@
module {
func @const_int_nonsplat(%arg0: memref<2x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 2]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: constant
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
@@ -124,7 +124,7 @@
module {
func @zero_element_1dtensor(%arg0 : memref<f32>, %arg1 : memref<4xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<f32>) : tensor<f32>
// CHECK: xla_hlo.broadcast_in_dim
// CHECK-SAME: iree.index_computation_info
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/concatenate.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/concatenate.mlir
index 0dd53db..7d14b70 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/concatenate.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/concatenate.mlir
@@ -7,24 +7,24 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0, d0)>]
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: memref<1x10xf32>
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (0, d1 - 64)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0, d0 - 64)>]
func @concatenate(%arg0: memref<1x64xf32>, %arg1 : memref<1x10xf32>, %arg2 : memref<1x74xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[1, 74]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<1x64xf32>) : tensor<1x64xf32>
%1 = iree.load_input(%arg1 : memref<1x10xf32>) : tensor<1x10xf32>
// CHECK: xla_hlo.concatenate
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (0, d1)>]
- // CHECK-SAME: [affine_map<(d0, d1) -> (0, d1 - 64)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0, d0 - 64)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0, d0)>]
%2 = "xla_hlo.concatenate"(%0, %1) {dimension = 1 : i64} : (tensor<1x64xf32>, tensor<1x10xf32>) -> tensor<1x74xf32>
iree.store_output(%2 : tensor<1x74xf32>, %arg2 : memref<1x74xf32>)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/copy.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/copy.mlir
index 93bb5f6..961c0da 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/copy.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/copy.mlir
@@ -7,22 +7,22 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
func @simple_load_store(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: iree.load_input
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
// CHECK: xla_hlo.copy
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.copy"(%0) : (tensor<12x42xi32>) -> tensor<12x42xi32>
iree.store_output(%1 : tensor<12x42xi32>, %arg1 : memref<12x42xi32>)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/gather.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/gather.mlir
index 66860da..98269d5 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/gather.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/gather.mlir
@@ -17,7 +17,7 @@
// CHECK-SAME: iree.symbol_number_info
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>, 0 : i32]
func @foo(%arg0: memref<5x1x10xf32>, %arg1: memref<i64>, %arg2: memref<1x10xf32>)
- attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5x1x10xf32>) : tensor<5x1x10xf32>
%1 = iree.load_input(%arg1 : memref<i64>) : tensor<i64>
%2 = "xla_hlo.gather"(%0, %1) {
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/pad.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/pad.mlir
index fde7bdf..248575e 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/pad.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/pad.mlir
@@ -7,18 +7,18 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1 - 4, d0 - 5)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 - 4, d0 - 5)>]
func @pad_zero_interior(%arg0 : memref<12x4xf32>, %arg1 : memref<18x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 18, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
%1 = constant dense<0.0> : tensor<f32>
// CHECK: xla_hlo.pad
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1 - 4, d0 - 5)>]
- // CHECK-SAME: [affine_map<(d0, d1) -> (0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 - 4, d0 - 5)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%2 = "xla_hlo.pad"(%0, %1) {edge_padding_high = dense<[2, 3]> : tensor<2xi64>, edge_padding_low = dense<[4, 5]> : tensor<2xi64>, interior_padding = dense<0> : tensor<2xi64>} : (tensor<12x4xf32>, tensor<f32>) -> tensor<18x12xf32>
iree.store_output(%2 : tensor<18x12xf32>, %arg1 : memref<18x12xf32>)
iree.return
@@ -33,18 +33,18 @@
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices = []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
func @pad_no_op(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[4, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
%1 = constant dense<0.0> : tensor<f32>
// CHECK: xla_hlo.pad
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
- // CHECK-SAME: [affine_map<(d0, d1) -> (0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%2 = "xla_hlo.pad"(%0, %1) {edge_padding_high = dense<[0, 0]> : tensor<2xi64>, edge_padding_low = dense<[0, 0]> : tensor<2xi64>, interior_padding = dense<0> : tensor<2xi64>} : (tensor<12x4xf32>, tensor<f32>) -> tensor<12x4xf32>
iree.store_output(%2 : tensor<12x4xf32>, %arg1 : memref<12x4xf32>)
iree.return
@@ -60,18 +60,18 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1 floordiv 2 - 2, (d0 - 5) floordiv 3)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 floordiv 2 - 2, (d0 - 5) floordiv 3)>]
func @pad_with_stride(%arg0 : memref<12x4xf32>, %arg1 : memref<29x18xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[18, 29, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
%1 = constant dense<0.0> : tensor<f32>
// CHECK: xla_hlo.pad
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1 floordiv 2 - 2, (d0 - 5) floordiv 3)>]
- // CHECK-SAME: [affine_map<(d0, d1) -> (0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 floordiv 2 - 2, (d0 - 5) floordiv 3)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%2 = "xla_hlo.pad"(%0, %1) {edge_padding_high = dense<[2, 3]> : tensor<2xi64>, edge_padding_low = dense<[4, 5]> : tensor<2xi64>, interior_padding = dense<[1, 2]> : tensor<2xi64>} : (tensor<12x4xf32>, tensor<f32>) -> tensor<29x18xf32>
iree.store_output(%2 : tensor<29x18xf32>, %arg1 : memref<29x18xf32>)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/reverse.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/reverse.mlir
index db34482..2c5cce7 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/reverse.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/reverse.mlir
@@ -7,16 +7,16 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (-d1 + 11, -d0 + 11)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (-d1 + 11, -d0 + 11)>]
func @reverse_2d(%arg0: memref<12x12xf32>, %arg1 : memref<12x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 12]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x12xf32>) : tensor<12x12xf32>
// CHECK: xla_hlo.reverse
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (-d1 + 11, -d0 + 11)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (-d1 + 11, -d0 + 11)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.reverse"(%0) {dimensions = dense<[1, 0]> : tensor<2xi64>} : (tensor<12x12xf32>) -> tensor<12x12xf32>
iree.store_output(%1 : tensor<12x12xf32>, %arg1 : memref<12x12xf32>)
iree.return
@@ -34,7 +34,7 @@
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (d2, -d1 + 2, d0)>]
func @reverse_3d(%arg0: memref<3x3x3xf32>, %arg1 : memref<3x3x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[3, 3, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<3x3x3xf32>) : tensor<3x3x3xf32>
// CHECK: xla_hlo.reverse
// CHECK-SAME: iree.index_computation_info
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/slice.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/slice.mlir
index ea5c501..e5cd2d1 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/slice.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/slice.mlir
@@ -7,16 +7,16 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3 + 2, d0 mod 3 + 1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 + 2, d0 + 1)>]
func @slice_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<6x6xf32>) : tensor<6x6xf32>
// CHECK: xla_hlo.slice
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3 + 2, d0 mod 3 + 1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 + 2, d0 + 1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3, d0 mod 3)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.slice"(%0) {start_indices = dense<[2, 1]> : tensor<2xi64>, limit_indices = dense<[4, 4]> : tensor<2xi64>, strides = dense<[1, 1]> : tensor<2xi64>} : (tensor<6x6xf32>) -> tensor<2x3xf32>
iree.store_output(%1 : tensor<2x3xf32>, %arg1 : memref<2x3xf32>)
iree.return
@@ -32,16 +32,16 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3 + 2, (d0 mod 3) * 2 + 1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 + 2, d0 * 2 + 1)>]
func @slice_non_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<6x6xf32>) : tensor<6x6xf32>
// CHECK: xla_hlo.slice
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3 + 2, (d0 mod 3) * 2 + 1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1 + 2, d0 * 2 + 1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0 floordiv 3, d0 mod 3)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.slice"(%0) {start_indices = dense<[2, 1]> : tensor<2xi64>, limit_indices = dense<[4, 6]> : tensor<2xi64>, strides = dense<[1, 2]> : tensor<2xi64>} : (tensor<6x6xf32>) -> tensor<2x3xf32>
iree.store_output(%1 : tensor<2x3xf32>, %arg1 : memref<2x3xf32>)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/store_reduce.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/store_reduce.mlir
index d9c0501..cbe345c 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/store_reduce.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/store_reduce.mlir
@@ -14,7 +14,7 @@
// CHECK-SAME: []
// CHECK-SAME: result_index
// CHECK-SAME: [affine_map<(d0, d1, d2) -> (0)>]
- func @reduction_entry(%arg0: memref<5xi32>, %arg1: memref<i32>, %arg2: memref<i32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[5, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ func @reduction_entry(%arg0: memref<5xi32>, %arg1: memref<i32>, %arg2: memref<i32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5xi32>) : tensor<5xi32>
iree.store_reduce(%0 : tensor<5xi32>, %arg2 : memref<i32>, @reduction_apply)
iree.return
@@ -30,14 +30,14 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9_]*]]: memref<4xi32>
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0)>]
- func @reduction_2D_dim0_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<4xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[4, 5, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d0)>]
+ func @reduction_2D_dim0_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<4xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5x4xi32>) : tensor<5x4xi32>
iree.store_reduce(%0 : tensor<5x4xi32>, %arg2 : memref<4xi32>, @reduction_apply)
iree.return
@@ -53,14 +53,14 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9_]*]]: memref<5xi32>
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1)>]
- func @reduction_2D_dim1_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<5xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 1 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[4, 5, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1)>]
+ func @reduction_2D_dim1_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<5xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 1 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5x4xi32>) : tensor<5x4xi32>
iree.store_reduce(%0 : tensor<5x4xi32>, %arg2 : memref<5xi32>, @reduction_apply)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/IndexComputation/test/transpose_add.mlir b/iree/compiler/Translation/SPIRV/IndexComputation/test/transpose_add.mlir
index 60c2e25..853a474 100644
--- a/iree/compiler/Translation/SPIRV/IndexComputation/test/transpose_add.mlir
+++ b/iree/compiler/Translation/SPIRV/IndexComputation/test/transpose_add.mlir
@@ -7,38 +7,38 @@
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: operand_indices
// CHECK-SAME: []
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d0, d1)>]
func @transpose_add(%arg0: memref<12x12xf32>, %arg1: memref<12x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: iree.load_input
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d0, d1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d0, d1)>]
%0 = iree.load_input(%arg0 : memref<12x12xf32>) : tensor<12x12xf32>
// CHECK: xla_hlo.transpose
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d0, d1)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d0, d1)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%1 = "xla_hlo.transpose"(%0) {permutation = dense<[1, 0]> : tensor<2xi64>} : (tensor<12x12xf32>) -> tensor<12x12xf32>
// CHECK: xla_hlo.add
// CHECK-SAME: iree.index_computation_info
// CHECK-SAME: operand_indices
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
// CHECK-SAME: result_index
- // CHECK-SAME: [affine_map<(d0, d1) -> (d1, d0)>]
+ // CHECK-SAME: [affine_map<(d0, d1, d2) -> (d1, d0)>]
%2 = xla_hlo.add %0, %1 : tensor<12x12xf32>
iree.store_output(%2 : tensor<12x12xf32>, %arg1 : memref<12x12xf32>)
iree.return
diff --git a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/BUILD b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/BUILD
index 9677f61..cf03821 100644
--- a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/BUILD
+++ b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/BUILD
@@ -29,8 +29,8 @@
],
deps = [
"//iree/compiler/Dialect/IREE/IR",
+ "//iree/compiler/Translation/CodegenUtils",
"//iree/compiler/Translation/XLAToLinalg:IREELinalgTensorToBuffer",
- "//iree/compiler/Utils",
"@llvm-project//llvm:support",
"@llvm-project//mlir:AffineOps",
"@llvm-project//mlir:AffineToStandardTransforms",
diff --git a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/CMakeLists.txt b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/CMakeLists.txt
index a18b6fa..83f9713 100644
--- a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/CMakeLists.txt
+++ b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/CMakeLists.txt
@@ -43,8 +43,8 @@
MLIRSupport
MLIRTransforms
iree::compiler::Dialect::IREE::IR
+ iree::compiler::Translation::CodegenUtils
iree::compiler::Translation::XLAToLinalg::IREELinalgTensorToBuffer
- iree::compiler::Utils
tensorflow::mlir_xla
ALWAYSLINK
PUBLIC
diff --git a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
index dfe6090..128a9e9 100644
--- a/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
+++ b/iree/compiler/Translation/SPIRV/LinalgToSPIRV/LowerToSPIRV.cpp
@@ -18,9 +18,9 @@
//
//===----------------------------------------------------------------------===//
+#include "iree/compiler/Translation/CodegenUtils/CodegenUtils.h"
#include "iree/compiler/Translation/SPIRV/LinalgToSPIRV/Passes.h"
#include "iree/compiler/Translation/XLAToLinalg/LinalgTensorToBuffer.h"
-#include "iree/compiler/Utils/IREECodegenUtils.h"
#include "mlir/Conversion/GPUToSPIRV/ConvertGPUToSPIRV.h"
#include "mlir/Conversion/LoopsToGPU/LoopsToGPU.h"
#include "mlir/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.h"
@@ -49,18 +49,6 @@
namespace mlir {
namespace iree_compiler {
-static ArrayRef<int64_t> dropTrailingOnes(ArrayRef<int64_t> vector) {
- if (vector.empty()) return vector;
- auto numTrailingOnes = 0;
- for (unsigned i = vector.size() - 1; i > 0; --i) {
- if (vector[i] != 1) {
- break;
- }
- numTrailingOnes++;
- }
- return vector.drop_back(numTrailingOnes);
-}
-
namespace {
/// These options are only for testing purposes. For actual execution with IREE,
/// these are computed by IREE/Backends automatically.
@@ -109,7 +97,7 @@
FuncOp funcOp = getFunction();
SmallVector<int64_t, 3> workGroupSizeVec;
workGroupSizeVec.reserve(3);
- if (failed(getLegacyWorkGroupSize(funcOp, workGroupSizeVec))) return;
+ if (failed(getWorkGroupSize(funcOp, workGroupSizeVec))) return;
ArrayRef<int64_t> workGroupSize = dropTrailingOnes(workGroupSizeVec);
OpBuilder builder(funcOp);
@@ -162,29 +150,16 @@
FuncOp funcOp = getFunction();
SmallVector<int64_t, 3> workGroupSizeVec;
workGroupSizeVec.reserve(3);
- if (failed(getLegacyWorkGroupSize(funcOp, workGroupSizeVec))) return;
+ if (failed(getWorkGroupSize(funcOp, workGroupSizeVec))) return;
ArrayRef<int64_t> workGroupSize = dropTrailingOnes(workGroupSizeVec);
- // While we can use any valid input for numWorkGroups, there might be
- // canonicalizations that can be used if the workgroup size is passed
- // accurately. For now compute the workgroup size based on the workload and
- // workgroup size.
- // TODO(ravishankarm): This assumes this is static for now. To handle
- // dynamic cases, generate the IR that corresponds to the operations here.
- SmallVector<int64_t, 3> workLoad;
- workLoad.reserve(3);
- if (failed(getLegacyLaunchSize(funcOp, workLoad))) {
- funcOp.emitError("unable to retrieve workload size in dispatch function");
- return signalPassFailure();
- }
- workLoad.resize(workGroupSize.size());
-
+ // For now just use number of workgroups to be [1, 1, 1]. The loop to GPU
+ // lowering doesnt use the value of number of workgroups in the codegen
+ // itself, but rather only uses this in the gpu.launch op which is
+ // irrelevant for IREE.
+ // TODO(ravishankarm): Fix the GPU lowering to allow not using gpu.launch at
+ // all.
SmallVector<int64_t, 3> numWorkGroups(workGroupSize.size(), 1);
- for (auto index : llvm::seq<unsigned>(0, workGroupSize.size())) {
- numWorkGroups[index] = workLoad[index] / workGroupSize[index];
- numWorkGroups[index] +=
- static_cast<bool>(workLoad[index] % workGroupSize[index]);
- }
SmallVector<Value, 3> numWorkGroupsVal, workGroupSizeVal;
numWorkGroupsVal.reserve(3);
@@ -210,7 +185,7 @@
ModuleOp moduleOp = getModule();
FuncOp funcOp = nullptr;
auto walkResult = moduleOp.walk([&funcOp](FuncOp fOp) -> WalkResult {
- if (fOp.getAttr("iree.executable.export")) {
+ if (isDispatchFunction(fOp)) {
if (funcOp) return WalkResult::interrupt();
funcOp = fOp;
}
@@ -258,7 +233,7 @@
SPIRVTypeConverter typeConverter;
OwningRewritePatternList patterns;
SmallVector<int32_t, 3> workGroupSize;
- if (failed(getLegacyWorkGroupSize(funcOp, workGroupSize))) return;
+ if (failed(getWorkGroupSize(funcOp, workGroupSize))) return;
// Set spv.entry_point_abi on each kernel functions to drive SPIR-V CodeGen.
// This is required because SPIR-V CodeGen's contract.
@@ -292,7 +267,7 @@
: workGroupSize(workGroupSize.begin(), workGroupSize.end()) {}
void runOnFunction() {
FuncOp funcOp = getFunction();
- if (!funcOp.getAttr("iree.executable.export")) return;
+ if (!isDispatchFunction(funcOp)) return;
if (workGroupSize.empty()) {
// By default look at the number of "parallel" loops in the generic op.
diff --git a/iree/compiler/Translation/SPIRV/ReductionCodegen/test/simple.mlir b/iree/compiler/Translation/SPIRV/ReductionCodegen/test/simple.mlir
index 985171d..903485a 100644
--- a/iree/compiler/Translation/SPIRV/ReductionCodegen/test/simple.mlir
+++ b/iree/compiler/Translation/SPIRV/ReductionCodegen/test/simple.mlir
@@ -6,7 +6,7 @@
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: memref<i32>
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9]*]]: memref<4xi32> {iree.executable.reduction.output}
// CHECK-SAME: iree.executable.reduction.apply = [[APPLYFN:@[a-zA-Z0-9_]*]]
- func @reduction_entry(memref<5x4xi32>, memref<i32>, memref<4xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[4, 5, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32}
+ func @reduction_entry(memref<5x4xi32>, memref<i32>, memref<4xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32}
// CHECK: [[TENSOR:%.*]] = iree.load_input([[ARG0]] : memref<5x4xi32>) : tensor<5x4xi32>
// CHECK: iree.store_reduce([[TENSOR]] : tensor<5x4xi32>, [[ARG2]] : memref<4xi32>, [[APPLYFN]])
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/BUILD b/iree/compiler/Translation/SPIRV/XLAToSPIRV/BUILD
index b11906a..200b440 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/BUILD
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/BUILD
@@ -34,10 +34,10 @@
],
deps = [
"//iree/compiler/Dialect/IREE/IR",
+ "//iree/compiler/Translation/CodegenUtils",
"//iree/compiler/Translation/SPIRV/IndexComputation",
"//iree/compiler/Translation/SPIRV/Passes",
"//iree/compiler/Translation/SPIRV/ReductionCodegen",
- "//iree/compiler/Utils",
"@llvm-project//llvm:support",
"@llvm-project//mlir:AffineOps",
"@llvm-project//mlir:AffineToStandardTransforms",
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/CMakeLists.txt b/iree/compiler/Translation/SPIRV/XLAToSPIRV/CMakeLists.txt
index 5199a5a..109844c 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/CMakeLists.txt
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/CMakeLists.txt
@@ -41,10 +41,10 @@
MLIRSupport
MLIRTransforms
iree::compiler::Dialect::IREE::IR
+ iree::compiler::Translation::CodegenUtils
iree::compiler::Translation::SPIRV::IndexComputation
iree::compiler::Translation::SPIRV::Passes
iree::compiler::Translation::SPIRV::ReductionCodegen
- iree::compiler::Utils
tensorflow::mlir_xla
ALWAYSLINK
PUBLIC
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.cpp b/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.cpp
index ca7bb66..692618e 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.cpp
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.cpp
@@ -139,7 +139,7 @@
entryFn.addEntryBlock();
SmallVector<int32_t, 3> workGroupSize;
- if (failed(getLegacyWorkGroupSize(fn, workGroupSize))) {
+ if (failed(getWorkGroupSize(fn, workGroupSize))) {
return failure();
}
auto entryFnAttr =
@@ -188,7 +188,7 @@
FuncOp fn) {
// First check that the global invocation id is in bounds.
SmallVector<int64_t, 3> launchSize;
- if (failed(getLegacyLaunchSize(fn, launchSize))) {
+ if (failed(getLaunchSize(fn, launchSize))) {
return failure();
}
auto loc = fn.getLoc();
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.h b/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.h
index 6a35a39..2d2ed7f 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.h
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/SPIRVLowering.h
@@ -20,9 +20,9 @@
#ifndef IREE_COMPILER_TRANSLATION_SPIRV_XLATOSPIRV_SPIRVLOWERING_H
#define IREE_COMPILER_TRANSLATION_SPIRV_XLATOSPIRV_SPIRVLOWERING_H
+#include "iree/compiler/Translation/CodegenUtils/CodegenUtils.h"
#include "iree/compiler/Translation/SPIRV/IndexComputation/IndexComputationAttribute.h"
#include "iree/compiler/Translation/SPIRV/XLAToSPIRV/TensorIndexToScalarValueMap.h"
-#include "iree/compiler/Utils/IREECodegenUtils.h"
#include "mlir/Dialect/SPIRV/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/SPIRVOps.h"
#include "mlir/Dialect/StandardOps/Ops.h"
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/arithmetic_ops.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/arithmetic_ops.mlir
index 30dd16f..c02f4de 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/arithmetic_ops.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/arithmetic_ops.mlir
@@ -6,7 +6,7 @@
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<4 x f32 [4]> [0]>, StorageBuffer>
func @mul_1D(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xf32>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]]
// CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]]
// CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32]
@@ -31,7 +31,7 @@
// -----
func @frem(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xf32>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FRem
@@ -43,7 +43,7 @@
// -----
func @srem(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi32>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SRem
@@ -55,7 +55,7 @@
// -----
func @srem(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi32>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SRem
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
index d8f1979..1d7360a 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast.mlir
@@ -8,7 +8,7 @@
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<1512 x i32 [4]> [0]>, StorageBuffer>
func @broadcast_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
@@ -29,7 +29,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<1512 x i32 [4]> [0]>, StorageBuffer>
func @broadcast_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
index 80378e9..00666f0 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/broadcast_in_dim.mlir
@@ -7,7 +7,7 @@
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1512 x i32 [4]> [0]>, StorageBuffer>
func @broadcast_in_dim_2D_3D(%arg0: memref<12x42xi32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[1, 2]> : tensor<2xi64>} : (tensor<12x42xi32>) -> tensor<3x12x42xi32>
iree.store_output(%1 : tensor<3x12x42xi32>, %arg1 : memref<3x12x42xi32>)
@@ -24,7 +24,7 @@
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<1512 x i32 [4]> [0]>, StorageBuffer>
func @broadcast_in_dim_scalar_3D(%arg0: memref<i32>, %arg1: memref<3x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<i32>) : tensor<i32>
@@ -40,7 +40,7 @@
module {
func @const_float_splat(%arg0: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: spv.constant 1.000000e+00 : f32
%0 = constant dense<1.0> : tensor<12xf32>
%1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[0]> : tensor<1xi64>} : (tensor<12xf32>) -> tensor<12x42xf32>
@@ -53,7 +53,7 @@
module {
func @const_int_splat(%arg0: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: spv.constant 42 : i32
%0 = constant dense<42> : tensor<12xi32>
%1 = "xla_hlo.broadcast_in_dim"(%0) {broadcast_dimensions = dense<[0]> : tensor<1xi64>} : (tensor<12xi32>) -> tensor<12x42xi32>
@@ -68,7 +68,7 @@
// CHECK: spv.func @const_int_nonsplat
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<1008 x i32 [4]> [0]>, StorageBuffer>
func @const_int_nonsplat(%arg0: memref<2x12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 2]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[CST:%.*]] = spv.constant dense<[42, 21]>
// CHECK: [[VAR:%.*]] = spv.Variable init([[CST]]) : !spv.ptr<!spv.array<2 x i32 [4]>, Function>
// CHECK: [[LOADPTR:%.*]] = spv.AccessChain [[VAR]]
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/compare.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/compare.mlir
index 73f842c..8f64a4b 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/compare.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/compare.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -iree-index-computation -simplify-spirv-affine-exprs=false -convert-iree-to-spirv -verify-diagnostics -o - %s | IreeFileCheck %s
func @ieq(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.IEqual
@@ -13,7 +13,7 @@
// -----
func @ineq(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.INotEqual
@@ -25,7 +25,7 @@
// -----
func @islt(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SLessThan
@@ -37,7 +37,7 @@
// -----
func @isle(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SLessThanEqual
@@ -49,7 +49,7 @@
// -----
func @isgt(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SGreaterThan
@@ -61,7 +61,7 @@
// -----
func @isge(%arg0: memref<4xi32>, %arg1: memref<4xi32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi32>) : tensor<4xi32>
%1 = iree.load_input(%arg1 : memref<4xi32>) : tensor<4xi32>
// CHECK: spv.SGreaterThanEqual
@@ -73,7 +73,7 @@
// -----
func @oeq(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FOrdEqual
@@ -85,7 +85,7 @@
// -----
func @oge(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FOrdGreaterThanEqual
@@ -97,7 +97,7 @@
// -----
func @ogt(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FOrdGreaterThan
@@ -109,7 +109,7 @@
// -----
func @ole(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FOrdLessThanEqual
@@ -121,7 +121,7 @@
// -----
func @olt(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FOrdLessThan
@@ -133,7 +133,7 @@
// -----
func @ueq(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FUnordEqual
@@ -145,7 +145,7 @@
// -----
func @uge(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FUnordGreaterThanEqual
@@ -157,7 +157,7 @@
// -----
func @ugt(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FUnordGreaterThan
@@ -169,7 +169,7 @@
// -----
func @ule(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FUnordLessThanEqual
@@ -181,7 +181,7 @@
// -----
func @ult(%arg0: memref<4xf32>, %arg1: memref<4xf32>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xf32>) : tensor<4xf32>
%1 = iree.load_input(%arg1 : memref<4xf32>) : tensor<4xf32>
// CHECK: spv.FUnordLessThan
@@ -193,7 +193,7 @@
// -----
func @beq(%arg0: memref<4xi1>, %arg1: memref<4xi1>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi1>) : tensor<4xi1>
%1 = iree.load_input(%arg1 : memref<4xi1>) : tensor<4xi1>
// CHECK: spv.LogicalEqual
@@ -205,7 +205,7 @@
// -----
func @bneq(%arg0: memref<4xi1>, %arg1: memref<4xi1>, %arg2: memref<4xi1>)
-attributes {iree.executable.export, iree.executable.workload = dense<[4, 1, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<4xi1>) : tensor<4xi1>
%1 = iree.load_input(%arg1 : memref<4xi1>) : tensor<4xi1>
// CHECK: spv.LogicalNotEqual
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/concatenate.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/concatenate.mlir
index 79510a0..6870aa3 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/concatenate.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/concatenate.mlir
@@ -6,17 +6,17 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<64 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<10 x f32 [4]> [0]>, StorageBuffer>
func @concatenate(%arg0: memref<1x64xf32>, %arg1 : memref<1x10xf32>, %arg2 : memref<1x74xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[1, 74]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]]
// CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]]
- // CHECK: [[GLOBALIDY:%.*]] = spv.CompositeExtract [[GLOBALID]][1 : i32]
+ // CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32]
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[INPUTVAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
// CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG1]]
// CHECK: [[INPUTVAL1:%.*]] = spv.Load "StorageBuffer" [[ARG1LOADPTR]] : f32
// CHECK: [[TRUE:%.*]] = spv.constant true
// CHECK: [[SIXTY_FOUR:%.*]] = spv.constant 64 : i32
- // CHECK: [[CHECK:%.*]] = spv.SGreaterThanEqual [[GLOBALIDY]], [[SIXTY_FOUR]] : i32
+ // CHECK: [[CHECK:%.*]] = spv.SGreaterThanEqual [[GLOBALIDX]], [[SIXTY_FOUR]] : i32
// CHECK: [[COND:%.*]] = spv.LogicalAnd [[TRUE]], [[CHECK]] : i1
// CHECK: [[RESULT:%.*]] = spv.Select [[COND]], [[INPUTVAL1]], [[INPUTVAL0]] : i1, f32
%0 = iree.load_input(%arg0 : memref<1x64xf32>) : tensor<1x64xf32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/constant.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/constant.mlir
index a94c6ce..9328f00 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/constant.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/constant.mlir
@@ -2,7 +2,7 @@
module {
func @const_f32(%arg0: memref<2x3xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[3, 2]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[CONST:%.*]] = spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32 [4]>
// CHECK: [[VAR:%.*]] = spv.Variable init([[CONST]])
// CHECK: [[NUMPTR:%.*]] = spv.AccessChain [[VAR]]
@@ -19,7 +19,7 @@
module {
func @splat_const_f32(%arg0: memref<2x3xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[3, 2]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: spv.constant 1.000000e+00 : f32
%0 = iree.load_input(%arg0 : memref<2x3xf32>) : tensor<2x3xf32>
%1 = "xla_hlo.constant"() {value = dense<1.0> : tensor<2x3xf32>} : () -> (tensor<2x3xf32>)
@@ -33,7 +33,7 @@
module {
func @const_i32(%arg0: memref<2x3xi32>, %arg1: memref<2x3xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[3, 2]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[CONST:%.*]] = spv.constant dense<[1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32 [4]>
// CHECK: [[VAR:%.*]] = spv.Variable init([[CONST]])
// CHECK: [[NUMPTR:%.*]] = spv.AccessChain [[VAR]]
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/convert.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/convert.mlir
index b7c12bb..9969849 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/convert.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/convert.mlir
@@ -2,7 +2,7 @@
module {
func @convert_f2f_nop(%arg0: memref<12xf32>, %arg1 : memref<12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xf32>) : tensor<12xf32>
// CHECK-NOT: spv.FConvert
%1 = "xla_hlo.convert"(%0) : (tensor<12xf32>) -> tensor<12xf32>
@@ -15,7 +15,7 @@
module {
func @convert_f2f(%arg0: memref<12xf32>, %arg1 : memref<12xf16>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xf32>) : tensor<12xf32>
// CHECK: spv.FConvert {{%.*}} f32 to f16
%1 = "xla_hlo.convert"(%0) : (tensor<12xf32>) -> tensor<12xf16>
@@ -28,7 +28,7 @@
module {
func @convert_i2i_nop(%arg0: memref<12xi32>, %arg1 : memref<12xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xi32>) : tensor<12xi32>
// CHECK-NOT: spv.SConvert
%1 = "xla_hlo.convert"(%0) : (tensor<12xi32>) -> tensor<12xi32>
@@ -41,7 +41,7 @@
module {
func @convert_i2i(%arg0: memref<12xi32>, %arg1 : memref<12xi16>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xi32>) : tensor<12xi32>
// CHECK: spv.SConvert {{%.*}} i32 to i16
%1 = "xla_hlo.convert"(%0) : (tensor<12xi32>) -> tensor<12xi16>
@@ -54,7 +54,7 @@
module {
func @convert_i2f(%arg0: memref<12xi32>, %arg1 : memref<12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xi32>) : tensor<12xi32>
// CHECK: spv.ConvertSToF
%1 = "xla_hlo.convert"(%0) : (tensor<12xi32>) -> tensor<12xf32>
@@ -67,7 +67,7 @@
module {
func @convert_f2i_nop(%arg0: memref<12xf32>, %arg1 : memref<12xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12xf32>) : tensor<12xf32>
// CHECK: spv.ConvertFToS
%1 = "xla_hlo.convert"(%0) : (tensor<12xf32>) -> tensor<12xi32>
@@ -80,7 +80,7 @@
module {
func @convert_b2i(%arg0: memref<12xi1>, %arg1 : memref<12xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12]> : tensor<1xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" %{{.*}} : i1
// CHECK: [[ZERO:%.*]] = spv.constant 0 : i32
// CHECK: [[ONE:%.*]] = spv.constant 1 : i32
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
index 5012f3e..425ca20 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/copy.mlir
@@ -8,7 +8,7 @@
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: {{spirv|spv}}.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}
func @simple_load_store(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[GLOBALIDPTR:%.*]] = spv._address_of [[GLOBALIDVAR]]
// CHECK: [[GLOBALID:%.*]] = spv.Load "Input" [[GLOBALIDPTR]]
// CHECK: [[GLOBALIDX:%.*]] = spv.CompositeExtract [[GLOBALID]][0 : i32]
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/exp_test.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/exp_test.mlir
index 8c32f38..ea51f14 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/exp_test.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/exp_test.mlir
@@ -2,7 +2,7 @@
module {
func @scalar_rgn_dispatch_0(%arg0: memref<f32>)
- attributes {iree.executable.export, iree.executable.workload = dense<1> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%cst = constant dense<1.000000e+00> : tensor<f32>
//CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32
%0 = "xla_hlo.exp"(%cst) : (tensor<f32>) -> tensor<f32>
@@ -15,7 +15,7 @@
module {
func @exp(%arg0: memref<12x42xf32>, %arg2 : memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: {{%.*}} = spv.GLSL.Exp {{%.*}} : f32
%2 = "xla_hlo.exp"(%0) : (tensor<12x42xf32>) -> tensor<12x42xf32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/extract_element.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/extract_element.mlir
index 23f0ec9..8ca5a78 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/extract_element.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/extract_element.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<i1 [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<i1 [0]>, StorageBuffer>
func @extract_element(%arg0: memref<i1>, %arg1: memref<i1>)
- attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<1> : tensor<3xi32>, iree.num_dims = 3 : i32, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.num_dims = 3 : i32, iree.ordinal = 0 : i32} {
%0 = "iree.load_input"(%arg0) : (memref<i1>) -> tensor<i1>
// CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
// CHECK: {{%.*}} = spv.AccessChain [[ARG0]]{{\[}}[[ZERO1]]{{\]}}
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/gather.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/gather.mlir
index 528c852..90ccb00 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/gather.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/gather.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<50 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<i64 [0]>, StorageBuffer>
func @foo(%arg0: memref<5x1x10xf32>, %arg1: memref<i64>, %arg2: memref<1x10xf32>)
- attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[10, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
// CHECK: [[ZERO1:%.*]] = spv.constant 0
// CHECK: [[LOAD_ADDRESS_ARG1:%.*]] = spv.AccessChain [[ARG1]]{{\[}}[[ZERO1]]{{\]}}
// CHECK: [[INDEXI64:%.*]] = spv.Load {{".*"}} [[LOAD_ADDRESS_ARG1]]
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/max.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/max.mlir
index cd9c8b2..f715e2d 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/max.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/max.mlir
@@ -2,7 +2,7 @@
module {
func @maxf(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2 : memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.GLSL.FMax [[VAL1:%.*]], [[VAL2:%.*]] : f32
@@ -16,7 +16,7 @@
module {
func @maxi(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2 : memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.GLSL.SMax [[VAL1:%.*]], [[VAL2:%.*]] : i32
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/pad.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/pad.mlir
index c8c9909..ab02c05 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/pad.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/pad.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<216 x f32 [4]> [0]>, StorageBuffer>
func @pad_zero_interior(%arg0 : memref<12x4xf32>, %arg1 : memref<18x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 18, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
%0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
@@ -30,7 +30,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer>
func @pad_no_op(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[4, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
@@ -54,7 +54,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<48 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<522 x f32 [4]> [0]>, StorageBuffer>
func @pad_zero_interior(%arg0 : memref<12x4xf32>, %arg1 : memref<29x18xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[18, 29, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[INPUTVAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
%0 = iree.load_input(%arg0 : memref<12x4xf32>) : tensor<12x4xf32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape.mlir
index 666c153..d94987d 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_2D_2D(%arg0: memref<24x21xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<24x21xi32>) : tensor<24x21xi32>
@@ -24,7 +24,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_3D_2D(%arg0: memref<4x6x21xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<4x6x21xi32>) : tensor<4x6x21xi32>
@@ -43,7 +43,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_2D_3D(%arg0: memref<24x21xi32>, %arg1: memref<12x6x7xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<24x21xi32>) : tensor<24x21xi32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape_dropdims.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape_dropdims.mlir
index 49f4935..a934a78 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape_dropdims.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reshape_dropdims.mlir
@@ -6,7 +6,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_4D_3D(%arg0: memref<12x42x1xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x42x1xi32>) : tensor<12x42x1xi32>
@@ -26,7 +26,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_4D_2D(%arg0: memref<12x42x1x1xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x42x1x1xi32>) : tensor<12x42x1x1xi32>
@@ -46,7 +46,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_2D_4D(%arg0: memref<12x42xi32>, %arg1: memref<12x42x1x1xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
@@ -66,7 +66,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_2D_4D(%arg0: memref<12x42xi32>, %arg1: memref<12x1x1x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
@@ -86,7 +86,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<504 x i32 [4]> [0]>, StorageBuffer>
func @reshape_2D_4D(%arg0: memref<12x1x1x42xi32>, %arg1: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
%0 = iree.load_input(%arg0 : memref<12x1x1x42xi32>) : tensor<12x1x1x42xi32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reverse.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reverse.mlir
index d5013c4..4721e58 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reverse.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/reverse.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<144 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<144 x f32 [4]> [0]>, StorageBuffer>
func @reverse_2d(%arg0: memref<12x12xf32>, %arg1 : memref<12x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 12]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
%0 = iree.load_input(%arg0 : memref<12x12xf32>) : tensor<12x12xf32>
@@ -24,7 +24,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<27 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<27 x f32 [4]> [0]>, StorageBuffer>
func @reverse_3d(%arg0: memref<3x3x3xf32>, %arg1 : memref<3x3x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[3, 3, 3]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
%0 = iree.load_input(%arg0 : memref<3x3x3xf32>) : tensor<3x3x3xf32>
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/select.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/select.mlir
index 8ac5038..cd612c9 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/select.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/select.mlir
@@ -2,7 +2,7 @@
module {
func @select_ford_ge(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdGreaterThanEqual {{%.*}}, {{%.*}}
@@ -18,7 +18,7 @@
module {
func @select_ford_eq(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdEqual {{%.*}}, {{%.*}}
@@ -33,7 +33,7 @@
module {
func @select_ford_gt(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdGreaterThan {{%.*}}, {{%.*}}
@@ -48,7 +48,7 @@
module {
func @select_ford_lt(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdLessThan {{%.*}}, {{%.*}}
@@ -63,7 +63,7 @@
module {
func @select_ford_le(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdLessThanEqual {{%.*}}, {{%.*}}
@@ -78,7 +78,7 @@
module {
func @select_ford_ne(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FOrdNotEqual {{%.*}}, {{%.*}}
@@ -93,7 +93,7 @@
module {
func @select_funord_eq(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordEqual {{%.*}}, {{%.*}}
@@ -108,7 +108,7 @@
module {
func @select_funord_ge(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordGreaterThanEqual {{%.*}}, {{%.*}}
@@ -123,7 +123,7 @@
module {
func @select_funord_gt(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordGreaterThan {{%.*}}, {{%.*}}
@@ -138,7 +138,7 @@
module {
func @select_funord_lt(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordLessThan {{%.*}}, {{%.*}}
@@ -153,7 +153,7 @@
module {
func @select_funord_le(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordLessThanEqual {{%.*}}, {{%.*}}
@@ -168,7 +168,7 @@
module {
func @select_funord_ne(%arg0: memref<12x42xf32>, %arg1: memref<12x42xf32>, %arg2: memref<12x42xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xf32>) : tensor<12x42xf32>
%1 = iree.load_input(%arg1 : memref<12x42xf32>) : tensor<12x42xf32>
//CHECK: [[COMPARE:%.*]] = spv.FUnordNotEqual {{%.*}}, {{%.*}}
@@ -183,7 +183,7 @@
module {
func @select_int_eq(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.IEqual {{%.*}}, {{%.*}}
@@ -198,7 +198,7 @@
module {
func @select_int_ne(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.INotEqual {{%.*}}, {{%.*}}
@@ -213,7 +213,7 @@
module {
func @select_int_lt(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.SLessThan {{%.*}}, {{%.*}}
@@ -228,7 +228,7 @@
module {
func @select_int_le(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.SLessThanEqual {{%.*}}, {{%.*}}
@@ -243,7 +243,7 @@
module {
func @select_int_ge(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.SGreaterThanEqual {{%.*}}, {{%.*}}
@@ -259,7 +259,7 @@
module {
func @select_int_gt(%arg0: memref<12x42xi32>, %arg1: memref<12x42xi32>, %arg2: memref<12x42xi32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[42, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<12x42xi32>) : tensor<12x42xi32>
%1 = iree.load_input(%arg1 : memref<12x42xi32>) : tensor<12x42xi32>
//CHECK: [[COMPARE:%.*]] = spv.SGreaterThan {{%.*}}, {{%.*}}
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/slice.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/slice.mlir
index a556c92..07b3b90 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/slice.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/slice.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<36 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<6 x f32 [4]> [0]>, StorageBuffer>
func @slice_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
// CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]
@@ -24,7 +24,7 @@
// CHECK-SAME: [[ARG0:%.*]]: !spv.ptr<!spv.struct<!spv.array<36 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%.*]]: !spv.ptr<!spv.struct<!spv.array<6 x f32 [4]> [0]>, StorageBuffer>
func @slice_non_unit_stride(%arg0: memref<6x6xf32>, %arg1: memref<2x3xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[6, 1]> : tensor<2xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL0:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]] : f32
// CHECK: [[ARG1STOREPTR:%.*]] = spv.AccessChain [[ARG1]]
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/store_reduce.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/store_reduce.mlir
index 5311af3..9654bb1 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/store_reduce.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/store_reduce.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<5 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
- func @reduction_entry(%arg0: memref<5xi32>, %arg1: memref<i32>, %arg2: memref<i32> {iree.executable.reduction.output}) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[5, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ func @reduction_entry(%arg0: memref<5xi32>, %arg1: memref<i32>, %arg2: memref<i32> {iree.executable.reduction.output}) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5xi32>) : tensor<5xi32>
// CHECK: [[LOADPTR:%[a-zA-Z0-9_]*]] = spv.AccessChain [[ARG0]]
// CHECK: [[LOADVAL:%[a-zA-Z0-9_]*]] = spv.Load "StorageBuffer" [[LOADPTR]]
@@ -27,7 +27,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<20 x i32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<i32 [0]>, StorageBuffer>
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<4 x i32 [4]> [0]>, StorageBuffer>
- func @reduction_2D_dim0_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<4xi32> {iree.executable.reduction.output}) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[4, 5, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ func @reduction_2D_dim0_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<4xi32> {iree.executable.reduction.output}) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 0 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
%0 = iree.load_input(%arg0 : memref<5x4xi32>) : tensor<5x4xi32>
// CHECK: [[LOADPTR:%[a-zA-Z0-9_]*]] = spv.AccessChain [[ARG0]]
// CHECK: [[LOADVAL:%[a-zA-Z0-9_]*]] = spv.Load "StorageBuffer" [[LOADPTR]]
@@ -47,7 +47,7 @@
module {
// CHECK: spv.func @reduction_2D_dim1_entry
// CHECK-SAME: [[ARG2:%[a-zA-Z0-9_]*]]: !spv.ptr<!spv.struct<!spv.array<5 x i32 [4]> [0]>, StorageBuffer>
- func @reduction_2D_dim1_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<5xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 1 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.executable.workload = dense<[4, 5, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ func @reduction_2D_dim1_entry(%arg0: memref<5x4xi32>, %arg1: memref<i32>, %arg2: memref<5xi32>) attributes {iree.executable.export, iree.executable.reduction, iree.executable.reduction.apply = @reduction_apply, iree.executable.reduction.dimension = 1 : i32, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi64>, iree.ordinal = 0 : i32} {
// CHECK: [[GLOBALIDPTR:%[a-zA-Z0-9_]*]] = spv._address_of @globalInvocationID
// CHECK: [[GLOBALID:%[a-zA-Z0-9_]*]] = spv.Load "Input" [[GLOBALIDPTR]] : vector<3xi32>
// CHECK: [[GLOBALIDY:%[a-zA-Z0-9_]*]] = spv.CompositeExtract [[GLOBALID]]{{\[}}1 : i32{{\]}}
diff --git a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/transpose_add.mlir b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/transpose_add.mlir
index a6cf7af..331490e 100644
--- a/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/transpose_add.mlir
+++ b/iree/compiler/Translation/SPIRV/XLAToSPIRV/test/transpose_add.mlir
@@ -5,7 +5,7 @@
// CHECK-SAME: [[ARG0:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<144 x f32 [4]> [0]>, StorageBuffer>
// CHECK-SAME: [[ARG1:%[a-zA-Z0-9]*]]: !spv.ptr<!spv.struct<!spv.array<144 x f32 [4]> [0]>, StorageBuffer>
func @transpose_add(%arg0: memref<12x12xf32>, %arg1: memref<12x12xf32>)
- attributes {iree.executable.export, iree.executable.workload = dense<[12, 12, 1]> : tensor<3xi32>, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
+ attributes {iree.executable.export, iree.executable.workgroup_size = dense<[32, 1, 1]> : tensor<3xi32>, iree.ordinal = 0 : i32} {
// CHECK: [[ARG0LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
// CHECK: [[VAL1:%.*]] = spv.Load "StorageBuffer" [[ARG0LOADPTR]]
// CHECK: [[ARG1LOADPTR:%.*]] = spv.AccessChain [[ARG0]]
diff --git a/iree/compiler/Utils/BUILD b/iree/compiler/Utils/BUILD
index a76ef85..dca2b51 100644
--- a/iree/compiler/Utils/BUILD
+++ b/iree/compiler/Utils/BUILD
@@ -23,12 +23,10 @@
name = "Utils",
srcs = [
"GraphUtils.cpp",
- "IREECodegenUtils.cpp",
"TypeConversionUtils.cpp",
],
hdrs = [
"GraphUtils.h",
- "IREECodegenUtils.h",
"TypeConversionUtils.h",
],
deps = [
diff --git a/iree/compiler/Utils/CMakeLists.txt b/iree/compiler/Utils/CMakeLists.txt
index 3094d47..aa87771 100644
--- a/iree/compiler/Utils/CMakeLists.txt
+++ b/iree/compiler/Utils/CMakeLists.txt
@@ -17,11 +17,9 @@
Utils
HDRS
"GraphUtils.h"
- "IREECodegenUtils.h"
"TypeConversionUtils.h"
SRCS
"GraphUtils.cpp"
- "IREECodegenUtils.cpp"
"TypeConversionUtils.cpp"
DEPS
LLVMSupport
diff --git a/iree/compiler/Utils/IREECodegenUtils.cpp b/iree/compiler/Utils/IREECodegenUtils.cpp
deleted file mode 100644
index 1b65cb9..0000000
--- a/iree/compiler/Utils/IREECodegenUtils.cpp
+++ /dev/null
@@ -1,78 +0,0 @@
-// Copyright 2019 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "iree/compiler/Utils/IREECodegenUtils.h"
-
-namespace mlir {
-namespace iree_compiler {
-
-/// Gets the launch size associated with the dispatch function.
-LogicalResult getLegacyLaunchSize(Operation *funcOp,
- SmallVectorImpl<int64_t> &launchSize) {
- if (!funcOp->getAttr("iree.executable.export")) {
- return funcOp->emitError(
- "expected operation to be in dispatch function to get launch size");
- }
- auto workloadAttr =
- funcOp->getAttrOfType<DenseElementsAttr>("iree.executable.workload");
- if (!workloadAttr) {
- return funcOp->emitError(
- "unable to find workload size, missing attribute "
- "iree.executable.workload in dispatch function");
- }
- launchSize.clear();
- for (auto value : workloadAttr.getValues<APInt>()) {
- launchSize.push_back(value.getSExtValue());
- }
- // Drop trailing ones.
- auto dropFrom = launchSize.size() - 1;
- while (dropFrom > 0 && launchSize[dropFrom] == 1) {
- --dropFrom;
- }
- if (dropFrom > 0) {
- launchSize.erase(std::next(launchSize.begin(), dropFrom + 1),
- launchSize.end());
- }
- return success();
-}
-
-/// Gets the workgroup size.
-template <typename intType>
-LogicalResult getLegacyWorkGroupSize(Operation *funcOp,
- SmallVectorImpl<intType> &workGroupSize) {
- if (!funcOp->getAttr("iree.executable.export")) {
- return funcOp->emitError(
- "expected operation to be in dispatch function to get launch size");
- }
- auto workGroupSizeAttr = funcOp->getAttrOfType<DenseElementsAttr>(
- "iree.executable.workgroup_size");
- if (!workGroupSizeAttr) {
- return funcOp->emitError(
- "unable to find workload size, missing attribute "
- "iree.executable.workload in dispatch function");
- }
- workGroupSize.clear();
- for (auto value : workGroupSizeAttr.getValues<APInt>()) {
- workGroupSize.push_back(value.getSExtValue());
- }
- return success();
-}
-
-template LogicalResult getLegacyWorkGroupSize<int32_t>(
- Operation *funcOp, SmallVectorImpl<int32_t> &workGroupSize);
-template LogicalResult getLegacyWorkGroupSize<int64_t>(
- Operation *funcOp, SmallVectorImpl<int64_t> &workGroupSize);
-
-} // namespace iree_compiler
-} // namespace mlir
diff --git a/iree/compiler/Utils/IREECodegenUtils.h b/iree/compiler/Utils/IREECodegenUtils.h
deleted file mode 100644
index 3f82e60..0000000
--- a/iree/compiler/Utils/IREECodegenUtils.h
+++ /dev/null
@@ -1,40 +0,0 @@
-// Copyright 2019 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef IREE_COMPILER_UTILS_IREECODEGENUTILS_H
-#define IREE_COMPILER_UTILS_IREECODEGENUTILS_H
-
-#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
-#include "mlir/IR/Function.h"
-#include "mlir/Support/LogicalResult.h"
-
-namespace mlir {
-namespace iree_compiler {
-
-// WARNING: this file is deprecated and will be removed soon. Do not use.
-
-// TODO(ravishankarm): remove this; it does not work with dynamic shapes.
-/// Gets the launch size associated with the dispatch function.
-LogicalResult getLegacyLaunchSize(Operation *funcOp,
- SmallVectorImpl<int64_t> &launchSize);
-
-/// Gets the workgroup size. Has to be a static constant.
-template <typename intType>
-LogicalResult getLegacyWorkGroupSize(Operation *funcOp,
- SmallVectorImpl<intType> &workGroupSize);
-
-} // namespace iree_compiler
-} // namespace mlir
-
-#endif // IREE_COMPILER_UTILS_IREECODEGENUTILS_H
