Move `LoweringConfigAttr` and `TranslationInfoAttr` to `AttrDef`. (#7385)
Previously these were defined using StructAttr which is limited in
its capabilities. Moving them to AttrDef allows this to be more
robust and verifiable. Also allows specification of custom builder
methods that remove a lot of the bioler plate.
In addition to the above, this change also adds a new attribute,
CompilationInfoAttr. This is the attribute that when added to
operations like linalg.matmul/linalg.*conv* (or their mhlo/tosa
couterparts) use the information from the specified attribute instead
of the default heuristics. This allows an external search to search
for better values than the defaults used. While there are some
preliminary verifies in place, additional verifiers are needed to
ensure that the user-provided attribute has all the information needed
(and in a consistent manner) to achieve correct compilation.
In addition, all these attributes are moved from the HAL dialect into
a new IREECodegenDialect added to compiler/Codegen/Dialect.
This change also brings in some code from D111594 to allow for this to
land without having to wait for that patch to land.
diff --git a/iree/compiler/Codegen/Common/BUILD b/iree/compiler/Codegen/Common/BUILD
index 6f71b73..0593b9b 100644
--- a/iree/compiler/Codegen/Common/BUILD
+++ b/iree/compiler/Codegen/Common/BUILD
@@ -46,6 +46,7 @@
deps = [
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common:FoldTensorExtractOpIncGen",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/Transforms",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Codegen/Common/CMakeLists.txt b/iree/compiler/Codegen/Common/CMakeLists.txt
index 5e77f08..ce29221 100644
--- a/iree/compiler/Codegen/Common/CMakeLists.txt
+++ b/iree/compiler/Codegen/Common/CMakeLists.txt
@@ -56,6 +56,7 @@
MLIRTransforms
MLIRVector
iree::compiler::Codegen::Common::FoldTensorExtractOpIncGen
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::Transforms
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
index 0f70605..06f3f51 100644
--- a/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
+++ b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
@@ -4,14 +4,13 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
@@ -90,14 +89,9 @@
if (!workloadPerWorkgroup.empty()) {
currWorkloadPerWorkgroup.assign(workloadPerWorkgroup.begin(),
workloadPerWorkgroup.end());
- } else if (IREE::HAL::TranslationInfo translationInfo =
+ } else if (IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(entryPointOp)) {
- if (ArrayAttr workloadPerWorkgroupAttr =
- translationInfo.workloadPerWorkgroup()) {
- currWorkloadPerWorkgroup = llvm::to_vector<4>(llvm::map_range(
- workloadPerWorkgroupAttr,
- [](Attribute attr) { return attr.cast<IntegerAttr>().getInt(); }));
- }
+ currWorkloadPerWorkgroup = translationInfo.getWorkloadPerWorkgroupVals();
}
if (!currWorkloadPerWorkgroup.empty()) {
diff --git a/iree/compiler/Codegen/Dialect/BUILD b/iree/compiler/Codegen/Dialect/BUILD
new file mode 100644
index 0000000..9c6008b
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/BUILD
@@ -0,0 +1,107 @@
+# Copyright 2019 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+load("@llvm-project//mlir:tblgen.bzl", "gentbl_cc_library", "td_library")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+exports_files([
+ "IREECodegenAttributes.td",
+ "IREECodegenDialect.td",
+ "LoweringConfig.td",
+])
+
+td_library(
+ name = "td_files",
+ srcs = enforce_glob(
+ [
+ "IREECodegenAttributes.td",
+ "IREECodegenDialect.td",
+ "LoweringConfig.td",
+ ],
+ include = ["*.td"],
+ ),
+ deps = [
+ "@llvm-project//mlir:OpBaseTdFiles",
+ ],
+)
+
+cc_library(
+ name = "IREECodegenDialect",
+ srcs = [
+ "IREECodegenDialect.cpp",
+ "LoweringConfig.cpp",
+ ],
+ hdrs = [
+ "IREECodegenDialect.h",
+ "LoweringConfig.h",
+ ],
+ textual_hdrs = [
+ "IREECodegenDialect.cpp.inc",
+ "IREECodegenDialect.h.inc",
+ "LoweringConfig.cpp.inc",
+ "LoweringConfig.h.inc",
+ "LoweringConfigEnums.cpp.inc",
+ "LoweringConfigEnums.h.inc",
+ ],
+ deps = [
+ ":IREECodegenDialectGen",
+ ":LoweringConfigGen",
+ "//iree/compiler/Codegen/Utils",
+ "@llvm-project//llvm:Support",
+ "@llvm-project//mlir:DialectUtils",
+ "@llvm-project//mlir:IR",
+ "@llvm-project//mlir:Parser",
+ "@llvm-project//mlir:StandardOps",
+ ],
+)
+
+gentbl_cc_library(
+ name = "IREECodegenDialectGen",
+ tbl_outs = [
+ (
+ ["-gen-dialect-decls"],
+ "IREECodegenDialect.h.inc",
+ ),
+ (
+ ["-gen-dialect-defs"],
+ "IREECodegenDialect.cpp.inc",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "IREECodegenAttributes.td",
+ deps = [":td_files"],
+)
+
+gentbl_cc_library(
+ name = "LoweringConfigGen",
+ tbl_outs = [
+ (
+ ["-gen-attrdef-decls"],
+ "LoweringConfig.h.inc",
+ ),
+ (
+ ["-gen-attrdef-defs"],
+ "LoweringConfig.cpp.inc",
+ ),
+ (
+ ["-gen-enum-decls"],
+ "LoweringConfigEnums.h.inc",
+ ),
+ (
+ ["-gen-enum-defs"],
+ "LoweringConfigEnums.cpp.inc",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "LoweringConfig.td",
+ deps = [":td_files"],
+)
diff --git a/iree/compiler/Codegen/Dialect/CMakeLists.txt b/iree/compiler/Codegen/Dialect/CMakeLists.txt
new file mode 100644
index 0000000..16f6826
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/CMakeLists.txt
@@ -0,0 +1,62 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# iree/compiler/Codegen/Dialect/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_cc_library(
+ NAME
+ IREECodegenDialect
+ HDRS
+ "IREECodegenDialect.h"
+ "LoweringConfig.h"
+ TEXTUAL_HDRS
+ "IREECodegenDialect.cpp.inc"
+ "IREECodegenDialect.h.inc"
+ "LoweringConfig.cpp.inc"
+ "LoweringConfig.h.inc"
+ "LoweringConfigEnums.cpp.inc"
+ "LoweringConfigEnums.h.inc"
+ SRCS
+ "IREECodegenDialect.cpp"
+ "LoweringConfig.cpp"
+ DEPS
+ ::IREECodegenDialectGen
+ ::LoweringConfigGen
+ LLVMSupport
+ MLIRIR
+ MLIRParser
+ MLIRStandard
+ iree::compiler::Codegen::Utils
+ PUBLIC
+)
+
+iree_tablegen_library(
+ NAME
+ IREECodegenDialectGen
+ TD_FILE
+ "IREECodegenAttributes.td"
+ OUTS
+ -gen-dialect-decls IREECodegenDialect.h.inc
+ -gen-dialect-defs IREECodegenDialect.cpp.inc
+)
+
+iree_tablegen_library(
+ NAME
+ LoweringConfigGen
+ TD_FILE
+ "LoweringConfig.td"
+ OUTS
+ -gen-attrdef-decls LoweringConfig.h.inc
+ -gen-attrdef-defs LoweringConfig.cpp.inc
+ -gen-enum-decls LoweringConfigEnums.h.inc
+ -gen-enum-defs LoweringConfigEnums.cpp.inc
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/iree/compiler/Codegen/Dialect/IREECodegenAttributes.td b/iree/compiler/Codegen/Dialect/IREECodegenAttributes.td
new file mode 100644
index 0000000..e5f8b26
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/IREECodegenAttributes.td
@@ -0,0 +1,14 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_CODEGEN_DIALECT_IREECODEGEN_ATTRIBUTES
+#define IREE_CODEGEN_DIALECT_IREECODEGEN_ATTRIBUTES
+
+include "iree/compiler/Codegen/Dialect/IREECodegenDialect.td"
+include "iree/compiler/Codegen/Dialect/LoweringConfig.td"
+
+
+#endif // IREE_CODEGEN_DIALECT_IREECODEGEN_ATTRIBUTES
diff --git a/iree/compiler/Codegen/Dialect/IREECodegenDialect.cpp b/iree/compiler/Codegen/Dialect/IREECodegenDialect.cpp
new file mode 100644
index 0000000..8f2cf47
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/IREECodegenDialect.cpp
@@ -0,0 +1,62 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
+
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.cpp.inc"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
+#include "mlir/IR/DialectImplementation.h"
+
+namespace mlir {
+namespace iree_compiler {
+namespace IREE {
+namespace Codegen {
+
+struct IREECodegenDialectOpAsmInterface : public OpAsmDialectInterface {
+ using OpAsmDialectInterface::OpAsmDialectInterface;
+ AliasResult getAlias(Attribute attr, raw_ostream &os) const override {
+ if (attr.isa<TranslationInfoAttr>()) {
+ os << "translation";
+ return AliasResult::OverridableAlias;
+ } else if (attr.isa<CompilationInfoAttr>()) {
+ os << "compilation";
+ return AliasResult::OverridableAlias;
+ } else if (attr.isa<LoweringConfigAttr>()) {
+ os << "config";
+ return AliasResult::OverridableAlias;
+ }
+ return AliasResult::NoAlias;
+ }
+};
+
+void IREECodegenDialect::initialize() {
+ initializeCodegenAttrs();
+ addInterfaces<IREECodegenDialectOpAsmInterface>();
+}
+
+Attribute IREECodegenDialect::parseAttribute(DialectAsmParser &parser,
+ Type type) const {
+ StringRef mnemonic;
+ if (failed(parser.parseKeyword(&mnemonic))) return {};
+ Attribute genAttr;
+ OptionalParseResult parseResult =
+ parseCodegenAttrs(parser, mnemonic, type, genAttr);
+ if (parseResult.hasValue()) return genAttr;
+ parser.emitError(parser.getNameLoc(), "unknown iree_codegen attribute");
+ return Attribute();
+}
+
+void IREECodegenDialect::printAttribute(Attribute attr,
+ DialectAsmPrinter &p) const {
+ if (failed(printCodegenAttrs(attr, p))) {
+ llvm_unreachable("unhandled iree_codegen attribute");
+ }
+}
+
+} // namespace Codegen
+} // namespace IREE
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Codegen/Dialect/IREECodegenDialect.h b/iree/compiler/Codegen/Dialect/IREECodegenDialect.h
new file mode 100644
index 0000000..bdb9473
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/IREECodegenDialect.h
@@ -0,0 +1,17 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_CODEGEN_DIALECT_IREECODEGEN_DIALECT_H_
+#define IREE_COMPILER_CODEGEN_DIALECT_IREECODEGEN_DIALECT_H_
+
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/OpDefinition.h"
+
+// clang-format off: must be included after all LLVM/MLIR eaders
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h.inc" // IWYU pragma: keep
+// clang-format on
+
+#endif // IREE_COMPILER_CODEGEN_DIALECT_IREECODEGEN_DIALECT_H_
diff --git a/iree/compiler/Codegen/Dialect/IREECodegenDialect.td b/iree/compiler/Codegen/Dialect/IREECodegenDialect.td
new file mode 100644
index 0000000..b1233e4
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/IREECodegenDialect.td
@@ -0,0 +1,44 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_CODEGEN_DIALECT_IREECODEGEN_DIALECT
+#define IREE_CODEGEN_DIALECT_IREECODEGEN_DIALECT
+
+include "mlir/IR/OpBase.td"
+
+//===----------------------------------------------------------------------===//
+// IREE Codegen dialect
+//===----------------------------------------------------------------------===//
+
+def IREECodegen_Dialect : Dialect {
+ let name = "iree_codegen";
+ let cppNamespace = "::mlir::iree_compiler::IREE::Codegen";
+
+ let summary = [{
+ A dialect representing attributes used by the IREE Code generation.
+ }];
+ let description = [{
+ This dialect is primarily meant to hold attributes that carry the
+ state of the compilation when lowered to scalar code for an
+ architecture. Typically, a backend starts by analysing the entry
+ point functions within the `hal.executable.variant` and deciding
+ which compilation pipeline to chose. During this, even the values
+ for parameters such as tile sizes, etc. are also decided. The rest
+ of the compilation flow does not make any heuristic decisions,
+ rather just looks at the values of the decision specified in
+ attributes that belong to this dialect. This allows an external
+ search to easily override the heuristics that are hard-coded
+ within a backend.
+ }];
+ let extraClassDeclaration = [{
+ void initializeCodegenAttrs();
+ OptionalParseResult parseCodegenAttrs(DialectAsmParser &parser,
+ StringRef mnemonic, Type type, Attribute &value) const;
+ LogicalResult printCodegenAttrs(Attribute attr, DialectAsmPrinter &p) const;
+ }];
+}
+
+#endif // IREE_CODEGEN_DIALECT_IREECODEGEN_DIALECT
\ No newline at end of file
diff --git a/iree/compiler/Codegen/Dialect/LoweringConfig.cpp b/iree/compiler/Codegen/Dialect/LoweringConfig.cpp
new file mode 100644
index 0000000..79990d6
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/LoweringConfig.cpp
@@ -0,0 +1,632 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
+
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
+#include "llvm/ADT/TypeSwitch.h"
+#include "mlir/Dialect/StandardOps/IR/Ops.h"
+#include "mlir/IR/DialectImplementation.h"
+
+#define GET_ATTRDEF_CLASSES
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.cpp.inc"
+#include "iree/compiler/Codegen/Dialect/LoweringConfigEnums.cpp.inc"
+
+static const char kConfigAttrName[] = "lowering.config";
+static const char kTranslationInfoAttrName[] = "translation.info";
+static const char kCompilationInfoAttrName[] = "compilation.info";
+
+namespace mlir {
+namespace iree_compiler {
+
+//===----------------------------------------------------------------------===//
+// Utility function for common code patterns.
+//===----------------------------------------------------------------------===//
+
+static bool checkIntegerArrayAttr(ArrayAttr arrayAttr) {
+ return !llvm::any_of(arrayAttr,
+ [](Attribute attr) { return !attr.isa<IntegerAttr>(); });
+}
+
+/// Returns an `ArrayAttr` where each element is an `IntegerAttr` of `IndexType`
+/// whose values is obtained from `values`.
+static ArrayAttr getIndexIntegerArrayAttr(MLIRContext *context,
+ ArrayRef<int64_t> values) {
+ auto attrs = llvm::to_vector<4>(
+ llvm::map_range(values, [&context](int64_t value) -> Attribute {
+ return IntegerAttr::get(IndexType::get(context), APInt(64, value));
+ }));
+ return ArrayAttr::get(context, attrs);
+}
+
+/// Returns an `ArrayAttr` where each element is an `IntegerAttr` of 64-bit
+/// integer type whose values is obtained from `values`.
+static ArrayAttr getI64IntegerArrayAttr(MLIRContext *context,
+ ArrayRef<int64_t> values) {
+ auto attrs = llvm::to_vector<4>(
+ llvm::map_range(values, [&context](int64_t value) -> Attribute {
+ return IntegerAttr::get(IntegerType::get(context, 64),
+ APInt(64, value));
+ }));
+ return ArrayAttr::get(context, attrs);
+}
+
+/// Assumes that `arrayAttr` is a list of `IntegerAttr`s and returns the values
+/// in these attributes as a vector.
+static SmallVector<int64_t> getIntegerVals(ArrayAttr arrayAttr) {
+ if (!arrayAttr) return {};
+ SmallVector<int64_t> values(arrayAttr.size());
+ for (auto attr : llvm::enumerate(arrayAttr)) {
+ values[attr.index()] = attr.value().cast<IntegerAttr>().getInt();
+ }
+ return values;
+}
+
+namespace IREE {
+namespace Codegen {
+
+namespace {
+
+// TODO(ravishankarm): The IREEFieldParser is part of the patch D111594 (where
+// it is called ::mlir::FieldParser). Remove this when the upstream change lands
+// in IREE.
+
+//===----------------------------------------------------------------------===//
+// Parse Fields
+//===----------------------------------------------------------------------===//
+
+/// Provide a template class that can be specialized by users to dispatch to
+/// parsers. Auto-generated parsers generate calls to
+/// `IREEFieldParser<T>::parse`, where `T` is the parameter storage type, to
+/// parse custom types.
+template <typename T, typename = T>
+struct IREEFieldParser;
+
+/// Parse an attribute.
+template <typename AttributeT>
+struct IREEFieldParser<
+ AttributeT, std::enable_if_t<std::is_base_of<Attribute, AttributeT>::value,
+ AttributeT>> {
+ static FailureOr<AttributeT> parse(DialectAsmParser &parser) {
+ AttributeT value;
+ if (parser.parseAttribute(value)) return failure();
+ return value;
+ }
+};
+
+/// Parse any integer.
+template <typename IntT>
+struct IREEFieldParser<IntT,
+ std::enable_if_t<std::is_integral<IntT>::value, IntT>> {
+ static FailureOr<IntT> parse(DialectAsmParser &parser) {
+ IntT value;
+ if (parser.parseInteger(value)) return failure();
+ return value;
+ }
+};
+
+/// Parse a string.
+template <>
+struct IREEFieldParser<std::string> {
+ static FailureOr<std::string> parse(DialectAsmParser &parser) {
+ std::string value;
+ if (parser.parseString(&value)) return failure();
+ return value;
+ }
+};
+
+/// Parse any container that supports back insertion as a list.
+template <typename ContainerT>
+struct IREEFieldParser<
+ ContainerT, std::enable_if_t<std::is_member_function_pointer<decltype(
+ &ContainerT::push_back)>::value,
+ ContainerT>> {
+ using ElementT = typename ContainerT::value_type;
+ static FailureOr<ContainerT> parse(DialectAsmParser &parser) {
+ ContainerT elements;
+ auto elementParser = [&]() {
+ auto element = IREEFieldParser<ElementT>::parse(parser);
+ if (failed(element)) return failure();
+ elements.push_back(element.getValue());
+ return success();
+ };
+ if (parser.parseCommaSeparatedList(elementParser)) return failure();
+ return elements;
+ }
+};
+} // namespace
+
+//===----------------------------------------------------------------------===//
+// iree_codegen.translation.info
+//===----------------------------------------------------------------------===//
+
+TranslationInfoAttr TranslationInfoAttr::get(
+ MLIRContext *context, DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workloadPerWorkgroup) {
+ auto pipelineAttr = StringAttr::get(context, stringifyEnum(passPipeline));
+ ArrayAttr workloadPerWorkgroupAttr =
+ getI64IntegerArrayAttr(context, workloadPerWorkgroup);
+ return get(context, pipelineAttr, workloadPerWorkgroupAttr);
+}
+
+DispatchLoweringPassPipeline
+TranslationInfoAttr::getDispatchLoweringPassPipeline() {
+ Optional<DispatchLoweringPassPipeline> passPipeline =
+ symbolizeEnum<DispatchLoweringPassPipeline>(getPassPipeline().getValue());
+ return passPipeline.getValue();
+}
+
+SmallVector<int64_t> TranslationInfoAttr::getWorkloadPerWorkgroupVals() {
+ return getIntegerVals(getWorkloadPerWorkgroup());
+}
+
+LogicalResult TranslationInfoAttr::verify(
+ function_ref<InFlightDiagnostic()> emitError, StringAttr passPipeline,
+ ArrayAttr workloadPerWorkgroup) {
+ if (!passPipeline) {
+ return emitError() << "missing pass pipeline specification";
+ }
+ auto passPipelineValue =
+ symbolizeEnum<IREE::Codegen::DispatchLoweringPassPipeline>(
+ passPipeline.getValue());
+ if (!passPipelineValue) {
+ return emitError() << "invalid pass pipeline value : "
+ << passPipeline.getValue();
+ }
+ if (!workloadPerWorkgroup) {
+ return emitError() << "expected workload_per_wg to be specified (even if "
+ "specified as empty)";
+ }
+ if (!checkIntegerArrayAttr(workloadPerWorkgroup)) {
+ return emitError() << "expected workload_per_wg to be an IntegerAttr list";
+ }
+ return success();
+}
+
+::mlir::Attribute TranslationInfoAttr::parse(::mlir::DialectAsmParser &parser,
+ ::mlir::Type attrType) {
+ ::mlir::FailureOr<StringAttr> _result_passPipeline;
+ ::mlir::FailureOr<ArrayAttr> _result_workloadPerWorkgroup;
+ // Parse literal '<'
+ if (parser.parseLess()) return {};
+ // Parse variable 'passPipeline'
+ _result_passPipeline = IREEFieldParser<StringAttr>::parse(parser);
+ if (failed(_result_passPipeline)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "failed to parse IREECodegen_TranslationInfoAttr "
+ "parameter 'passPipeline' which is to be a `StringAttr`");
+ return {};
+ }
+ // Parse literal ','
+ if (parser.parseComma()) return {};
+ // Parse literal 'workload_per_wg'
+ if (parser.parseKeyword("workload_per_wg")) return {};
+ // Parse literal '='
+ if (parser.parseEqual()) return {};
+ // Parse variable 'workloadPerWorkgroup'
+ _result_workloadPerWorkgroup = IREEFieldParser<ArrayAttr>::parse(parser);
+ if (failed(_result_workloadPerWorkgroup)) {
+ parser.emitError(
+ parser.getCurrentLocation(),
+ "failed to parse IREECodegen_TranslationInfoAttr parameter "
+ "'workloadPerWorkgroup' which is to be a `ArrayAttr`");
+ return {};
+ }
+ // Parse literal '>'
+ if (parser.parseGreater()) return {};
+ return TranslationInfoAttr::get(parser.getContext(),
+ _result_passPipeline.getValue(),
+ _result_workloadPerWorkgroup.getValue());
+}
+
+void TranslationInfoAttr::print(::mlir::DialectAsmPrinter &printer) const {
+ printer << "translation.info";
+ printer << "<";
+ printer << getPassPipeline();
+ printer << ",";
+ printer << ' ' << "workload_per_wg";
+ printer << ' ' << "=";
+ printer << ' ';
+ printer << getWorkloadPerWorkgroup();
+ printer << ">";
+}
+
+//===----------------------------------------------------------------------===//
+// iree_codegen.lowering.config
+//===----------------------------------------------------------------------===//
+
+LoweringConfigAttr LoweringConfigAttr::get(MLIRContext *context,
+ TileSizesListTypeRef tileSizes,
+ ArrayRef<int64_t> nativeVectorSize) {
+ auto attrList = llvm::to_vector<4>(
+ llvm::map_range(tileSizes, [&](ArrayRef<int64_t> sizes) -> Attribute {
+ return getI64IntegerArrayAttr(context, sizes);
+ }));
+ ArrayAttr tileSizesAttr = ArrayAttr::get(context, attrList);
+ ArrayAttr nativeVectorSizeAttr =
+ getI64IntegerArrayAttr(context, nativeVectorSize);
+ return get(context, tileSizesAttr, nativeVectorSizeAttr);
+}
+
+TileSizesListType LoweringConfigAttr::getTileSizeVals() {
+ auto tileSizesAttr = getTileSizes();
+ if (!tileSizesAttr) return {};
+ TileSizesListType tileSizes;
+ for (auto attr : tileSizesAttr) {
+ auto vals = getIntegerVals(attr.cast<ArrayAttr>());
+ tileSizes.emplace_back(std::move(vals));
+ }
+ return tileSizes;
+}
+
+SmallVector<int64_t> LoweringConfigAttr::getTileSizeVals(unsigned level) {
+ ArrayAttr tileSizesAttr = getTileSizes();
+ if (!tileSizesAttr || tileSizesAttr.size() <= level) return {};
+ return getIntegerVals(tileSizesAttr[level].cast<ArrayAttr>());
+}
+
+SmallVector<int64_t> LoweringConfigAttr::getNativeVectorSizeVals() {
+ ArrayAttr nativeVectorSizeAttr = getNativeVectorSize();
+ if (!nativeVectorSizeAttr) return {};
+ return getIntegerVals(nativeVectorSizeAttr);
+}
+
+LogicalResult LoweringConfigAttr::verify(
+ function_ref<InFlightDiagnostic()> emitError, ArrayAttr tileSizes,
+ ArrayAttr nativeVectorSize) {
+ if (!tileSizes) {
+ return emitError() << "expected tile_sizes to be specified (even is "
+ "specified as empty)";
+ }
+ if (llvm::any_of(tileSizes, [](Attribute attr) {
+ auto arrayAttr = attr.dyn_cast<ArrayAttr>();
+ return !arrayAttr || !checkIntegerArrayAttr(arrayAttr);
+ })) {
+ return emitError()
+ << "expected all elements of tile_sizes to be a list of integers";
+ }
+ if (!nativeVectorSize) {
+ return emitError() << "expected native_vector_size to be specified (even "
+ "if specified as empty)";
+ }
+ if (!checkIntegerArrayAttr(nativeVectorSize)) {
+ return emitError()
+ << "expected native_vector_size to be a list of integer values";
+ }
+ return success();
+}
+
+::mlir::Attribute LoweringConfigAttr::parse(::mlir::DialectAsmParser &parser,
+ ::mlir::Type attrType) {
+ ::mlir::FailureOr<ArrayAttr> _result_tileSizes;
+ ::mlir::FailureOr<ArrayAttr> _result_nativeVectorSize;
+ // Parse literal '<'
+ if (parser.parseLess()) return {};
+ // Parse literal 'tile_sizes'
+ if (parser.parseKeyword("tile_sizes")) return {};
+ // Parse literal '='
+ if (parser.parseEqual()) return {};
+ // Parse variable 'tileSizes'
+ _result_tileSizes = IREEFieldParser<ArrayAttr>::parse(parser);
+ if (failed(_result_tileSizes)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "failed to parse IREECodegen_LoweringConfigAttr parameter "
+ "'tileSizes' which is to be a `ArrayAttr`");
+ return {};
+ }
+ // Parse literal ','
+ if (parser.parseComma()) return {};
+ // Parse literal 'native_vector_size'
+ if (parser.parseKeyword("native_vector_size")) return {};
+ // Parse literal '='
+ if (parser.parseEqual()) return {};
+ // Parse variable 'nativeVectorSize'
+ _result_nativeVectorSize = IREEFieldParser<ArrayAttr>::parse(parser);
+ if (failed(_result_nativeVectorSize)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "failed to parse IREECodegen_LoweringConfigAttr parameter "
+ "'nativeVectorSize' which is to be a `ArrayAttr`");
+ return {};
+ }
+ // Parse literal '>'
+ if (parser.parseGreater()) return {};
+ return LoweringConfigAttr::get(parser.getContext(),
+ _result_tileSizes.getValue(),
+ _result_nativeVectorSize.getValue());
+}
+
+void LoweringConfigAttr::print(::mlir::DialectAsmPrinter &printer) const {
+ printer << "lowering.config";
+ printer << "<";
+ printer << "tile_sizes";
+ printer << ' ' << "=";
+ printer << ' ';
+ printer << getTileSizes();
+ printer << ",";
+ printer << ' ' << "native_vector_size";
+ printer << ' ' << "=";
+ printer << ' ';
+ printer << getNativeVectorSize();
+ printer << ">";
+}
+
+//===----------------------------------------------------------------------===//
+// iree.compilation.info
+//===----------------------------------------------------------------------===//
+
+CompilationInfoAttr CompilationInfoAttr::get(MLIRContext *context,
+ TileSizesListTypeRef tileSizes,
+ ArrayRef<int64_t> nativeVectorSize,
+ ArrayRef<int64_t> workgroupSize) {
+ LoweringConfigAttr configAttr =
+ LoweringConfigAttr::get(context, tileSizes, nativeVectorSize);
+ TranslationInfoAttr translationInfo =
+ TranslationInfoAttr::get(context, DispatchLoweringPassPipeline::None);
+ ArrayAttr workgroupSizeAttr = getI64IntegerArrayAttr(context, workgroupSize);
+ return get(context, configAttr, translationInfo, workgroupSizeAttr);
+}
+
+CompilationInfoAttr CompilationInfoAttr::get(
+ MLIRContext *context, TileSizesListTypeRef tileSizes,
+ ArrayRef<int64_t> nativeVectorSize,
+ DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workloadPerWorkgroup, ArrayRef<int64_t> workgroupSize) {
+ LoweringConfigAttr configAttr =
+ LoweringConfigAttr::get(context, tileSizes, nativeVectorSize);
+ TranslationInfoAttr translationInfoAttr =
+ TranslationInfoAttr::get(context, passPipeline, workloadPerWorkgroup);
+ ArrayAttr workgroupSizeAttr = getI64IntegerArrayAttr(context, workgroupSize);
+ return get(context, configAttr, translationInfoAttr, workgroupSizeAttr);
+}
+
+LogicalResult CompilationInfoAttr::verify(
+ function_ref<InFlightDiagnostic()> emitError,
+ LoweringConfigAttr loweringConfig, TranslationInfoAttr translationInfo,
+ ArrayAttr workgroupSize) {
+ if (!loweringConfig) {
+ return emitError() << "missing lowering config";
+ }
+ if (failed(
+ LoweringConfigAttr::verify(emitError, loweringConfig.getTileSizes(),
+ loweringConfig.getNativeVectorSize()))) {
+ return failure();
+ }
+ if (!translationInfo) {
+ return emitError() << "missing translation info";
+ }
+ if (failed(TranslationInfoAttr::verify(
+ emitError, translationInfo.getPassPipeline(),
+ translationInfo.getWorkloadPerWorkgroup()))) {
+ return failure();
+ }
+ if (!workgroupSize) {
+ return emitError() << "expected workgroup_size to be specified (even if "
+ "specified empty)";
+ }
+ if (!checkIntegerArrayAttr(workgroupSize)) {
+ return emitError() << "expected workgroup_size to be a list of integers";
+ }
+ return success();
+}
+
+/// Parser method that is copied from the auto-generated using `assemblyFormat`
+/// available with patch D111594. Replace after that change is in IREE.
+::mlir::Attribute CompilationInfoAttr::parse(::mlir::DialectAsmParser &parser,
+ ::mlir::Type attrType) {
+ ::mlir::FailureOr<LoweringConfigAttr> _result_loweringConfig;
+ ::mlir::FailureOr<TranslationInfoAttr> _result_translationInfo;
+ ::mlir::FailureOr<ArrayAttr> _result_workgroupSize;
+ // Parse literal '<'
+ if (parser.parseLess()) return {};
+ // Parse variable 'loweringConfig'
+ _result_loweringConfig = IREEFieldParser<LoweringConfigAttr>::parse(parser);
+ if (failed(_result_loweringConfig)) {
+ parser.emitError(
+ parser.getCurrentLocation(),
+ "failed to parse IREECodegen_CompilationInfoAttr parameter "
+ "'loweringConfig' which is to be a `LoweringConfigAttr`");
+ return {};
+ }
+ // Parse literal ','
+ if (parser.parseComma()) return {};
+ // Parse variable 'translationInfo'
+ _result_translationInfo = IREEFieldParser<TranslationInfoAttr>::parse(parser);
+ if (failed(_result_translationInfo)) {
+ parser.emitError(
+ parser.getCurrentLocation(),
+ "failed to parse IREECodegen_CompilationInfoAttr parameter "
+ "'translationInfo' which is to be a `TranslationInfoAttr`");
+ return {};
+ }
+ // Parse literal ','
+ if (parser.parseComma()) return {};
+ // Parse literal 'workgroup_size'
+ if (parser.parseKeyword("workgroup_size")) return {};
+ // Parse literal '='
+ if (parser.parseEqual()) return {};
+ // Parse variable 'workgroupSize'
+ _result_workgroupSize = IREEFieldParser<ArrayAttr>::parse(parser);
+ if (failed(_result_workgroupSize)) {
+ parser.emitError(parser.getCurrentLocation(),
+ "failed to parse IREECodegen_CompilationInfoAttr "
+ "parameter 'workgroupSize' which is to be a `ArrayAttr`");
+ return {};
+ }
+ // Parse literal '>'
+ if (parser.parseGreater()) return {};
+ return CompilationInfoAttr::get(
+ parser.getContext(), _result_loweringConfig.getValue(),
+ _result_translationInfo.getValue(), _result_workgroupSize.getValue());
+}
+
+/// Printer method that is copied from the auto-generated using `assemblyFormat`
+/// available with patch D111594. Replace after that change is in IREE.
+void CompilationInfoAttr::print(::mlir::DialectAsmPrinter &printer) const {
+ printer << "compilation.info";
+ printer << "<";
+ printer << getLoweringConfig();
+ printer << ",";
+ printer << ' ';
+ printer << getTranslationInfo();
+ printer << ",";
+ printer << ' ' << "workgroup_size";
+ printer << ' ' << "=";
+ printer << ' ';
+ printer << getWorkgroupSize();
+ printer << ">";
+}
+
+SmallVector<int64_t> CompilationInfoAttr::getWorkgroupSizeVals() {
+ ArrayAttr workgroupSizeAttr = getWorkgroupSize();
+ if (!workgroupSizeAttr) return {};
+ return getIntegerVals(workgroupSizeAttr);
+}
+
+//===----------------------------------------------------------------------===//
+// Initialize attributes
+//===----------------------------------------------------------------------===//
+
+void IREECodegenDialect::initializeCodegenAttrs() {
+ addAttributes<
+#define GET_ATTRDEF_LIST
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.cpp.inc" // IWYU pragma: keeep
+ >();
+}
+
+OptionalParseResult IREECodegenDialect::parseCodegenAttrs(
+ DialectAsmParser &parser, StringRef mnemonic, Type type,
+ Attribute &value) const {
+ return generatedAttributeParser(parser, mnemonic, type, value);
+}
+
+LogicalResult IREECodegenDialect::printCodegenAttrs(
+ Attribute attr, DialectAsmPrinter &p) const {
+ return generatedAttributePrinter(attr, p);
+}
+
+} // namespace Codegen
+} // namespace IREE
+
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting iree_codegen.translation.info attribute on the
+// `hal.executable.entry_point`
+// ===----------------------------------------------------------------------===//
+
+IREE::Codegen::TranslationInfoAttr getTranslationInfo(
+ IREE::HAL::ExecutableEntryPointOp entryPointOp) {
+ return entryPointOp->getAttrOfType<IREE::Codegen::TranslationInfoAttr>(
+ kTranslationInfoAttrName);
+}
+
+SmallVector<int64_t> getWorkgroupSize(
+ IREE::HAL::ExecutableEntryPointOp entryPointOp) {
+ if (Optional<ArrayAttr> workgroupSizeAttrList =
+ entryPointOp.workgroup_size()) {
+ return getIntegerVals(*workgroupSizeAttrList);
+ }
+ return {};
+}
+
+void setTranslationInfo(IREE::HAL::ExecutableEntryPointOp entryPointOp,
+ IREE::Codegen::TranslationInfoAttr translationInfo,
+ ArrayRef<int64_t> workgroupSize) {
+ entryPointOp->setAttr(kTranslationInfoAttrName, translationInfo);
+ // The workgroup size is set on the entry point op directly.
+ if (!workgroupSize.empty()) {
+ MLIRContext *context = entryPointOp->getContext();
+ auto attrs = getIndexIntegerArrayAttr(context, workgroupSize);
+ entryPointOp.workgroup_sizeAttr(attrs);
+ }
+}
+
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting `iree_codegen.lowering.config` attribute on root
+// operations.
+// ===----------------------------------------------------------------------===//
+
+IREE::Codegen::LoweringConfigAttr getLoweringConfig(Operation *op) {
+ return op->getAttrOfType<IREE::Codegen::LoweringConfigAttr>(kConfigAttrName);
+}
+
+SmallVector<int64_t> getTileSizes(Operation *op, unsigned level) {
+ IREE::Codegen::LoweringConfigAttr configAttr = getLoweringConfig(op);
+ if (!configAttr) return {};
+ return configAttr.getTileSizeVals(level);
+}
+SmallVector<Value, 4> getTileSizes(OpBuilder &b, Operation *op,
+ unsigned level) {
+ return llvm::to_vector<4>(
+ llvm::map_range(getTileSizes(op, level), [&](int64_t t) -> Value {
+ return b.create<arith::ConstantIndexOp>(op->getLoc(), t);
+ }));
+}
+
+void setLoweringConfig(Operation *op,
+ IREE::Codegen::LoweringConfigAttr config) {
+ op->setAttr(kConfigAttrName, config);
+}
+
+LogicalResult setOpConfigAndEntryPointFnTranslation(
+ FuncOp entryPointFn, Operation *op,
+ IREE::Codegen::LoweringConfigAttr config,
+ IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workgroupSize) {
+ auto partitionedLoops = getPartitionedLoops(op);
+ SmallVector<int64_t, 3> workloadPerWorkgroup;
+ auto tileSizes = config.getTileSizeVals(0);
+ if (!tileSizes.empty() && !partitionedLoops.empty()) {
+ for (unsigned depth : partitionedLoops) {
+ if (depth >= tileSizes.size()) {
+ return op->emitOpError(
+ "illegal configuration for lowering op, expect first level "
+ "tile size to contain at least ")
+ << partitionedLoops.back() << " elements";
+ }
+ if (tileSizes[depth] == 0) {
+ return op->emitOpError("illegal to set tilesize of loop ")
+ << depth
+ << " to zero since it is set to be partitioned at the flow "
+ "level";
+ }
+ workloadPerWorkgroup.push_back(tileSizes[depth]);
+ }
+ if (!workloadPerWorkgroup.empty()) {
+ workloadPerWorkgroup =
+ llvm::to_vector<3>(llvm::reverse(workloadPerWorkgroup));
+ }
+ }
+ auto entryPointOp = getEntryPoint(entryPointFn);
+ if (!entryPointOp) {
+ return entryPointFn.emitOpError(
+ "unable to find entry point op for entry point function");
+ }
+ auto translationInfo = IREE::Codegen::TranslationInfoAttr::get(
+ entryPointOp->getContext(), passPipeline, workloadPerWorkgroup);
+ setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting `iree_codegen.compilation.info` attribute on root
+// operations to override IREEs default compilation.
+// ===----------------------------------------------------------------------===//
+
+IREE::Codegen::CompilationInfoAttr getCompilationInfo(Operation *op) {
+ return op->getAttrOfType<IREE::Codegen::CompilationInfoAttr>(
+ kCompilationInfoAttrName);
+}
+
+void setCompilationInfo(Operation *op,
+ IREE::Codegen::CompilationInfoAttr config) {
+ op->setAttr(kCompilationInfoAttrName, config);
+}
+
+void eraseCompilationInfo(Operation *op) {
+ op->removeAttr(kCompilationInfoAttrName);
+}
+
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Codegen/Dialect/LoweringConfig.h b/iree/compiler/Codegen/Dialect/LoweringConfig.h
new file mode 100644
index 0000000..6d99215
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/LoweringConfig.h
@@ -0,0 +1,152 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+//===- LoweringConfig.h - Declares configuration for lowering Linalg ops --===//
+//
+// This file declares an attribute that drives how a dispatch region containing
+// a set of operations are lowered. The attribute itself is attached to Linalg
+// operations, and help converting a Linalg operation into "scalar code".
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
+#define IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
+
+#include "iree/compiler/Codegen/Utils/Utils.h"
+#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
+#include "mlir/IR/Builders.h"
+#include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+
+namespace mlir {
+namespace iree_compiler {
+/// Typedef for tile sizes to use at different levels of tiling.
+using TileSizesListType = SmallVector<SmallVector<int64_t>>;
+using TileSizesListTypeRef = ArrayRef<SmallVector<int64_t>>;
+} // namespace iree_compiler
+} // namespace mlir
+
+// clang-format off
+#include "iree/compiler/Codegen/Dialect/LoweringConfigEnums.h.inc"
+#define GET_ATTRDEF_CLASSES
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h.inc"
+// clang-format on
+
+namespace mlir {
+namespace iree_compiler {
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting iree_codegen.translation.info attribute on the
+// `hal.executable.entry_point`
+// ===----------------------------------------------------------------------===//
+
+/// Gets the translate executable info attribute value associated with
+/// `entryPointOp`. It expects that the attribute is stored using the identifier
+/// `translation.info`.
+IREE::Codegen::TranslationInfoAttr getTranslationInfo(
+ IREE::HAL::ExecutableEntryPointOp entryPointOp);
+/// Returns the translation info for the `funcOp` (by looking at the entry
+/// point). Returns `nullptr` on failure.
+inline IREE::Codegen::TranslationInfoAttr getTranslationInfo(FuncOp funcOp) {
+ auto entryPointOp = getEntryPoint(funcOp);
+ if (!entryPointOp) return nullptr;
+ return getTranslationInfo(entryPointOp);
+}
+
+/// Returns the workgroup size specified on the `entryPointOp`.
+SmallVector<int64_t> getWorkgroupSize(
+ IREE::HAL::ExecutableEntryPointOp entryPointOp);
+
+/// Set the translate executable info with the entry point op. Overwrites the
+/// existing attributes.
+// TODO(ravishankarm, benvanik): Eventually all the information needed for the
+// lowering will be consolidated into a single attribute with richer
+// information.
+void setTranslationInfo(IREE::HAL::ExecutableEntryPointOp entryPointOp,
+ IREE::Codegen::TranslationInfoAttr translationInfo,
+ ArrayRef<int64_t> workgroupSize = {});
+inline void setTranslationInfo(
+ FuncOp entryPointFn, IREE::Codegen::TranslationInfoAttr translationInfo,
+ ArrayRef<int64_t> workgroupSize = {}) {
+ auto entryPointOp = getEntryPoint(entryPointFn);
+ return setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
+}
+
+/// Sets the translation info on the `hal.executable.entry_point` op
+/// corresponding to the `entryPointFn`. Returns failure if a translation info
+/// is already set on the entry point op and is incompatible with what is being
+/// set.
+inline void setTranslationInfo(
+ FuncOp entryPointFn,
+ IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workloadPerWorkgroup, ArrayRef<int64_t> workgroupSize) {
+ auto entryPointOp = getEntryPoint(entryPointFn);
+ MLIRContext *context = entryPointFn.getContext();
+ auto translationInfo = IREE::Codegen::TranslationInfoAttr::get(
+ context, passPipeline, workloadPerWorkgroup);
+ setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
+}
+
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting `iree_codegen.lowering.config` attribute on root
+// operations.
+// ===----------------------------------------------------------------------===//
+
+/// Returns the lowering configuration set for an operation. Returns `nullptr`
+/// if no value is set. It expects that the attribute is stored using the
+/// identifier `lowering.config`.
+IREE::Codegen::LoweringConfigAttr getLoweringConfig(Operation *op);
+
+/// Returns the tile sizes for a particular operation if the
+/// `iree_codegen.lowering.config` attribute is set on it.
+SmallVector<int64_t> getTileSizes(Operation *op, unsigned level);
+SmallVector<Value, 4> getTileSizes(OpBuilder &b, Operation *op, unsigned level);
+
+/// Sets the lowering configuration, overwriting existing attribute values.
+void setLoweringConfig(Operation *op, IREE::Codegen::LoweringConfigAttr config);
+
+/// Sets translation for the entry-point function based on op configuration.
+LogicalResult setOpConfigAndEntryPointFnTranslation(
+ FuncOp entryPointFn, Operation *op,
+ IREE::Codegen::LoweringConfigAttr config,
+ IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workgroupSize = {});
+inline LogicalResult setOpConfigAndEntryPointFnTranslation(
+ FuncOp entryPointFn, Operation *op, TileSizesListTypeRef tileSizes,
+ ArrayRef<int64_t> nativeVectorSize,
+ IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
+ ArrayRef<int64_t> workgroupSize = {}) {
+ MLIRContext *context = entryPointFn.getContext();
+ auto config = IREE::Codegen::LoweringConfigAttr::get(context, tileSizes,
+ nativeVectorSize);
+ setLoweringConfig(op, config);
+ return setOpConfigAndEntryPointFnTranslation(entryPointFn, op, config,
+ passPipeline, workgroupSize);
+}
+
+//===----------------------------------------------------------------------===//
+// Helpers for getting/setting `iree_codegen.compilation.info` attribute on root
+// operations to override IREEs default compilation.
+// ===----------------------------------------------------------------------===//
+
+/// Returns the `#iree_codegen.compilation.info` set on the operation. Assumes
+/// that the identifier used is `compilation.info`.
+IREE::Codegen::CompilationInfoAttr getCompilationInfo(Operation *op);
+
+/// Sets the `config` to use for compiling the operation. If `op` is the root
+/// operation of the dispatch region, overrides the default configuration that
+/// is used for compilation.
+void setCompilationInfo(Operation *op,
+ IREE::Codegen::CompilationInfoAttr config);
+
+/// Removes the `#iree_codegen.compilation.info` attribute that is set on the
+/// operation.
+void eraseCompilationInfo(Operation *op);
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
diff --git a/iree/compiler/Codegen/Dialect/LoweringConfig.td b/iree/compiler/Codegen/Dialect/LoweringConfig.td
new file mode 100644
index 0000000..13489da
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/LoweringConfig.td
@@ -0,0 +1,192 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_CODEGEN_DIALECT_LOWERINGCONFIG
+#define IREE_COMPILER_CODEGEN_DIALECT_LOWERINGCONFIG
+
+include "iree/compiler/Codegen/Dialect/IREECodegenDialect.td"
+
+// List of pre-existing pipelines for translating executables.
+def CPU_Default
+ : StrEnumAttrCase<"CPUDefault">;
+def CPU_Vectorization
+ : StrEnumAttrCase<"CPUVectorization">;
+def CPU_TensorToVectors
+ : StrEnumAttrCase<"CPUTensorToVectors">;
+
+def LLVMGPU_SimpleDistribute
+ : StrEnumAttrCase<"LLVMGPUDistribute">;
+def LLVMGPU_Vectorize
+ : StrEnumAttrCase<"LLVMGPUVectorize">;
+def LLVMGPU_MatmulSimt
+ : StrEnumAttrCase<"LLVMGPUMatmulSimt">;
+
+def SPIRV_SimpleDistribute
+ : StrEnumAttrCase<"SPIRVDistribute">;
+def SPIRV_Vectorize
+ : StrEnumAttrCase<"SPIRVVectorize">;
+def SPIRV_DistributeToGlobalID
+ : StrEnumAttrCase<"SPIRVDistributeToGlobalID">;
+def SPIRV_VectorizeToCooperativeOps
+ : StrEnumAttrCase<"SPIRVVectorizeToCooperativeOps">;
+
+def None
+ : StrEnumAttrCase<"None">;
+
+// EnumAttrCase for all known lowerings for ops within dispatch region
+// to scalar/native-vector code.
+def DispatchLoweringPassPipelineEnum : StrEnumAttr<
+ "DispatchLoweringPassPipeline",
+ "identifier for pass pipeline use to lower dispatch region",
+ [CPU_Default, CPU_TensorToVectors, CPU_Vectorization,
+ LLVMGPU_SimpleDistribute, LLVMGPU_Vectorize, LLVMGPU_MatmulSimt,
+ SPIRV_SimpleDistribute, SPIRV_DistributeToGlobalID,
+ SPIRV_Vectorize, SPIRV_VectorizeToCooperativeOps, None]> {
+ let cppNamespace = "::mlir::iree_compiler::IREE::Codegen";
+}
+
+def IREECodegen_TranslationInfoAttr :
+ AttrDef<IREECodegen_Dialect, "TranslationInfo", []> {
+ let mnemonic = "translation.info";
+ let summary = [{drive dispatch entry point lowering}];
+ let description = [{
+ Specifies the information that is used to drive the translation of
+ an entry point function using Linalg based structured-op
+ lowering.. During executable translation this is attached to the
+ `hal.executable.entry_point` operation.
+
+ If this operation is already set on the root operation (as part of
+ `iree_codegen.compilation.info`) that drives the compilation of a
+ dispatch region (like `linalg.matmul`/`linalg.*conv*`), this
+ attribute gets propagated to the entry point function.
+
+ The fields are
+ - `passPipeline` : The pass pipeline to use.
+ - `workloadPerWorkgroup` : Specifies how much of the original
+ `workload` is handled by a workgroup along `x`, `y` and `z`. If
+ left empty it implies that that there is a single workgroup that
+ does the entire `workload`.
+
+ }];
+
+ // TODO(ravishankarm): Commented out till patch D111594 lands.
+ // let assemblyFormat = [{
+ // `<` $passPipeline `,` `workload_per_wg` `=` $workloadPerWorkgroup `>`
+ // }];
+
+ let parameters = (ins
+ AttrParameter<"StringAttr", "">:$passPipeline,
+ AttrParameter<"ArrayAttr", "">:$workloadPerWorkgroup
+ );
+ let builders = [
+ AttrBuilder<(ins "DispatchLoweringPassPipeline":$passPipeline,
+ CArg<"ArrayRef<int64_t>", "{}">:$workloadPerWorkgroup)>
+ ];
+ let extraClassDeclaration = [{
+ // Returns the lowering pass pipeline set.
+ DispatchLoweringPassPipeline getDispatchLoweringPassPipeline();
+
+ // Returns values of the workloadPerWorkgroup field if set.
+ SmallVector<int64_t> getWorkloadPerWorkgroupVals();
+ }];
+ let genVerifyDecl = 1;
+}
+
+def IREECodegen_LoweringConfigAttr :
+ AttrDef<IREECodegen_Dialect, "LoweringConfig", []> {
+ let mnemonic = "lowering.config";
+ let summary = [{drive lowering of an operation within dispatch region}];
+ let description = [{
+ Specifies the information that is used by backend compiler to
+ translate an operation to scalar code. The way the information is
+ used is specific to each backend (indeed specific to the pass
+ pipeline used) to compile that operation.
+
+ TODO: Currently there is no verification that the configuration
+ specifies everything needed for a pass-pipeline. The values to set
+ for these parameters is dependent on the pass-pipeline
+ implementation. In future, each pass pipeline could verify that
+ the lowering configuration has all the necessary attributes for
+ the pipeline.
+
+ }];
+
+ // TODO(ravishankarm): Commented out till patch D111594 lands.
+ // let assemblyFormat = [{
+ // `<` `tile_sizes` `=` $tileSizes `,` `native_vector_size` `=` $nativeVectorSize `>`
+ // }];
+
+ let parameters = (ins
+ AttrParameter<"ArrayAttr", "">:$tileSizes,
+ AttrParameter<"ArrayAttr", "">:$nativeVectorSize
+ );
+ let builders = [
+ AttrBuilder<(ins "TileSizesListTypeRef":$tileSizes,
+ CArg<"ArrayRef<int64_t>", "{}">:$nativeVectorSize)>
+ ];
+ let extraClassDeclaration = [{
+ // Returns the tile sizes for all levels set for the op.
+ TileSizesListType getTileSizeVals();
+
+ // Returns the tile sizes for a level set for the op.
+ SmallVector<int64_t> getTileSizeVals(unsigned level = 0);
+
+ // Returns the native vector size to use.
+ SmallVector<int64_t> getNativeVectorSizeVals();
+ }];
+ let genVerifyDecl = 1;
+}
+
+def IREECodegen_CompilationInfoAttr :
+ AttrDef<IREECodegen_Dialect, "CompilationInfo", []> {
+ let mnemonic = "compilation.info";
+ let summary = [{drive lowering of an operation from input dialect}];
+ let description = [{
+ Specifies the information that allows controlling the compilation
+ of operations like `linalg.matmul`/`linalg.*conv` within
+ IREE. This information is used to override the defaults used by
+ the IREE compiler. Currently it is only valid to set this on
+ `linalg.matmul`/`linalg.*conv*` operations.
+
+ TODO: It is expected that the `TranslationInfoAttr` and the
+ `LoweringConfigAttr` are specified. Currently there is no
+ verification that the values of the `LoweringConfigAttr` fully
+ specifies the behaviour of the compilation path chosen with
+ `TranslationInfoAttr`. This could be added in the future. Note:
+ Typically the values used for the first-level tiling in
+ `LoweringConfigAttr` and `workload_per_wg` value in the
+ `TranslationInfoAttr` are the same since the first-level of tile +
+ distribute is already done at the `Flow` level. This verification
+ is also a TODO.
+ }];
+ let parameters = (ins
+ AttrParameter<"LoweringConfigAttr", "">:$loweringConfig,
+ AttrParameter<"TranslationInfoAttr", "">:$translationInfo,
+ AttrParameter<"ArrayAttr", "">:$workgroupSize
+ );
+
+ // TODO(ravishankarm): Commented out till patch D111594 lands.
+ // let assemblyFormat = [{
+ // `<` $loweringConfig `,` $translationInfo `,` `workgroup_size` `=` $workgroupSize `>`
+ // }];
+
+ let builders = [
+ AttrBuilder<(ins "TileSizesListTypeRef":$tileSizes,
+ "ArrayRef<int64_t>":$nativeVectorSize,
+ CArg<"ArrayRef<int64_t>", "{}">:$workgroupSize)>,
+ AttrBuilder<(ins "TileSizesListTypeRef":$tileSizes,
+ "ArrayRef<int64_t>":$nativeVectorSize,
+ "DispatchLoweringPassPipeline":$passPipeline,
+ "ArrayRef<int64_t>":$workloadPerWorkgroup,
+ CArg<"ArrayRef<int64_t>", "{}">:$workgroupSize)>,
+ ];
+ let extraClassDeclaration = [{
+ SmallVector<int64_t> getWorkgroupSizeVals();
+ }];
+ let genVerifyDecl = 1;
+}
+
+#endif // IREE_COMPILER_CODEGEN_DIALECT_LOWERINGCONFIG
diff --git a/iree/compiler/Codegen/Dialect/test/BUILD b/iree/compiler/Codegen/Dialect/test/BUILD
new file mode 100644
index 0000000..704286c
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/test/BUILD
@@ -0,0 +1,30 @@
+# Copyright 2021 The IREE Authors
+#
+# Licensed under the Apache License v2.0 with LLVM Exceptions.
+# See https://llvm.org/LICENSE.txt for license information.
+# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+# Tests for common transforms.
+
+load("//iree:lit_test.bzl", "iree_lit_test_suite")
+load("//build_tools/bazel:enforce_glob.bzl", "enforce_glob")
+
+package(
+ default_visibility = ["//visibility:public"],
+ features = ["layering_check"],
+ licenses = ["notice"], # Apache 2.0
+)
+
+iree_lit_test_suite(
+ name = "lit",
+ srcs = enforce_glob(
+ [
+ "lowering_config_attr.mlir",
+ ],
+ include = ["*.mlir"],
+ ),
+ data = [
+ "//iree/tools:IreeFileCheck",
+ "//iree/tools:iree-opt",
+ ],
+)
diff --git a/iree/compiler/Codegen/Dialect/test/CMakeLists.txt b/iree/compiler/Codegen/Dialect/test/CMakeLists.txt
new file mode 100644
index 0000000..4de932a
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/test/CMakeLists.txt
@@ -0,0 +1,23 @@
+################################################################################
+# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
+# iree/compiler/Codegen/Dialect/test/BUILD #
+# #
+# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
+# CMake-only content. #
+# #
+# To disable autogeneration for this file entirely, delete this header. #
+################################################################################
+
+iree_add_all_subdirs()
+
+iree_lit_test_suite(
+ NAME
+ lit
+ SRCS
+ "lowering_config_attr.mlir"
+ DATA
+ iree::tools::IreeFileCheck
+ iree::tools::iree-opt
+)
+
+### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir b/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
new file mode 100644
index 0000000..3639996
--- /dev/null
+++ b/iree/compiler/Codegen/Dialect/test/lowering_config_attr.mlir
@@ -0,0 +1,37 @@
+// RUN: iree-opt -split-input-file %s | IreeFileCheck %s
+
+module attributes {
+ translation.info = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [32, 42]>
+} { }
+// CHECK: #translation = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [32, 42]>
+
+// -----
+
+module attributes {
+ translation.info = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = []>
+} { }
+// CHECK: #translation = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = []>
+
+// -----
+
+module attributes {
+ lowering.config = #iree_codegen.lowering.config<tile_sizes = [], native_vector_size = []>
+} { }
+// CHECK: #config = #iree_codegen.lowering.config<tile_sizes = [], native_vector_size = []>
+
+// -----
+
+module attributes {
+ lowering.config = #iree_codegen.lowering.config<tile_sizes = [[], [10]], native_vector_size = [32, 32]>
+} { }
+// CHECK: #config = #iree_codegen.lowering.config<tile_sizes = {{\[}}[], [10]{{\]}}, native_vector_size = [32, 32]>
+
+// -----
+
+module attributes {
+ compilation.info = #iree_codegen.compilation.info<
+ #iree_codegen.lowering.config<tile_sizes = [], native_vector_size = []>,
+ #iree_codegen.translation.info<"CPUDefault", workload_per_wg = []>,
+ workgroup_size = []>
+} { }
+// CHECK: #compilation = #iree_codegen.compilation.info<#iree_codegen.lowering.config<tile_sizes = [], native_vector_size = []>, #iree_codegen.translation.info<"CPUDefault", workload_per_wg = []>, workgroup_size = []>
\ No newline at end of file
diff --git a/iree/compiler/Codegen/LLVMCPU/BUILD b/iree/compiler/Codegen/LLVMCPU/BUILD
index 0d41055..21e2802 100644
--- a/iree/compiler/Codegen/LLVMCPU/BUILD
+++ b/iree/compiler/Codegen/LLVMCPU/BUILD
@@ -29,6 +29,7 @@
deps = [
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/Transforms",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
index 50207fd..44d43b1 100644
--- a/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
@@ -57,6 +57,7 @@
MLIRVectorToLLVM
MLIRVectorToSCF
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::Transforms
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 3504b5b..dbbdcca 100644
--- a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -8,9 +8,7 @@
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/CommandLine.h"
@@ -227,8 +225,8 @@
getDefaultWorkloadPerWorkgroup(tiledLoops, nativeVectorSizeInElements);
setTranslationInfo(
- entryPointFn, IREE::HAL::DispatchLoweringPassPipeline::CPUDefault,
- /*workgroupSize =*/ArrayRef<int64_t>{}, workloadPerWorkgroup);
+ entryPointFn, IREE::Codegen::DispatchLoweringPassPipeline::CPUDefault,
+ workloadPerWorkgroup, /*workgroupSize =*/ArrayRef<int64_t>{});
return success();
}
@@ -299,8 +297,9 @@
vectorSizeVals[i]);
}
setTranslationInfo(
- entryPointFn, IREE::HAL::DispatchLoweringPassPipeline::CPUTensorToVectors,
- /*workgroupSize =*/ArrayRef<int64_t>{}, workloadPerWorkgroup);
+ entryPointFn,
+ IREE::Codegen::DispatchLoweringPassPipeline::CPUTensorToVectors,
+ workloadPerWorkgroup, /*workgroupSize =*/ArrayRef<int64_t>{});
SmallVector<int64_t, 4> l1TileSizes, vectorTileSizes;
if (isBatchMatmul) {
@@ -321,8 +320,8 @@
// level tiling.
tileSizes.emplace_back(std::move(l1TileSizes));
tileSizes.emplace_back(std::move(vectorTileSizes));
- IREE::HAL::LoweringConfig config =
- buildConfigAttr(tileSizes, vectorSizeVals, entryPointFn.getContext());
+ auto config = IREE::Codegen::LoweringConfigAttr::get(
+ entryPointFn.getContext(), tileSizes, vectorSizeVals);
setLoweringConfig(contractionOp, config);
return success();
}
@@ -368,14 +367,14 @@
return {1, 1, 1, M0, N0, K0};
};
- SmallVector<int64_t, 4> nativeVectorSize = getVectorSizes();
+ SmallVector<int64_t> nativeVectorSize = getVectorSizes();
TileSizesListType tileSizes = {getWorkgroupTileSizes(), getL1TileSizes(),
nativeVectorSize};
return setOpConfigAndEntryPointFnTranslation(
entryPointFn, mmt4dOp, tileSizes, nativeVectorSize,
- IREE::HAL::DispatchLoweringPassPipeline::CPUVectorization);
+ IREE::Codegen::DispatchLoweringPassPipeline::CPUVectorization);
}
/// Sets the lowering configuration for dispatch region for linalg_ext.fft
@@ -384,8 +383,7 @@
ArrayRef<TiledLoopInfo> tiledLoops) {
auto partitionedLoops = getPartitionedLoops(fftOp);
unsigned maxDepth = partitionedLoops.back() + 1;
- SmallVector<int64_t, 4> workgroupTileSizes(maxDepth,
- defaultWorkgroupTileSize);
+ SmallVector<int64_t> workgroupTileSizes(maxDepth, defaultWorkgroupTileSize);
llvm::DenseSet<unsigned> partitionedLoopsSet(partitionedLoops.begin(),
partitionedLoops.end());
for (auto dim : llvm::seq<int64_t>(0, workgroupTileSizes.size())) {
@@ -412,7 +410,7 @@
return setOpConfigAndEntryPointFnTranslation(
entryPointFn, fftOp, tileSizes,
/*nativeVectorSizes=*/ArrayRef<int64_t>{},
- IREE::HAL::DispatchLoweringPassPipeline::CPUDefault);
+ IREE::Codegen::DispatchLoweringPassPipeline::CPUDefault);
}
/// Finds the root operation in the given list of linalg operations and sets
@@ -454,11 +452,8 @@
for (auto computeOp : computeOps) {
if (!hasMarker(computeOp, getWorkgroupMarker())) continue;
- if (auto config = getLoweringConfig(computeOp)) {
- // Check if the op has a preset pipeline.
- auto passPipeline = getLoweringPassPipeline(config);
- if (!passPipeline) continue;
-
+ if (IREE::Codegen::CompilationInfoAttr compilationInfo =
+ getCompilationInfo(computeOp)) {
// If the function already has a translation, error out.
if (auto translationInfo = getTranslationInfo(entryPointFn)) {
return computeOp->emitOpError(
@@ -466,17 +461,12 @@
"info");
}
- SmallVector<int64_t, 4> workgroupSize;
- if (auto workgroupSizeAttr = config.workgroupSize()) {
- workgroupSize = llvm::to_vector<4>(
- llvm::map_range(workgroupSizeAttr, [](Attribute intAttr) {
- return intAttr.cast<IntegerAttr>().getInt();
- }));
- }
- if (failed(setOpConfigAndEntryPointFnTranslation(
- entryPointFn, computeOp, config, *passPipeline, workgroupSize))) {
- return failure();
- }
+ SmallVector<int64_t> workgroupSize =
+ compilationInfo.getWorkgroupSizeVals();
+ setTranslationInfo(entryPointFn, compilationInfo.getTranslationInfo(),
+ workgroupSize);
+ setLoweringConfig(computeOp, compilationInfo.getLoweringConfig());
+ eraseCompilationInfo(computeOp);
}
}
diff --git a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
index dbacd1a..afb616e 100644
--- a/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
+++ b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
@@ -7,7 +7,7 @@
#ifndef IREE_COMPILER_CODEGEN_LLVMCPU_KERNELDISPATCH_H_
#define IREE_COMPILER_CODEGEN_LLVMCPU_KERNELDISPATCH_H_
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "mlir/IR/BuiltinOps.h"
namespace mlir {
diff --git a/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
index 87bd94b..0732fbb 100644
--- a/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -30,8 +30,9 @@
LLVMCPULowerExecutableTargetPass> {
public:
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<IREE::HAL::HALDialect, linalg::LinalgDialect,
- LLVM::LLVMDialect, vector::VectorDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, IREE::HAL::HALDialect,
+ linalg::LinalgDialect, LLVM::LLVMDialect,
+ vector::VectorDialect>();
}
LLVMCPULowerExecutableTargetPass(bool vectorize = true)
@@ -123,16 +124,15 @@
// is fine.
llvm::StringMap<IREE::HAL::ExecutableEntryPointOp> entryPoints =
getAllEntryPoints(moduleOp);
- Optional<IREE::HAL::DispatchLoweringPassPipeline> passPipeline;
+ Optional<IREE::Codegen::DispatchLoweringPassPipeline> passPipeline;
for (auto &it : entryPoints) {
auto entryPointOp = it.second;
- if (IREE::HAL::TranslationInfo translationInfo =
+ if (IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(entryPointOp)) {
- Optional<IREE::HAL::DispatchLoweringPassPipeline> currPipeline =
- getLoweringPassPipeline(translationInfo);
- if (!currPipeline) continue;
+ IREE::Codegen::DispatchLoweringPassPipeline currPipeline =
+ translationInfo.getDispatchLoweringPassPipeline();
if (passPipeline) {
- if (currPipeline.getValue() != passPipeline.getValue()) {
+ if (currPipeline != passPipeline.getValue()) {
moduleOp.emitError(
"unhandled compilation of entry point function with different "
"pass pipelines within a module");
@@ -150,14 +150,14 @@
OpPassManager &nestedModulePM =
executableLoweringPipeline.nest<ModuleOp>();
switch (passPipeline.getValue()) {
- case IREE::HAL::DispatchLoweringPassPipeline::CPUDefault:
- case IREE::HAL::DispatchLoweringPassPipeline::None:
+ case IREE::Codegen::DispatchLoweringPassPipeline::CPUDefault:
+ case IREE::Codegen::DispatchLoweringPassPipeline::None:
addCPUDefaultPassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::CPUVectorization:
+ case IREE::Codegen::DispatchLoweringPassPipeline::CPUVectorization:
addCPUVectorizationPassPipeline(nestedModulePM, lowerToVectors);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::CPUTensorToVectors:
+ case IREE::Codegen::DispatchLoweringPassPipeline::CPUTensorToVectors:
addTensorToVectorsPassPipeline(nestedModulePM, lowerToVectors);
break;
default:
diff --git a/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndVectorizeLinalgTensorOps.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndVectorizeLinalgTensorOps.cpp
index aa1a3c6..5a110e7 100644
--- a/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndVectorizeLinalgTensorOps.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTileAndVectorizeLinalgTensorOps.cpp
@@ -4,6 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
@@ -78,9 +79,8 @@
l1patterns.insert<TileWorkgroups>(
context,
linalg::LinalgTilingOptions().setTileSizeComputationFunction(
- [](OpBuilder &builder,
- Operation *operation) -> SmallVector<Value, 4> {
- return getTileSizes(builder, operation,
+ [](OpBuilder &builder, Operation *op) -> SmallVector<Value, 4> {
+ return getTileSizes(builder, op,
static_cast<unsigned>(TilingLevel::L1Tiles));
}),
linalg::LinalgTransformationFilter(
@@ -112,11 +112,9 @@
l2patterns.insert<TileWorkgroups>(
context,
linalg::LinalgTilingOptions().setTileSizeComputationFunction(
- [](OpBuilder &builder,
- Operation *operation) -> SmallVector<Value, 4> {
+ [](OpBuilder &builder, Operation *op) -> SmallVector<Value, 4> {
return getTileSizes(
- builder, operation,
- static_cast<unsigned>(TilingLevel::VectorTiles));
+ builder, op, static_cast<unsigned>(TilingLevel::VectorTiles));
}),
linalg::LinalgTransformationFilter(
Identifier::get(getWorkgroupL1TileMarker(), context),
diff --git a/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
index 2fc0185..53092ae 100644
--- a/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
@@ -167,9 +167,8 @@
l1patterns.insert<TileWorkgroups>(
context,
linalg::LinalgTilingOptions().setTileSizeComputationFunction(
- [](OpBuilder &builder,
- Operation *operation) -> SmallVector<Value, 4> {
- return getTileSizes(builder, operation,
+ [](OpBuilder &builder, Operation *op) -> SmallVector<Value, 4> {
+ return getTileSizes(builder, op,
static_cast<unsigned>(TilingLevel::L1Tiles));
}),
linalg::LinalgTransformationFilter(
@@ -188,11 +187,9 @@
l2patterns.insert<TileWorkgroups>(
context,
linalg::LinalgTilingOptions().setTileSizeComputationFunction(
- [](OpBuilder &builder,
- Operation *operation) -> SmallVector<Value, 4> {
+ [](OpBuilder &builder, Operation *op) -> SmallVector<Value, 4> {
return getTileSizes(
- builder, operation,
- static_cast<unsigned>(TilingLevel::VectorTiles));
+ builder, op, static_cast<unsigned>(TilingLevel::VectorTiles));
}),
linalg::LinalgTransformationFilter(
Identifier::get(getWorkgroupL1TileMarker(), context),
diff --git a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
index 91728ce..c307c9b 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -59,10 +59,11 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {nativeVectorSize = [4, 4, 4], tileSizes = {{\[}}{{\[}}{{\]}}, [32, 32, 32], [4, 4, 4]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[], [32, 32, 32], [4, 4, 4]{{\]}}, native_vector_size = [4, 4, 4]>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUTensorToVectors", workload_per_wg = [64, 64]>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK: hal.executable.entry_point public @matmul_tensors
-// CHECK-SAME: translation.info = {passPipeline = "CPUTensorToVectors", workloadPerWorkgroup = [64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -118,11 +119,10 @@
}
}
}
-
-// CHECK: #[[CONFIG:[a-zA-Z]+]] = {passPipeline = "CPUDefault"}
-// CHECK: hal.executable private @add_no_config
-// CHECK: hal.executable.entry_point public @add_no_config
-// CHECK-SAME: translation.info = #[[CONFIG]]
+// CHECK: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = []>
+// CHECK: hal.executable private @add_no_config
+// CHECK: hal.executable.entry_point public @add_no_config
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// -----
@@ -192,9 +192,10 @@
}
}
}
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64]>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK: hal.executable.entry_point public @add
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -296,8 +297,9 @@
}
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable.entry_point public @add4D
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG2:[a-zA-Z0-9_]+]]: index)
@@ -378,8 +380,9 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {nativeVectorSize = [1, 4, 4, 4], tileSizes = {{\[}}[], [1, 32, 32, 32], [1, 4, 4, 4]{{\]}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[], [1, 32, 32, 32], [1, 4, 4, 4]{{\]}}, native_vector_size = [1, 4, 4, 4]>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUTensorToVectors", workload_per_wg = [64, 64]>
// CHECK: hal.executable.entry_point public @batch_matmul_tensors
// CHECK-NEXT: (%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9]+]]: index
@@ -393,6 +396,10 @@
// -----
+#compilation = #iree_codegen.compilation.info<
+ #iree_codegen.lowering.config<tile_sizes = [[32, 32, 32]], native_vector_size = []>,
+ #iree_codegen.translation.info<"CPUVectorization", workload_per_wg = [32, 32]>,
+ workgroup_size = []>
hal.executable private @preset_config_matmul_tensors {
hal.executable.variant @system_elf_x86_64, target = #hal.executable.target<"llvm", "system-elf-x86_64"> {
hal.executable.entry_point @preset_config attributes {interface = @io, ordinal = 0 : index}
@@ -427,7 +434,11 @@
%14 = affine.min affine_map<(d0)[s0] -> (-d0 + 512, s0)>(%arg1)[%workgroup_size_x]
%15 = linalg.init_tensor [%13, %14] : tensor<?x?xf32>
%16 = linalg.fill(%cst, %15) : f32, tensor<?x?xf32> -> tensor<?x?xf32>
- %17 = linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = {passPipeline = "CPUVectorization", tileSizes = [[32, 32, 32]]}} ins(%8, %10 : tensor<?x256xf32>, tensor<256x?xf32>) outs(%16 : tensor<?x?xf32>) -> tensor<?x?xf32>
+ %17 = linalg.matmul {
+ __internal_linalg_transform__ = "workgroup",
+ compilation.info = #compilation}
+ ins(%8, %10 : tensor<?x256xf32>, tensor<256x?xf32>)
+ outs(%16 : tensor<?x?xf32>) -> tensor<?x?xf32>
flow.dispatch.tensor.store %17, %2, offsets = [%arg0, %arg1], sizes = [%11, %12], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:128x512xf32>
}
}
@@ -441,11 +452,12 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {passPipeline = "CPUVectorization", tileSizes = {{\[}}[32, 32, 32]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[32, 32, 32]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 * 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUVectorization", workload_per_wg = [32, 32]>
// CHECK: hal.executable.entry_point
-// CHECK-SAME: translation.info = {passPipeline = "CPUVectorization", workloadPerWorkgroup = [32, 32]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[NWG_X:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
@@ -511,9 +523,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64]>
// CHECK: hal.executable.entry_point public @tensor_insert_slice
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -548,12 +561,11 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[64]]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[64]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @static_1d_fft_stage2
-// CHECK-SAME: translation.info = {
-// CHECK-SAME: passPipeline = "CPUDefault"
-// CHECK-SAME: workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %[[ARG1:.+]]: index, %[[ARG2:.+]]: index):
// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
// CHECK-NEXT: %[[T0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
@@ -620,12 +632,11 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[64, 64, 64]]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[64, 64, 64]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable.entry_point public @static_3d_fft_stage3
-// CHECK-SAME: translation.info = {
-// CHECK-SAME: passPipeline = "CPUDefault"
-// CHECK-SAME: workloadPerWorkgroup = [64, 64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %[[ARG1:.+]]: index, %[[ARG2:.+]]: index):
// CHECK-NEXT: %[[T0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
// CHECK-NEXT: %[[T1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]]
@@ -700,8 +711,9 @@
}
}
}
+// CHECK: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64]>
// CHECK: hal.executable.entry_point public @outs_fusion_fn
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// -----
@@ -768,9 +780,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64, 64]>
// CHECK: hal.executable.entry_point public @conv attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64, 64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index)
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]
// CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]
@@ -844,8 +857,9 @@
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 64, 32]>
// CHECK: hal.executable.entry_point public @conv_static attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 64, 32]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index)
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]
// CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP0]]()[%[[ARG1]]
@@ -902,8 +916,9 @@
}
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [32, 8]>
// CHECK: hal.executable.entry_point public @generic_static attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [32, 8]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index)
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]
@@ -960,11 +975,12 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {nativeVectorSize = [4, 4, 4], tileSizes = {{\[}}[], [28, 8, 24], [4, 4, 4]{{\]}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[], [28, 8, 24], [4, 4, 4]{{\]}}, native_vector_size = [4, 4, 4]>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 28)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUTensorToVectors", workload_per_wg = [8, 28]>
// CHECK: hal.executable.entry_point public @matmul_static attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUTensorToVectors", workloadPerWorkgroup = [8, 28]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index)
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
@@ -1035,8 +1051,9 @@
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
// CHECK-DAG: #[[MAP2:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64, 8, 4]>
// CHECK: hal.executable.entry_point public @restrict_num_workgroups attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64, 8, 4]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index, %[[ARG2:[a-zA-Z0-9]+]]: index)
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
// CHECK-DAG: %[[D1:.+]] = affine.apply #[[MAP1]]()[%[[ARG1]]]
@@ -1074,9 +1091,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_0 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
@@ -1113,9 +1131,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECk-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_1 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
@@ -1152,9 +1171,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_2 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
@@ -1191,9 +1211,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_3 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
@@ -1230,9 +1251,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_4 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
@@ -1269,9 +1291,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"CPUDefault", workload_per_wg = [64]>
// CHECK: hal.executable.entry_point public @test_exp_5 attributes
-// CHECK-SAME: translation.info = {passPipeline = "CPUDefault", workloadPerWorkgroup = [64]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK-DAG: %[[D0:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
diff --git a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
index e491217..21e0e9a 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.variant(iree-llvmcpu-lower-executable-target{use-lowering-pipeline='builtin.func(iree-llvmcpu-vectorization)'}))" -split-input-file %s | IreeFileCheck %s
// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.variant(iree-llvmcpu-lower-executable-target{use-lowering-pipeline='builtin.func(iree-llvmcpu-vectorization{promote-workgroup-to-full-tiles}),cse'}))" -split-input-file %s | IreeFileCheck %s -check-prefix=CHECK-PROMOTED
-#config = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
+#config = #iree_codegen.lowering.config<tile_sizes = [[64, 64], [32, 32, 32], [4, 4, 4]], native_vector_size = [4, 4, 4]>
hal.executable private @dynamic_matmul {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -114,7 +114,7 @@
// -----
-#config = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
+#config = #iree_codegen.lowering.config<tile_sizes = [[64, 64], [32, 32, 32], [4, 4, 4]], native_vector_size = [4, 4, 4]>
hal.executable private @matmul_i8_i8_i32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
diff --git a/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir
index ddc4b8f..4a51e1a 100644
--- a/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/tile_and_vectorize.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt %s -cse -iree-llvmcpu-tile-and-vectorize -cse -canonicalize -split-input-file | IreeFileCheck %s
-#config0 = {tileSizes = [[64, 64]]}
-#config1 = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
+#config0 = #iree_codegen.lowering.config<tile_sizes = [[64, 64]], native_vector_size = []>
+#config1 = #iree_codegen.lowering.config<tile_sizes = [[64, 64], [32, 32, 32], [4, 4, 4]], native_vector_size = [4, 4, 4]>
#map0 = affine_map<()[s0] -> (s0 * 64)>
#map1 = affine_map<(d0) -> (64, -d0 + 383)>
#map2 = affine_map<(d0) -> (64, -d0 + 513)>
diff --git a/iree/compiler/Codegen/LLVMGPU/BUILD b/iree/compiler/Codegen/LLVMGPU/BUILD
index e9cc75d..51adc6e 100644
--- a/iree/compiler/Codegen/LLVMGPU/BUILD
+++ b/iree/compiler/Codegen/LLVMGPU/BUILD
@@ -34,6 +34,7 @@
deps = [
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/Transforms",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
index cf63b73..f3c058e 100644
--- a/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
@@ -61,6 +61,7 @@
MLIRVectorToLLVM
MLIRVectorToSCF
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::Transforms
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
index fa009cd..0d1ae4e 100644
--- a/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -8,7 +8,7 @@
#include <numeric>
-#include "iree/compiler/Codegen/Utils/Utils.h"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
#include "llvm/Support/Debug.h"
@@ -100,7 +100,7 @@
}
// Currently just a basic tile size to enable tiling and vectorization.
// TODO: pick a more efficient tile size and tile at subgroup level.
- SmallVector<int64_t, 4> ts;
+ SmallVector<int64_t> ts;
// Tile all the higher parallel dimension with a size of 1 and the 2 most
// inner dimension with the tileX/tileY size.
ts.append(op.getNumParallelLoops() - 2, 1);
@@ -110,14 +110,14 @@
tileSizes.push_back(ts); // Workgroup level.
return setOpConfigAndEntryPointFnTranslation(
entryPoint, op, tileSizes, /*nativeVectorSizes=*/ArrayRef<int64_t>{},
- IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUMatmulSimt,
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUMatmulSimt,
workgroupSize);
}
static LogicalResult setFftConfig(FuncOp entryPoint, linalg_ext::FftOp op) {
auto partitionedLoops = getPartitionedLoops(op);
unsigned loopDepth = partitionedLoops.back() + 1;
- SmallVector<int64_t, 4> workgroupTileSize(loopDepth, 0);
+ SmallVector<int64_t> workgroupTileSize(loopDepth, 0);
SmallVector<int64_t, 3> workgroupSize = {cudaWarpSize, 1, 1};
// Tiling along partitioned loops with size 1.
@@ -137,14 +137,14 @@
TileSizesListType tileSizes = {workgroupTileSize};
return setOpConfigAndEntryPointFnTranslation(
entryPoint, op, tileSizes, /*nativeVectorSizes=*/ArrayRef<int64_t>{},
- IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUDistribute,
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUDistribute,
workgroupSize);
}
// Basic default properties for linalg ops that haven't been tuned.
static LogicalResult setRootDefaultConfig(FuncOp entryPoint, Operation *op) {
- IREE::HAL::DispatchLoweringPassPipeline passPipeline =
- IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUDistribute;
+ IREE::Codegen::DispatchLoweringPassPipeline passPipeline =
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUDistribute;
TileSizesListType tileSizes;
SmallVector<unsigned> partitionedLoops = getPartitionedLoops(op);
if (partitionedLoops.empty()) {
@@ -210,43 +210,34 @@
tileSizes.emplace_back(std::move(workgroupTileSizes)); // Workgroup level
return setOpConfigAndEntryPointFnTranslation(
entryPoint, op, tileSizes, /*nativeVectorSizes=*/ArrayRef<int64_t>{},
- IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUVectorize, workgroupSize);
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUVectorize,
+ workgroupSize);
}
/// Propagate the configuration annotated in the incoming IR.
-static LogicalResult setUserConfig(FuncOp entryPointFn, Operation *computeOp,
- IREE::HAL::LoweringConfig config) {
- IREE::HAL::DispatchLoweringPassPipeline passPipeline =
- IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUVectorize;
- if (auto setPassPipeline = getLoweringPassPipeline(config)) {
- passPipeline = setPassPipeline.getValue();
+static LogicalResult setUserConfig(
+ FuncOp entryPointFn, Operation *computeOp,
+ IREE::Codegen::CompilationInfoAttr compilationInfo) {
+ if (auto translationInfo = getTranslationInfo(entryPointFn)) {
+ return computeOp->emitOpError(
+ "multiple ops within dispatch trying to set the translation "
+ "info");
}
- SmallVector<int64_t, 4> workgroupSize;
- if (auto workgroupSizeAttr = config.workgroupSize()) {
- workgroupSize = llvm::to_vector<4>(
- llvm::map_range(workgroupSizeAttr, [](Attribute intAttr) {
- return intAttr.cast<IntegerAttr>().getInt();
- }));
- }
- if (failed(setOpConfigAndEntryPointFnTranslation(
- entryPointFn, computeOp, config, passPipeline, workgroupSize))) {
- return failure();
- }
- // Reset the op configuration to drop the pass-pipeline and workgroup size
- // info. The op does not carry that information anymore.
- auto resetConfig = IREE::HAL::LoweringConfig::get(
- config.tileSizes(), config.nativeVectorSize(),
- /*passPipeline =*/nullptr,
- /*workgroupSize =*/nullptr, computeOp->getContext());
- setLoweringConfig(computeOp, resetConfig);
+
+ SmallVector<int64_t> workgroupSize = compilationInfo.getWorkgroupSizeVals();
+ setTranslationInfo(entryPointFn, compilationInfo.getTranslationInfo(),
+ workgroupSize);
+ setLoweringConfig(computeOp, compilationInfo.getLoweringConfig());
+ eraseCompilationInfo(computeOp);
return success();
}
static LogicalResult setRootConfig(FuncOp entryPointFn, Operation *computeOp) {
- if (IREE::HAL::LoweringConfig config = getLoweringConfig(computeOp)) {
+ if (IREE::Codegen::CompilationInfoAttr compilationInfo =
+ getCompilationInfo(computeOp)) {
// If the op already has a lowering config coming from the IR use this and
// bypass the heuristic.
- return setUserConfig(entryPointFn, computeOp, config);
+ return setUserConfig(entryPointFn, computeOp, compilationInfo);
}
if (auto linalgOp = dyn_cast<linalg::LinalgOp>(computeOp)) {
if (linalg::isaContractionOpInterface(linalgOp) &&
@@ -294,8 +285,9 @@
// anything. Without any compute ops, this shouldnt be using tile and
// distribute.
setTranslationInfo(
- funcOp, IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUDistribute,
- workgroupSize, workloadPerWorkgroup);
+ funcOp,
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUDistribute,
+ workloadPerWorkgroup, workgroupSize);
continue;
}
@@ -330,8 +322,9 @@
// anything. Without any compute ops, this shouldnt be using tile and
// distribute.
setTranslationInfo(
- funcOp, IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUDistribute,
- {1, 1, 1}, /*workloadPerWorkgroup=*/{});
+ funcOp,
+ IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUDistribute,
+ /*workloadPerWorkgroup=*/{}, {1, 1, 1});
continue;
}
if (failed(setRootConfig(funcOp, rootOperation))) continue;
@@ -342,7 +335,7 @@
// and distributed. The rest of the compilation must be structured to either
// use `TileAndFuse` or they are independent configurations that are
// determined based on the op.
- IREE::HAL::LoweringConfig config = getLoweringConfig(rootOperation);
+ IREE::Codegen::LoweringConfigAttr config = getLoweringConfig(rootOperation);
for (auto op : computeOps) {
if (op == rootOperation) continue;
setLoweringConfig(op, config);
diff --git a/iree/compiler/Codegen/LLVMGPU/KernelConfig.h b/iree/compiler/Codegen/LLVMGPU/KernelConfig.h
index 2717b90..d085245 100644
--- a/iree/compiler/Codegen/LLVMGPU/KernelConfig.h
+++ b/iree/compiler/Codegen/LLVMGPU/KernelConfig.h
@@ -7,7 +7,6 @@
#ifndef IREE_COMPILER_CODEGEN_LLVMGPU_KERNELCONFIG_H_
#define IREE_COMPILER_CODEGEN_LLVMGPU_KERNELCONFIG_H_
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/IR/BuiltinOps.h"
namespace mlir {
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp
index a07d472..5dfd252 100644
--- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUDistributeSharedMemoryCopy.cpp
@@ -7,6 +7,7 @@
#include <algorithm>
#include <numeric>
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
index 371d33c..5ef1239 100644
--- a/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
@@ -4,10 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/LLVMGPU/KernelConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h"
@@ -32,9 +33,9 @@
LLVMGPULowerExecutableTargetPass> {
public:
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<IREE::HAL::HALDialect, linalg::LinalgDialect,
- linalg_ext::LinalgExtDialect, vector::VectorDialect,
- gpu::GPUDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, IREE::HAL::HALDialect,
+ linalg::LinalgDialect, linalg_ext::LinalgExtDialect,
+ vector::VectorDialect, gpu::GPUDialect>();
}
LLVMGPULowerExecutableTargetPass() = default;
@@ -72,16 +73,15 @@
// is fine.
llvm::StringMap<IREE::HAL::ExecutableEntryPointOp> entryPoints =
getAllEntryPoints(moduleOp);
- Optional<IREE::HAL::DispatchLoweringPassPipeline> passPipeline;
+ Optional<IREE::Codegen::DispatchLoweringPassPipeline> passPipeline;
for (auto &it : entryPoints) {
auto entryPointOp = it.second;
- if (IREE::HAL::TranslationInfo translationInfo =
+ if (IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(entryPointOp)) {
- Optional<IREE::HAL::DispatchLoweringPassPipeline> currPipeline =
- getLoweringPassPipeline(translationInfo);
- if (!currPipeline) continue;
+ IREE::Codegen::DispatchLoweringPassPipeline currPipeline =
+ translationInfo.getDispatchLoweringPassPipeline();
if (passPipeline) {
- if (currPipeline.getValue() != passPipeline.getValue()) {
+ if (currPipeline != passPipeline.getValue()) {
moduleOp.emitError(
"unhandled compilation of entry point function with different "
"pass pipelines within a module");
@@ -98,13 +98,13 @@
if (!testLoweringConfiguration && passPipeline.hasValue()) {
OpPassManager &nestedModulePM = executableLoweringPipeline.nest<ModuleOp>();
switch (*passPipeline) {
- case IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUDistribute:
+ case IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUDistribute:
addGPUSimpleDistributePassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUVectorize:
+ case IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUVectorize:
addGPUVectorizationPassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::LLVMGPUMatmulSimt:
+ case IREE::Codegen::DispatchLoweringPassPipeline::LLVMGPUMatmulSimt:
addGPUMatmulSimtPassPipeline(nestedModulePM);
break;
default:
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
index 383b74d..df0e622 100644
--- a/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
@@ -4,6 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
@@ -77,11 +78,8 @@
auto translationInfo = getTranslationInfo(entryPointOp);
if (!translationInfo) return SmallVector<int64_t>();
- ArrayAttr workloadPerWorkgroupAttr = translationInfo.workloadPerWorkgroup();
- if (!workloadPerWorkgroupAttr) return SmallVector<int64_t>();
- auto workloadPerWorkgroup = llvm::to_vector<4>(llvm::map_range(
- workloadPerWorkgroupAttr,
- [](Attribute attr) { return attr.cast<IntegerAttr>().getInt(); }));
+ SmallVector<int64_t> workloadPerWorkgroup =
+ translationInfo.getWorkloadPerWorkgroupVals();
if (workloadSize.size() != workloadPerWorkgroup.size())
return SmallVector<int64_t>();
SmallVector<int64_t> numWorkgroups;
diff --git a/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
index 4409c33..2a701f8 100644
--- a/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
@@ -4,14 +4,13 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/LLVMGPU/KernelConfig.h"
#include "iree/compiler/Codegen/LLVMGPU/LLVMGPUUtils.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
#include "iree/compiler/Dialect/LinalgExt/Transforms/Transforms.h"
#include "iree/compiler/Dialect/Util/IR/UtilOps.h"
@@ -38,19 +37,14 @@
auto tileSizesFn = [&](OpBuilder &builder,
Operation *op) -> SmallVector<Value, 4> {
SmallVector<unsigned> partitionedLoops = getPartitionedLoops(op);
- SmallVector<int64_t, 4> tileSizes = getTileSizes(op, 0);
- Location loc = op->getLoc();
- auto tileSizesVal =
- llvm::to_vector<4>(llvm::map_range(tileSizes, [&](int64_t v) -> Value {
- return builder.create<arith::ConstantIndexOp>(loc, v);
- }));
- auto zero = builder.create<arith::ConstantIndexOp>(loc, 0);
+ SmallVector<Value, 4> tileSizes = getTileSizes(builder, op, 0);
+ auto zero = builder.create<arith::ConstantIndexOp>(op->getLoc(), 0);
for (unsigned depth : partitionedLoops) {
- if (depth < tileSizesVal.size()) {
- tileSizesVal[depth] = zero;
+ if (depth < tileSizes.size()) {
+ tileSizes[depth] = zero;
}
}
- return tileSizesVal;
+ return tileSizes;
};
auto tilingOptions = linalg::LinalgTilingOptions()
@@ -69,8 +63,8 @@
/// Patterns for thread level tiling.
static void populateTilingToInvocationPatterns(
MLIRContext *context, OwningRewritePatternList &patterns,
- SmallVector<int64_t, 4> &workgroupSize,
- SmallVector<int64_t, 4> &workloadPerWorkgroup) {
+ SmallVectorImpl<int64_t> &workgroupSize,
+ SmallVectorImpl<int64_t> &workloadPerWorkgroup) {
linalg::TileSizeComputationFunction getInnerTileSizeFn =
[&](OpBuilder &builder, Operation *operation) {
SmallVector<Value, 4> tileSizesVal;
@@ -95,7 +89,7 @@
return tileSizesVal;
};
- auto getThreadProcInfoFn = [workgroupSize](
+ auto getThreadProcInfoFn = [&workgroupSize](
OpBuilder &builder, Location loc,
ArrayRef<Range> parallelLoopRanges) {
return getGPUThreadIdsAndCounts(builder, loc, parallelLoopRanges.size(),
@@ -240,11 +234,8 @@
auto workgroupSize = llvm::to_vector<4>(llvm::map_range(
getEntryPoint(funcOp).workgroup_size().getValue(),
[&](Attribute attr) { return attr.cast<IntegerAttr>().getInt(); }));
- auto workloadPerWorkgroup = llvm::to_vector<4>(llvm::map_range(
- getTranslationInfo(getEntryPoint(funcOp))
- .workloadPerWorkgroup()
- .getValue(),
- [&](Attribute attr) { return attr.cast<IntegerAttr>().getInt(); }));
+ auto workloadPerWorkgroup =
+ getTranslationInfo(getEntryPoint(funcOp)).getWorkloadPerWorkgroupVals();
int64_t flatWorkgroupSize =
workgroupSize[0] * workgroupSize[1] * workgroupSize[2];
diff --git a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
index 2b7e6da..e7aaf8d 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -1,6 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-llvmgpu-tile-and-distribute))))' %s | IreeFileCheck %s
-#config = {tileSizes = [[2, 256, 4]]}
+#config = #iree_codegen.lowering.config<tile_sizes = [[2, 256, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"LLVMGPUMatmulSimt", workload_per_wg = [256, 2]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#map0 = affine_map<()[s0] -> (s0 * 2)>
#map1 = affine_map<()[s0] -> (s0 * 256)>
@@ -12,7 +13,7 @@
hal.executable.entry_point @dot_dispatch_0 attributes {
interface = @legacy_io,
ordinal = 0 : index,
- translation.info = {passPipeline = "LLVMGPUMatmulSimt" : i32, workloadPerWorkgroup = [256, 2]},
+ translation.info = #translation,
workgroup_size = [64 : index, 1 : index, 1 : index]}
builtin.module {
builtin.func @dot_dispatch_0() {
@@ -86,14 +87,15 @@
// -----
-#config = {tileSizes = [[]]}
+#config = #iree_codegen.lowering.config<tile_sizes = [[]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"LLVMGPUVectorize", workload_per_wg = []>
// Pure reducion case, skip tiling.
hal.executable @reduction_dispatch {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @predict_dispatch_153 attributes {
interface = @io,
ordinal = 0 : index,
- translation.info = {passPipeline = "LLVMGPUVectorize" : i32},
+ translation.info = #translation,
workgroup_size = [1: index, 1: index, 1: index]}
builtin.module {
builtin.func @predict_dispatch_153() {
@@ -120,7 +122,7 @@
}
}
}
-// CHECK: #[[CONFIG:.+]] = {tileSizes = {{\[}}[]{{\]}}}
+// CHECK: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[]{{\]}}, native_vector_size = []>
// CHECK: hal.executable public @reduction_dispatch
// CHECK: linalg.fill
// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
index b67847a..8690542 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/gpu_set_num_workgroups.mlir
@@ -33,11 +33,11 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[256]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[256]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 256)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUVectorize", workload_per_wg = [256]>
// CHECK: hal.executable.entry_point public @add_dispatch_0
-// CHECK-SAME: passPipeline = "LLVMGPUVectorize"
-// CHECK-SAME: workloadPerWorkgroup = [256]
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [64 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index,
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -92,12 +92,12 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[4, 2, 4]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[4, 2, 4]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 2)>
// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUMatmulSimt", workload_per_wg = [2, 4]>
// CHECK: hal.executable.entry_point public @dot_dispatch_1
-// CHECK-SAME: passPipeline = "LLVMGPUMatmulSimt"
-// CHECK-SAME: workloadPerWorkgroup = [2, 4]
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [2 : index, 4 : index, 1 : index]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index, %[[ARG1:[a-zA-Z0-9]+]]: index,
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -143,18 +143,18 @@
}
}
-// CHECK-DAG: #[[CONFIG0:.+]] = {passPipeline = "LLVMGPUDistribute"}
-// CHECK-DAG: #[[CONFIG1:.+]] = {tileSizes = {{\[}}[]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = []>
// CHECK: hal.executable.entry_point public @predict_dispatch_153
-// CHECK-SAME: translation.info = #[[CONFIG0]]
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^bb0(%[[ARG0:[a-zA-Z0-9]+]]: index,
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
// CHECK: hal.return %[[C1]], %[[C1]], %[[C1]]
// CHECK: linalg.fill
-// CHECK-SAME: lowering.config = #[[CONFIG1]]
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// CHECK: linalg.generic
-// CHECK-SAME: lowering.config = #[[CONFIG1]]
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -199,9 +199,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = [128, 1]>
// CHECK: hal.executable.entry_point public @tensor_insert_slice
-// CHECK-SAME: translation.info = {passPipeline = "LLVMGPUDistribute", workloadPerWorkgroup = [128, 1]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -246,10 +247,11 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[1, 256]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 256]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 256)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUVectorize", workload_per_wg = [256, 1]>
// CHECK: hal.executable.entry_point public @tensor_insert_slice
-// CHECK-SAME: translation.info = {passPipeline = "LLVMGPUVectorize", workloadPerWorkgroup = [256, 1]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -286,11 +288,11 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[4]]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[4]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = [4]>
// CHECK: hal.executable.entry_point public @static_1d_fft_stage2
-// CHECK-SAME: translation.info = {passPipeline = "LLVMGPUDistribute"
-// CHECK-SAME: workloadPerWorkgroup = [4]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %{{.+}}: index, %{{.+}}: index):
// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
@@ -351,11 +353,11 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[1, 1, 8]]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 1, 8]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = [8, 1, 1]>
// CHECK: hal.executable.entry_point public @static_3d_fft_stage3
-// CHECK-SAME: translation.info = {passPipeline = "LLVMGPUDistribute"
-// CHECK-SAME: workloadPerWorkgroup = [8, 1, 1]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %[[ARG1:.+]]: index, %[[ARG2:.+]]: index):
// CHECK-NEXT: %[[T:.+]] = affine.apply #[[MAP0]]()[%[[ARG0]]]
@@ -367,6 +369,10 @@
// -----
+#compilation = #iree_codegen.compilation.info<
+ #iree_codegen.lowering.config<tile_sizes = [[32, 256, 64]], native_vector_size = []>,
+ #iree_codegen.translation.info<"LLVMGPUMatmulSimt", workload_per_wg = [256, 32]>,
+ workgroup_size = [16, 8, 1]>
hal.executable @user_config {
hal.executable.variant public @cuda_nvptx_fb, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point public @_lowering_config_test_dispatch_1 attributes {interface = @io, ordinal = 0 : index}
@@ -401,7 +407,7 @@
%14 = affine.min affine_map<(d0)[s0] -> (-d0 + 1024, s0)>(%arg1)[%workgroup_size_x]
%15 = linalg.init_tensor [%13, %14] : tensor<?x?xf32>
%16 = linalg.fill(%cst, %15) : f32, tensor<?x?xf32> -> tensor<?x?xf32>
- %17 = linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = {passPipeline = "LLVMGPUMatmulSimt", tileSizes = [[32, 256, 64]], workgroupSize = [16, 8, 1]}} ins(%8, %10 : tensor<?x256xf32>, tensor<256x?xf32>) outs(%16 : tensor<?x?xf32>) -> tensor<?x?xf32>
+ %17 = linalg.matmul {__internal_linalg_transform__ = "workgroup", compilation.info = #compilation} ins(%8, %10 : tensor<?x256xf32>, tensor<256x?xf32>) outs(%16 : tensor<?x?xf32>) -> tensor<?x?xf32>
flow.dispatch.tensor.store %17, %2, offsets = [%arg0, %arg1], sizes = [%11, %12], strides = [1, 1] : tensor<?x?xf32> -> !flow.dispatch.tensor<writeonly:128x1024xf32>
}
}
@@ -416,10 +422,10 @@
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {{{.*}}tileSizes = {{\[}}[32, 256, 64]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[32, 256, 64]{{\]}}
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"LLVMGPUMatmulSimt", workload_per_wg = [256, 32]>
// CHECK: hal.executable.entry_point public @_lowering_config_test_dispatch_1
-// CHECK-SAME: passPipeline = "LLVMGPUMatmulSimt"
-// CHECK-SAME: workloadPerWorkgroup = [256, 32]
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [16 : index, 8 : index, 1 : index]
// CHECK: func @_lowering_config_test_dispatch_1
// CHECK: linalg.fill
diff --git a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
index 46b6931..bed9ee3 100644
--- a/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
@@ -42,12 +42,13 @@
// -----
// CHECK-LABEL: func @workgroup_tile_loop()
+#translation = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = [32]>
hal.executable private @workgroup_tile_loop {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @workgroup_tile_loop attributes {
interface = @io,
ordinal = 0 : index,
- translation.info = {passPipeline = "LLVMGPUDistribute", workloadPerWorkgroup = [32]}
+ translation.info = #translation
}
builtin.module {
builtin.func @workgroup_tile_loop() {
@@ -71,12 +72,13 @@
// -----
// CHECK-LABEL: func @workgroup_tile_loop_negative()
+#translation = #iree_codegen.translation.info<"LLVMGPUDistribute", workload_per_wg = [16]>
hal.executable private @workgroup_tile_loop_negative {
hal.executable.variant @cuda, target = #hal.executable.target<"cuda", "cuda-nvptx-fb"> {
hal.executable.entry_point @workgroup_tile_loop_negative attributes {
interface = @io,
ordinal = 0 : index,
- translation.info = {passPipeline = "LLVMGPUDistribute", workloadPerWorkgroup = [16]}
+ translation.info = #translation
}
builtin.module {
builtin.func @workgroup_tile_loop_negative() {
diff --git a/iree/compiler/Codegen/Passes.h b/iree/compiler/Codegen/Passes.h
index b513b99..84b4fb5 100644
--- a/iree/compiler/Codegen/Passes.h
+++ b/iree/compiler/Codegen/Passes.h
@@ -10,7 +10,6 @@
#include <memory>
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Pass/PassOptions.h"
diff --git a/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp b/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
index 38240a8..877cac7 100644
--- a/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/AdrenoConfig.cpp
@@ -13,7 +13,9 @@
#include <array>
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
+#include "llvm/ADT/TypeSwitch.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
+#include "mlir/IR/BuiltinOps.h"
namespace mlir {
namespace iree_compiler {
diff --git a/iree/compiler/Codegen/SPIRV/BUILD b/iree/compiler/Codegen/SPIRV/BUILD
index a63331f..bc50e24 100644
--- a/iree/compiler/Codegen/SPIRV/BUILD
+++ b/iree/compiler/Codegen/SPIRV/BUILD
@@ -39,6 +39,7 @@
deps = [
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/Transforms",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Codegen/SPIRV/CMakeLists.txt b/iree/compiler/Codegen/SPIRV/CMakeLists.txt
index f6bbb28..9929b63 100644
--- a/iree/compiler/Codegen/SPIRV/CMakeLists.txt
+++ b/iree/compiler/Codegen/SPIRV/CMakeLists.txt
@@ -71,6 +71,7 @@
MLIRVectorInterfaces
MLIRVectorToSPIRV
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::Transforms
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
index a956c04..4f9253a 100644
--- a/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
@@ -15,11 +15,11 @@
#include <tuple>
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/SPIRV/Utils.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/Util/IR/UtilOps.h"
#include "llvm/ADT/DenseMapInfo.h"
diff --git a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
index 612947b..e4439d6 100644
--- a/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/KernelConfig.cpp
@@ -6,11 +6,10 @@
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/SPIRV/Utils.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/MathExtras.h"
@@ -18,6 +17,7 @@
#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h"
+#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/Matchers.h"
#define DEBUG_TYPE "iree-spirv-kernel-config"
@@ -34,12 +34,11 @@
// TODO(ravishankarm): Remove this when that pipeline is deprecated.
static LogicalResult setTranslationUsingDistributeToGlobalId(
FuncOp funcOp, ArrayRef<int64_t> workgroupSize) {
- auto entryPointOp = getEntryPoint(funcOp);
- MLIRContext *context = entryPointOp.getContext();
- auto translationInfo = buildTranslationInfo(
- IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistributeToGlobalID,
- /*workloadPerWorkgroup =*/{}, context);
- setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
+ setTranslationInfo(
+ funcOp,
+ IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistributeToGlobalID,
+ /*workloadPerWorkgroup=*/{}, workgroupSize);
+ MLIRContext *context = funcOp.getContext();
OpBuilder builder(context);
int64_t workgroupSizeX = workgroupSize[0];
auto numWorkgroupsFn = [workgroupSizeX](OpBuilder &b, Location loc,
@@ -107,9 +106,9 @@
int64_t residualThreads = subgroupSize;
int64_t residualTilingFactor = bestTilingFactor;
- SmallVector<int64_t, 3> workgroupSize(3, 1); // (X, Y, Z)
- SmallVector<int64_t, 4> workgroupTileSizes(4, 0); // (N, OH, OW, OC)
- SmallVector<int64_t, 4> invocationTileSizes(4, 0); // (N, OH, OW, OC)
+ SmallVector<int64_t, 3> workgroupSize(3, 1); // (X, Y, Z)
+ SmallVector<int64_t> workgroupTileSizes(4, 0); // (N, OH, OW, OC)
+ SmallVector<int64_t> invocationTileSizes(4, 0); // (N, OH, OW, OC)
// Deduce the configuration for the OC dimension.
for (int64_t x = residualThreads; x >= 2; x >>= 1) {
@@ -181,7 +180,7 @@
}
}
- auto pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorize;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize;
TileSizesListType tileSizes;
tileSizes.push_back(workgroupTileSizes);
tileSizes.push_back(invocationTileSizes);
@@ -244,10 +243,10 @@
int64_t residualThreads = bestX * bestY;
int64_t residualTilingFactor = (bestThreadM + bestThreadK) * bestThreadN;
- SmallVector<int64_t, 3> workgroupSize(3, 1); // (X, Y, Z)
- SmallVector<int64_t, 4> workgroupTileSizes(2 + isBM, 0); // (B, M, N)
- SmallVector<int64_t, 4> invocationTileSizes(2 + isBM, 0); // (B, M, N)
- SmallVector<int64_t, 4> reductionTileSizes(3 + isBM, 0); // (B, M, N, K)
+ SmallVector<int64_t, 3> workgroupSize(3, 1); // (X, Y, Z)
+ SmallVector<int64_t> workgroupTileSizes(2 + isBM, 0); // (B, M, N, K)
+ SmallVector<int64_t> invocationTileSizes(2 + isBM, 0); // (B, M, N, K)
+ SmallVector<int64_t> reductionTileSizes(3 + isBM, 0); // (B, M, N, K)
if (isBM) workgroupTileSizes[0] = invocationTileSizes[0] = 1;
@@ -302,7 +301,7 @@
}
if (reductionTileSizes[2 + isBM] == 0) return success();
- auto pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorize;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize;
TileSizesListType tileSizes;
tileSizes.push_back(workgroupTileSizes);
tileSizes.push_back(invocationTileSizes);
@@ -321,13 +320,13 @@
static LogicalResult setOpConfig(spirv::ResourceLimitsAttr limits,
linalg_ext::FftOp op) {
const int64_t subgroupSize = limits.subgroup_size().getValue().getSExtValue();
- auto pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistribute;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistribute;
std::array<int64_t, 3> workgroupSize = {subgroupSize, 1, 1};
auto partitionedLoops = getPartitionedLoops(op);
unsigned loopDepth = partitionedLoops.back() + 1;
- SmallVector<int64_t, 4> workgroupTileSize(loopDepth, 0);
+ SmallVector<int64_t> workgroupTileSize(loopDepth, 0);
// Tiling along partitioned loops with size 1.
for (int64_t loopIndex : partitionedLoops) {
@@ -357,7 +356,7 @@
Operation *op) {
auto partitionedLoops = getPartitionedLoops(op);
if (partitionedLoops.empty()) {
- auto pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorize;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize;
std::array<int64_t, 3> workgroupSize = {1, 1, 1};
auto funcOp = op->getParentOfType<FuncOp>();
return setOpConfigAndEntryPointFnTranslation(funcOp, op, {}, {}, pipeline,
@@ -367,7 +366,7 @@
const int64_t subgroupSize = limits.subgroup_size().getValue().getSExtValue();
int64_t numElementsPerWorkgroup = subgroupSize;
int64_t numElementsPerThread = 1;
- auto pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistribute;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistribute;
// Returns true if the given `operand` has 32-bit element type.
auto has32BitElementType = [](Value operand) {
@@ -415,15 +414,15 @@
if (vectorize) {
numElementsPerThread = numElementsPerWorkgroup / subgroupSize;
- pipeline = IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorize;
+ pipeline = IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize;
}
}
std::array<int64_t, 3> workgroupSize = {subgroupSize, 1, 1};
unsigned loopDepth = partitionedLoops.back() + 1;
- SmallVector<int64_t, 4> workgroupTileSize(loopDepth, 0);
- SmallVector<int64_t, 4> threadTileSize(loopDepth, 0);
+ SmallVector<int64_t> workgroupTileSize(loopDepth, 0);
+ SmallVector<int64_t> threadTileSize(loopDepth, 0);
// Tiling along partitioned loops with size 1.
for (int64_t loopIndex : partitionedLoops) {
@@ -597,8 +596,9 @@
SmallVector<int64_t> workloadPerWorkgroup(tiledLoops.size(), 1);
workloadPerWorkgroup.front() = subgroupSize * 4;
setTranslationInfo(
- funcOp, IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistribute,
- workgroupSize, workloadPerWorkgroup);
+ funcOp,
+ IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistribute,
+ workloadPerWorkgroup, workgroupSize);
return success();
}
return funcOp.emitError("contains no root Linalg operation");
@@ -610,7 +610,7 @@
// and distributed. The rest of the compilation must be structured to either
// use `TileAndFuse` or they are independent configurations that are
// determined based on the op.
- IREE::HAL::LoweringConfig config = getLoweringConfig(rootOperation);
+ IREE::Codegen::LoweringConfigAttr config = getLoweringConfig(rootOperation);
for (auto op : computeOps) {
if (op == rootOperation) continue;
setLoweringConfig(op, config);
diff --git a/iree/compiler/Codegen/SPIRV/KernelConfig.h b/iree/compiler/Codegen/SPIRV/KernelConfig.h
index c0d4f31..8185894 100644
--- a/iree/compiler/Codegen/SPIRV/KernelConfig.h
+++ b/iree/compiler/Codegen/SPIRV/KernelConfig.h
@@ -17,9 +17,9 @@
#include <array>
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/Dialect/Linalg/IR/LinalgInterfaces.h"
#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h"
+#include "mlir/IR/BuiltinOps.h"
namespace mlir {
namespace iree_compiler {
diff --git a/iree/compiler/Codegen/SPIRV/MaliConfig.cpp b/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
index 1c19955..9577d43 100644
--- a/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/MaliConfig.cpp
@@ -13,7 +13,9 @@
#include <array>
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
+#include "llvm/ADT/TypeSwitch.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
+#include "mlir/IR/BuiltinOps.h"
namespace mlir {
namespace iree_compiler {
diff --git a/iree/compiler/Codegen/SPIRV/NVIDIAConfig.cpp b/iree/compiler/Codegen/SPIRV/NVIDIAConfig.cpp
index 3d07b75..dc1e447 100644
--- a/iree/compiler/Codegen/SPIRV/NVIDIAConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/NVIDIAConfig.cpp
@@ -10,10 +10,12 @@
//
//===----------------------------------------------------------------------===//
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
+#include "mlir/IR/BuiltinOps.h"
#define DEBUG_TYPE "iree-spirv-nvidia-config"
@@ -80,8 +82,8 @@
getElementType(init), lhsShape[0], rhsShape[1], lhsShape[1]);
if (!coopMatSize) return success();
- auto pipeline =
- IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorizeToCooperativeOps;
+ auto pipeline = IREE::Codegen::DispatchLoweringPassPipeline::
+ SPIRVVectorizeToCooperativeOps;
// For now only support one subgroup per workgroup because in the above
// configuration deduction step we only consider whether the input workload is
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTargetPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTargetPass.cpp
index fbc151e..f559ee4 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTargetPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTargetPass.cpp
@@ -4,10 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h"
@@ -36,7 +37,8 @@
SPIRVLowerExecutableTargetPass(const SPIRVLowerExecutableTargetPass &pass) {}
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<AffineDialect, gpu::GPUDialect, IREE::HAL::HALDialect,
+ registry.insert<IREE::Codegen::IREECodegenDialect, AffineDialect,
+ gpu::GPUDialect, IREE::HAL::HALDialect,
linalg::LinalgDialect, linalg_ext::LinalgExtDialect,
memref::MemRefDialect, scf::SCFDialect, ShapeDialect,
spirv::SPIRVDialect, vector::VectorDialect>();
@@ -71,16 +73,15 @@
// is fine.
llvm::StringMap<IREE::HAL::ExecutableEntryPointOp> entryPoints =
getAllEntryPoints(moduleOp);
- Optional<IREE::HAL::DispatchLoweringPassPipeline> passPipeline;
+ Optional<IREE::Codegen::DispatchLoweringPassPipeline> passPipeline;
for (auto &it : entryPoints) {
auto entryPointOp = it.second;
- if (IREE::HAL::TranslationInfo translationInfo =
+ if (IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(entryPointOp)) {
- Optional<IREE::HAL::DispatchLoweringPassPipeline> currPipeline =
- getLoweringPassPipeline(translationInfo);
- if (!currPipeline) continue;
+ IREE::Codegen::DispatchLoweringPassPipeline currPipeline =
+ translationInfo.getDispatchLoweringPassPipeline();
if (passPipeline) {
- if (currPipeline.getValue() != passPipeline.getValue()) {
+ if (currPipeline != passPipeline.getValue()) {
moduleOp.emitError(
"unhandled compilation of entry point function with different "
"pass pipelines within a module");
@@ -97,16 +98,17 @@
if (!testLoweringConfiguration && passPipeline.hasValue()) {
OpPassManager &nestedModulePM = executableLoweringPipeline.nest<ModuleOp>();
switch (*passPipeline) {
- case IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistribute:
+ case IREE::Codegen::DispatchLoweringPassPipeline::SPIRVDistribute:
addSPIRVTileAndDistributePassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::SPIRVDistributeToGlobalID:
+ case IREE::Codegen::DispatchLoweringPassPipeline::
+ SPIRVDistributeToGlobalID:
addSPIRVDistributeToGlobalIDPassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::SPIRVVectorize:
+ case IREE::Codegen::DispatchLoweringPassPipeline::SPIRVVectorize:
addSPIRVTileAndVectorizePassPipeline(nestedModulePM);
break;
- case IREE::HAL::DispatchLoweringPassPipeline::
+ case IREE::Codegen::DispatchLoweringPassPipeline::
SPIRVVectorizeToCooperativeOps:
addSPIRVTileAndVectorizeToCooperativeOpsPassPipeline(nestedModulePM);
break;
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVRemoveOneTripTiledLoops.cpp b/iree/compiler/Codegen/SPIRV/SPIRVRemoveOneTripTiledLoops.cpp
index 079f895..21b7e62 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVRemoveOneTripTiledLoops.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVRemoveOneTripTiledLoops.cpp
@@ -4,6 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
@@ -112,11 +113,7 @@
auto translationInfo = getTranslationInfo(entryPointOp);
if (!translationInfo) return;
- ArrayAttr workloadPerWorkgroupAttr = translationInfo.workloadPerWorkgroup();
- if (!workloadPerWorkgroupAttr) return;
- auto workloadPerWorkgroup = llvm::to_vector<4>(llvm::map_range(
- workloadPerWorkgroupAttr,
- [](Attribute attr) { return attr.cast<IntegerAttr>().getInt(); }));
+ auto workloadPerWorkgroup = translationInfo.getWorkloadPerWorkgroupVals();
MLIRContext *context = &getContext();
removeOneTripTiledLoops(context, funcOp, cast<linalg::LinalgOp>(rootOp[0]),
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp b/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp
index 07b3a62..15dc4fe 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVTileAndDistribute.cpp
@@ -11,12 +11,12 @@
//
//===----------------------------------------------------------------------===//
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/SPIRV/Utils.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
-#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
#include "iree/compiler/Dialect/LinalgExt/Transforms/Transforms.h"
#include "llvm/ADT/STLExtras.h"
@@ -79,11 +79,7 @@
RewritePatternSet &patterns) {
linalg::TileSizeComputationFunction getInnerTileSizeFn =
[&](OpBuilder &builder, Operation *op) {
- SmallVector<int64_t> tileSizes = getTileSizes(op, 1);
- return llvm::to_vector<4>(
- llvm::map_range(tileSizes, [&](int64_t v) -> Value {
- return builder.create<arith::ConstantIndexOp>(op->getLoc(), v);
- }));
+ return getTileSizes(builder, op, 1);
};
auto getThreadProcInfoFn = [](OpBuilder &builder, Location loc,
@@ -161,11 +157,7 @@
MLIRContext *context, RewritePatternSet &patterns,
linalg::LinalgTransformationFilter marker) {
auto getTileSizeFn = [&](OpBuilder &builder, Operation *op) {
- SmallVector<int64_t> tileSizes = getTileSizes(op, 2);
- return llvm::to_vector<4>(
- llvm::map_range(tileSizes, [&](int64_t v) -> Value {
- return builder.create<arith::ConstantIndexOp>(op->getLoc(), v);
- }));
+ return getTileSizes(builder, op, 2);
};
auto tilingOptions = linalg::LinalgTilingOptions()
diff --git a/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorizeToCooperativeOps.cpp b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorizeToCooperativeOps.cpp
index dc611de..3a3c10b 100644
--- a/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorizeToCooperativeOps.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorizeToCooperativeOps.cpp
@@ -13,6 +13,7 @@
#include <algorithm>
+#include "iree/compiler/Codegen/Dialect/LoweringConfig.h"
#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Codegen/SPIRV/KernelConfig.h"
@@ -20,7 +21,6 @@
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/GPU/GPUDialect.h"
diff --git a/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir b/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir
index aadabbc..c178156 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_adreno_conv.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
// Conv - large OC - distribute to only one workgroup dimension.
@@ -74,18 +74,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_112x112x512
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [256, 8, 1]}
-// CHECK-SAME: workgroup_size = [64 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index
-// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index
-// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
-// CHECK-NEXT: hal.return %[[C2]], %[[C14]], %[[C112]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 1, 8, 256], [0, 1, 8, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [256, 8, 1]>
+// CHECK: hal.executable.entry_point public @conv_112x112x512
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [64 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index
+// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index
+// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
+// CHECK-NEXT: hal.return %[[C2]], %[[C14]], %[[C112]]
-// CHECK: func @conv_112x112x512()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 1, 8, 256], [0, 1, 8, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_112x112x512()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -163,18 +165,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_112x112x32
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 16, 4]}
-// CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
-// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C7]], %[[C28]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 4, 16, 32], [0, 4, 2, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 16, 4]>
+// CHECK: hal.executable.entry_point public @conv_112x112x32
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
+// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C7]], %[[C28]]
-// CHECK: func @conv_112x112x32()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 4, 16, 32], [0, 4, 2, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_112x112x32()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -251,17 +255,19 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_16x16x16
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 8, 8]}
-// CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C2]], %[[C2]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 8, 8, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 8, 8]>
+// CHECK: hal.executable.entry_point public @conv_16x16x16
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[C2:.+]] = arith.constant 2 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C2]], %[[C2]]
-// CHECK: func @conv_16x16x16()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 8, 8, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_16x16x16()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -340,17 +346,19 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @dwconv_28x28x144
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]}
-// CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index
-// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
-// CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 4, 4, 16], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
+// CHECK: hal.executable.entry_point public @dwconv_28x28x144
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 4 : index, 4 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index
+// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
+// CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]]
-// CHECK: func @dwconv_28x28x144()
-// CHECK: linalg.depthwise_conv2D_nhw
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 4, 4, 16], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]]}
+// CHECK: func @dwconv_28x28x144()
+// CHECK: linalg.depthwise_conv2D_nhw
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -428,14 +436,15 @@
}
}
}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 4, 4, 8], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 4, 4]>
+// CHECK: hal.executable.entry_point public @dwconv_4x4x8
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 4 : index, 4 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]]
-// CHECK-LABEL: hal.executable.entry_point public @dwconv_4x4x8
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 4, 4]}
-// CHECK-SAME: workgroup_size = [2 : index, 4 : index, 4 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]]
-
-// CHECK: func @dwconv_4x4x8()
-// CHECK: linalg.depthwise_conv2D_nhw
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 4, 4, 8], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]]}
+// CHECK: func @dwconv_4x4x8()
+// CHECK: linalg.depthwise_conv2D_nhw
+// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir
index 8412a28..251f21f 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_adreno_matmul.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
// Large matmul that can match the best tiling scheme.
@@ -62,18 +62,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_1024x2048x512
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [128, 32]}
-// CHECK-SAME: workgroup_size = [32 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 128)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[32, 128], [16, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [128, 32]>
+// CHECK: hal.executable.entry_point public @matmul_1024x2048x512
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [32 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_1024x2048x512()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[32, 128], [16, 4], [0, 0, 4]]}
+// CHECK: func @matmul_1024x2048x512()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -139,18 +143,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_3136x24x96
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 448]}
-// CHECK-SAME: workgroup_size = [2 : index, 32 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 448)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[448, 8], [14, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 448)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 448]>
+// CHECK: hal.executable.entry_point public @matmul_3136x24x96
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 32 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_3136x24x96()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[448, 8], [14, 4], [0, 0, 4]]}
+// CHECK: func @matmul_3136x24x96()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -216,18 +224,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_196x64x192
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 28]}
-// CHECK-SAME: workgroup_size = [16 : index, 4 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 64)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 28)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[28, 64], [7, 4], [0, 0, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 28)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 28]>
+// CHECK: hal.executable.entry_point public @matmul_196x64x192
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [16 : index, 4 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_196x64x192()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[28, 64], [7, 4], [0, 0, 8]]}
+// CHECK: func @matmul_196x64x192()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -288,18 +300,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_12544x96x16
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 128]}
-// CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 128)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[128, 32], [16, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 128]>
+// CHECK: hal.executable.entry_point public @matmul_12544x96x16
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 8 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_12544x96x16()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[128, 32], [16, 4], [0, 0, 4]]}
+// CHECK: func @matmul_12544x96x16()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -365,18 +381,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_49x160x576
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 7]}
-// CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 7)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[7, 32], [7, 4], [0, 0, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 7)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 7]>
+// CHECK: hal.executable.entry_point public @matmul_49x160x576
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_49x160x576()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[7, 32], [7, 4], [0, 0, 8]]}
+// CHECK: func @matmul_49x160x576()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -453,17 +473,21 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @batch_matmul_4x384x384
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [128, 32, 1]}
-// CHECK-SAME: workgroup_size = [32 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 128)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 32, 128], [1, 16, 4], [0, 0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 128)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [128, 32, 1]>
+// CHECK: hal.executable.entry_point public @batch_matmul_4x384x384
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [32 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
-// CHECK: func @batch_matmul_4x384x384()
-// CHECK: linalg.batch_matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 32, 128], [1, 16, 4], [0, 0, 0, 4]]}
+// CHECK: func @batch_matmul_4x384x384()
+// CHECK: linalg.batch_matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -540,14 +564,17 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @batch_matmul_4x8x8
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 8, 1]}
-// CHECK-SAME: workgroup_size = [2 : index, 8 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 8, 8], [1, 1, 4], [0, 0, 0, 16]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 8, 1]>
+// CHECK: hal.executable.entry_point public @batch_matmul_4x8x8
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 8 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
-// CHECK: func @batch_matmul_4x8x8()
-// CHECK: linalg.batch_matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 8, 8], [1, 1, 4], [0, 0, 0, 16]]}
+// CHECK: func @batch_matmul_4x8x8()
+// CHECK: linalg.batch_matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
\ No newline at end of file
diff --git a/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir
index 0b584bf..10dc64f 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_default_matmul.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
// Odd K that forbids vectorization.
@@ -74,16 +74,19 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @batch_matmul_1x3x32
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [4, 1, 1]}
-// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[X]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[Z]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 1, 4], [1, 1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [4, 1, 1]>
+// CHECK: hal.executable.entry_point public @batch_matmul_1x3x32
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[Z]]
-// CHECK: func @batch_matmul_1x3x32()
-// CHECK: linalg.batch_matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 1, 4], [1, 1, 1]]}
+// CHECK: func @batch_matmul_1x3x32()
+// CHECK: linalg.batch_matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -149,14 +152,17 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_64x16
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [4, 1]}
-// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[X]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 4], [1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [4, 1]>
+// CHECK: hal.executable.entry_point public @matmul_64x16
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]]
-// CHECK: func @matmul_64x16()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 4], [1, 1]]}
+// CHECK: func @matmul_64x16()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir
index ae68820..57ea3d8 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_linalg_ext_ops.mlir
@@ -1,5 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
-
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
hal.executable private @static_1d_sort {
hal.interface @io {
hal.interface.binding @s0b0_rw_external, set=0, binding=0, type="StorageBuffer", access="Read|Write"
@@ -34,8 +33,10 @@
// Check that the workgroup count and size are (1, 1, 1) for serializing the computation.
-// CHECK-LABEL: hal.executable.entry_point public @static_1d_sort
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize"}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = [], native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = []>
+// CHECK: hal.executable.entry_point public @static_1d_sort
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [1 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%{{.+}}: index, %{{.+}}: index, %{{.+}}: index):
// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
@@ -43,7 +44,7 @@
// CHECK: func @static_1d_sort()
// CHECK: linalg_ext.sort
-// CHECK-SAME: lowering.config = {}
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -98,17 +99,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @static_3d_sort
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [16, 1]}
-// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[DIV:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 16)>()[%[[X]]]
-// CHECK-NEXT: hal.return %[[DIV]], %[[Y]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 0, 16], [1, 0, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 16)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [16, 1]>
+// CHECK: hal.executable.entry_point public @static_3d_sort
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[DIV:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: hal.return %[[DIV]], %[[Y]], %[[ONE]]
-// CHECK: func @static_3d_sort()
-// CHECK: linalg_ext.sort
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 0, 16], [1, 0, 1]]}
+// CHECK: func @static_3d_sort()
+// CHECK: linalg_ext.sort
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -144,18 +148,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @static_1d_fft_stage2
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute"
-// CHECK-SAME: workloadPerWorkgroup = [4]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [4]>
+// CHECK: hal.executable.entry_point public @static_1d_fft_stage2
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %{{.+}}: index, %{{.+}}: index):
// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[T:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[ARG0]]]
+// CHECK-NEXT: %[[T:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
// CHECK-NEXT: hal.return %[[T]], %[[ONE]], %[[ONE]]
// CHECK: func @static_1d_fft_stage2()
// CHECK: linalg_ext.fft
-// CHECK-SAME: lowering.config = {tileSizes = {{\[}}[4]]}
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -214,14 +220,16 @@
}
-// CHECK-LABEL: hal.executable.entry_point public @static_3d_fft_stage3
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute"
-// CHECK-SAME: workloadPerWorkgroup = [8, 1, 1]}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 1, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [8, 1, 1]>
+// CHECK: hal.executable.entry_point public @static_3d_fft_stage3
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
// CHECK-NEXT: ^{{.+}}(%[[ARG0:.+]]: index, %[[ARG1:.+]]: index, %[[ARG2:.+]]: index):
-// CHECK-NEXT: %[[T:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[ARG0]]]
+// CHECK-NEXT: %[[T:.+]] = affine.apply #[[MAP]]()[%[[ARG0]]]
// CHECK-NEXT: hal.return %[[T]], %[[ARG1]], %[[ARG2]]
// CHECK: func @static_3d_fft_stage3()
// CHECK: linalg_ext.fft
-// CHECK-SAME: lowering.config = {tileSizes = {{\[}}[1, 1, 8]]}
+// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir
index 0cb6a62..4ed7441 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_linalg_ops.mlir
@@ -47,9 +47,10 @@
}
}
}
-// CHECK: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 64)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [64, 1]>
// CHECK: hal.executable.entry_point public @tensor_insert_slice
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [64, 1]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
@@ -100,10 +101,11 @@
}
}
}
-// CHECK-DAG: #[[CONFIG:.+]] = {tileSizes = {{\[}}[1, 16], [1, 1]{{\]}}}
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 16], [1, 1]{{\]}}, native_vector_size = []>
// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 16)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [16, 1]>
// CHECK: hal.executable.entry_point public @tensor_insert_slice
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVDistribute", workloadPerWorkgroup = [16, 1]}
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
// CHECK-NEXT: %[[ARG0:[a-zA-Z0-9_]+]]: index
// CHECK-SAME: %[[ARG1:[a-zA-Z0-9_]+]]: index
// CHECK-DAG: %[[C1:.+]] = arith.constant 1 : index
diff --git a/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir b/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir
index 02000f2..644ab49 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_mali_conv.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
// Conv - large OC - distribute to only one workgroup dimension.
@@ -74,18 +74,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_112x112x512
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 4, 1]}
-// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C8:.+]] = arith.constant 8 : index
-// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index
-// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
-// CHECK-NEXT: hal.return %[[C8]], %[[C28]], %[[C112]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 1, 4, 64], [0, 1, 4, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 4, 1]>
+// CHECK: hal.executable.entry_point public @conv_112x112x512
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [16 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C8:.+]] = arith.constant 8 : index
+// CHECK-NEXT: %[[C28:.+]] = arith.constant 28 : index
+// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
+// CHECK-NEXT: hal.return %[[C8]], %[[C28]], %[[C112]]
-// CHECK: func @conv_112x112x512()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 1, 4, 64], [0, 1, 4, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_112x112x512()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -163,18 +165,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_112x112x32
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8, 1]}
-// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index
-// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C14]], %[[C112]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 1, 8, 32], [0, 1, 4, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 8, 1]>
+// CHECK: hal.executable.entry_point public @conv_112x112x32
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[C14:.+]] = arith.constant 14 : index
+// CHECK-NEXT: %[[C112:.+]] = arith.constant 112 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C14]], %[[C112]]
-// CHECK: func @conv_112x112x32()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 1, 8, 32], [0, 1, 4, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_112x112x32()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -251,17 +255,19 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @conv_16x16x16
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]}
-// CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[C4:.+]] = arith.constant 4 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C4]], %[[C4]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 4, 4, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
+// CHECK: hal.executable.entry_point public @conv_16x16x16
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[C4:.+]] = arith.constant 4 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C4]], %[[C4]]
-// CHECK: func @conv_16x16x16()
-// CHECK: linalg.conv_2d_nhwc_hwcf
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 4, 4, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1, 4]]}
+// CHECK: func @conv_16x16x16()
+// CHECK: linalg.conv_2d_nhwc_hwcf
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -340,17 +346,19 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @dwconv_28x28x144
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]}
-// CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index
-// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
-// CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 4, 4, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
+// CHECK: hal.executable.entry_point public @dwconv_28x28x144
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [4 : index, 2 : index, 2 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C9:.+]] = arith.constant 9 : index
+// CHECK-NEXT: %[[C7:.+]] = arith.constant 7 : index
+// CHECK-NEXT: hal.return %[[C9]], %[[C7]], %[[C7]]
-// CHECK: func @dwconv_28x28x144()
-// CHECK: linalg.depthwise_conv2D_nhw
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 4, 4, 16], [0, 2, 2, 4], [0, 0, 0, 0, 1, 1]]}
+// CHECK: func @dwconv_28x28x144()
+// CHECK: linalg.depthwise_conv2D_nhw
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -430,14 +438,16 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @dwconv_1x2x8
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 2, 1]}
-// CHECK-SAME: workgroup_size = [2 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[0, 1, 2, 8], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 2, 1]>
+// CHECK: hal.executable.entry_point public @dwconv_1x2x8
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: hal.return %[[C1]], %[[C1]], %[[C1]]
-// CHECK: func @dwconv_1x2x8()
-// CHECK: linalg.depthwise_conv2D_nhw
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[0, 1, 2, 8], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]]}
+// CHECK: func @dwconv_1x2x8()
+// CHECK: linalg.depthwise_conv2D_nhw
+// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir
index 4d72dc4..fe03afa 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_mali_matmul.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
// Large matmul that can match the best tiling scheme.
@@ -62,18 +62,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_1024x2048x512
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8]}
-// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[8, 32], [4, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 8]>
+// CHECK: hal.executable.entry_point public @matmul_1024x2048x512
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_1024x2048x512()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[8, 32], [4, 4], [0, 0, 4]]}
+// CHECK: func @matmul_1024x2048x512()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -139,18 +143,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_3136x24x96
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 32]}
-// CHECK-SAME: workgroup_size = [2 : index, 8 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[32, 8], [4, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 32]>
+// CHECK: hal.executable.entry_point public @matmul_3136x24x96
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 8 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_3136x24x96()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[32, 8], [4, 4], [0, 0, 4]]}
+// CHECK: func @matmul_3136x24x96()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -216,18 +224,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_196x64x192
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 4]}
-// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 4)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[4, 32], [2, 4], [0, 0, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 4)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 4]>
+// CHECK: hal.executable.entry_point public @matmul_196x64x192
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_196x64x192()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[4, 32], [2, 4], [0, 0, 8]]}
+// CHECK: func @matmul_196x64x192()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -288,18 +300,22 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_12544x96x16
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 8]}
-// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[8, 32], [4, 4], [0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 8]>
+// CHECK: hal.executable.entry_point public @matmul_12544x96x16
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[ONE]]
-// CHECK: func @matmul_12544x96x16()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[8, 32], [4, 4], [0, 0, 4]]}
+// CHECK: func @matmul_12544x96x16()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -365,17 +381,20 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_49x160x576
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 1]}
-// CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 32], [1, 4], [0, 0, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 1]>
+// CHECK: hal.executable.entry_point public @matmul_49x160x576
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[ONE:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y]], %[[ONE]]
-// CHECK: func @matmul_49x160x576()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 32], [1, 4], [0, 0, 8]]}
+// CHECK: func @matmul_49x160x576()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -452,17 +471,21 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @batch_matmul_4x384x384
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [32, 12, 1]}
-// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 32)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 12)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 12, 32], [1, 6, 4], [0, 0, 0, 4]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 32)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 12)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 12, 1]>
+// CHECK: hal.executable.entry_point public @batch_matmul_4x384x384
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [8 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
-// CHECK: func @batch_matmul_4x384x384()
-// CHECK: linalg.batch_matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 12, 32], [1, 6, 4], [0, 0, 0, 4]]}
+// CHECK: func @batch_matmul_4x384x384()
+// CHECK: linalg.batch_matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -540,14 +563,18 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @batch_matmul_4x2x8
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [8, 2, 1]}
-// CHECK-SAME: workgroup_size = [2 : index, 2 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 8)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 2)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[1, 2, 8], [1, 1, 4], [0, 0, 0, 8]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP0:.+]] = affine_map<()[s0] -> (s0 ceildiv 8)>
+// CHECK-DAG: #[[MAP1:.+]] = affine_map<()[s0] -> (s0 ceildiv 2)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 2, 1]>
+// CHECK: hal.executable.entry_point public @batch_matmul_4x2x8
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [2 : index, 2 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %[[Z:.+]]: index):
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP0]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP1]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[Z]]
-// CHECK: func @batch_matmul_4x2x8()
-// CHECK: linalg.batch_matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[1, 2, 8], [1, 1, 4], [0, 0, 0, 8]]}
+// CHECK: func @batch_matmul_4x2x8()
+// CHECK: linalg.batch_matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
diff --git a/iree/compiler/Codegen/SPIRV/test/config_nvidia_matmul_cooperative_ops.mlir b/iree/compiler/Codegen/SPIRV/test/config_nvidia_matmul_cooperative_ops.mlir
index 3968d72..1097c3c 100644
--- a/iree/compiler/Codegen/SPIRV/test/config_nvidia_matmul_cooperative_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/config_nvidia_matmul_cooperative_ops.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -mlir-print-local-scope -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-spirv-lower-executable-target-pass{test-lowering-configuration=true}))' %s | IreeFileCheck %s
#map0 = affine_map<()[s0, s1] -> (s0 * s1)>
#map1 = affine_map<(d0)[s0] -> (s0, -d0 + 256)>
@@ -100,18 +100,21 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_256x1024x128_div_sub
-// CHECK-SAME: translation.info = {passPipeline = "SPIRVVectorizeToCooperativeOps", workloadPerWorkgroup = [16, 16]}
-// CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index]
-// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
-// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
-// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 16)>()[%[[X]]]
-// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply affine_map<()[s0] -> (s0 ceildiv 16)>()[%[[Y]]]
-// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[C1]]
+// CHECK-DAG: #[[CONFIG:.+]] = #iree_codegen.lowering.config<tile_sizes = {{\[}}[16, 16, 16], [16, 16, 16]{{\]}}, native_vector_size = []>
+// CHECK-DAG: #[[MAP:.+]] = affine_map<()[s0] -> (s0 ceildiv 16)>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorizeToCooperativeOps", workload_per_wg = [16, 16]>
+// CHECK: hal.executable.entry_point public @matmul_256x1024x128_div_sub
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
+// CHECK-SAME: workgroup_size = [32 : index, 1 : index, 1 : index]
+// CHECK-NEXT: ^{{.+}}(%[[X:.+]]: index, %[[Y:.+]]: index, %{{.+}}: index):
+// CHECK-NEXT: %[[C1:.+]] = arith.constant 1 : index
+// CHECK-NEXT: %[[X_COUNT:.+]] = affine.apply #[[MAP]]()[%[[X]]]
+// CHECK-NEXT: %[[Y_COUNT:.+]] = affine.apply #[[MAP]]()[%[[Y]]]
+// CHECK-NEXT: hal.return %[[X_COUNT]], %[[Y_COUNT]], %[[C1]]
-// CHECK: func @matmul_256x1024x128_div_sub()
-// CHECK: linalg.matmul
-// CHECK-SAME{LITERAL}: lowering.config = {tileSizes = [[16, 16, 16], [16, 16, 16]]}
+// CHECK: func @matmul_256x1024x128_div_sub()
+// CHECK: linalg.matmul
+// CHECK-SAME: lowering.config = #[[CONFIG]]
// -----
@@ -194,5 +197,6 @@
}
}
-// CHECK-LABEL: hal.executable.entry_point public @matmul_256x1024x8
-// CHECK-SAME: passPipeline = "SPIRVVectorize"
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation.info<"SPIRVVectorize"
+// CHECK: hal.executable.entry_point public @matmul_256x1024x8
+// CHECK-SAME: translation.info = #[[TRANSLATION]]
diff --git a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
index fc1058e..3315774 100644
--- a/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-codegen-linalg-to-spirv-pipeline))' %s | IreeFileCheck %s
-#config = {tileSizes = [[8, 64], [8, 4], [0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[8, 64], [8, 4], [0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8]>
hal.executable private @fuse_and_vectorize_fill_matmul {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -13,7 +13,7 @@
hal.executable.entry_point @fuse_and_vectorize_fill_matmul attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8]}
+ translation.info = #translation
}
builtin.module {
func @fuse_and_vectorize_fill_matmul() {
@@ -70,8 +70,8 @@
// -----
-#config = {tileSizes = [[8, 64], [8, 4], [0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[8, 64], [8, 4], [0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8]>
hal.executable private @fuse_and_vectorize_matmul_add {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -83,7 +83,7 @@
hal.executable.entry_point @fuse_and_vectorize_matmul_add attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8]}
+ translation.info = #translation
}
builtin.module {
func @fuse_and_vectorize_matmul_add() {
diff --git a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
index 2dd847e..6b8b7bc 100644
--- a/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/remove_one_trip_tiled_loop.mlir
@@ -1,5 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-remove-one-trip-tiled-loop))))' %s | IreeFileCheck %s
+#config = #iree_codegen.lowering.config<tile_sizes = [[0, 4, 4, 16], [], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
hal.executable private @static_shaped_conv {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -9,7 +11,7 @@
hal.executable.variant @vulkan_spirv_fb, target = #hal.executable.target<"vulkan", "vulkan-spirv-fb"> {
hal.executable.entry_point @static_shaped_conv attributes {
interface = @io, ordinal = 0 : index,
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]},
+ translation.info = #translation,
workgroup_size = [4 : index, 4 : index, 1 : index]
}
builtin.module {
@@ -46,8 +48,8 @@
%16 = affine.min affine_map<(d0) -> (4, -d0 + 112)>(%arg0)
%17 = affine.min affine_map<(d0) -> (4, -d0 + 112)>(%arg1)
%18 = memref.subview %2[0, %arg0, %arg1, %arg2] [1, %16, %17, %14] [1, 1, 1, 1] : memref<1x112x112x32xf32> to memref<1x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 401408 + s0 + d1 * 3584 + d2 * 32 + d3)>>
- linalg.fill(%cst, %18) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[0, 4, 4, 16], [], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]]}} : f32, memref<1x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 401408 + s0 + d1 * 3584 + d2 * 32 + d3)>>
- linalg.conv_2d_nhwc_hwcf {__internal_linalg_transform__ = "workgroup", dilations = dense<1> : tensor<2xi64>, lowering.config = {tileSizes = [[0, 4, 4, 16], [], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]]}, strides = dense<2> : tensor<2xi64>}
+ linalg.fill(%cst, %18) {__internal_linalg_transform__ = "workgroup", lowering.config = #config} : f32, memref<1x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 401408 + s0 + d1 * 3584 + d2 * 32 + d3)>>
+ linalg.conv_2d_nhwc_hwcf {__internal_linalg_transform__ = "workgroup", dilations = dense<1> : tensor<2xi64>, lowering.config = #config, strides = dense<2> : tensor<2xi64>}
ins(%13, %15 : memref<1x?x?x3xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 151875 + s0 + d1 * 675 + d2 * 3 + d3)>>, memref<3x3x3x?xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 288 + s0 + d1 * 96 + d2 * 32 + d3)>>)
outs(%18 : memref<1x?x?x?xf32, affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 401408 + s0 + d1 * 3584 + d2 * 32 + d3)>>)
}
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
index bdd9652..b12a099 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_scatter.mlir
@@ -1,5 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute))))' %s | IreeFileCheck %s
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 16], [1, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [16, 1]>
hal.executable private @static_scatter_update_slice {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -10,7 +12,7 @@
hal.executable.variant @vulkan_spirv_fb, target = #hal.executable.target<"vulkan", "vulkan-spirv-fb"> {
hal.executable.entry_point @static_scatter_update_slice attributes {
interface = @io, ordinal = 0 : index,
- translation.info = {passPipeline = 5 : i32, workloadPerWorkgroup = [16, 1]},
+ translation.info = #translation,
workgroup_size = [16 : index, 1 : index, 1 : index]
}
@@ -36,7 +38,7 @@
%8 = memref.subview %1[%arg0, 0] [1, 1] [1, 1] : memref<40x1xi32> to memref<1x1xi32, affine_map<(d0, d1)[s0] -> (d0 + s0 + d1)>>
%9 = memref.cast %8 : memref<1x1xi32, affine_map<(d0, d1)[s0] -> (d0 + s0 + d1)>> to memref<?x1xi32, affine_map<(d0, d1)[s0] -> (d0 + s0 + d1)>>
%10 = memref.subview %2[0, %arg1] [100, %5] [1, 1] : memref<100x500xi32> to memref<100x?xi32, affine_map<(d0, d1)[s0] -> (d0 * 500 + s0 + d1)>>
- linalg_ext.scatter {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[1, 16], [1, 1]]}} ins(%7, %9 : memref<?x?xi32, affine_map<(d0, d1)[s0] -> (d0 * 500 + s0 + d1)>>, memref<?x1xi32, affine_map<(d0, d1)[s0] -> (d0 + s0 + d1)>>) outs(%10 : memref<100x?xi32, affine_map<(d0, d1)[s0] -> (d0 * 500 + s0 + d1)>>) {
+ linalg_ext.scatter {__internal_linalg_transform__ = "workgroup", lowering.config = #config} ins(%7, %9 : memref<?x?xi32, affine_map<(d0, d1)[s0] -> (d0 * 500 + s0 + d1)>>, memref<?x1xi32, affine_map<(d0, d1)[s0] -> (d0 + s0 + d1)>>) outs(%10 : memref<100x?xi32, affine_map<(d0, d1)[s0] -> (d0 * 500 + s0 + d1)>>) {
^bb0(%arg2: i32, %arg3: i32): // no predecessors
linalg_ext.yield %arg2 : i32
}
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
index 201f344..345c731 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_distribute_sort.mlir
@@ -1,5 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-distribute, cse))))' %s | IreeFileCheck %s
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 0, 16], [1, 0, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVDistribute", workload_per_wg = [16, 1]>
hal.executable private @static_3d_sort {
hal.interface @io {
hal.interface.binding @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -8,7 +10,7 @@
hal.executable.variant @vulkan_spirv_fb, target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb"> {
hal.executable.entry_point @static_3d_sort attributes {
interface = @io, ordinal = 0 : index,
- translation.info = {passPipeline = 5 : i32, workloadPerWorkgroup = [16, 1]},
+ translation.info = #translation,
workgroup_size = [16 : index, 1 : index, 1 : index]
}
builtin.module {
@@ -30,8 +32,8 @@
%5 = memref.cast %4 : memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>> to memref<?x?x?xi32>
%6 = memref.subview %1[%arg0, 0, %arg1] [1, 32, 16] [1, 1, 1] : memref<64x32x128xi32> to memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>
%7 = memref.cast %6 : memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>> to memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>
- linalg.copy(%5, %6) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[1, 0, 16], [1, 0, 1]]}} : memref<?x?x?xi32>, memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>
- linalg_ext.sort dimension(1) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[1, 0, 16], [1, 0, 1]]}} outs(%7 : memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>) {
+ linalg.copy(%5, %6) {__internal_linalg_transform__ = "workgroup", lowering.config = #config} : memref<?x?x?xi32>, memref<1x32x16xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>
+ linalg_ext.sort dimension(1) {__internal_linalg_transform__ = "workgroup", lowering.config = #config} outs(%7 : memref<?x32x?xi32, affine_map<(d0, d1, d2)[s0] -> (d0 * 4096 + s0 + d1 * 128 + d2)>>) {
^bb0(%arg2: i32, %arg3: i32): // no predecessors
%8 = arith.cmpi slt, %arg2, %arg3 : i32
linalg_ext.yield %8 : i1
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
index 8bf2d36..958207c 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
@@ -8,7 +8,8 @@
#map5 = affine_map<(d0, d1, d2) -> (d2, d1)>
#map6 = affine_map<(d0, d1, d2) -> (d0, d1)>
-#config = {tileSizes = [[8, 16], [1, 1], [0, 0, 1]]}
+#config = #iree_codegen.lowering.config<tile_sizes = [[8, 16], [1, 1], [0, 0, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [8, 16]>
hal.executable private @matmul {
hal.interface @io {
@@ -20,7 +21,7 @@
hal.executable.entry_point @matmul attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 8: index, 1: index],
- translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [8, 16]}
+ translation.info = #translation
}
builtin.module {
func @matmul() {
@@ -82,8 +83,8 @@
// -----
-#config = {tileSizes = [[1, 4, 32], [1, 1, 1]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 4, 32], [1, 1, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 4, 1]>
hal.executable private @conv_1d {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -94,7 +95,7 @@
hal.executable.entry_point @conv_1d attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [32: index, 4: index, 1: index],
- translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]}
+ translation.info = #translation
}
builtin.module {
func @conv_1d() {
@@ -165,8 +166,8 @@
#map6 = affine_map<(d0)[s0] -> (4, -d0 + s0)>
#map7 = affine_map<(d0)[s0] -> (32, -d0 + s0)>
-#config = {tileSizes = [[0, 1, 4, 32], [0, 1, 1, 1], [0, 0, 0, 0, 1, 1, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[0, 1, 4, 32], [0, 1, 1, 1], [0, 0, 0, 0, 1, 1, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 4, 1]>
hal.executable private @conv_no_padding {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -177,7 +178,7 @@
hal.executable.entry_point @conv_no_padding attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [32: index, 4: index, 1: index],
- translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]}
+ translation.info = #translation
}
builtin.module {
func @conv_no_padding() {
@@ -292,8 +293,8 @@
// -----
-#config = {tileSizes = [[0, 0, 1, 4, 32], [0, 0, 1, 1, 1]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[0, 0, 1, 4, 32], [0, 0, 1, 1, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 4, 1]>
hal.executable private @conv_3d {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -304,7 +305,7 @@
hal.executable.entry_point @conv_3d attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [32: index, 4: index, 1: index],
- translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]}
+ translation.info = #translation
}
builtin.module {
func @conv_3d() {
@@ -365,8 +366,8 @@
#map6 = affine_map<()[s0] -> (32, s0 * -32 + 13)>
#map7 = affine_map<(d0, d1, d2, d3)[s0] -> (d0 * 1092 + s0 + d1 * 78 + d2 * 6 + d3)>
-#config = {tileSizes = [[1, 4, 32], [1, 1, 1]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 4, 32], [1, 1, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [32, 4, 1]>
module {
hal.executable private @pooling_nhwc_max {
hal.interface @io {
@@ -378,7 +379,7 @@
hal.executable.entry_point @pooling_nhwc_max attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [32: index, 4: index, 1: index],
- translation.info = {passPipeline = 6 : i32, workloadPerWorkgroup = [32, 4, 1]}
+ translation.info = #translation
}
builtin.module {
func @pooling_nhwc_max() {
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
index 2cd1b62..6018c6b 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(iree-spirv-tile-and-distribute,iree-spirv-vectorize))))' -canonicalize -cse %s | IreeFileCheck %s
-#config = {tileSizes = [[1, 8, 64], [1, 8, 4], [0, 0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 8, 64], [1, 8, 4], [0, 0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8, 1]>
hal.executable private @batch_matmul_static_shape {
hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -12,7 +12,7 @@
hal.executable.entry_point @batch_matmul_static_shape attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8, 1]}
+ translation.info = #translation
}
builtin.module {
func @batch_matmul_static_shape() {
@@ -370,8 +370,8 @@
// -----
-#config = {tileSizes = [[1, 8, 64], [1, 8, 4], [0, 0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 8, 64], [1, 8, 4], [0, 0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8, 1]>
hal.executable private @fused_fill_batch_matmul {
hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -382,7 +382,7 @@
hal.executable.entry_point @fused_fill_batch_matmul attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8, 1]}
+ translation.info = #translation
}
builtin.module {
func @fused_fill_batch_matmul() {
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
index ffd3b8f..757dd46 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(canonicalize,iree-spirv-remove-one-trip-tiled-loop,iree-spirv-tile-and-distribute,iree-spirv-vectorize))))' -canonicalize -cse %s | IreeFileCheck %s
-#config = {tileSizes = [[0, 4, 4, 16], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[0, 4, 4, 16], [0, 4, 1, 4], [0, 0, 0, 0, 1, 1, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
hal.executable private @conv_static_shape_f32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -13,7 +13,7 @@
interface = @io,
ordinal = 0 : index,
workgroup_size = [4: index, 4: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]}
+ translation.info = #translation
} {
^bb0(%arg0 : index, %arg1 : index, %arg2 : index):
%x = arith.constant 2: index
@@ -99,8 +99,8 @@
// -----
-#config = {tileSizes = [[0, 4, 4, 16], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[0, 4, 4, 16], [0, 1, 1, 4], [0, 0, 0, 0, 1, 1]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [16, 4, 4]>
hal.executable private @depthwise_conv_static_shape_f32 {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -112,7 +112,7 @@
interface = @io,
ordinal = 0 : index,
workgroup_size = [4: index, 4: index, 4: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [16, 4, 4]}
+ translation.info = #translation
} {
^bb0(%arg0 : index, %arg1 : index, %arg2 : index):
%x = arith.constant 6: index
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
index 0906c26..a3c5db1 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(iree-set-num-workgroups,builtin.module(builtin.func(iree-spirv-tile-and-distribute,iree-spirv-vectorize))))' -canonicalize -cse %s | IreeFileCheck %s
-#config = {tileSizes = [[8, 64], [8, 4], [0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[8, 64], [8, 4], [0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8]>
hal.executable private @matmul_static_shape_f16 {
hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -12,7 +12,7 @@
hal.executable.entry_point @matmul_static_shape_f16 attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8]}
+ translation.info = #translation
}
builtin.module {
func @matmul_static_shape_f16() {
@@ -66,8 +66,8 @@
// -----
-#config = {tileSizes = [[8, 64], [8, 4], [0, 0, 4]]}
-
+#config = #iree_codegen.lowering.config<tile_sizes = [[8, 64], [8, 4], [0, 0, 4]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorize", workload_per_wg = [64, 8]>
hal.executable private @matmul_static_shape_f32 {
hal.interface private @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -78,7 +78,7 @@
hal.executable.entry_point @matmul_static_shape_f32 attributes {
interface = @io, ordinal = 0 : index,
workgroup_size = [16: index, 1: index, 1: index],
- translation.info = {passPipeline = "SPIRVVectorize", workloadPerWorkgroup = [64, 8]}
+ translation.info = #translation
}
builtin.module {
func @matmul_static_shape_f32() {
diff --git a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir
index e63bd03..42e13de 100644
--- a/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_to_cooperative_ops.mlir
@@ -1,5 +1,7 @@
// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.variant(builtin.module(builtin.func(iree-spirv-tile-and-vectorize-to-cooperative-ops))))' %s | IreeFileCheck %s
+#config = #iree_codegen.lowering.config<tile_sizes = [[16, 16, 16], [16, 16, 16]], native_vector_size = []>
+#translation = #iree_codegen.translation.info<"SPIRVVectorizeToCooperativeOps", workload_per_wg = [16, 16]>
hal.executable public @matmul_256x1024x128_div_sub {
hal.interface public @io {
hal.interface.binding public @s0b0_ro_external, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -28,7 +30,7 @@
subgroup_size = 32 : i32}>}> {
hal.executable.entry_point public @matmul_256x1024x128_div_sub attributes {
interface = @io, ordinal = 0 : index,
- translation.info = {passPipeline = "SPIRVVectorizeToCooperativeOps", workloadPerWorkgroup = [16, 16]},
+ translation.info = #translation,
workgroup_size = [32 : index, 1 : index, 1 : index]
} {
^bb0(%arg0: index, %arg1: index, %arg2: index): // no predecessors
@@ -63,14 +65,14 @@
%11 = memref.subview %2[%arg0, 0] [16, 128] [1, 1] : memref<256x128xf16> to memref<16x128xf16, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>
%12 = memref.subview %3[0, %arg1] [128, 16] [1, 1] : memref<128x1024xf16> to memref<128x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>
%13 = memref.subview %4[%arg0, %arg1] [16, 16] [1, 1] : memref<256x1024xf16> to memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>
- linalg.fill(%cst, %13) {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[16, 16, 16], [16, 16, 16]]}} : f16, memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>
- linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[16, 16, 16], [16, 16, 16]]}}
+ linalg.fill(%cst, %13) {__internal_linalg_transform__ = "workgroup", lowering.config = #config} : f16, memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>
+ linalg.matmul {__internal_linalg_transform__ = "workgroup", lowering.config = #config}
ins(%11, %12 : memref<16x128xf16, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>, memref<128x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>)
outs(%13 : memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>)
linalg.generic {indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>, affine_map<(d0, d1) -> (d0, d1)>], iterator_types = ["parallel", "parallel"]}
ins(%13, %9, %10 : memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>, memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>, memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>)
outs(%13 : memref<16x16xf16, affine_map<(d0, d1)[s0] -> (d0 * 1024 + s0 + d1)>>)
- attrs = {__internal_linalg_transform__ = "workgroup", lowering.config = {tileSizes = [[16, 16, 16], [16, 16, 16]]}} {
+ attrs = {__internal_linalg_transform__ = "workgroup", lowering.config = #config} {
^bb0(%arg2: f16, %arg3: f16, %arg4: f16, %arg5: f16): // no predecessors
%14 = arith.divf %arg2, %arg3 : f16
%15 = arith.subf %14, %arg4 : f16
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
index 667dab3..3785956 100644
--- a/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
@@ -5,6 +5,7 @@
// CHECK: vector.transfer_read %{{.+}}[%c0], {{.+}} memref<4xf32, #{{.+}}>, vector<4xf32>
// CHECK: addf %{{.*}}, %{{.*}} : vector<4xf32>
// CHECK: vector.transfer_write {{.*}} : vector<4xf32>, memref<4xf32
+#config = #iree_codegen.lowering.config<tile_sizes = [[128], [4]], native_vector_size = []>
hal.executable private @elementwise_static_shape {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -24,7 +25,7 @@
%ret0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128xf32>
linalg.generic {
__internal_linalg_transform__ = "workgroup",
- lowering.config = {tileSizes = [[128], [4]]},
+ lowering.config = #config,
indexing_maps = [affine_map<(i) -> (i)>,
affine_map<(i) -> (i)>,
affine_map<(i) -> (i)>],
@@ -54,6 +55,7 @@
// CHECK-NOT: vector.transfer_read
// CHECK: scf.for
// CHECK: scf.for
+#config = #iree_codegen.lowering.config<tile_sizes = [[1, 32], [1, 1]], native_vector_size = []>
hal.executable private @elementwise_transpose {
hal.interface @io {
hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
@@ -73,7 +75,7 @@
%ret0 = hal.interface.binding.subspan @io::@ret0[%c0] : memref<128x8xf32>
linalg.generic {
__internal_linalg_transform__ = "workgroup",
- lowering.config = {tileSizes = [[1, 32], [1, 1]]},
+ lowering.config = #config,
indexing_maps = [affine_map<(d0, d1) -> (d0, d1)>,
affine_map<(d0, d1) -> (d0)>,
affine_map<(d0, d1) -> (d0, d1)>],
diff --git a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
index bd16b5d..618ec1f 100644
--- a/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
@@ -1,5 +1,6 @@
// RUN: iree-opt -split-input-file -iree-spirv-vectorize %s | IreeFileCheck %s
+#config = #iree_codegen.lowering.config<tile_sizes = [[2, 128], [], [1, 4], [0, 0, 4]], native_vector_size = []>
func @matmul_2x128x4() {
%c0 = arith.constant 0 : index
%c128 = arith.constant 128 : index
@@ -25,10 +26,10 @@
%11 = "gpu.thread_id"() {dimension = "y"} : () -> index
%12 = affine.apply affine_map<()[s0] -> (s0 * 4)>()[%10]
%13 = memref.subview %9[%11, %12] [1, 4] [1, 1] : memref<2x128xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>> to memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>
- linalg.fill(%cst, %13) {__internal_linalg_transform__ = "vectorize", lowering.config = {tileSizes = [[2, 128], [], [1, 4], [0, 0, 4]]}} : f32, memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>
+ linalg.fill(%cst, %13) {__internal_linalg_transform__ = "vectorize", lowering.config = #config} : f32, memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>
%17 = memref.subview %7[%11, 0] [1, 4] [1, 1] : memref<2x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 4 + s0 + d1)>> to memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 4 + s0 + d1)>>
%18 = memref.subview %8[0, %12] [4, 4] [1, 1] : memref<4x128xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>> to memref<4x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>
- linalg.matmul {__internal_linalg_transform__ = "vectorize", lowering.config = {tileSizes = [[2, 128], [], [1, 4], [0, 0, 4]]}}
+ linalg.matmul {__internal_linalg_transform__ = "vectorize", lowering.config = #config}
ins(%17, %18 : memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 4 + s0 + d1)>>, memref<4x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>)
outs(%13 : memref<1x4xf32, affine_map<(d0, d1)[s0] -> (d0 * 128 + s0 + d1)>>)
}
diff --git a/iree/compiler/Codegen/Utils/Utils.cpp b/iree/compiler/Codegen/Utils/Utils.cpp
index 4f55f59..58a1e04 100644
--- a/iree/compiler/Codegen/Utils/Utils.cpp
+++ b/iree/compiler/Codegen/Utils/Utils.cpp
@@ -19,6 +19,10 @@
namespace mlir {
namespace iree_compiler {
+//===----------------------------------------------------------------------===//
+// Utility functions to get entry point(s)
+//===----------------------------------------------------------------------===//
+
bool isEntryPoint(FuncOp func) { return func.isPublic(); }
IREE::HAL::ExecutableEntryPointOp getEntryPoint(FuncOp funcOp) {
@@ -41,21 +45,9 @@
return entryPointOps;
}
-IREE::HAL::TranslationInfo getTranslationInfo(FuncOp funcOp) {
- auto entryPointOp = getEntryPoint(funcOp);
- if (!entryPointOp) return nullptr;
- return getTranslationInfo(entryPointOp);
-}
-
-void setTranslationInfo(FuncOp entryPointFn,
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize,
- ArrayRef<int64_t> workloadPerWorkgroup) {
- auto entryPointOp = getEntryPoint(entryPointFn);
- auto translationInfo = buildTranslationInfo(
- passPipeline, workloadPerWorkgroup, entryPointFn.getContext());
- setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
-}
+//===----------------------------------------------------------------------===//
+// Utility functions used in setting default configurations.
+//===----------------------------------------------------------------------===//
SmallVector<unsigned> getPartitionedLoops(Operation *op) {
if (auto mmt4dOp = dyn_cast<linalg::Mmt4DOp>(op)) {
@@ -80,45 +72,6 @@
return {};
}
-LogicalResult setOpConfigAndEntryPointFnTranslation(
- FuncOp entryPointFn, Operation *op, IREE::HAL::LoweringConfig config,
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize) {
- auto partitionedLoops = getPartitionedLoops(op);
- SmallVector<int64_t, 3> workloadPerWorkgroup;
- auto tileSizes = getTileSizes(config, 0);
- if (!tileSizes.empty() && !partitionedLoops.empty()) {
- for (unsigned depth : partitionedLoops) {
- if (depth >= tileSizes.size()) {
- return op->emitOpError(
- "illegal configuration for lowering op, expect first level "
- "tile size to contain at least ")
- << partitionedLoops.back() << " elements";
- }
- if (tileSizes[depth] == 0) {
- return op->emitOpError("illegal to set tilesize of loop ")
- << depth
- << " to zero since it is set to be partitioned at the flow "
- "level";
- }
- workloadPerWorkgroup.push_back(tileSizes[depth]);
- }
- if (!workloadPerWorkgroup.empty()) {
- workloadPerWorkgroup =
- llvm::to_vector<3>(llvm::reverse(workloadPerWorkgroup));
- }
- }
- auto entryPointOp = getEntryPoint(entryPointFn);
- if (!entryPointOp) {
- return entryPointFn.emitOpError(
- "unable to find entry point op for entry point function");
- }
- IREE::HAL::TranslationInfo translationInfo = buildTranslationInfo(
- passPipeline, workloadPerWorkgroup, entryPointOp->getContext());
- setTranslationInfo(entryPointOp, translationInfo, workgroupSize);
- return success();
-}
-
/// Walk up the defs of the view, to get the untiled value. Either walks up
/// `ViewOpInterface` op-chains or the `subtensor` op-chains.
static Value getViewSource(Value view) {
diff --git a/iree/compiler/Codegen/Utils/Utils.h b/iree/compiler/Codegen/Utils/Utils.h
index 120e2c4..f0c563a 100644
--- a/iree/compiler/Codegen/Utils/Utils.h
+++ b/iree/compiler/Codegen/Utils/Utils.h
@@ -8,7 +8,6 @@
#define IREE_COMPILER_CODEGEN_UTILS_UTILS_H_
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "llvm/ADT/StringMap.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
#include "mlir/IR/BuiltinOps.h"
@@ -18,6 +17,10 @@
static constexpr unsigned kNumMaxParallelDims = 3;
+//===----------------------------------------------------------------------===//
+// Utility functions to get entry point(s)
+//===----------------------------------------------------------------------===//
+
/// Returns true if the given `func` is a kernel dispatch entry point.
bool isEntryPoint(FuncOp func);
@@ -28,18 +31,9 @@
/// Returns the entry point op for the `funcOp`. Returns `nullptr` on failure.
IREE::HAL::ExecutableEntryPointOp getEntryPoint(FuncOp funcOp);
-/// Returns the translation info for the `funcOp` (by looking at the entry
-/// point). Returns `nullptr` on failure.
-IREE::HAL::TranslationInfo getTranslationInfo(FuncOp funcOp);
-
-/// Sets the translation info on the `hal.executable.entry_point` op
-/// corresponding to the `entryPointFn`. Returns failure if a translation info
-/// is already set on the entry point op and is incompatible with what is being
-/// set.
-void setTranslationInfo(FuncOp entryPointFn,
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize,
- ArrayRef<int64_t> workloadPerWorkgroup);
+//===----------------------------------------------------------------------===//
+// Utility functions used in setting default configurations.
+//===----------------------------------------------------------------------===//
/// Returns the loops that are partitioned during dispatch region formations, in
/// order, i.e. starting from the outer-most to innermost.
@@ -47,23 +41,6 @@
/// formation to tile and distribute the ops.
SmallVector<unsigned> getPartitionedLoops(Operation *op);
-/// Sets translation for the entry-point function based on op configuration.
-LogicalResult setOpConfigAndEntryPointFnTranslation(
- FuncOp entryPointFn, Operation *op, IREE::HAL::LoweringConfig config,
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize = {});
-inline LogicalResult setOpConfigAndEntryPointFnTranslation(
- FuncOp entryPointFn, Operation *op, TileSizesListTypeRef tileSizes,
- ArrayRef<int64_t> nativeVectorSize,
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize = {}) {
- IREE::HAL::LoweringConfig config =
- buildConfigAttr(tileSizes, nativeVectorSize, op->getContext());
- setLoweringConfig(op, config);
- return setOpConfigAndEntryPointFnTranslation(entryPointFn, op, config,
- passPipeline, workgroupSize);
-}
-
/// Returns the untiled type of a tiled view for both tensor and memref
/// types. Either walks the `ViewOpInterface` chain (for memrefs) or the
/// `subtensor` op chain (for tensors).
diff --git a/iree/compiler/Dialect/HAL/IR/BUILD b/iree/compiler/Dialect/HAL/IR/BUILD
index 9586bfc..c0cdf75 100644
--- a/iree/compiler/Dialect/HAL/IR/BUILD
+++ b/iree/compiler/Dialect/HAL/IR/BUILD
@@ -27,7 +27,6 @@
"HALDialect.td",
"HALInterfaces.td",
"HALOps.td",
- "LoweringConfig.td",
],
include = ["*.td"],
),
@@ -46,14 +45,12 @@
"HALOpFolders.cpp",
"HALOps.cpp",
"HALTypes.cpp",
- "LoweringConfig.cpp",
],
hdrs = [
"HALDialect.h",
"HALOps.h",
"HALTraits.h",
"HALTypes.h",
- "LoweringConfig.h",
],
textual_hdrs = [
"HALAttrs.cpp.inc",
@@ -70,18 +67,12 @@
"HALStructs.h.inc",
"HALTypeInterfaces.cpp.inc",
"HALTypeInterfaces.h.inc",
- "LoweringConfig.h.inc",
- "LoweringConfig.cpp.inc",
- "LoweringConfigEnums.h.inc",
- "LoweringConfigEnums.cpp.inc",
],
deps = [
":HALInterfacesGen",
":HALOpsGen",
":HALStructsGen",
":HALTypesGen",
- ":LoweringConfigEnumGen",
- ":LoweringConfigGen",
"//iree/compiler/Dialect/Shape/IR",
"//iree/compiler/Dialect/Util/IR",
"@llvm-project//llvm:Support",
@@ -221,37 +212,3 @@
td_file = "HALOps.td",
deps = [":td_files"],
)
-
-gentbl_cc_library(
- name = "LoweringConfigGen",
- tbl_outs = [
- (
- ["-gen-struct-attr-decls"],
- "LoweringConfig.h.inc",
- ),
- (
- ["-gen-struct-attr-defs"],
- "LoweringConfig.cpp.inc",
- ),
- ],
- tblgen = "@llvm-project//mlir:mlir-tblgen",
- td_file = "LoweringConfig.td",
- deps = [":td_files"],
-)
-
-gentbl_cc_library(
- name = "LoweringConfigEnumGen",
- tbl_outs = [
- (
- ["-gen-enum-decls"],
- "LoweringConfigEnums.h.inc",
- ),
- (
- ["-gen-enum-defs"],
- "LoweringConfigEnums.cpp.inc",
- ),
- ],
- tblgen = "@llvm-project//mlir:mlir-tblgen",
- td_file = "LoweringConfig.td",
- deps = [":td_files"],
-)
diff --git a/iree/compiler/Dialect/HAL/IR/CMakeLists.txt b/iree/compiler/Dialect/HAL/IR/CMakeLists.txt
index 341d7ea..8d2316a 100644
--- a/iree/compiler/Dialect/HAL/IR/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/IR/CMakeLists.txt
@@ -18,7 +18,6 @@
"HALOps.h"
"HALTraits.h"
"HALTypes.h"
- "LoweringConfig.h"
TEXTUAL_HDRS
"HALAttrInterfaces.cpp.inc"
"HALAttrInterfaces.h.inc"
@@ -34,22 +33,15 @@
"HALStructs.h.inc"
"HALTypeInterfaces.cpp.inc"
"HALTypeInterfaces.h.inc"
- "LoweringConfig.cpp.inc"
- "LoweringConfig.h.inc"
- "LoweringConfigEnums.cpp.inc"
- "LoweringConfigEnums.h.inc"
SRCS
"HALOpFolders.cpp"
"HALOps.cpp"
"HALTypes.cpp"
- "LoweringConfig.cpp"
DEPS
::HALInterfacesGen
::HALOpsGen
::HALStructsGen
::HALTypesGen
- ::LoweringConfigEnumGen
- ::LoweringConfigGen
LLVMSupport
MLIRIR
MLIRMemRef
@@ -145,24 +137,4 @@
-gen-dialect-doc HALDialect.md
)
-iree_tablegen_library(
- NAME
- LoweringConfigGen
- TD_FILE
- "LoweringConfig.td"
- OUTS
- -gen-struct-attr-decls LoweringConfig.h.inc
- -gen-struct-attr-defs LoweringConfig.cpp.inc
-)
-
-iree_tablegen_library(
- NAME
- LoweringConfigEnumGen
- TD_FILE
- "LoweringConfig.td"
- OUTS
- -gen-enum-decls LoweringConfigEnums.h.inc
- -gen-enum-defs LoweringConfigEnums.cpp.inc
-)
-
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
diff --git a/iree/compiler/Dialect/HAL/IR/HALDialect.cpp b/iree/compiler/Dialect/HAL/IR/HALDialect.cpp
index 7b38e4a..a3e30eb 100644
--- a/iree/compiler/Dialect/HAL/IR/HALDialect.cpp
+++ b/iree/compiler/Dialect/HAL/IR/HALDialect.cpp
@@ -10,7 +10,6 @@
#include "iree/compiler/Dialect/HAL/Conversion/HALToVM/ConvertHALToVM.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALTypes.h"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "iree/compiler/Dialect/HAL/hal.imports.h"
#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
#include "iree/compiler/Dialect/VM/Conversion/ConversionDialectInterface.h"
@@ -44,9 +43,6 @@
} else if (auto targetAttr = attr.dyn_cast<ExecutableTargetAttr>()) {
os << "executable_target_" << targetAttr.getSymbolNameFragment();
return AliasResult::OverridableAlias;
- } else if (attr.isa<LoweringConfig>()) {
- os << "config";
- return AliasResult::OverridableAlias;
}
return AliasResult::NoAlias;
}
diff --git a/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp b/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp
deleted file mode 100644
index 4bc7e4e..0000000
--- a/iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp
+++ /dev/null
@@ -1,153 +0,0 @@
-// Copyright 2021 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
-
-#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "mlir/Dialect/StandardOps/IR/Ops.h"
-
-static const char kConfigAttrName[] = "lowering.config";
-static const char kTranslationInfoAttrName[] = "translation.info";
-
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.cpp.inc"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfigEnums.cpp.inc"
-
-namespace mlir {
-namespace iree_compiler {
-
-//===----------------------------------------------------------------------===//
-// Helpers for getting/setting information needed to lower an executable. These
-// are information that are stored as attributes on the
-// `hal.executable.entry_point`
-//===----------------------------------------------------------------------===//
-
-IREE::HAL::TranslationInfo buildTranslationInfo(
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workloadPerWorkgroup, MLIRContext *context) {
- OpBuilder builder(context);
- auto pipelineAttr = StringAttr::get(context, stringifyEnum(passPipeline));
- ArrayAttr workloadPerWorkgroupAttr = nullptr;
- if (!workloadPerWorkgroup.empty()) {
- workloadPerWorkgroupAttr = builder.getI64ArrayAttr(workloadPerWorkgroup);
- }
- return IREE::HAL::TranslationInfo::get(pipelineAttr, workloadPerWorkgroupAttr,
- context);
-}
-
-IREE::HAL::TranslationInfo getTranslationInfo(
- IREE::HAL::ExecutableEntryPointOp entryPointOp) {
- return entryPointOp->getAttrOfType<IREE::HAL::TranslationInfo>(
- kTranslationInfoAttrName);
-}
-
-SmallVector<int64_t> getWorkgroupSize(
- IREE::HAL::ExecutableEntryPointOp entryPointOp) {
- SmallVector<int64_t> workgroupSize;
- if (Optional<ArrayAttr> workgroupSizeAttrList =
- entryPointOp.workgroup_size()) {
- workgroupSize.resize(workgroupSizeAttrList->size());
- for (auto attr : llvm::enumerate(workgroupSizeAttrList.getValue())) {
- workgroupSize[attr.index()] = attr.value().cast<IntegerAttr>().getInt();
- }
- }
- return workgroupSize;
-}
-
-void setTranslationInfo(IREE::HAL::ExecutableEntryPointOp entryPointOp,
- IREE::HAL::TranslationInfo translationInfo,
- ArrayRef<int64_t> workgroupSize) {
- entryPointOp->setAttr(kTranslationInfoAttrName, translationInfo);
- // The workgroup size is set on the entry point op directly.
- if (!workgroupSize.empty()) {
- MLIRContext *context = entryPointOp->getContext();
- auto indexType = IndexType::get(context);
- auto attrs = llvm::to_vector<4>(
- llvm::map_range(workgroupSize, [&](int64_t v) -> Attribute {
- return IntegerAttr::get(indexType, v);
- }));
- entryPointOp.workgroup_sizeAttr(ArrayAttr::get(context, attrs));
- }
-}
-
-//===----------------------------------------------------------------------===//
-// Helpers for getting/setting the `hal.lowering.*` attributes that drive the
-// linalg-based lowering.
-// ===----------------------------------------------------------------------===//
-
-IREE::HAL::LoweringConfig getLoweringConfig(Operation *op) {
- return op->getAttrOfType<IREE::HAL::LoweringConfig>(kConfigAttrName);
-}
-
-void setLoweringConfig(Operation *op, IREE::HAL::LoweringConfig config) {
- op->setAttr(kConfigAttrName, config);
-}
-
-void eraseLoweringConfig(Operation *op) { op->removeAttr(kConfigAttrName); }
-
-//===----------------------------------------------------------------------===//
-// Helpers for accessing values from the LoweringConfig attribute.
-//===----------------------------------------------------------------------===//
-
-IREE::HAL::LoweringConfig buildConfigAttr(TileSizesListTypeRef tileSizes,
- ArrayRef<int64_t> nativeVectorSize,
- MLIRContext *context) {
- OpBuilder builder(context);
- ArrayAttr tileSizesAttr = nullptr;
- if (!tileSizes.empty()) {
- auto attrList = llvm::to_vector<4>(
- llvm::map_range(tileSizes, [&](ArrayRef<int64_t> sizes) -> Attribute {
- return builder.getI64ArrayAttr(sizes);
- }));
- tileSizesAttr = builder.getArrayAttr(attrList);
- }
- ArrayAttr nativeVectorSizeAttr = nullptr;
- if (!nativeVectorSize.empty()) {
- nativeVectorSizeAttr = builder.getI64ArrayAttr(nativeVectorSize);
- }
- return IREE::HAL::LoweringConfig::get(tileSizesAttr, nativeVectorSizeAttr,
- /*passPipeline = */ nullptr,
- /*workgroupSize = */ nullptr, context);
-}
-
-TileSizesListType getTileSizes(IREE::HAL::LoweringConfig config) {
- auto tileSizesAttr = config.tileSizes();
- if (!tileSizesAttr) return {};
- return llvm::to_vector<1>(llvm::map_range(
- tileSizesAttr, [&](Attribute attr) -> SmallVector<int64_t, 4> {
- return llvm::to_vector<4>(
- llvm::map_range(attr.cast<ArrayAttr>(), [&](Attribute intAttr) {
- return intAttr.cast<IntegerAttr>().getInt();
- }));
- }));
-}
-
-SmallVector<int64_t, 4> getTileSizes(IREE::HAL::LoweringConfig config,
- unsigned level) {
- ArrayAttr tileSizesAttr = config.tileSizes();
- if (!tileSizesAttr || tileSizesAttr.size() <= level) return {};
- return llvm::to_vector<4>(llvm::map_range(
- tileSizesAttr.getValue()[level].cast<ArrayAttr>(),
- [&](Attribute intAttr) { return intAttr.cast<IntegerAttr>().getInt(); }));
-}
-
-SmallVector<Value, 4> getTileSizes(OpBuilder &b, Operation *op,
- unsigned level) {
- return llvm::to_vector<4>(
- llvm::map_range(getTileSizes(op, level), [&](int64_t t) -> Value {
- return b.create<arith::ConstantIndexOp>(op->getLoc(), t);
- }));
-}
-
-SmallVector<int64_t, 4> getNativeVectorSize(IREE::HAL::LoweringConfig config) {
- ArrayAttr nativeVectorSizeAttr = config.nativeVectorSize();
- if (!nativeVectorSizeAttr) return {};
- return llvm::to_vector<4>(llvm::map_range(
- nativeVectorSizeAttr,
- [&](Attribute intAttr) { return intAttr.cast<IntegerAttr>().getInt(); }));
-}
-
-} // namespace iree_compiler
-} // namespace mlir
diff --git a/iree/compiler/Dialect/HAL/IR/LoweringConfig.h b/iree/compiler/Dialect/HAL/IR/LoweringConfig.h
deleted file mode 100644
index a48d60a..0000000
--- a/iree/compiler/Dialect/HAL/IR/LoweringConfig.h
+++ /dev/null
@@ -1,163 +0,0 @@
-// Copyright 2021 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-//===- LoweringConfig.h - Declares configuration for lowering Linalg ops --===//
-//
-// This file declares an attribute that drives how a dispatch region containing
-// a set of operations are lowered. The attribute itself is attached to Linalg
-// operations, and help converting a Linalg operation into "scalar code".
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
-#define IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
-
-#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "mlir/IR/Builders.h"
-#include "mlir/IR/BuiltinAttributes.h"
-#include "mlir/IR/BuiltinTypes.h"
-
-// clang-format off
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfigEnums.h.inc"
-#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h.inc"
-// clang-format on
-
-namespace mlir {
-namespace iree_compiler {
-
-namespace IREE {
-namespace HAL {
-
-inline bool operator==(const TranslationInfo &lhs, const TranslationInfo &rhs) {
- return lhs.passPipeline() == rhs.passPipeline() &&
- lhs.workloadPerWorkgroup() == rhs.workloadPerWorkgroup();
-}
-
-inline bool operator!=(const TranslationInfo &lhs, const TranslationInfo &rhs) {
- return !(lhs == rhs);
-}
-
-} // namespace HAL
-} // namespace IREE
-
-//===----------------------------------------------------------------------===//
-// Helpers for getting/setting information needed to lower an executable. These
-// are information that are stored as attributes on the
-// `hal.executable.entry_point`
-//===----------------------------------------------------------------------===//
-
-/// Builder method for IREE::HAL::TranslationInfoAttr.
-IREE::HAL::TranslationInfo buildTranslationInfo(
- IREE::HAL::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workloadPerWorkgroup, MLIRContext *context);
-
-/// Gets the translate executable info attribute value associated with
-/// `entryPointOp`.
-IREE::HAL::TranslationInfo getTranslationInfo(
- IREE::HAL::ExecutableEntryPointOp entryPointOp);
-
-/// Get the pass pipeline specified in the `translationInfo`
-inline Optional<IREE::HAL::DispatchLoweringPassPipeline>
-getLoweringPassPipeline(IREE::HAL::TranslationInfo translationInfo) {
- return IREE::HAL::symbolizeDispatchLoweringPassPipeline(
- translationInfo.passPipeline().getValue());
-}
-
-/// Returns the workgroup size specified on the `entryPointOp`.
-SmallVector<int64_t> getWorkgroupSize(
- IREE::HAL::ExecutableEntryPointOp entryPointOp);
-
-/// Set the translate executable info with the entry point op. Overwrites the
-/// existing attributes.
-// TODO(ravishankarm, benvanik): Eventually all the information needed for the
-// lowering will be consolidated into a single attribute with richer
-// information.
-void setTranslationInfo(IREE::HAL::ExecutableEntryPointOp entryPointOp,
- IREE::HAL::TranslationInfo translationInfo,
- ArrayRef<int64_t> workgroupSize = {});
-
-//===----------------------------------------------------------------------===//
-// Helpers for getting/setting the `hal.lowering.*` attributes that drive the
-// linalg-based lowering.
-// ===----------------------------------------------------------------------===//
-
-/// Returns the lowering configuration set for an operation.
-IREE::HAL::LoweringConfig getLoweringConfig(Operation *op);
-
-/// Sets the lowering configuration, overwriting existing attribute values.
-void setLoweringConfig(Operation *op, IREE::HAL::LoweringConfig config);
-
-/// Removes the lowering configuration on the operation if it exists.
-void eraseLoweringConfig(Operation *op);
-
-//===----------------------------------------------------------------------===//
-// Helpers for accessing values from the LoweringConfig attribute.
-//===----------------------------------------------------------------------===//
-
-// TODO(ravishankarm): Struct attributes dont have a way of defining extra class
-// methods. When they do, these could all be moved into the attribute definition
-// itself.
-
-/// Stores the tile sizes to use at different levels of tiling as a vector of
-/// vectors.
-/// - First level tiling maps to workgroups.
-/// - Second level tiling maps to subgroups.
-/// - Third level tiling maps to invocations.
-using TileSizesListType = SmallVector<SmallVector<int64_t, 4>, 1>;
-using TileSizesListTypeRef = ArrayRef<SmallVector<int64_t, 4>>;
-
-/// Construct a lowering configuration.
-IREE::HAL::LoweringConfig buildConfigAttr(TileSizesListTypeRef tileSizes,
- ArrayRef<int64_t> nativeVectorSize,
- MLIRContext *context);
-
-/// Get the tile sizes for all levels.
-TileSizesListType getTileSizes(IREE::HAL::LoweringConfig config);
-
-/// Get the tile sizes for all levels for an operation if the lowering
-/// configuration is set.
-inline TileSizesListType getTileSizes(Operation *op) {
- auto configAttr = getLoweringConfig(op);
- if (!configAttr) return {};
- return getTileSizes(configAttr);
-}
-
-/// Get the tile sizes for level `level`, if it is defined. Returns {} if tile
-/// sizes are not set for that level.
-SmallVector<int64_t, 4> getTileSizes(IREE::HAL::LoweringConfig config,
- unsigned level);
-
-/// Get the tile sizes for level `level` for an operation if the lowering
-/// configuration for the operation is set, and tile sizes are defined for that
-/// level.
-inline SmallVector<int64_t, 4> getTileSizes(Operation *op, unsigned level) {
- auto configAttr = getLoweringConfig(op);
- if (!configAttr) return {};
- return getTileSizes(configAttr, level);
-}
-SmallVector<Value, 4> getTileSizes(OpBuilder &b, Operation *op, unsigned level);
-
-/// Gets the native vector size defined in the lowering configuration.
-SmallVector<int64_t, 4> getNativeVectorSize(IREE::HAL::LoweringConfig config);
-
-/// Gets the native vector size defined for lowering an operation, if the
-/// lowering configuration is defined. If not returns empty vector.
-inline SmallVector<int64_t, 4> getNativeVectorSize(Operation *op) {
- auto configAttr = getLoweringConfig(op);
- if (!configAttr) return {};
- return getNativeVectorSize(configAttr);
-}
-
-/// Get the pass pipeline specified in the `loweringConfig`
-inline Optional<IREE::HAL::DispatchLoweringPassPipeline>
-getLoweringPassPipeline(IREE::HAL::LoweringConfig config) {
- return IREE::HAL::symbolizeDispatchLoweringPassPipeline(
- config.passPipeline().getValue());
-}
-
-} // namespace iree_compiler
-} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_COMMON_LOWERINGCONFIG_H_
diff --git a/iree/compiler/Dialect/HAL/IR/LoweringConfig.td b/iree/compiler/Dialect/HAL/IR/LoweringConfig.td
deleted file mode 100644
index 9e520ce..0000000
--- a/iree/compiler/Dialect/HAL/IR/LoweringConfig.td
+++ /dev/null
@@ -1,79 +0,0 @@
-// Copyright 2021 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_DIALECT_HAL_IR_LOWERINGCONFIG
-#define IREE_COMPILER_DIALECT_HAL_IR_LOWERINGCONFIG
-
-// Putting this in HAL dialect for now.
-include "iree/compiler/Dialect/HAL/IR/HALDialect.td"
-
-// List of pre-existing pipelines for translating executables.
-def CPU_Default
- : StrEnumAttrCase<"CPUDefault">;
-def CPU_Vectorization
- : StrEnumAttrCase<"CPUVectorization">;
-def CPU_TensorToVectors
- : StrEnumAttrCase<"CPUTensorToVectors">;
-
-def LLVMGPU_SimpleDistribute
- : StrEnumAttrCase<"LLVMGPUDistribute">;
-def LLVMGPU_Vectorize
- : StrEnumAttrCase<"LLVMGPUVectorize">;
-def LLVMGPU_MatmulSimt
- : StrEnumAttrCase<"LLVMGPUMatmulSimt">;
-
-def SPIRV_SimpleDistribute
- : StrEnumAttrCase<"SPIRVDistribute">;
-def SPIRV_DistributeToGlobalID
- : StrEnumAttrCase<"SPIRVDistributeToGlobalID">;
-def SPIRV_Vectorize
- : StrEnumAttrCase<"SPIRVVectorize">;
-def SPIRV_VectorizeToCooperativeOps
- : StrEnumAttrCase<"SPIRVVectorizeToCooperativeOps">;
-def None
- : StrEnumAttrCase<"None">;
-
-// EnumAttrCase for all known lowerings for ops within dispatch region
-// to scalar/native-vector code.
-def DispatchLoweringPassPipelineEnum : StrEnumAttr<
- "DispatchLoweringPassPipeline",
- "identifier for pass pipeline use to lower dispatch region",
- [CPU_Default, CPU_TensorToVectors, CPU_Vectorization,
- LLVMGPU_SimpleDistribute, LLVMGPU_Vectorize, LLVMGPU_MatmulSimt,
- SPIRV_SimpleDistribute, SPIRV_DistributeToGlobalID,
- SPIRV_Vectorize, SPIRV_VectorizeToCooperativeOps,
- None]> {
- let cppNamespace = "::mlir::iree_compiler::IREE::HAL";
-}
-
-def TileSizesListAttr :
- TypedArrayAttrBase<I64ArrayAttr,
- "list of tile sizes for all levels"> { }
-
-// Attribute that captures information needed for translating the executables.
-def TranslationInfoAttr :
- StructAttr<"TranslationInfo", HAL_Dialect, [
- StructFieldAttr<"passPipeline", DispatchLoweringPassPipelineEnum>,
- StructFieldAttr<"workloadPerWorkgroup",
- DefaultValuedAttr<I64ArrayAttr, "{}">>,
- ]>;
-
-// Attribute that carries information needed to perform
-// tiling/vectorization, etc.
-def HAL_LoweringConfigAttr :
- StructAttr<"LoweringConfig", HAL_Dialect, [
- StructFieldAttr<"tileSizes",
- DefaultValuedAttr<TileSizesListAttr, "{}">>,
- StructFieldAttr<"nativeVectorSize",
- DefaultValuedAttr<I64ArrayAttr, "{}">>,
- StructFieldAttr<"passPipeline",
- DefaultValuedAttr<
- DispatchLoweringPassPipelineEnum,
- "\"IREE::HAL::DispatchLoweringPassPipeline::None\"">>,
- StructFieldAttr<"workgroupSize",
- DefaultValuedAttr<I64ArrayAttr, "{}">>
- ]>;
-
-#endif // IREE_COMPILER_DIALECT_HAL_IR_LOWERINGCONFIG
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/BUILD b/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
index 0be0ecd..bf84382 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
@@ -42,6 +42,7 @@
deps = [
":cuda_libdevice",
"//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/LLVMGPU",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Utils",
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
index 51210f0..4974fa9 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
@@ -6,6 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.h"
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/CUDA/LLVMPasses.h"
#include "iree/compiler/Dialect/HAL/Target/CUDA/libdevice.h"
@@ -150,7 +151,7 @@
std::string name() const override { return "cuda"; }
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<gpu::GPUDialect>();
+ registry.insert<gpu::GPUDialect, IREE::Codegen::IREECodegenDialect>();
mlir::registerLLVMDialectTranslation(registry);
mlir::registerNVVMDialectTranslation(registry);
}
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
index 398b138..710eb0f 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
@@ -37,6 +37,7 @@
":StaticLibraryGenerator",
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/LLVMCPU",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/HAL/Target",
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
index 7c24113..f6ea0e4 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
@@ -46,6 +46,7 @@
MLIRLLVMToLLVMIRTranslation
MLIRTargetLLVMIRExport
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::LLVMCPU
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
index 4cf5cd7..5275258 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
@@ -8,6 +8,7 @@
#include <cstdlib>
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h"
#include "iree/compiler/Dialect/HAL/Target/LLVM/LibraryBuilder.h"
@@ -110,6 +111,7 @@
void getDependentDialects(DialectRegistry ®istry) const override {
mlir::registerLLVMDialectTranslation(registry);
+ registry.insert<IREE::Codegen::IREECodegenDialect>();
}
IREE::HAL::DeviceTargetAttr getDefaultDeviceTarget(
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
index 01f3323..6fd21ab 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
@@ -28,6 +28,7 @@
":SPIRVToMSL",
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/SPIRV",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/HAL/Target",
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
index d5fecf0..d01431a 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
@@ -30,6 +30,7 @@
MLIRSPIRVSerialization
MLIRVector
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::SPIRV
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/MetalSPIRVTarget.cpp b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/MetalSPIRVTarget.cpp
index ed74c70..6d8ec73 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/MetalSPIRVTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/MetalSPIRVTarget.cpp
@@ -6,6 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/MetalSPIRV/MetalSPIRVTarget.h"
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/MetalSPIRV/SPIRVToMSL.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
@@ -44,7 +45,8 @@
std::string name() const override { return "metal"; }
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<spirv::SPIRVDialect, gpu::GPUDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, spirv::SPIRVDialect,
+ gpu::GPUDialect>();
}
IREE::HAL::DeviceTargetAttr getDefaultDeviceTarget(
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/BUILD b/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
index 67ffe56..6c1f412 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
@@ -31,6 +31,7 @@
],
deps = [
"//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/LLVMGPU",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Utils",
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
index a57a8b8..041ca47 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
@@ -38,6 +38,7 @@
MLIRROCDLToLLVMIRTranslation
MLIRSupport
MLIRTargetLLVMIRExport
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::LLVMGPU
iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::HAL::Target
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp b/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
index 08f2643..90aaeb3 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
@@ -8,6 +8,7 @@
#include <mutex>
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
@@ -77,13 +78,13 @@
void getDependentDialects(DialectRegistry ®istry) const override {
mlir::registerLLVMDialectTranslation(registry);
mlir::registerROCDLDialectTranslation(registry);
+ registry.insert<IREE::Codegen::IREECodegenDialect>();
}
IREE::HAL::DeviceTargetAttr getDefaultDeviceTarget(
MLIRContext *context) const override {
Builder b(context);
SmallVector<NamedAttribute> configItems;
- ;
configItems.emplace_back(b.getIdentifier("executable_targets"),
getExecutableTargets(context));
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/BUILD b/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
index fdf19ce..4582466 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
@@ -30,6 +30,7 @@
],
deps = [
"//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/Modules/VMVX/IR:VMVXDialect",
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
index f21e2c0..757b3fd 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
@@ -26,6 +26,7 @@
MLIRIR
MLIRPass
MLIRSupport
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/VMVXTarget.cpp b/iree/compiler/Dialect/HAL/Target/VMVX/VMVXTarget.cpp
index 9206b75..8b3fa52 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/VMVXTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/VMVXTarget.cpp
@@ -6,6 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/VMVX/VMVXTarget.h"
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
#include "iree/compiler/Dialect/Modules/VMVX/IR/VMVXDialect.h"
@@ -35,7 +36,8 @@
std::string name() const override { return "vmvx"; }
void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<VM::VMDialect, VMVX::VMVXDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, VM::VMDialect,
+ VMVX::VMVXDialect>();
}
IREE::HAL::DeviceTargetAttr getDefaultDeviceTarget(
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
index 0345228..2d398cc 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
@@ -31,6 +31,7 @@
deps = [
"//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Codegen/SPIRV",
"//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
index d7f4590..d653bfe 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
@@ -34,6 +34,7 @@
MLIRSupport
MLIRVector
iree::compiler::Codegen::Common
+ iree::compiler::Codegen::Dialect::IREECodegenDialect
iree::compiler::Codegen::PassHeaders
iree::compiler::Codegen::SPIRV
iree::compiler::Codegen::Utils
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
index 3afdcba..9af89bc 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -6,6 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h"
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
@@ -119,8 +120,8 @@
std::string name() const override { return "vulkan"; }
void getDependentDialects(DialectRegistry ®istry) const override {
- registry
- .insert<Vulkan::VulkanDialect, spirv::SPIRVDialect, gpu::GPUDialect>();
+ registry.insert<IREE::Codegen::IREECodegenDialect, Vulkan::VulkanDialect,
+ spirv::SPIRVDialect, gpu::GPUDialect>();
}
IREE::HAL::DeviceTargetAttr getDefaultDeviceTarget(
diff --git a/iree/test/e2e/regression/lowering_config.mlir b/iree/test/e2e/regression/lowering_config.mlir
index a70c491..17b401f 100644
--- a/iree/test/e2e/regression/lowering_config.mlir
+++ b/iree/test/e2e/regression/lowering_config.mlir
@@ -1,11 +1,17 @@
-#config1 = {tileSizes = [[32, 32, 32]], passPipeline = 1 : i32}
-#config2 = {tileSizes = [[64, 64, 64]], passPipeline = 1 : i32}
+#compilation0 = #iree_codegen.compilation.info<
+ #iree_codegen.lowering.config<tile_sizes = [[32, 32, 32]], native_vector_size = []>,
+ #iree_codegen.translation.info<"CPUVectorization", workload_per_wg = [32, 32]>,
+ workgroup_size = []>
+#compilation1 = #iree_codegen.compilation.info<
+ #iree_codegen.lowering.config<tile_sizes = [[64, 64, 64]], native_vector_size = []>,
+ #iree_codegen.translation.info<"CPUVectorization", workload_per_wg = [64, 64]>,
+ workgroup_size = []>
func @lowering_config_test() {
%a = util.unfoldable_constant dense<1.0> : tensor<128x256xf32>
%b = util.unfoldable_constant dense<2.0> : tensor<256x512xf32>
%c = util.unfoldable_constant dense<2.0> : tensor<256x1024xf32>
- %d = "mhlo.dot"(%a, %b) {lowering.config = #config1} : (tensor<128x256xf32>, tensor<256x512xf32>) -> tensor<128x512xf32>
- %e = "mhlo.dot"(%a, %c) {lowering.config = #config2} : (tensor<128x256xf32>, tensor<256x1024xf32>) -> tensor<128x1024xf32>
+ %d = "mhlo.dot"(%a, %b) {compilation.info = #compilation0} : (tensor<128x256xf32>, tensor<256x512xf32>) -> tensor<128x512xf32>
+ %e = "mhlo.dot"(%a, %c) {compilation.info = #compilation1} : (tensor<128x256xf32>, tensor<256x1024xf32>) -> tensor<128x1024xf32>
check.expect_almost_eq_const(%d, dense<512.0> : tensor<128x512xf32>) : tensor<128x512xf32>
check.expect_almost_eq_const(%e, dense<512.0> : tensor<128x1024xf32>) : tensor<128x1024xf32>
return
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index 2723b38..9a68ab2 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -100,6 +100,7 @@
deps = [
"//iree/compiler/Bindings/Native/Transforms",
"//iree/compiler/Bindings/TFLite/Transforms",
+ "//iree/compiler/Codegen/Dialect:IREECodegenDialect",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/Flow/Transforms",
"//iree/compiler/Dialect/HAL/IR:HALDialect",
diff --git a/iree/tools/init_iree_dialects.h b/iree/tools/init_iree_dialects.h
index b374934..184af77 100644
--- a/iree/tools/init_iree_dialects.h
+++ b/iree/tools/init_iree_dialects.h
@@ -14,6 +14,7 @@
#include "iree-dialects/Dialect/IREE/IREEDialect.h"
#include "iree-dialects/Dialect/IREEPyDM/IR/Dialect.h"
+#include "iree/compiler/Codegen/Dialect/IREECodegenDialect.h"
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h"
@@ -32,7 +33,8 @@
// Add all the IREE dialects to the provided registry.
inline void registerIreeDialects(DialectRegistry ®istry) {
// clang-format off
- registry.insert<IREE::Flow::FlowDialect,
+ registry.insert<IREE::Codegen::IREECodegenDialect,
+ IREE::Flow::FlowDialect,
IREE::HAL::HALDialect,
ShapeDialect,
IREE::Stream::StreamDialect,
