[Codegen] Move remaining pipelines to `iree_codegen` attrs (#24398)
Follow-up to the LLVMGPU, SPIRV, and LLVMCPU pipeline migrations
(#23816, #23851, #23864). Same goal: get backend-specific pipeline
identifiers out of the global `DispatchLoweringPassPipeline` enum and
behind `PipelineAttrInterface`.
This moves the remaining non-CPU/GPU pipeline cases to `iree_codegen`:
`#iree_codegen.vmvx_pipeline`,
`#iree_codegen.transform_dialect_codegen`, and
`#iree_codegen.no_pipeline`. It keeps custom pass pipelines represented
as `#iree_codegen.pass_pipeline<...>`.
These attrs intentionally stay in `iree_codegen`. `vmvx_pipeline` does
not move to a backend dialect because IREE does not have a VMVX codegen
dialect today; it can become a backend-owned enum attr if such a dialect
is introduced later. `transform_dialect_codegen` and `no_pipeline` are
also codegen-level sentinels, so modeling the remaining cases as unit
attrs avoids creating dialect surface just for pipeline spelling.
After this, `TranslationInfoAttr` and `iree_codegen.smt.constraints`
parse pipeline attributes directly through the interface. The old enum,
keyword compatibility parser, helper predicates, and textual parsing in
Python dialect tests can be removed.
Issue: https://github.com/iree-org/iree/issues/23535
Assisted-by: codex
diff --git a/compiler/bindings/c/iree/compiler/dialects/iree_codegen.h b/compiler/bindings/c/iree/compiler/dialects/iree_codegen.h
index c236024..40948f8 100644
--- a/compiler/bindings/c/iree/compiler/dialects/iree_codegen.h
+++ b/compiler/bindings/c/iree/compiler/dialects/iree_codegen.h
@@ -19,17 +19,40 @@
// It mirrors the IREE Codegen Dialect which is not stable itself.
MLIR_CAPI_EXPORTED bool
-ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr(MlirAttribute attr);
+ireeAttributeIsACodegenVMVXPipelineAttr(MlirAttribute attr);
+
+MLIR_CAPI_EXPORTED MlirTypeID ireeCodegenVMVXPipelineAttrGetTypeID(void);
+
+MLIR_CAPI_EXPORTED MlirAttribute
+ireeCodegenVMVXPipelineAttrGet(MlirContext mlirCtx);
+
+MLIR_CAPI_EXPORTED bool
+ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr(MlirAttribute attr);
MLIR_CAPI_EXPORTED MlirTypeID
-ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID(void);
+ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID(void);
-MLIR_CAPI_EXPORTED MlirAttribute ireeCodegenDispatchLoweringPassPipelineAttrGet(
- MlirContext mlirCtx, uint32_t value);
+MLIR_CAPI_EXPORTED MlirAttribute
+ireeCodegenTransformDialectCodegenPipelineAttrGet(MlirContext mlirCtx);
-MLIR_CAPI_EXPORTED
-uint32_t
-ireeCodegenDispatchLoweringPassPipelineAttrGetValue(MlirAttribute attr);
+MLIR_CAPI_EXPORTED bool
+ireeAttributeIsACodegenNoPipelineAttr(MlirAttribute attr);
+
+MLIR_CAPI_EXPORTED MlirTypeID ireeCodegenNoPipelineAttrGetTypeID(void);
+
+MLIR_CAPI_EXPORTED MlirAttribute
+ireeCodegenNoPipelineAttrGet(MlirContext mlirCtx);
+
+MLIR_CAPI_EXPORTED bool
+ireeAttributeIsACodegenPassPipelineAttr(MlirAttribute attr);
+
+MLIR_CAPI_EXPORTED MlirTypeID ireeCodegenPassPipelineAttrGetTypeID(void);
+
+MLIR_CAPI_EXPORTED MlirAttribute
+ireeCodegenPassPipelineAttrGet(MlirContext mlirCtx, MlirStringRef pipeline);
+
+MLIR_CAPI_EXPORTED MlirStringRef
+ireeCodegenPassPipelineAttrGetPipeline(MlirAttribute attr);
MLIR_CAPI_EXPORTED bool
ireeAttributeIsACodegenTranslationInfoAttr(MlirAttribute attr);
@@ -37,8 +60,8 @@
MLIR_CAPI_EXPORTED MlirTypeID ireeCodegenTranslationInfoAttrGetTypeID(void);
struct ireeCodegenTranslationInfoParameters {
- // DispatchLoweringPassPipelineAttr or any attribute implementing
- // PipelineAttrInterface (e.g., #iree_gpu.pipeline<...>).
+ // Attribute implementing PipelineAttrInterface (e.g.,
+ // #iree_gpu.pipeline<...>).
MlirAttribute passPipeline;
MlirAttribute codegenSpec; // Optional SymbolRefAttr.
const int64_t *workgroupSize; // Optional ArrayRef<int64_t>.
diff --git a/compiler/bindings/python/IREECompilerDialectsModule.cpp b/compiler/bindings/python/IREECompilerDialectsModule.cpp
index 3a9570e..c435144 100644
--- a/compiler/bindings/python/IREECompilerDialectsModule.cpp
+++ b/compiler/bindings/python/IREECompilerDialectsModule.cpp
@@ -74,29 +74,57 @@
m.def_submodule("iree_codegen", "iree_codegen dialect bindings");
//===-------------------------------------------------------------------===//
- // CodegenDispatchLoweringPassPipelineAttr
+ // Codegen pipeline attrs
//===-------------------------------------------------------------------===//
- mlir_attribute_subclass(
- iree_codegen_module, "DispatchLoweringPassPipelineAttr",
- ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr,
- ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID)
+ mlir_attribute_subclass(iree_codegen_module, "VMVXPipelineAttr",
+ ireeAttributeIsACodegenVMVXPipelineAttr,
+ ireeCodegenVMVXPipelineAttrGetTypeID)
.def_classmethod(
"get",
- [](const py::object &, uint32_t value, MlirContext ctx) {
- return ireeCodegenDispatchLoweringPassPipelineAttrGet(ctx, value);
+ [](const py::object &, MlirContext ctx) {
+ return ireeCodegenVMVXPipelineAttrGet(ctx);
},
- "cls"_a, "value"_a, "ctx"_a = py::none(),
- "Gets an #iree_codegen.dispatch_lowering_pass_pipeline from "
- "parameters.")
- .def_property_readonly(
- "raw_value", ireeCodegenDispatchLoweringPassPipelineAttrGetValue)
- .def_property_readonly("value", [](MlirAttribute self) -> py::object {
- uint32_t rawValue =
- ireeCodegenDispatchLoweringPassPipelineAttrGetValue(self);
- return py::module_::import_(kCodegenModuleImportPath)
- .attr("DispatchLoweringPassPipeline")(rawValue);
- });
+ "cls"_a, "ctx"_a = py::none(),
+ "Gets an #iree_codegen.vmvx_pipeline attribute.");
+
+ mlir_attribute_subclass(
+ iree_codegen_module, "TransformDialectCodegenPipelineAttr",
+ ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr,
+ ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID)
+ .def_classmethod(
+ "get",
+ [](const py::object &, MlirContext ctx) {
+ return ireeCodegenTransformDialectCodegenPipelineAttrGet(ctx);
+ },
+ "cls"_a, "ctx"_a = py::none(),
+ "Gets an #iree_codegen.transform_dialect_codegen attribute.");
+
+ mlir_attribute_subclass(iree_codegen_module, "NoPipelineAttr",
+ ireeAttributeIsACodegenNoPipelineAttr,
+ ireeCodegenNoPipelineAttrGetTypeID)
+ .def_classmethod(
+ "get",
+ [](const py::object &, MlirContext ctx) {
+ return ireeCodegenNoPipelineAttrGet(ctx);
+ },
+ "cls"_a, "ctx"_a = py::none(),
+ "Gets an #iree_codegen.no_pipeline attribute.");
+
+ mlir_attribute_subclass(iree_codegen_module, "PassPipelineAttr",
+ ireeAttributeIsACodegenPassPipelineAttr,
+ ireeCodegenPassPipelineAttrGetTypeID)
+ .def_classmethod(
+ "get",
+ [](const py::object &, std::string pipeline, MlirContext ctx) {
+ MlirStringRef pipelineRef =
+ mlirStringRefCreate(pipeline.data(), pipeline.size());
+ return ireeCodegenPassPipelineAttrGet(ctx, pipelineRef);
+ },
+ "cls"_a, "pipeline"_a, "ctx"_a = py::none(),
+ "Gets an #iree_codegen.pass_pipeline attribute.")
+ .def_property_readonly("pipeline",
+ ireeCodegenPassPipelineAttrGetPipeline);
//===-------------------------------------------------------------------===//
// CodegenTranslationInfoAttr
diff --git a/compiler/bindings/python/test/ir/dialects_test.py b/compiler/bindings/python/test/ir/dialects_test.py
index a5664f9..2b80a3b 100644
--- a/compiler/bindings/python/test/ir/dialects_test.py
+++ b/compiler/bindings/python/test/ir/dialects_test.py
@@ -74,16 +74,38 @@
@run
-def codegen_dispatch_lowering_pass_pipeline():
- pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
- iree_codegen.DispatchLoweringPassPipeline.VMVXDefault
- )
+def codegen_vmvx_pipeline_attr():
+ pipeline_attr = iree_codegen.VMVXPipelineAttr.get()
assert pipeline_attr is not None
- assert pipeline_attr.value == iree_codegen.DispatchLoweringPassPipeline.VMVXDefault
- assert pipeline_attr.raw_value == int(
- iree_codegen.DispatchLoweringPassPipeline.VMVXDefault
+ assert isinstance(pipeline_attr, iree_codegen.VMVXPipelineAttr)
+ assert str(pipeline_attr) == "#iree_codegen.vmvx_pipeline"
+
+
+@run
+def codegen_transform_dialect_codegen_pipeline_attr():
+ pipeline_attr = iree_codegen.TransformDialectCodegenPipelineAttr.get()
+ assert pipeline_attr is not None
+ assert isinstance(pipeline_attr, iree_codegen.TransformDialectCodegenPipelineAttr)
+ assert str(pipeline_attr) == "#iree_codegen.transform_dialect_codegen"
+
+
+@run
+def codegen_no_pipeline_attr():
+ pipeline_attr = iree_codegen.NoPipelineAttr.get()
+ assert pipeline_attr is not None
+ assert isinstance(pipeline_attr, iree_codegen.NoPipelineAttr)
+ assert str(pipeline_attr) == "#iree_codegen.no_pipeline"
+
+
+@run
+def codegen_pass_pipeline_attr():
+ pipeline_attr = iree_codegen.PassPipelineAttr.get("func.func(canonicalize)")
+ assert pipeline_attr is not None
+ assert isinstance(pipeline_attr, iree_codegen.PassPipelineAttr)
+ assert pipeline_attr.pipeline == "func.func(canonicalize)"
+ assert (
+ str(pipeline_attr) == '#iree_codegen.pass_pipeline<"func.func(canonicalize)">'
)
- assert "VMVXDefault" in str(pipeline_attr)
@run
@@ -97,12 +119,13 @@
@run
def codegen_translation_info_minimal():
- pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
- iree_codegen.DispatchLoweringPassPipeline.None_
- )
+ pipeline_attr = iree_codegen.NoPipelineAttr.get()
translation_info = iree_codegen.TranslationInfoAttr.get(pipeline_attr)
assert translation_info is not None
- assert str(translation_info) == "#iree_codegen.translation_info<pipeline = None>"
+ assert (
+ str(translation_info)
+ == "#iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline>"
+ )
assert translation_info.pass_pipeline == pipeline_attr
assert translation_info.codegen_spec is None
assert translation_info.workgroup_size == []
@@ -112,14 +135,13 @@
@run
def codegen_translation_info_with_sizes():
- pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
- iree_codegen.DispatchLoweringPassPipeline.Custom
- )
+ pipeline_attr = iree_codegen.PassPipelineAttr.get("canonicalize")
translation_info = iree_codegen.TranslationInfoAttr.get(
pipeline_attr, None, [64, 4, 1], 32
)
assert translation_info is not None
assert translation_info.pass_pipeline == pipeline_attr
+ assert pipeline_attr.pipeline == "canonicalize"
assert translation_info.codegen_spec is None
assert translation_info.workgroup_size == [64, 4, 1]
assert translation_info.subgroup_size == 32
@@ -128,9 +150,7 @@
@run
def codegen_translation_info_full():
- pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
- iree_codegen.DispatchLoweringPassPipeline.TransformDialectCodegen
- )
+ pipeline_attr = iree_codegen.TransformDialectCodegenPipelineAttr.get()
foo_symbol = ir.SymbolRefAttr.get(["foo"])
configuration = ir.DictAttr.get({"A": ir.IntegerAttr.get(ir.IndexType.get(), 42)})
translation_info = iree_codegen.TranslationInfoAttr.get(
@@ -449,9 +469,7 @@
def compilation_info():
attributes = ir.DictAttr.get({"reduction": get_index_array_attr([])})
lowering_config = iree_gpu.LoweringConfigAttr.get(attributes)
- pipeline_attr = iree_codegen.DispatchLoweringPassPipelineAttr.get(
- iree_codegen.DispatchLoweringPassPipeline.None_
- )
+ pipeline_attr = iree_codegen.NoPipelineAttr.get()
translation_info = iree_codegen.TranslationInfoAttr.get(pipeline_attr)
compilation_info = iree_codegen.CompilationInfoAttr.get(
diff --git a/compiler/src/iree/compiler/API/Internal/IREECodegenDialectCAPI.cpp b/compiler/src/iree/compiler/API/Internal/IREECodegenDialectCAPI.cpp
index 75c723b..79c8819 100644
--- a/compiler/src/iree/compiler/API/Internal/IREECodegenDialectCAPI.cpp
+++ b/compiler/src/iree/compiler/API/Internal/IREECodegenDialectCAPI.cpp
@@ -7,7 +7,7 @@
#include <cassert>
#include <cstdint>
#include <optional>
-#include <type_traits>
+
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenInterfaces.h"
@@ -35,41 +35,71 @@
using mlir::iree_compiler::IREE::Codegen::CompilationInfoAttr;
using mlir::iree_compiler::IREE::Codegen::ConstraintsOp;
-using mlir::iree_compiler::IREE::Codegen::DispatchLoweringPassPipeline;
-using mlir::iree_compiler::IREE::Codegen::DispatchLoweringPassPipelineAttr;
using mlir::iree_compiler::IREE::Codegen::IntKnobAttr;
using mlir::iree_compiler::IREE::Codegen::LoweringConfigAttrInterface;
+using mlir::iree_compiler::IREE::Codegen::NoPipelineAttr;
using mlir::iree_compiler::IREE::Codegen::OneOfKnobAttr;
+using mlir::iree_compiler::IREE::Codegen::PassPipelineAttr;
+using mlir::iree_compiler::IREE::Codegen::PipelineAttrInterface;
using mlir::iree_compiler::IREE::Codegen::RootOpAttr;
+using mlir::iree_compiler::IREE::Codegen::TransformDialectCodegenPipelineAttr;
using mlir::iree_compiler::IREE::Codegen::TranslationInfoAttr;
+using mlir::iree_compiler::IREE::Codegen::VMVXPipelineAttr;
using mlir::iree_compiler::IREE::HAL::ExecutableVariantOp;
-bool ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr(
+bool ireeAttributeIsACodegenVMVXPipelineAttr(MlirAttribute attr) {
+ return llvm::isa<VMVXPipelineAttr>(unwrap(attr));
+}
+
+MlirTypeID ireeCodegenVMVXPipelineAttrGetTypeID() {
+ return wrap(VMVXPipelineAttr::getTypeID());
+}
+
+MlirAttribute ireeCodegenVMVXPipelineAttrGet(MlirContext mlirCtx) {
+ return wrap(VMVXPipelineAttr::get(unwrap(mlirCtx)));
+}
+
+bool ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr(
MlirAttribute attr) {
- return llvm::isa<DispatchLoweringPassPipelineAttr>(unwrap(attr));
+ return llvm::isa<TransformDialectCodegenPipelineAttr>(unwrap(attr));
}
-MlirTypeID ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID() {
- return wrap(DispatchLoweringPassPipelineAttr::getTypeID());
+MlirTypeID ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID() {
+ return wrap(TransformDialectCodegenPipelineAttr::getTypeID());
}
-static_assert(
- std::is_same_v<uint32_t,
- std::underlying_type_t<DispatchLoweringPassPipeline>>,
- "Enum type changed");
-
MlirAttribute
-ireeCodegenDispatchLoweringPassPipelineAttrGet(MlirContext mlirCtx,
- uint32_t value) {
- mlir::MLIRContext *ctx = unwrap(mlirCtx);
- return wrap(DispatchLoweringPassPipelineAttr::get(
- ctx, static_cast<DispatchLoweringPassPipeline>(value)));
+ireeCodegenTransformDialectCodegenPipelineAttrGet(MlirContext mlirCtx) {
+ return wrap(TransformDialectCodegenPipelineAttr::get(unwrap(mlirCtx)));
}
-uint32_t
-ireeCodegenDispatchLoweringPassPipelineAttrGetValue(MlirAttribute attr) {
- return static_cast<uint32_t>(
- llvm::cast<DispatchLoweringPassPipelineAttr>(unwrap(attr)).getValue());
+bool ireeAttributeIsACodegenNoPipelineAttr(MlirAttribute attr) {
+ return llvm::isa<NoPipelineAttr>(unwrap(attr));
+}
+
+MlirTypeID ireeCodegenNoPipelineAttrGetTypeID() {
+ return wrap(NoPipelineAttr::getTypeID());
+}
+
+MlirAttribute ireeCodegenNoPipelineAttrGet(MlirContext mlirCtx) {
+ return wrap(NoPipelineAttr::get(unwrap(mlirCtx)));
+}
+
+bool ireeAttributeIsACodegenPassPipelineAttr(MlirAttribute attr) {
+ return llvm::isa<PassPipelineAttr>(unwrap(attr));
+}
+
+MlirTypeID ireeCodegenPassPipelineAttrGetTypeID() {
+ return wrap(PassPipelineAttr::getTypeID());
+}
+
+MlirAttribute ireeCodegenPassPipelineAttrGet(MlirContext mlirCtx,
+ MlirStringRef pipeline) {
+ return wrap(PassPipelineAttr::get(unwrap(mlirCtx), unwrap(pipeline)));
+}
+
+MlirStringRef ireeCodegenPassPipelineAttrGetPipeline(MlirAttribute attr) {
+ return wrap(llvm::cast<PassPipelineAttr>(unwrap(attr)).getPipeline());
}
bool ireeAttributeIsACodegenTranslationInfoAttr(MlirAttribute attr) {
@@ -85,11 +115,10 @@
assert(!mlirAttributeIsNull(parameters.passPipeline) &&
"Invalid pass pipeline attr: cannot be null");
+ mlir::Attribute pipelineAttr = unwrap(parameters.passPipeline);
assert(
- (ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr(
- parameters.passPipeline) ||
- ireeAttributeIsAGPUPipelineAttr(parameters.passPipeline)) &&
- "passPipeline must be DispatchLoweringPassPipelineAttr or PipelineAttr");
+ pipelineAttr.hasPromiseOrImplementsInterface<PipelineAttrInterface>() &&
+ "passPipeline must implement PipelineAttrInterface");
assert((mlirAttributeIsNull(parameters.codegenSpec) ||
mlirAttributeIsASymbolRef(parameters.codegenSpec)) &&
@@ -99,7 +128,6 @@
mlirAttributeIsADictionary(parameters.configuration)) &&
"Invalid configuration attr");
- mlir::Attribute pipelineAttr = unwrap(parameters.passPipeline);
auto codegenSpec = llvm::cast_if_present<mlir::SymbolRefAttr>(
unwrap(parameters.codegenSpec));
diff --git a/compiler/src/iree/compiler/API/api_exports.c b/compiler/src/iree/compiler/API/api_exports.c
index cad40ac..1521ba2 100644
--- a/compiler/src/iree/compiler/API/api_exports.c
+++ b/compiler/src/iree/compiler/API/api_exports.c
@@ -11,11 +11,14 @@
#include <stdint.h>
extern void ireeAttributeIsACodegenCompilationInfoAttr();
-extern void ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr();
extern void ireeAttributeIsACodegenIntKnobAttr();
+extern void ireeAttributeIsACodegenNoPipelineAttr();
extern void ireeAttributeIsACodegenOneOfKnobAttr();
+extern void ireeAttributeIsACodegenPassPipelineAttr();
extern void ireeAttributeIsACodegenRootOpAttr();
+extern void ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr();
extern void ireeAttributeIsACodegenTranslationInfoAttr();
+extern void ireeAttributeIsACodegenVMVXPipelineAttr();
extern void ireeAttributeIsAGPULoweringConfigAttr();
extern void ireeAttributeIsAGPUMMAAttr();
extern void ireeAttributeIsAGPUMMAIntrinsicAttr();
@@ -28,9 +31,6 @@
extern void ireeCodegenCompilationInfoAttrGetParameters();
extern void ireeCodegenCompilationInfoAttrGetTypeID();
extern void ireeCodegenConvertConstraintsOpToSMTLIB();
-extern void ireeCodegenDispatchLoweringPassPipelineAttrGet();
-extern void ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID();
-extern void ireeCodegenDispatchLoweringPassPipelineAttrGetValue();
extern void ireeCodegenGetAttentionOpDetail();
extern void ireeCodegenGetExecutableVariantOps();
extern void ireeCodegenGetIGEMMGenericConvDetails();
@@ -41,15 +41,24 @@
extern void ireeCodegenIntKnobAttrGetTypeID();
extern void ireeCodegenMlirOperationIsACodegenAttentionOp();
extern void ireeCodegenMlirOperationIsAScaledContractionOp();
+extern void ireeCodegenNoPipelineAttrGet();
+extern void ireeCodegenNoPipelineAttrGetTypeID();
extern void ireeCodegenOneOfKnobAttrGetName();
extern void ireeCodegenOneOfKnobAttrGetOptions();
extern void ireeCodegenOneOfKnobAttrGetTypeID();
+extern void ireeCodegenPassPipelineAttrGet();
+extern void ireeCodegenPassPipelineAttrGetPipeline();
+extern void ireeCodegenPassPipelineAttrGetTypeID();
extern void ireeCodegenRootOpAttrGet();
extern void ireeCodegenRootOpAttrGetSet();
extern void ireeCodegenRootOpAttrGetTypeID();
+extern void ireeCodegenTransformDialectCodegenPipelineAttrGet();
+extern void ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID();
extern void ireeCodegenTranslationInfoAttrGet();
extern void ireeCodegenTranslationInfoAttrGetParameters();
extern void ireeCodegenTranslationInfoAttrGetTypeID();
+extern void ireeCodegenVMVXPipelineAttrGet();
+extern void ireeCodegenVMVXPipelineAttrGetTypeID();
extern void ireeCompilerEnumeratePlugins();
extern void ireeCompilerEnumerateRegisteredHALTargetBackends();
extern void ireeCompilerErrorDestroy();
@@ -1179,11 +1188,14 @@
uintptr_t __iree_compiler_hidden_force_extern() {
uintptr_t x = 0;
x += (uintptr_t)&ireeAttributeIsACodegenCompilationInfoAttr;
- x += (uintptr_t)&ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr;
x += (uintptr_t)&ireeAttributeIsACodegenIntKnobAttr;
+ x += (uintptr_t)&ireeAttributeIsACodegenNoPipelineAttr;
x += (uintptr_t)&ireeAttributeIsACodegenOneOfKnobAttr;
+ x += (uintptr_t)&ireeAttributeIsACodegenPassPipelineAttr;
x += (uintptr_t)&ireeAttributeIsACodegenRootOpAttr;
+ x += (uintptr_t)&ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr;
x += (uintptr_t)&ireeAttributeIsACodegenTranslationInfoAttr;
+ x += (uintptr_t)&ireeAttributeIsACodegenVMVXPipelineAttr;
x += (uintptr_t)&ireeAttributeIsAGPULoweringConfigAttr;
x += (uintptr_t)&ireeAttributeIsAGPUMMAAttr;
x += (uintptr_t)&ireeAttributeIsAGPUMMAIntrinsicAttr;
@@ -1196,9 +1208,6 @@
x += (uintptr_t)&ireeCodegenCompilationInfoAttrGetParameters;
x += (uintptr_t)&ireeCodegenCompilationInfoAttrGetTypeID;
x += (uintptr_t)&ireeCodegenConvertConstraintsOpToSMTLIB;
- x += (uintptr_t)&ireeCodegenDispatchLoweringPassPipelineAttrGet;
- x += (uintptr_t)&ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID;
- x += (uintptr_t)&ireeCodegenDispatchLoweringPassPipelineAttrGetValue;
x += (uintptr_t)&ireeCodegenGetAttentionOpDetail;
x += (uintptr_t)&ireeCodegenGetExecutableVariantOps;
x += (uintptr_t)&ireeCodegenGetIGEMMGenericConvDetails;
@@ -1209,15 +1218,24 @@
x += (uintptr_t)&ireeCodegenIntKnobAttrGetTypeID;
x += (uintptr_t)&ireeCodegenMlirOperationIsACodegenAttentionOp;
x += (uintptr_t)&ireeCodegenMlirOperationIsAScaledContractionOp;
+ x += (uintptr_t)&ireeCodegenNoPipelineAttrGet;
+ x += (uintptr_t)&ireeCodegenNoPipelineAttrGetTypeID;
x += (uintptr_t)&ireeCodegenOneOfKnobAttrGetName;
x += (uintptr_t)&ireeCodegenOneOfKnobAttrGetOptions;
x += (uintptr_t)&ireeCodegenOneOfKnobAttrGetTypeID;
+ x += (uintptr_t)&ireeCodegenPassPipelineAttrGet;
+ x += (uintptr_t)&ireeCodegenPassPipelineAttrGetPipeline;
+ x += (uintptr_t)&ireeCodegenPassPipelineAttrGetTypeID;
x += (uintptr_t)&ireeCodegenRootOpAttrGet;
x += (uintptr_t)&ireeCodegenRootOpAttrGetSet;
x += (uintptr_t)&ireeCodegenRootOpAttrGetTypeID;
+ x += (uintptr_t)&ireeCodegenTransformDialectCodegenPipelineAttrGet;
+ x += (uintptr_t)&ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID;
x += (uintptr_t)&ireeCodegenTranslationInfoAttrGet;
x += (uintptr_t)&ireeCodegenTranslationInfoAttrGetParameters;
x += (uintptr_t)&ireeCodegenTranslationInfoAttrGetTypeID;
+ x += (uintptr_t)&ireeCodegenVMVXPipelineAttrGet;
+ x += (uintptr_t)&ireeCodegenVMVXPipelineAttrGetTypeID;
x += (uintptr_t)&ireeCompilerEnumeratePlugins;
x += (uintptr_t)&ireeCompilerEnumerateRegisteredHALTargetBackends;
x += (uintptr_t)&ireeCompilerErrorDestroy;
diff --git a/compiler/src/iree/compiler/API/api_exports.def b/compiler/src/iree/compiler/API/api_exports.def
index f1851a7..5892229 100644
--- a/compiler/src/iree/compiler/API/api_exports.def
+++ b/compiler/src/iree/compiler/API/api_exports.def
@@ -1,11 +1,14 @@
; Generated by generate_exports.py: Do not edit.
EXPORTS
ireeAttributeIsACodegenCompilationInfoAttr
- ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr
ireeAttributeIsACodegenIntKnobAttr
+ ireeAttributeIsACodegenNoPipelineAttr
ireeAttributeIsACodegenOneOfKnobAttr
+ ireeAttributeIsACodegenPassPipelineAttr
ireeAttributeIsACodegenRootOpAttr
+ ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr
ireeAttributeIsACodegenTranslationInfoAttr
+ ireeAttributeIsACodegenVMVXPipelineAttr
ireeAttributeIsAGPULoweringConfigAttr
ireeAttributeIsAGPUMMAAttr
ireeAttributeIsAGPUMMAIntrinsicAttr
@@ -18,9 +21,6 @@
ireeCodegenCompilationInfoAttrGetParameters
ireeCodegenCompilationInfoAttrGetTypeID
ireeCodegenConvertConstraintsOpToSMTLIB
- ireeCodegenDispatchLoweringPassPipelineAttrGet
- ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID
- ireeCodegenDispatchLoweringPassPipelineAttrGetValue
ireeCodegenGetAttentionOpDetail
ireeCodegenGetExecutableVariantOps
ireeCodegenGetIGEMMGenericConvDetails
@@ -31,15 +31,24 @@
ireeCodegenIntKnobAttrGetTypeID
ireeCodegenMlirOperationIsACodegenAttentionOp
ireeCodegenMlirOperationIsAScaledContractionOp
+ ireeCodegenNoPipelineAttrGet
+ ireeCodegenNoPipelineAttrGetTypeID
ireeCodegenOneOfKnobAttrGetName
ireeCodegenOneOfKnobAttrGetOptions
ireeCodegenOneOfKnobAttrGetTypeID
+ ireeCodegenPassPipelineAttrGet
+ ireeCodegenPassPipelineAttrGetPipeline
+ ireeCodegenPassPipelineAttrGetTypeID
ireeCodegenRootOpAttrGet
ireeCodegenRootOpAttrGetSet
ireeCodegenRootOpAttrGetTypeID
+ ireeCodegenTransformDialectCodegenPipelineAttrGet
+ ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID
ireeCodegenTranslationInfoAttrGet
ireeCodegenTranslationInfoAttrGetParameters
ireeCodegenTranslationInfoAttrGetTypeID
+ ireeCodegenVMVXPipelineAttrGet
+ ireeCodegenVMVXPipelineAttrGetTypeID
ireeCompilerEnumeratePlugins
ireeCompilerEnumerateRegisteredHALTargetBackends
ireeCompilerErrorDestroy
diff --git a/compiler/src/iree/compiler/API/api_exports.ld b/compiler/src/iree/compiler/API/api_exports.ld
index 8234447..4040ec7 100644
--- a/compiler/src/iree/compiler/API/api_exports.ld
+++ b/compiler/src/iree/compiler/API/api_exports.ld
@@ -2,11 +2,14 @@
VER_0 {
global:
ireeAttributeIsACodegenCompilationInfoAttr;
- ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr;
ireeAttributeIsACodegenIntKnobAttr;
+ ireeAttributeIsACodegenNoPipelineAttr;
ireeAttributeIsACodegenOneOfKnobAttr;
+ ireeAttributeIsACodegenPassPipelineAttr;
ireeAttributeIsACodegenRootOpAttr;
+ ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr;
ireeAttributeIsACodegenTranslationInfoAttr;
+ ireeAttributeIsACodegenVMVXPipelineAttr;
ireeAttributeIsAGPULoweringConfigAttr;
ireeAttributeIsAGPUMMAAttr;
ireeAttributeIsAGPUMMAIntrinsicAttr;
@@ -19,9 +22,6 @@
ireeCodegenCompilationInfoAttrGetParameters;
ireeCodegenCompilationInfoAttrGetTypeID;
ireeCodegenConvertConstraintsOpToSMTLIB;
- ireeCodegenDispatchLoweringPassPipelineAttrGet;
- ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID;
- ireeCodegenDispatchLoweringPassPipelineAttrGetValue;
ireeCodegenGetAttentionOpDetail;
ireeCodegenGetExecutableVariantOps;
ireeCodegenGetIGEMMGenericConvDetails;
@@ -32,15 +32,24 @@
ireeCodegenIntKnobAttrGetTypeID;
ireeCodegenMlirOperationIsACodegenAttentionOp;
ireeCodegenMlirOperationIsAScaledContractionOp;
+ ireeCodegenNoPipelineAttrGet;
+ ireeCodegenNoPipelineAttrGetTypeID;
ireeCodegenOneOfKnobAttrGetName;
ireeCodegenOneOfKnobAttrGetOptions;
ireeCodegenOneOfKnobAttrGetTypeID;
+ ireeCodegenPassPipelineAttrGet;
+ ireeCodegenPassPipelineAttrGetPipeline;
+ ireeCodegenPassPipelineAttrGetTypeID;
ireeCodegenRootOpAttrGet;
ireeCodegenRootOpAttrGetSet;
ireeCodegenRootOpAttrGetTypeID;
+ ireeCodegenTransformDialectCodegenPipelineAttrGet;
+ ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID;
ireeCodegenTranslationInfoAttrGet;
ireeCodegenTranslationInfoAttrGetParameters;
ireeCodegenTranslationInfoAttrGetTypeID;
+ ireeCodegenVMVXPipelineAttrGet;
+ ireeCodegenVMVXPipelineAttrGetTypeID;
ireeCompilerEnumeratePlugins;
ireeCompilerEnumerateRegisteredHALTargetBackends;
ireeCompilerErrorDestroy;
diff --git a/compiler/src/iree/compiler/API/api_exports.macos.lst b/compiler/src/iree/compiler/API/api_exports.macos.lst
index 0fb74bc..f9e2a0d 100644
--- a/compiler/src/iree/compiler/API/api_exports.macos.lst
+++ b/compiler/src/iree/compiler/API/api_exports.macos.lst
@@ -1,10 +1,13 @@
# Generated by generate_exports.py: Do not edit.
_ireeAttributeIsACodegenCompilationInfoAttr
-_ireeAttributeIsACodegenDispatchLoweringPassPipelineAttr
_ireeAttributeIsACodegenIntKnobAttr
+_ireeAttributeIsACodegenNoPipelineAttr
_ireeAttributeIsACodegenOneOfKnobAttr
+_ireeAttributeIsACodegenPassPipelineAttr
_ireeAttributeIsACodegenRootOpAttr
+_ireeAttributeIsACodegenTransformDialectCodegenPipelineAttr
_ireeAttributeIsACodegenTranslationInfoAttr
+_ireeAttributeIsACodegenVMVXPipelineAttr
_ireeAttributeIsAGPULoweringConfigAttr
_ireeAttributeIsAGPUMMAAttr
_ireeAttributeIsAGPUMMAIntrinsicAttr
@@ -17,9 +20,6 @@
_ireeCodegenCompilationInfoAttrGetParameters
_ireeCodegenCompilationInfoAttrGetTypeID
_ireeCodegenConvertConstraintsOpToSMTLIB
-_ireeCodegenDispatchLoweringPassPipelineAttrGet
-_ireeCodegenDispatchLoweringPassPipelineAttrGetTypeID
-_ireeCodegenDispatchLoweringPassPipelineAttrGetValue
_ireeCodegenGetAttentionOpDetail
_ireeCodegenGetExecutableVariantOps
_ireeCodegenGetIGEMMGenericConvDetails
@@ -30,15 +30,24 @@
_ireeCodegenIntKnobAttrGetTypeID
_ireeCodegenMlirOperationIsACodegenAttentionOp
_ireeCodegenMlirOperationIsAScaledContractionOp
+_ireeCodegenNoPipelineAttrGet
+_ireeCodegenNoPipelineAttrGetTypeID
_ireeCodegenOneOfKnobAttrGetName
_ireeCodegenOneOfKnobAttrGetOptions
_ireeCodegenOneOfKnobAttrGetTypeID
+_ireeCodegenPassPipelineAttrGet
+_ireeCodegenPassPipelineAttrGetPipeline
+_ireeCodegenPassPipelineAttrGetTypeID
_ireeCodegenRootOpAttrGet
_ireeCodegenRootOpAttrGetSet
_ireeCodegenRootOpAttrGetTypeID
+_ireeCodegenTransformDialectCodegenPipelineAttrGet
+_ireeCodegenTransformDialectCodegenPipelineAttrGetTypeID
_ireeCodegenTranslationInfoAttrGet
_ireeCodegenTranslationInfoAttrGetParameters
_ireeCodegenTranslationInfoAttrGetTypeID
+_ireeCodegenVMVXPipelineAttrGet
+_ireeCodegenVMVXPipelineAttrGetTypeID
_ireeCompilerEnumeratePlugins
_ireeCompilerEnumerateRegisteredHALTargetBackends
_ireeCompilerErrorDestroy
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/expand_gpu_ops_scan.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/expand_gpu_ops_scan.mlir
index d19fa03..0817216 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/expand_gpu_ops_scan.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/expand_gpu_ops_scan.mlir
@@ -4,7 +4,7 @@
// Expects 2 shuffle steps for inclusive scan, then total shuffle from last
// lane, then exclusive shift shuffle + select identity at lane 0.
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
func.func @scan_add_f32(%x: f32, %zero: f32) -> (f32, f32) attributes {translation_info = #translation_info} {
%scan, %total = iree_gpu.subgroup_scan(%x, identity = %zero : f32) cluster(size = 4) {
^bb0(%lhs: f32, %rhs: f32):
@@ -59,7 +59,7 @@
// Clustered scan with cluster_size=4, stride=16.
// Expects shuffle offsets of 16, 32. lanePos = (laneId / 16) % 4.
-#translation_info2 = #iree_codegen.translation_info<pipeline = None workgroup_size = [64, 1, 1] subgroup_size = 64>
+#translation_info2 = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [64, 1, 1] subgroup_size = 64>
func.func @scan_add_f32_stride16(%x: f32, %zero: f32) -> (f32, f32) attributes {translation_info = #translation_info2} {
%scan, %total = iree_gpu.subgroup_scan(%x, identity = %zero : f32) cluster(size = 4, stride = 16) {
^bb0(%lhs: f32, %rhs: f32):
@@ -116,7 +116,7 @@
// i32 scan: no bitcast needed.
-#translation_info3 = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info3 = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
func.func @scan_add_i32(%x: i32, %zero: i32) -> (i32, i32) attributes {translation_info = #translation_info3} {
%scan, %total = iree_gpu.subgroup_scan(%x, identity = %zero : i32) cluster(size = 4) {
^bb0(%lhs: i32, %rhs: i32):
@@ -148,7 +148,7 @@
// f16 scan: requires bitcast to i32 for shuffle.
-#translation_info4 = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info4 = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
func.func @scan_add_f16(%x: f16, %zero: f16) -> (f16, f16) attributes {translation_info = #translation_info4} {
%scan, %total = iree_gpu.subgroup_scan(%x, identity = %zero : f16) cluster(size = 4) {
^bb0(%lhs: f16, %rhs: f16):
@@ -187,7 +187,7 @@
// Expects 5 shuffle steps for inclusive scan (offsets 1, 2, 4, 8, 16),
// then total shuffle + exclusive shift.
-#translation_info5 = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info5 = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
func.func @scan_full_subgroup(%x: f32, %zero: f32) -> (f32, f32) attributes {translation_info = #translation_info5} {
%scan, %total = iree_gpu.subgroup_scan(%x, identity = %zero : f32) {
^bb0(%lhs: f32, %rhs: f32):
@@ -222,7 +222,7 @@
// Expects 2 shuffle steps for inclusive scan, then total shuffle from last
// lane. No exclusive shift shuffle or identity select.
-#translation_info6 = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info6 = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
func.func @scan_inclusive_add_f32(%x: f32) -> (f32, f32) attributes {translation_info = #translation_info6} {
%scan, %total = iree_gpu.subgroup_scan inclusive (%x) cluster(size = 4) {
^bb0(%lhs: f32, %rhs: f32):
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_distribute_shared_memory.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_distribute_shared_memory.mlir
index 7c487f6..c086e4f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_distribute_shared_memory.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/gpu_distribute_shared_memory.mlir
@@ -1,7 +1,7 @@
// RUN: iree-opt --split-input-file --pass-pipeline='builtin.module(func.func(iree-codegen-gpu-distribute-shared-memory-copy, fold-memref-alias-ops, canonicalize, cse))' %s | FileCheck %s
#executable_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 4, 1]>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 4, 1]>
module {
memref.global "private" @__shared_memory___1 : memref<3x512xf32, 3>
memref.global "private" @__shared_memory___0 : memref<256x4xf32, 3>
@@ -81,7 +81,7 @@
// -----
#executable_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 8, 1]>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 8, 1]>
module {
func.func @unaligned_shared_memory_copy(
@@ -127,7 +127,7 @@
// -----
#executable_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 8, 1]>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 8, 1]>
module {
func.func @zero_dim_shared_memory_copy(%global : memref<f32>, %shared : memref<f32>)
attributes {hal.executable.target = #executable_target, translation_info = #translation_info} {
@@ -153,7 +153,7 @@
// -----
#executable_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 8, 1]>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 8, 1]>
module {
func.func @dequant_shared_memory_copy(%A: memref<1x32x128xi4>, %B: memref<1x128xf32>, %C: memref<1x128xi4>,
%SM: memref<1x32x128xf32, #gpu.address_space<workgroup>>)
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_distribute_shared_memory.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_distribute_shared_memory.mlir
index 0b041d0..6cc129a 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_distribute_shared_memory.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/transform_gpu_distribute_shared_memory.mlir
@@ -2,7 +2,7 @@
#executable_target = #hal.executable.target<"cuda", "cuda-nvptx-fb">
#map1 = affine_map<(d0, d1) -> (d0, d1)>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 4, 1]>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 4, 1]>
module attributes {transform.with_named_sequence} {
memref.global "private" @__shared_memory__ : memref<64x16xf32, #gpu.address_space<workgroup>>
diff --git a/compiler/src/iree/compiler/Codegen/Common/GPU/test/vector_reduction_to_gpu.mlir b/compiler/src/iree/compiler/Codegen/Common/GPU/test/vector_reduction_to_gpu.mlir
index 6b88a7b..d99a2cd 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GPU/test/vector_reduction_to_gpu.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/GPU/test/vector_reduction_to_gpu.mlir
@@ -6,7 +6,7 @@
#hal.pipeline.binding<storage_buffer>
]>
#map = affine_map<()[s0, s1] -> (s1 * 2 + s0 floordiv 32)>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
module {
func.func @simple_reduce() attributes {translation_info = #translation_info} {
%c0 = arith.constant 0 : index
@@ -79,7 +79,7 @@
#hal.pipeline.binding<storage_buffer>,
#hal.pipeline.binding<storage_buffer>
]>
-#translation_info_forall = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info_forall = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
module {
func.func @reduce_forall() attributes {translation_info = #translation_info_forall} {
%c0 = arith.constant 0 : index
@@ -133,7 +133,7 @@
#hal.pipeline.binding<storage_buffer>,
#hal.pipeline.binding<uniform_buffer>
]>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
#map = affine_map<()[s0, s1] -> (s1 * 2 + s0 floordiv 32)>
module {
func.func @reduce_uniform_buffer_offset() attributes {translation_info = #translation_info} {
@@ -192,7 +192,7 @@
#hal.pipeline.binding<storage_buffer>
]>
#map = affine_map<()[s0, s1] -> (s1 * 2 + s0 floordiv 32)>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
module {
func.func @reduce_storage_buffer_offset() attributes {translation_info = #translation_info} {
%c0 = arith.constant 0 : index
@@ -246,7 +246,7 @@
#hal.pipeline.binding<storage_buffer>,
#hal.pipeline.binding<storage_buffer>
]>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
module {
func.func @shared_memory_copy() attributes {translation_info = #translation_info} {
%c0 = arith.constant 0 : index
@@ -282,7 +282,7 @@
#hal.pipeline.binding<storage_buffer>,
#hal.pipeline.binding<storage_buffer>
]>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [64, 1, 1] subgroup_size = 64>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [64, 1, 1] subgroup_size = 64>
#map = affine_map<()[s0] -> (s0 * 4)>
#map1 = affine_map<(d0, d1) -> (0, d1)>
module {
@@ -333,7 +333,7 @@
#hal.pipeline.binding<storage_buffer>,
#hal.pipeline.binding<storage_buffer>
]>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [32, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [32, 1, 1] subgroup_size = 32>
module {
func.func @simple_nd_write() attributes {translation_info = #translation_info} {
%c0 = arith.constant 0 : index
diff --git a/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp b/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
index 4dc3e8e..6f48179 100644
--- a/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
@@ -34,9 +34,9 @@
auto funcOp = *funcOps.begin();
IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(funcOp);
- if (!translationInfo || translationInfo.getDispatchLoweringPassPipeline() !=
- IREE::Codegen::DispatchLoweringPassPipeline::
- TransformDialectCodegen) {
+ if (!translationInfo ||
+ !isa<IREE::Codegen::TransformDialectCodegenPipelineAttr>(
+ translationInfo.getPassPipeline())) {
return;
}
@@ -53,14 +53,14 @@
return signalPassFailure();
}
- // Make sure that the translation info is set to `None` to avoid using
+ // Make sure that the translation info is set to `no_pipeline` to avoid using
// other pass pipelines.
auto translationInfoModified = getTranslationInfo(funcOp);
if (!translationInfoModified ||
- translationInfoModified.getDispatchLoweringPassPipeline() !=
- IREE::Codegen::DispatchLoweringPassPipeline::None) {
+ !isa<IREE::Codegen::NoPipelineAttr>(
+ translationInfoModified.getPassPipeline())) {
funcOp->emitOpError("expected transform dialect lowering to set the "
- "translation_info to use None");
+ "translation_info to use no_pipeline");
return signalPassFailure();
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp b/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
index a12a574..6d53e38 100644
--- a/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
@@ -197,7 +197,8 @@
// of the strategy, the variant needs to be annotated with:
// ```mlir
// "translation_info" =
- // #iree_codegen.translation_info<pipeline = None>
+ // #iree_codegen.translation_info<
+ // pipeline = #iree_codegen.no_pipeline>
// ```
LDBG() << "MaterializeUserConfigsPass on function: " << funcOp;
if (succeeded(userTransformLibrary)) {
@@ -252,9 +253,8 @@
/// We only need to resolve symbols for transform dialect based
/// strategies.
if (!translationInfo ||
- translationInfo.getDispatchLoweringPassPipeline() !=
- IREE::Codegen::DispatchLoweringPassPipeline::
- TransformDialectCodegen) {
+ !isa<IREE::Codegen::TransformDialectCodegenPipelineAttr>(
+ translationInfo.getPassPipeline())) {
continue;
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.cpp b/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.cpp
index bac339a..1a72010 100644
--- a/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.cpp
@@ -67,11 +67,10 @@
}
}
-ConstraintsOpShell
-createConstraintsOpShell(OpBuilder &builder, Operation *rootOp,
- IREE::Codegen::RootOpAttr rootOpAttr,
- Attribute pipelineAttr, DictionaryAttr knobs,
- unsigned numLoops, ArrayRef<AffineMap> indexingMaps) {
+ConstraintsOpShell createConstraintsOpShell(
+ OpBuilder &builder, Operation *rootOp, IREE::Codegen::RootOpAttr rootOpAttr,
+ IREE::Codegen::PipelineAttrInterface pipelineAttr, DictionaryAttr knobs,
+ unsigned numLoops, ArrayRef<AffineMap> indexingMaps) {
MLIRContext *ctx = rootOp->getContext();
Location loc = rootOp->getLoc();
builder.setInsertionPointAfter(rootOp);
diff --git a/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.h b/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.h
index 8a08e19..75a6c5f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.h
+++ b/compiler/src/iree/compiler/Codegen/Common/SMTConstraintUtils.h
@@ -30,11 +30,10 @@
/// op with block args, emits static dim constraints. Returns with builder
/// positioned at the end of the block, ready for pipeline-specific
/// constraints.
-ConstraintsOpShell
-createConstraintsOpShell(OpBuilder &builder, Operation *rootOp,
- IREE::Codegen::RootOpAttr rootOpAttr,
- Attribute pipelineAttr, DictionaryAttr knobs,
- unsigned numLoops, ArrayRef<AffineMap> indexingMaps);
+ConstraintsOpShell createConstraintsOpShell(
+ OpBuilder &builder, Operation *rootOp, IREE::Codegen::RootOpAttr rootOpAttr,
+ IREE::Codegen::PipelineAttrInterface pipelineAttr, DictionaryAttr knobs,
+ unsigned numLoops, ArrayRef<AffineMap> indexingMaps);
/// Helper to create an SMT integer constant.
Value mkIntConst(OpBuilder &builder, Location loc, int64_t v);
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info.mlir b/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info.mlir
index 6c6e648..398ee71 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info.mlir
@@ -8,15 +8,15 @@
hal.executable.variant public @reconcile_workgroup_size target(#hal.executable.target<"", "", {}>) {
hal.executable.export public @entry_point layout(#pipeline_layout)
builtin.module {
- func.func @entry_point() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @entry_point() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
func.call @fn1() : () -> ()
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
return
}
}
@@ -40,7 +40,7 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
return
}
func.func @fn2() {
@@ -68,10 +68,10 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4, 2]>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4, 2]>} {
return
}
}
@@ -93,10 +93,10 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4]>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4]>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline>} {
return
}
}
@@ -117,10 +117,10 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
return
}
}
@@ -145,10 +145,10 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 64>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 64>} {
return
}
}
@@ -173,10 +173,10 @@
func.call @fn2() : () -> ()
return
}
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None, {llvm_func_attrs = {"some-llvm-attr" = "2"}}>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline, {llvm_func_attrs = {"some-llvm-attr" = "2"}}>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None, {llvm_func_attrs = {"some-llvm-attr" = "4"}}>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline, {llvm_func_attrs = {"some-llvm-attr" = "4"}}>} {
return
}
}
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info_pure.mlir b/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info_pure.mlir
index d7b0ed2..bc79d0f 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info_pure.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/reconcile_translation_info_pure.mlir
@@ -14,7 +14,7 @@
// CHECK-NEXT: %[[X:.+]], %[[Y:.+]], %[[Z:.+]] = iree_tensor_ext.dispatch.workgroup_count_from_slice()
// CHECK-NEXT: iree_codegen.yield %[[X]], %[[Y]], %[[Z]]
func.func @no_workload_ordinals() attributes {
- translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [64]>
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [64]>
} {
return
}
@@ -34,7 +34,7 @@
// CHECK-NEXT: %[[X:.+]], %[[Y:.+]], %[[Z:.+]] = iree_tensor_ext.dispatch.workgroup_count_from_slice(%[[A0]], %[[A1]], %[[A2]], %[[A3]])
// CHECK-NEXT: iree_codegen.yield %[[X]], %[[Y]], %[[Z]]
func.func @with_workload_ordinals() attributes {
- translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [128]>
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [128]>
} {
%p0 = hal.interface.constant.load layout(#pipeline_layout) ordinal(0) : i32
%p1 = hal.interface.constant.load layout(#pipeline_layout) ordinal(1) : i32
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/strip_compilation_info.mlir b/compiler/src/iree/compiler/Codegen/Common/test/strip_compilation_info.mlir
index fad5cd5..719f283 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/strip_compilation_info.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/strip_compilation_info.mlir
@@ -18,10 +18,10 @@
hal.executable.variant public @strip_main target(#hal.executable.target<"", "">) {
hal.executable.export public @entry_point layout(#pipeline_layout)
builtin.module {
- func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+ func.func @fn1() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
return
}
- func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+ func.func @fn2() attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
return
}
}
@@ -37,7 +37,7 @@
// -----
#config = #iree_codegen.lowering_config<tile_sizes = [[128, 256], [16, 16]]>
-#translation = #iree_codegen.translation_info<pipeline = None workgroup_size = [16, 8, 1] subgroup_size = 64>
+#translation = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [16, 8, 1] subgroup_size = 64>
#compilation = #iree_codegen.compilation_info<lowering_config = #config, translation_info = #translation>
func.func @matmul_128x1024x256(%lhs : tensor<128x256xf32>, %rhs: tensor<256x1024xf32>, %init: tensor<128x1024xf32>) -> tensor<128x1024xf32> {
%result = linalg.matmul {compilation_info = #compilation} ins(%lhs, %rhs : tensor<128x256xf32>, tensor<256x1024xf32>) outs(%init : tensor<128x1024xf32>) -> tensor<128x1024xf32>
@@ -68,7 +68,7 @@
#map3 = affine_map<(d0, d1, d2, d3, d4, d5) -> ()>
#map4 = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3)>
#config = #iree_codegen.lowering_config<tile_sizes = [[128, 256], [16, 16]]>
-func.func @attention(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+func.func @attention(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
%init = tensor.empty() : tensor<2x10x6x4xf16>
%result = iree_linalg_ext.attention {decomposition_config = {pv_attrs = {x}, qk_attrs = {y}, z}, indexing_maps = [#map, #map1, #map2, #map3, #map4], lowering_config = #config} ins(%arg0, %arg1, %arg2, %arg3 : tensor<2x10x6x4xf16>, tensor<2x10x4x4xf16>, tensor<2x10x4x4xf16>, f16) outs(%init : tensor<2x10x6x4xf16>) {
^bb0(%arg: f32):
@@ -94,9 +94,9 @@
#map3 = affine_map<(d0, d1, d2, d3, d4, d5) -> ()>
#map4 = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3)>
#config = #iree_codegen.lowering_config<tile_sizes = [[128, 256], [16, 16]]>
-#translation = #iree_codegen.translation_info<pipeline = None workgroup_size = [16, 8, 1] subgroup_size = 64>
+#translation = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [16, 8, 1] subgroup_size = 64>
#compilation = #iree_codegen.compilation_info<lowering_config = #config, translation_info = #translation>
-func.func @attention_1(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+func.func @attention_1(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
%init = tensor.empty() : tensor<2x10x6x4xf16>
%result = iree_linalg_ext.attention {decomposition_config = {pv_attrs = {x}, qk_attrs = {y}, use_exp2}, indexing_maps = [#map, #map1, #map2, #map3, #map4], compilation_info = #compilation} ins(%arg0, %arg1, %arg2, %arg3 : tensor<2x10x6x4xf16>, tensor<2x10x4x4xf16>, tensor<2x10x4x4xf16>, f16) outs(%init : tensor<2x10x6x4xf16>) {
^bb0(%arg: f32):
@@ -124,9 +124,9 @@
#map3 = affine_map<(d0, d1, d2, d3, d4, d5) -> ()>
#map4 = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3)>
#config = #iree_codegen.lowering_config<tile_sizes = [[128, 256], [16, 16]]>
-#translation = #iree_codegen.translation_info<pipeline = None workgroup_size = [16, 8, 1] subgroup_size = 64>
+#translation = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [16, 8, 1] subgroup_size = 64>
#compilation = #iree_codegen.compilation_info<lowering_config = #config, translation_info = #translation>
-func.func @attention_2(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+func.func @attention_2(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> tensor<2x10x6x4xf16> attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
%init = tensor.empty() : tensor<2x10x6x4xf16>
%result = iree_linalg_ext.attention {decomposition_config = {pv_attrs = {x}, qk_attrs = {y}}, indexing_maps = [#map, #map1, #map2, #map3, #map4], compilation_info = #compilation} ins(%arg0, %arg1, %arg2, %arg3 : tensor<2x10x6x4xf16>, tensor<2x10x4x4xf16>, tensor<2x10x4x4xf16>, f16) outs(%init : tensor<2x10x6x4xf16>) {
^bb0(%arg: f32):
@@ -154,7 +154,7 @@
#map4 = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2, d3)>
#map5 = affine_map<(d0, d1, d2, d3, d4, d5) -> (d0, d1, d2)>
#config = #iree_codegen.lowering_config<tile_sizes = [[128, 256], [16, 16]]>
-func.func @online_attention(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> (tensor<2x10x6x4xf32>, tensor<2x10x6xf32>, tensor<2x10x6xf32>) attributes {translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = 32>} {
+func.func @online_attention(%arg0: tensor<2x10x6x4xf16>, %arg1 : tensor<2x10x4x4xf16>, %arg2 : tensor<2x10x4x4xf16>, %arg3 : f16) -> (tensor<2x10x6x4xf32>, tensor<2x10x6xf32>, tensor<2x10x6xf32>) attributes {translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = 32>} {
%acc = tensor.empty() : tensor<2x10x6x4xf32>
%max = tensor.empty() : tensor<2x10x6xf32>
%sum = tensor.empty() : tensor<2x10x6xf32>
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp
index 9b0bade..0243ab2 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp
@@ -21,49 +21,6 @@
#include "mlir/IR/StorageUniquerSupport.h"
#include "mlir/Pass/PassRegistry.h"
-// Custom parse/print directives for TranslationInfoAttr's pipeline field.
-// These must be defined before the generated .cpp.inc is included because
-// the ODS-generated parse/print methods call them.
-namespace mlir::iree_compiler::IREE::Codegen {
-
-/// Parses either a DispatchLoweringPassPipeline enum keyword (e.g.,
-/// `CPUDefault`) or a generic attribute implementing PipelineAttrInterface
-/// (e.g., `#iree_codegen.pass_pipeline<"canonicalize">`).
-ParseResult parsePipelineAttr(AsmParser &parser, Attribute &result) {
- StringRef keyword;
- SMLoc loc = parser.getCurrentLocation();
- if (succeeded(parser.parseOptionalKeyword(&keyword))) {
- std::optional<DispatchLoweringPassPipeline> pipeline =
- symbolizeDispatchLoweringPassPipeline(keyword);
- if (!pipeline) {
- parser.emitError(loc, "unknown pipeline keyword: ") << keyword;
- return failure();
- }
- result =
- DispatchLoweringPassPipelineAttr::get(parser.getContext(), *pipeline);
- return success();
- }
- Attribute attr;
- if (parser.parseAttribute(attr)) {
- return failure();
- }
- result = attr;
- return success();
-}
-
-/// Prints DispatchLoweringPassPipelineAttr as a bare keyword and other
-/// attributes (e.g., PipelineAttrInterface impls) via the generic printer.
-void printPipelineAttr(AsmPrinter &printer, Attribute pipelineAttr) {
- if (auto enumAttr =
- dyn_cast<DispatchLoweringPassPipelineAttr>(pipelineAttr)) {
- printer << stringifyEnum(enumAttr.getValue());
- return;
- }
- printer.printAttribute(pipelineAttr);
-}
-
-} // namespace mlir::iree_compiler::IREE::Codegen
-
#define GET_ATTRDEF_CLASSES
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.cpp.inc"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenEnums.cpp.inc"
@@ -116,6 +73,58 @@
}
//===----------------------------------------------------------------------===//
+// iree_codegen.vmvx_pipeline
+//===----------------------------------------------------------------------===//
+
+static VMVXPipelineBuilder &getVMVXPipelineBuilderStorage() {
+ static VMVXPipelineBuilder builder = nullptr;
+ return builder;
+}
+
+void registerVMVXPipelineBuilder(VMVXPipelineBuilder builder) {
+ // Expected to be called exactly once during global init, so thread safety is
+ // not a concern.
+ [[maybe_unused]] static bool registered = false;
+ assert(!registered && "VMVX pipeline builder registered more than once");
+ registered = true;
+ getVMVXPipelineBuilderStorage() = builder;
+}
+
+LogicalResult
+VMVXPipelineAttr::buildPipeline(OpPassManager &pm,
+ const CodegenPipelineOptions *options) const {
+ VMVXPipelineBuilder builder = getVMVXPipelineBuilderStorage();
+ if (!builder) {
+ return emitError(UnknownLoc::get(getContext()))
+ << "no VMVX pipeline builder registered; ensure "
+ "registerCodegenVMVXPasses() was called";
+ }
+ return builder(*this, pm, options);
+}
+
+//===----------------------------------------------------------------------===//
+// iree_codegen.transform_dialect_codegen
+//===----------------------------------------------------------------------===//
+
+LogicalResult TransformDialectCodegenPipelineAttr::buildPipeline(
+ OpPassManager &, const CodegenPipelineOptions *) const {
+ return emitError(UnknownLoc::get(getContext()))
+ << "transform dialect codegen pipeline should be consumed by "
+ "LowerExecutableUsingTransformDialect and replaced with "
+ "#iree_codegen.no_pipeline before executable target lowering";
+}
+
+//===----------------------------------------------------------------------===//
+// iree_codegen.no_pipeline
+//===----------------------------------------------------------------------===//
+
+LogicalResult
+NoPipelineAttr::buildPipeline(OpPassManager &,
+ const CodegenPipelineOptions *) const {
+ return success();
+}
+
+//===----------------------------------------------------------------------===//
// iree_codegen.pass_pipeline
//===----------------------------------------------------------------------===//
@@ -144,35 +153,6 @@
// iree_codegen.translation_info
//===----------------------------------------------------------------------===//
-TranslationInfoAttr TranslationInfoAttr::get(
- MLIRContext *context, DispatchLoweringPassPipeline passPipeline,
- SymbolRefAttr codegenSpec, ArrayRef<int64_t> workgroupSize,
- std::optional<int64_t> subgroupSize, DictionaryAttr configuration) {
- Attribute pipelineAttr =
- DispatchLoweringPassPipelineAttr::get(context, passPipeline);
- return get(context, pipelineAttr, codegenSpec, workgroupSize,
- subgroupSize.value_or(int64_t()), configuration);
-}
-
-TranslationInfoAttr TranslationInfoAttr::get(
- MLIRContext *context, DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize, std::optional<int64_t> subgroupSize,
- DictionaryAttr configuration) {
- Attribute pipelineAttr =
- DispatchLoweringPassPipelineAttr::get(context, passPipeline);
- return get(context, pipelineAttr, /*codegenSpec=*/SymbolRefAttr(),
- workgroupSize, subgroupSize.value_or(int64_t()), configuration);
-}
-
-DispatchLoweringPassPipeline
-TranslationInfoAttr::getDispatchLoweringPassPipeline() {
- if (auto enumAttr =
- dyn_cast<DispatchLoweringPassPipelineAttr>(getPassPipeline())) {
- return enumAttr.getValue();
- }
- return DispatchLoweringPassPipeline::None;
-}
-
LogicalResult TranslationInfoAttr::verify(
function_ref<InFlightDiagnostic()> emitError, Attribute passPipeline,
SymbolRefAttr codegenSpec, ArrayRef<int64_t> workgroupSize,
@@ -180,25 +160,12 @@
if (!passPipeline) {
return emitError() << "missing pass pipeline specification";
}
- if (auto enumAttr =
- dyn_cast<DispatchLoweringPassPipelineAttr>(passPipeline)) {
- DispatchLoweringPassPipeline passPipelineValue = enumAttr.getValue();
- if (passPipelineValue > IREE::Codegen::DispatchLoweringPassPipeline::None) {
- return emitError() << "invalid pass pipeline value : "
- << stringifyEnum(passPipelineValue);
- }
- DispatchLoweringPassPipeline tdPassPipeline =
- IREE::Codegen::DispatchLoweringPassPipeline::TransformDialectCodegen;
- if (codegenSpec && passPipelineValue != tdPassPipeline) {
- return emitError()
- << "transform dialect codegen spec requires pass pipeline : "
- << stringifyEnum(tdPassPipeline);
- }
- } else if (!passPipeline
- .hasPromiseOrImplementsInterface<PipelineAttrInterface>()) {
- return emitError()
- << "pass pipeline must be a DispatchLoweringPassPipelineAttr or "
- "implement PipelineAttrInterface";
+ if (!passPipeline.hasPromiseOrImplementsInterface<PipelineAttrInterface>()) {
+ return emitError() << "pass pipeline must implement PipelineAttrInterface";
+ }
+ if (codegenSpec && !isa<TransformDialectCodegenPipelineAttr>(passPipeline)) {
+ return emitError() << "transform dialect codegen spec requires transform "
+ "dialect codegen pipeline";
}
if (workgroupSize.size() > 3) {
return emitError() << "workgroup size cannot have more than 3 entries";
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h
index 90ea994..1f3575d 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h
@@ -40,20 +40,14 @@
namespace mlir::iree_compiler::IREE::Codegen {
-/// Parses either a DispatchLoweringPassPipeline enum keyword or a generic
-/// attribute implementing PipelineAttrInterface.
-ParseResult parsePipelineAttr(AsmParser &parser, Attribute &result);
-inline ParseResult parsePipelineAttr(OpAsmParser &parser, Attribute &result) {
- return parsePipelineAttr(static_cast<AsmParser &>(parser), result);
-}
+/// Callback type for VMVX pipeline builders.
+using VMVXPipelineBuilder =
+ LogicalResult (*)(Attribute pipelineAttr, OpPassManager &pm,
+ const CodegenPipelineOptions *options);
-/// Prints DispatchLoweringPassPipelineAttr as a bare keyword and other
-/// attributes via the generic printer.
-void printPipelineAttr(AsmPrinter &printer, Attribute pipelineAttr);
-inline void printPipelineAttr(OpAsmPrinter &printer, Operation *,
- Attribute pipelineAttr) {
- printPipelineAttr(printer, pipelineAttr);
-}
+/// Registers the VMVX pipeline builder callback. Must be called before
+/// any compilation that uses #iree_codegen.vmvx_pipeline.
+void registerVMVXPipelineBuilder(VMVXPipelineBuilder builder);
} // namespace mlir::iree_compiler::IREE::Codegen
@@ -178,58 +172,6 @@
return success();
}
-/// Convenience function that sets the lowering configuration on the operation
-/// and translation info for a generic lowering config, lowering pipeline,
-/// and optional workgroup/subgroup size.
-inline LogicalResult setOpConfigAndEntryPointFnTranslation(
- mlir::FunctionOpInterface entryPointFn, Operation *op,
- IREE::Codegen::LoweringConfigAttrInterface config,
- IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize = {},
- std::optional<int64_t> subgroupSize = {},
- DictionaryAttr pipelineConfig = DictionaryAttr()) {
- MLIRContext *context = entryPointFn.getContext();
- auto translationInfo = IREE::Codegen::TranslationInfoAttr::get(
- context, passPipeline, SymbolRefAttr(), workgroupSize, subgroupSize,
- pipelineConfig);
- return setOpConfigAndEntryPointFnTranslation(entryPointFn, op, config,
- translationInfo);
-}
-
-/// Convenience function that sets the lowering configuration on the operation
-/// and translation info on the entry point op for the common case of specifying
-/// tile sizes to use for the operation, and pass pipeline to use for the
-/// translation.
-inline LogicalResult setOpConfigAndEntryPointFnTranslation(
- mlir::FunctionOpInterface entryPointFn, Operation *op,
- TileSizesListTypeRef tileSizes,
- ScalableTileFlagsListTypeRef scalableTileFlags,
- IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize = {},
- std::optional<int64_t> subgroupSize = {},
- DictionaryAttr pipelineConfig = DictionaryAttr()) {
- MLIRContext *context = entryPointFn.getContext();
- auto config = IREE::Codegen::LoweringConfigAttr::get(context, tileSizes,
- scalableTileFlags);
- return setOpConfigAndEntryPointFnTranslation(entryPointFn, op, config,
- passPipeline, workgroupSize,
- subgroupSize, pipelineConfig);
-}
-
-/// Overload of setOpConfigAndEntryPointFnTranslation() for the "no scalable
-/// flags" case.
-inline LogicalResult setOpConfigAndEntryPointFnTranslation(
- mlir::FunctionOpInterface entryPointFn, Operation *op,
- TileSizesListTypeRef tileSizes,
- IREE::Codegen::DispatchLoweringPassPipeline passPipeline,
- ArrayRef<int64_t> workgroupSize = {},
- std::optional<int64_t> subgroupSize = {},
- DictionaryAttr pipelineConfig = DictionaryAttr()) {
- return setOpConfigAndEntryPointFnTranslation(entryPointFn, op, tileSizes, {},
- passPipeline, workgroupSize,
- subgroupSize, pipelineConfig);
-}
-
/// Function to erase lowering configs that are set on an operation.
void eraseLoweringConfig(Operation *op);
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td
index f692a93..417ec35 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.td
@@ -13,45 +13,6 @@
include "mlir/IR/EnumAttr.td"
//===---------------------------------------------------------------------===//
-// Pass pipelines
-//===---------------------------------------------------------------------===//
-
-// List of pre-existing pipelines for translating executables.
-def VMVX_Default : I32EnumAttrCase<"VMVXDefault", 300>;
-
-def Linalg_TransformDialectCodegen
- : I32EnumAttrCase<"TransformDialectCodegen", 1000>;
-def Custom
- : I32EnumAttrCase<"Custom", 1001>;
-
-def None : I32EnumAttrCase<"None", 0xffff>;
-
-// EnumAttrCase for all known lowerings for ops within dispatch region
-// to scalar/native-vector code.
-def DispatchLoweringPassPipelineEnum : I32EnumAttr<
- "DispatchLoweringPassPipeline",
- "identifier for pass pipeline use to lower dispatch region", [
- VMVX_Default,
-
- // Transform dialect based codegen
- Linalg_TransformDialectCodegen,
-
- // For out of tree pass-pipelines
- Custom,
-
- // None to specify no in-built pipelines to use.
- None
- ]> {
- let cppNamespace = "::mlir::iree_compiler::IREE::Codegen";
- // Don't generate a C++ class! We want to use the AttrDef
- let genSpecializedAttr = 0;
-}
-
-// Define the AttrDef
-def DispatchLoweringPassPipelineAttr :
- EnumAttr<IREECodegen_Dialect, DispatchLoweringPassPipelineEnum, "">;
-
-//===---------------------------------------------------------------------===//
// IREE Codegen workgroup mapping attributes
//===---------------------------------------------------------------------===//
@@ -190,9 +151,35 @@
//===---------------------------------------------------------------------===//
-// iree_codegen.pass_pipeline
+// iree_codegen pipeline attributes
//===---------------------------------------------------------------------===//
+class IREECodegen_UnitPipelineAttr<string name, string mnemonicStr,
+ string summaryStr> :
+ AttrDef<IREECodegen_Dialect, name, [
+ DeclareAttrInterfaceMethods<IREECodegen_PipelineAttrInterface, [
+ "buildPipeline"
+ ]>
+ ]> {
+ let mnemonic = mnemonicStr;
+ let summary = summaryStr;
+ let parameters = (ins);
+ let assemblyFormat = [{}];
+}
+
+def IREECodegen_VMVXPipelineAttr :
+ IREECodegen_UnitPipelineAttr<"VMVXPipeline", "vmvx_pipeline",
+ "The VMVX lowering pipeline.">;
+
+def IREECodegen_TransformDialectCodegenPipelineAttr :
+ IREECodegen_UnitPipelineAttr<
+ "TransformDialectCodegenPipeline", "transform_dialect_codegen",
+ "Transform dialect driven codegen pipeline marker.">;
+
+def IREECodegen_NoPipelineAttr :
+ IREECodegen_UnitPipelineAttr<"NoPipeline", "no_pipeline",
+ "No-op codegen pipeline marker.">;
+
def IREECodegen_PassPipelineAttr :
AttrDef<IREECodegen_Dialect, "PassPipeline", [
DeclareAttrInterfaceMethods<IREECodegen_PipelineAttrInterface, [
@@ -234,10 +221,9 @@
dispatch region (like `linalg.matmul`/`linalg.*conv*`), this
attribute gets propagated to the entry point function.
- The `passPipeline` field can be either:
- - A `DispatchLoweringPassPipelineAttr` (enum keyword like `CPUDefault`).
- - Any attribute implementing `PipelineAttrInterface` (e.g.,
- `#iree_codegen.pass_pipeline<"...">`).
+ The `passPipeline` field must be an attribute implementing
+ `PipelineAttrInterface` (e.g. `#iree_codegen.no_pipeline` or
+ `#iree_codegen.pass_pipeline<"...">`).
}];
let parameters = (ins
@@ -250,25 +236,9 @@
OptionalParameter<"DictionaryAttr",
"Pipeline specific configuration">:$configuration
);
- let builders = [
- AttrBuilder<(ins "DispatchLoweringPassPipeline":$passPipeline,
- CArg<"SymbolRefAttr", "{}">:$codegenSpec,
- CArg<"ArrayRef<int64_t>", "{}">:$workgroupSize,
- CArg<"std::optional<int64_t>", "std::nullopt">:$subgroupSize,
- CArg<"DictionaryAttr", "{}">:$configuration)>,
- AttrBuilder<(ins "DispatchLoweringPassPipeline":$passPipeline,
- "ArrayRef<int64_t>":$workgroupSize,
- CArg<"std::optional<int64_t>", "std::nullopt">:$subgroupSize,
- CArg<"DictionaryAttr", "{}">:$configuration)>
- ];
- let extraClassDeclaration = [{
- // Returns the lowering pass pipeline enum value. Returns None if the
- // pipeline is not a DispatchLoweringPassPipelineAttr.
- DispatchLoweringPassPipeline getDispatchLoweringPassPipeline();
- }];
let assemblyFormat = [{
- `<` `pipeline` `=` custom<PipelineAttr>($passPipeline)
+ `<` `pipeline` `=` $passPipeline
(`codegen_spec` `=` $codegenSpec^)?
(`workgroup_size` `=` `[` $workgroupSize^ `]`)?
(`subgroup_size` `=` $subgroupSize^)?
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenOps.td b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenOps.td
index 383bc15..b451594 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenOps.td
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenOps.td
@@ -609,15 +609,14 @@
}];
let arguments = (ins
IREECodegen_RootOpAttr:$target,
- AnyAttrOf<[DispatchLoweringPassPipelineAttr,
- IREECodegen_PipelineAttrInterface]>:$pipeline,
+ IREECodegen_PipelineAttrInterface:$pipeline,
DictionaryAttr:$knobs,
Variadic<Index>:$problem_dims
);
let regions = (region SizedRegion<1>:$body);
let results = (outs);
let assemblyFormat = [{
- `target` `=` $target `,` `pipeline` `=` custom<PipelineAttr>($pipeline) `,`
+ `target` `=` $target `,` `pipeline` `=` $pipeline `,`
custom<KnobsDictionary>($knobs)
`dims` `(` $problem_dims `)` attr-dict-with-keyword
$body
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/invalid.mlir b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/invalid.mlir
index 9d215a6..b1f0d88 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/invalid.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/invalid.mlir
@@ -46,7 +46,7 @@
// Constraints op: block arg wrong type.
func.func @constraints_block_arg_wrong_type(%arg0: index) {
// expected-error @+1 {{'iree_codegen.smt.constraints' op block argument #0 must be !smt.int but got 'index'}}
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {}
dims(%arg0) {
^bb0(%m: index):
@@ -68,7 +68,7 @@
// Constraints op: block arg count mismatch with problem_dims.
func.func @constraints_block_arg_mismatch(%arg0: index) {
// expected-error @+1 {{'iree_codegen.smt.constraints' op expected 1 block arguments but got 2}}
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {}
dims(%arg0) {
^bb0(%m: !smt.int, %extra: !smt.int):
@@ -99,7 +99,7 @@
// Knob op: duplicate knob name.
func.func @duplicate_knob_name(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {workgroup = [#iree_codegen.smt.int_knob<"wg_m">]}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -116,7 +116,7 @@
// Constraints op: too few block args for problem_dims.
func.func @constraints_block_arg_too_few(%arg0: index, %arg1: index) {
// expected-error @+1 {{'iree_codegen.smt.constraints' op expected 2 block arguments but got 1}}
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {}
dims(%arg0, %arg1) {
^bb0(%m: !smt.int):
@@ -136,7 +136,7 @@
// Knob op: knob name not found in knobs dict.
func.func @knob_name_not_found(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {workgroup = [#iree_codegen.smt.int_knob<"wg_m">]}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -150,7 +150,7 @@
// Knob op: bare string in knobs dict does not satisfy knob lookup.
func.func @string_attr_not_a_knob(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {name = "wg_m"}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -171,10 +171,10 @@
// -----
-// Constraints op: pipeline attr must be DispatchLoweringPassPipelineAttr or
-// implement PipelineAttrInterface -- a plain string attr is neither.
+// Constraints op: pipeline attr must implement PipelineAttrInterface -- a
+// plain string attr does not.
func.func @constraints_invalid_pipeline(%arg0: index) {
- // expected-error @+1 {{'iree_codegen.smt.constraints' op attribute 'pipeline' failed to satisfy constraint}}
+ // expected-error @+1 {{custom op 'iree_codegen.smt.constraints' invalid kind of attribute specified}}
iree_codegen.smt.constraints target = <set = 0>, pipeline = "not_a_pipeline",
knobs = {}
dims(%arg0) {
@@ -185,9 +185,27 @@
// -----
+func.func @translation_info_invalid_pipeline() attributes {
+ // expected-error @+1 {{pass pipeline must implement PipelineAttrInterface}}
+ translation_info = #iree_codegen.translation_info<pipeline = "not_a_pipeline">
+} {
+ return
+}
+
+// -----
+
+func.func @translation_info_spec_requires_transform_pipeline() attributes {
+ // expected-error @+1 {{transform dialect codegen spec requires transform dialect codegen pipeline}}
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline codegen_spec = @foo>
+} {
+ return
+}
+
+// -----
+
// LookupOp: keys and values size mismatch.
func.func @smt_lookup_size_mismatch(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -202,7 +220,7 @@
// LookupOp: empty table.
func.func @smt_lookup_empty(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -217,7 +235,7 @@
// LookupOp: duplicate keys.
func.func @smt_lookup_duplicate_keys(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -241,7 +259,7 @@
// AssertOp: too few args for format string placeholders.
func.func @assert_too_few_args(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -258,7 +276,7 @@
// AssertOp: too many args for format string placeholders.
func.func @assert_too_many_args(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -275,7 +293,7 @@
// OneOfKnobAttr: knob name not found in knobs dict.
func.func @one_of_knob_name_not_found(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {mma = #iree_codegen.smt.one_of_knob<"mma_idx", ["a", "b"]>}
dims(%arg0) {
^bb0(%m: !smt.int):
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir
index abd76c6..11f5599 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/lowering_config_attr.mlir
@@ -96,7 +96,7 @@
/// translation info cannot have more than 3 entries for workgroup size
func.func @workgroup_size_more_than_3_err() attributes {
// expected-error @+1 {{workgroup size cannot have more than 3 entries}}
- translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4, 1, 1, 1]> {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4, 1, 1, 1]> {
return
}
}
@@ -108,7 +108,7 @@
/// translation info workgroup_size values needs to have non-negative values.
func.func @workgroup_size_neg_err() attributes {
// expected-error @+1 {{workgroup size value has to be greater than zero}}
- translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [4, -1, 1]> {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [4, -1, 1]> {
return
}
}
@@ -120,7 +120,7 @@
/// translation info workgroup_size values needs to have non-negative values.
func.func @subgroup_size_neg_err() attributes {
// expected-error @+1 {{subgroup size value cannot be negative}}
- translation_info = #iree_codegen.translation_info<pipeline = None subgroup_size = -1> {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline subgroup_size = -1> {
return
}
}
@@ -151,9 +151,11 @@
// -----
module {
- /// Invalid pass pipeline string should be caught at verify time.
+ /// Invalid pass pipeline string should be caught while parsing the nested
+ /// pass pipeline attr.
func.func @invalid_pass_pipeline() attributes {
- // expected-error @+1 {{invalid pass pipeline specification: 'not_a_real_pass'}}
+ // expected-error @+2 {{invalid pass pipeline specification: 'not_a_real_pass'}}
+ // expected-error @+1 {{failed to parse IREECodegen_TranslationInfoAttr parameter 'passPipeline'}}
translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.pass_pipeline<"not_a_real_pass">>
} {
return
diff --git a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/roundtrip.mlir b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/roundtrip.mlir
index a1d77f1..2b5e042 100644
--- a/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/roundtrip.mlir
+++ b/compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR/test/roundtrip.mlir
@@ -114,6 +114,36 @@
// -----
+// Test no pipeline attr inside translation_info.
+func.func private @translation_info_no_pipeline() attributes {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline>
+}
+// CHECK: #translation = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline>
+// CHECK-LABEL: func.func private @translation_info_no_pipeline()
+// CHECK-SAME: translation_info = #translation
+
+// -----
+
+// Test transform dialect codegen pipeline attr inside translation_info.
+func.func private @translation_info_transform_dialect_codegen() attributes {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen codegen_spec = @__kernel_config workgroup_size = [64, 1, 1] subgroup_size = 32>
+}
+// CHECK: #translation = #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen codegen_spec = @__kernel_config workgroup_size = [64, 1, 1] subgroup_size = 32>
+// CHECK-LABEL: func.func private @translation_info_transform_dialect_codegen()
+// CHECK-SAME: translation_info = #translation
+
+// -----
+
+// Test VMVX pipeline attr inside translation_info.
+func.func private @translation_info_vmvx_pipeline() attributes {
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
+}
+// CHECK: #translation = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
+// CHECK-LABEL: func.func private @translation_info_vmvx_pipeline()
+// CHECK-SAME: translation_info = #translation
+
+// -----
+
// Test constraints op with knobs and dims.
func.func @constraints_op(%arg0: index, %arg1: index) {
iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_gpu.pipeline<VectorDistribute>,
@@ -170,7 +200,7 @@
// Test assert op with static message (no format args).
func.func @assert_static_message(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {x = #iree_codegen.smt.int_knob<"x">}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -210,7 +240,7 @@
// Test constraints op with empty dims.
func.func @constraints_op_empty_dims() {
- iree_codegen.smt.constraints target = <set = 1>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 1>, pipeline = #iree_codegen.no_pipeline,
knobs = {}
dims() {
^bb0:
@@ -218,7 +248,7 @@
return
}
// CHECK-LABEL: func.func @constraints_op_empty_dims(
-// CHECK: iree_codegen.smt.constraints target = <set = 1>, pipeline = None,
+// CHECK: iree_codegen.smt.constraints target = <set = 1>, pipeline = #iree_codegen.no_pipeline,
// CHECK: knobs = {}
// CHECK: dims()
@@ -257,7 +287,7 @@
// Test OneOfKnobAttr in constraints op knobs dict.
func.func @one_of_knob_attr(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {mma = #iree_codegen.smt.one_of_knob<"mma_idx", ["option_a", "option_b", "option_c"]>}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -272,7 +302,7 @@
// Test OneOfKnobAttr with heterogeneous options (integer attrs).
func.func @one_of_knob_int_options(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {size = #iree_codegen.smt.one_of_knob<"size_idx", [16 : i64, 32 : i64, 64 : i64]>}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -287,7 +317,7 @@
// Test smt.lookup op roundtrip.
func.func @smt_lookup(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {mma = #iree_codegen.smt.one_of_knob<"mma_idx", ["a", "b", "c"]>}
dims(%arg0) {
^bb0(%m: !smt.int):
@@ -304,7 +334,7 @@
// Test smt.lookup with non-contiguous keys not starting at 0.
func.func @smt_lookup_sparse(%arg0: index) {
- iree_codegen.smt.constraints target = <set = 0>, pipeline = None,
+ iree_codegen.smt.constraints target = <set = 0>, pipeline = #iree_codegen.no_pipeline,
knobs = {mma = #iree_codegen.smt.one_of_knob<"mma_idx", ["a", "b", "c"]>}
dims(%arg0) {
^bb0(%m: !smt.int):
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 5f22ca3..4ea30f3 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -138,8 +138,6 @@
"set distConfig.maxTileSizes[i] to 2 * distConfig.minTileSizes[i]."),
llvm::cl::init(false));
-using IREE::Codegen::DispatchLoweringPassPipeline;
-
// Encodes the pre-processing strategy to be applied on a Linalg operation
// before vectorization.
enum class VectorPreProcStrategy {
@@ -4049,9 +4047,11 @@
// The transform dialect codegen has different logics and codegen flow.
// Ignore the tile sizes adjustment.
- DispatchLoweringPassPipeline pipeline =
- getTranslationInfo(entryPointFn).getDispatchLoweringPassPipeline();
- if (pipeline != DispatchLoweringPassPipeline::TransformDialectCodegen) {
+ IREE::Codegen::TranslationInfoAttr translationInfo =
+ getTranslationInfo(entryPointFn);
+ if (!translationInfo ||
+ !isa<IREE::Codegen::TransformDialectCodegenPipelineAttr>(
+ translationInfo.getPassPipeline())) {
if (failed(adjustTileSizesForRootUnPackOp(entryPointFn, rootOperation))) {
return failure();
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
index 5be301f..5ca501e 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
@@ -171,13 +171,15 @@
OpPassManager passManager(func::FuncOp::getOperationName());
Attribute pipelineAttr = translationInfo.getPassPipeline();
+ // No pipeline specified, nothing to do.
+ if (isa<IREE::Codegen::NoPipelineAttr>(pipelineAttr)) {
+ return;
+ }
+
auto pipelineIface =
dyn_cast<IREE::Codegen::PipelineAttrInterface>(pipelineAttr);
if (!pipelineIface) {
- if (translationInfo.getDispatchLoweringPassPipeline() ==
- IREE::Codegen::DispatchLoweringPassPipeline::None) {
- return;
- }
+ // Not an interface implementor -- reject any remaining legacy pipeline.
funcOp.emitOpError("Unsupported pipeline on CPU target.");
return signalPassFailure();
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.td b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.td
index 4d59fe4..85a0430 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.td
@@ -81,7 +81,7 @@
def LLVMCPULowerExecutableTargetPass :
InterfacePass<"iree-llvmcpu-lower-executable-target", "mlir::FunctionOpInterface"> {
let summary =
- "Lower executable target using an IREE::HAL::DispatchLoweringPassPipeline";
+ "Lower executable target using a pipeline attribute";
let description = [{
Pass to lower the module an hal.executable.variant operation to external
dialect. Currently this pass lowers to LLVM dialect, but could be
@@ -113,7 +113,7 @@
def LLVMCPUSelectLoweringStrategyPass :
Pass<"iree-llvmcpu-select-lowering-strategy", "ModuleOp"> {
let summary =
- "Select a IREE::HAL::DispatchLoweringPassPipeline for lowering the variant";
+ "Select a pipeline attribute for lowering the variant";
let description = [{
Pass to select a lowering strategy for a hal.executable.variant operation.
The variant is annotated with the selected strategies, which are
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
index 8a5f83d..da14779 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -2666,9 +2666,12 @@
}
if (!rootOperation) {
- // No root operation found, set it to none.
+ // No root operation found, set it to no_pipeline.
+ MLIRContext *context = funcOp.getContext();
auto translationInfo = IREE::Codegen::TranslationInfoAttr::get(
- funcOp.getContext(), IREE::Codegen::DispatchLoweringPassPipeline::None);
+ context, IREE::Codegen::NoPipelineAttr::get(context),
+ /*codegenSpec=*/SymbolRefAttr(), /*workgroupSize=*/ArrayRef<int64_t>(),
+ /*subgroupSize=*/int64_t(), /*configuration=*/DictionaryAttr());
if (failed(setTranslationInfo(funcOp, translationInfo))) {
return failure();
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUConstraintGenerator.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUConstraintGenerator.cpp
index bbe16ed..84ed850 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUConstraintGenerator.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPUConstraintGenerator.cpp
@@ -68,8 +68,9 @@
}
/// Emit constraints for a single root op under VectorDistribute pipeline.
-static LogicalResult emitConstraintsForOp(Operation *rootOp,
- Attribute pipelineAttr) {
+static LogicalResult
+emitConstraintsForOp(Operation *rootOp,
+ IREE::Codegen::PipelineAttrInterface pipelineAttr) {
auto linalgOp = dyn_cast<linalg::LinalgOp>(rootOp);
if (!linalgOp) {
return success();
@@ -89,10 +90,12 @@
LogicalResult emitLLVMGPUConstraints(Attribute attr,
ArrayRef<Operation *> rootOps) {
- auto pipelineAttr = cast<IREE::GPU::PipelineAttr>(attr);
+ auto gpuPipelineAttr = cast<IREE::GPU::PipelineAttr>(attr);
+ auto pipelineAttr =
+ cast<IREE::Codegen::PipelineAttrInterface>(gpuPipelineAttr);
// Only VectorDistribute has constraint generation today.
- if (pipelineAttr.getValue() !=
+ if (gpuPipelineAttr.getValue() !=
IREE::GPU::LoweringPipeline::VectorDistribute) {
return success();
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
index 2d42324..afeceed 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
@@ -88,6 +88,10 @@
OpPassManager &pipeline = maybePipeline.value();
Attribute pipelineAttr = translationInfo.getPassPipeline();
+ // No pipeline specified, nothing to do.
+ if (isa<IREE::Codegen::NoPipelineAttr>(pipelineAttr)) {
+ return;
+ }
// Check for PipelineAttrInterface first (covers GPU::PipelineAttr via
// external model and any custom pipeline attrs).
@@ -95,11 +99,7 @@
dyn_cast<IREE::Codegen::PipelineAttrInterface>(pipelineAttr);
if (!pipelineIface) {
- // Not an interface implementor -- check for the legacy None pipeline.
- if (translationInfo.getDispatchLoweringPassPipeline() ==
- IREE::Codegen::DispatchLoweringPassPipeline::None) {
- return;
- }
+ // Not an interface implementor -- reject any remaining legacy pipeline.
funcOp.emitOpError("unsupported pipeline on GPU target");
return signalPassFailure();
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.td b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.td
index 0277b27..64ba6a5 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.td
@@ -107,7 +107,7 @@
def LLVMGPULowerExecutableTargetPass :
InterfacePass<"iree-llvmgpu-lower-executable-target", "mlir::FunctionOpInterface"> {
- let summary = "Perform lowering of executable target using one of the IREE::HAL::DispatchLoweringPassPipeline";
+ let summary = "Perform lowering of executable target using a pipeline attribute";
let options = [
Option<"forROCDL", "for-rocdl", "bool",
/*default=*/"false",
@@ -138,7 +138,7 @@
def LLVMGPUSelectLoweringStrategyPass :
Pass<"iree-llvmgpu-select-lowering-strategy", "ModuleOp"> {
- let summary = "Select a IREE::HAL::DispatchLoweringPassPipeline for lowering the target variant";
+ let summary = "Select a pipeline attribute for lowering the target variant";
let options = [
Option<"gpuOptions", "gpu-options", "GPUCodegenOptions",
/*default=*/"GPUCodegenOptions()",
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
index a357f61..36a59ec 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensions.cpp
@@ -99,12 +99,14 @@
IREE::Codegen::TranslationInfoAttr updatedTranslationInfo =
IREE::Codegen::TranslationInfoAttr::get(
rewriter.getContext(),
- IREE::Codegen::DispatchLoweringPassPipeline::None, getWorkgroupDims(),
- getSubgroupSize());
+ IREE::Codegen::NoPipelineAttr::get(rewriter.getContext()),
+ /*codegenSpec=*/SymbolRefAttr(), getWorkgroupDims(),
+ static_cast<int64_t>(getSubgroupSize()),
+ /*configuration=*/DictionaryAttr());
// Set config dictionary.
// Transform Dialect pipeline requires translation_info pass pipeline to
- // be set to None here.
+ // be set to no_pipeline here.
if (translationInfo) {
updatedTranslationInfo = IREE::Codegen::TranslationInfoAttr::get(
rewriter.getContext(), updatedTranslationInfo.getPassPipeline(),
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/Verifiers.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/Verifiers.cpp
index 11b0117..81b002e 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/Verifiers.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/Verifiers.cpp
@@ -12,8 +12,6 @@
namespace mlir::iree_compiler {
-using CodeGenPipeline = IREE::Codegen::DispatchLoweringPassPipeline;
-
/// Verifies pipelines that use iree_gpu.lowering_config attributes.
LogicalResult verifyLLVMGPUVectorDistributePipeline(
Operation *op, IREE::GPU::LoweringConfigAttr loweringConfig) {
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_bufferize_spec.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_bufferize_spec.mlir
index 24da799..22e16dd 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_bufferize_spec.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_codegen_bufferize_spec.mlir
@@ -10,7 +10,7 @@
} : !transform.op<"func.func">
// Annotate the exported function as already translated.
- %none = transform.param.constant #iree_codegen.translation_info<pipeline = None> -> !transform.any_param
+ %none = transform.param.constant #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline> -> !transform.any_param
transform.annotate %func_op_bufferized "translation_info" = %none : !transform.op<"func.func">, !transform.any_param
transform.yield
}
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_distribution.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_distribution.mlir
index 9863a4f..83f8afa 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_distribution.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_dialect_vector_distribution.mlir
@@ -9,7 +9,7 @@
#pipeline_layout = #hal.pipeline.layout<bindings = [
#hal.pipeline.binding<storage_buffer>
]>
-#translation_info = #iree_codegen.translation_info<pipeline = None workgroup_size = [64, 1, 1] subgroup_size = 32>
+#translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [64, 1, 1] subgroup_size = 32>
func.func @reduce_dispatch_0() attributes {translation_info = #translation_info} {
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_distribute_forall.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_distribute_forall.mlir
index feb2a50..afcd7a2 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_distribute_forall.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/transform_distribute_forall.mlir
@@ -4,7 +4,7 @@
#hal.pipeline.binding<storage_buffer>
]>
#executable_target_cuda_nvptx_fb = #hal.executable.target<"cuda", "cuda-nvptx-fb">
-#translation = #iree_codegen.translation_info<pipeline = TransformDialectCodegen, { config_test = "config_test" }>
+#translation = #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen, { config_test = "config_test" }>
module {
func.func @distribute() attributes {hal.executable.target = #executable_target_cuda_nvptx_fb, translation_info = #translation} {
%cst = arith.constant dense<0.000000e+00> : vector<1xf16>
@@ -43,7 +43,7 @@
}
// CHECK-DAG: #[[DIV32:.*]] = affine_map<()[s0] -> (s0 floordiv 32)>
-// CHECK-DAG: #[[TRANSLATION_INFO:.*]] = #iree_codegen.translation_info<pipeline = None workgroup_size = [256, 1, 1] subgroup_size = 32, {config_test = "config_test"}>
+// CHECK-DAG: #[[TRANSLATION_INFO:.*]] = #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline workgroup_size = [256, 1, 1] subgroup_size = 32, {config_test = "config_test"}>
// CHECK: func.func @distribute()
// CHECK-SAME: translation_info = #[[TRANSLATION_INFO]]
// CHECK: %[[TX:.+]] = gpu.thread_id x
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
index 546e506..4d71d1e 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.td
@@ -62,8 +62,7 @@
def SPIRVLowerExecutableTargetPass :
InterfacePass<"iree-spirv-lower-executable-target-pass",
"mlir::FunctionOpInterface"> {
- let summary = "Lower the executable target to SPIR-V using one of the "
- "IREE::HAL::DispatchLoweringPassPipeline";
+ let summary = "Lower the executable target to SPIR-V using a pipeline attribute";
let description = [{
Main pass to lower executables to scalar + vector code on SPIR-V path.
Invokes one of the pass pipelines that translate the executable to
@@ -96,8 +95,7 @@
def SPIRVSelectLoweringStrategyPass :
Pass<"iree-spirv-select-lowering-strategy-pass", "ModuleOp"> {
- let summary = "Select the IREE::HAL::DispatchLoweringPassPipeline for lowering"
- "to SPIR-V";
+ let summary = "Select the pipeline attribute for lowering to SPIR-V";
}
def SPIRVTileAndDistributePass : InterfacePass<"iree-spirv-tile-and-distribute", "mlir::FunctionOpInterface"> {
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTarget.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTarget.cpp
index 702e70b..6a146f0 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTarget.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableTarget.cpp
@@ -83,6 +83,10 @@
OpPassManager &pipeline = maybePipeline.value();
Attribute pipelineAttr = translationInfo.getPassPipeline();
+ // No pipeline specified, nothing to do.
+ if (isa<IREE::Codegen::NoPipelineAttr>(pipelineAttr)) {
+ return;
+ }
// Check for PipelineAttrInterface first (covers GPU::SPIRVPipelineAttr via
// external model and any custom pipeline attrs).
@@ -90,11 +94,7 @@
dyn_cast<IREE::Codegen::PipelineAttrInterface>(pipelineAttr);
if (!pipelineIface) {
- // Not an interface implementor -- check for the legacy None pipeline.
- if (translationInfo.getDispatchLoweringPassPipeline() ==
- IREE::Codegen::DispatchLoweringPassPipeline::None) {
- return;
- }
+ // Not an interface implementor -- reject any remaining legacy pipeline.
funcOp.emitOpError("unsupported pipeline on SPIR-V target");
return signalPassFailure();
}
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableUsingTransformDialect.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableUsingTransformDialect.cpp
index dc0b844..c9d5d4f 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableUsingTransformDialect.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVLowerExecutableUsingTransformDialect.cpp
@@ -35,9 +35,9 @@
auto funcOp = *funcOps.begin();
IREE::Codegen::TranslationInfoAttr translationInfo =
getTranslationInfo(funcOp);
- if (!translationInfo || translationInfo.getDispatchLoweringPassPipeline() !=
- IREE::Codegen::DispatchLoweringPassPipeline::
- TransformDialectCodegen) {
+ if (!translationInfo ||
+ !isa<IREE::Codegen::TransformDialectCodegenPipelineAttr>(
+ translationInfo.getPassPipeline())) {
return;
}
@@ -59,14 +59,14 @@
return signalPassFailure();
}
- // Make sure that the translation info is set to `None` to avoid using
+ // Make sure that the translation info is set to `no_pipeline` to avoid using
// other pass pipelines.
auto translationInfoModified = getTranslationInfo(funcOp);
if (!translationInfoModified ||
- translationInfoModified.getDispatchLoweringPassPipeline() !=
- IREE::Codegen::DispatchLoweringPassPipeline::None) {
+ !isa<IREE::Codegen::NoPipelineAttr>(
+ translationInfoModified.getPassPipeline())) {
funcOp->emitOpError("expected transform dialect lowering to set the "
- "translation_info to use None");
+ "translation_info to use no_pipeline");
return signalPassFailure();
}
}
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVSelectLoweringStrategy.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVSelectLoweringStrategy.cpp
index f3a4e9d..fdba309 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVSelectLoweringStrategy.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/SPIRVSelectLoweringStrategy.cpp
@@ -66,13 +66,8 @@
Attribute pipelineAttr = translationInfo.getPassPipeline();
// Transform dialect encodes configuration into the schedule directly.
- if (auto enumPipeline =
- dyn_cast<IREE::Codegen::DispatchLoweringPassPipelineAttr>(
- pipelineAttr)) {
- if (enumPipeline.getValue() ==
- IREE::Codegen::DispatchLoweringPassPipeline::TransformDialectCodegen) {
- return success();
- }
+ if (isa<IREE::Codegen::TransformDialectCodegenPipelineAttr>(pipelineAttr)) {
+ return success();
}
// Only SPIRV pipelines have additional verification. Other pipeline types
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/test/custom_pass_pipeline.mlir b/compiler/src/iree/compiler/Codegen/SPIRV/test/custom_pass_pipeline.mlir
index b26f43b..de561c7 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/test/custom_pass_pipeline.mlir
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/test/custom_pass_pipeline.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt --pass-pipeline='builtin.module(func.func(iree-spirv-lower-executable-target-pass))' %s | FileCheck %s
+// RUN: iree-opt --pass-pipeline='builtin.module(func.func(iree-spirv-lower-executable-target-pass))' --verify-diagnostics --split-input-file %s | FileCheck %s
// Verify that a custom pass pipeline specified via #iree_codegen.pass_pipeline
// attribute is executed by the SPIRV lower executable target pass.
@@ -17,3 +17,16 @@
// CHECK-LABEL: func.func @test_custom_pipeline
// CHECK-SAME: (%[[ARG0:.+]]: index)
// CHECK-NEXT: return %[[ARG0]]
+
+// -----
+
+#executable_target = #hal.executable.target<"vulkan-spirv", "vulkan-spirv-fb">
+
+// expected-error@unknown {{VMVX pipeline expects VMVXCodegenPipelineOptions}}
+// expected-error@+1 {{'func.func' op failed to build pass pipeline}}
+func.func @vmvx_pipeline_on_spirv_target() attributes {
+ hal.executable.target = #executable_target,
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
+} {
+ return
+}
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/BUILD.bazel b/compiler/src/iree/compiler/Codegen/VMVX/BUILD.bazel
index 89755e4..a958d42 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/BUILD.bazel
+++ b/compiler/src/iree/compiler/Codegen/VMVX/BUILD.bazel
@@ -39,6 +39,7 @@
deps = [
":PassesIncGen",
"//compiler/src/iree/compiler/Codegen/Dialect/Codegen/IR:IREECodegenDialect",
+ "//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/HAL/IR",
"//compiler/src/iree/compiler/Utils",
"@llvm-project//mlir:LinalgTransforms",
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/VMVX/CMakeLists.txt
index 48832f3..1c37af5 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Codegen/VMVX/CMakeLists.txt
@@ -33,6 +33,7 @@
MLIRTransformUtils
MLIRTransforms
iree::compiler::Codegen::Dialect::Codegen::IR::IREECodegenDialect
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::IR
iree::compiler::Utils
PUBLIC
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/KernelDispatch.cpp b/compiler/src/iree/compiler/Codegen/VMVX/KernelDispatch.cpp
index 791f784..fa60c75 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/KernelDispatch.cpp
+++ b/compiler/src/iree/compiler/Codegen/VMVX/KernelDispatch.cpp
@@ -19,6 +19,13 @@
constexpr int kDefaultDistTileSize = 64;
+static IREE::Codegen::TranslationInfoAttr
+getVMVXDefaultTranslationInfo(MLIRContext *context) {
+ return IREE::Codegen::TranslationInfoAttr::get(
+ context, IREE::Codegen::VMVXPipelineAttr::get(context), SymbolRefAttr(),
+ /*workgroupSize=*/{}, /*subgroupSize=*/0, DictionaryAttr());
+}
+
/// Returns true if the operation is nested inside a tiled and distributed loop.
static bool isInsideDistributedLoop(Operation *op) {
for (Operation *parent = op->getParentOp(); parent;
@@ -85,7 +92,7 @@
getLoweringConfigWithDistributionTiles(fftOp.getContext(), distTileSizes);
return setOpConfigAndEntryPointFnTranslation(
entryPointFn, fftOp, loweringConfig,
- IREE::Codegen::DispatchLoweringPassPipeline::VMVXDefault);
+ getVMVXDefaultTranslationInfo(fftOp.getContext()));
}
static LogicalResult setRootConfig(mlir::FunctionOpInterface entryPointFn,
@@ -100,7 +107,7 @@
distTileSizes);
return setOpConfigAndEntryPointFnTranslation(
entryPointFn, tilingInterfaceOp, loweringConfig,
- IREE::Codegen::DispatchLoweringPassPipeline::VMVXDefault);
+ getVMVXDefaultTranslationInfo(tilingInterfaceOp.getContext()));
}
static LogicalResult
@@ -119,10 +126,7 @@
static LogicalResult
lowerUsingVMVXDefaultPipeline(mlir::FunctionOpInterface op) {
- auto translationInfo = IREE::Codegen::TranslationInfoAttr::get(
- op.getContext(),
- IREE::Codegen::DispatchLoweringPassPipeline::VMVXDefault);
- return setTranslationInfo(op, translationInfo);
+ return setTranslationInfo(op, getVMVXDefaultTranslationInfo(op.getContext()));
}
/// Sets the translation information to use for a dispatch region.
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/Passes.cpp b/compiler/src/iree/compiler/Codegen/VMVX/Passes.cpp
index 9044583..c5b95f7 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/VMVX/Passes.cpp
@@ -9,8 +9,10 @@
#include "iree/compiler/Codegen/Common/CPU/Passes.h"
#include "iree/compiler/Codegen/Common/PassUtils.h"
#include "iree/compiler/Codegen/Common/Passes.h"
+#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
#include "iree/compiler/Codegen/VMVX/Passes.h"
#include "iree/compiler/Dialect/LinalgExt/Transforms/Passes.h"
+#include "llvm/Support/Casting.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Pass/PassManager.h"
@@ -79,6 +81,24 @@
}
}
+static LogicalResult buildVMVXPipeline(Attribute pipelineAttr,
+ OpPassManager &pm,
+ const CodegenPipelineOptions *options) {
+ assert(isa<IREE::Codegen::VMVXPipelineAttr>(pipelineAttr) &&
+ "unexpected VMVX pipeline attr");
+ if (!options) {
+ return emitError(UnknownLoc::get(pipelineAttr.getContext()))
+ << "VMVX pipeline expects VMVXCodegenPipelineOptions";
+ }
+ auto vmvxOptions = dyn_cast<VMVXCodegenPipelineOptions>(options);
+ if (!vmvxOptions) {
+ return emitError(UnknownLoc::get(pipelineAttr.getContext()))
+ << "VMVX pipeline expects VMVXCodegenPipelineOptions";
+ }
+ addVMVXDefaultPassPipeline(pm, vmvxOptions->enableUKernels);
+ return success();
+}
+
void buildVMVXConfigurationPassPipeline(OpPassManager &modulePassManager) {
{
FunctionLikeNest funcPassManager(modulePassManager);
@@ -128,6 +148,7 @@
void registerCodegenVMVXPasses() {
// Generated.
registerPasses();
+ IREE::Codegen::registerVMVXPipelineBuilder(buildVMVXPipeline);
static PassPipelineRegistration<> VMVXConfigPipeline(
"iree-codegen-vmvx-configuration-pipeline",
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/Passes.h b/compiler/src/iree/compiler/Codegen/VMVX/Passes.h
index b7beac1..cf12f05 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/Passes.h
+++ b/compiler/src/iree/compiler/Codegen/VMVX/Passes.h
@@ -12,6 +12,7 @@
#ifndef IREE_COMPILER_CODEGEN_VMVX_PASSES_H_
#define IREE_COMPILER_CODEGEN_VMVX_PASSES_H_
+#include "iree/compiler/Codegen/Utils/CodegenPipelineOptions.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Interfaces/FunctionInterfaces.h"
#include "mlir/Pass/Pass.h"
@@ -27,6 +28,16 @@
void addVMVXDefaultPassPipeline(OpPassManager &funcPassManager,
bool enableUKernels);
+/// Wraps VMVX pipeline options for passing through
+/// PipelineAttrInterface::buildPipeline.
+struct VMVXCodegenPipelineOptions final
+ : CodegenPipelineOptionsBase<VMVXCodegenPipelineOptions> {
+ explicit VMVXCodegenPipelineOptions(bool enableUKernels)
+ : enableUKernels(enableUKernels) {}
+
+ bool enableUKernels = false;
+};
+
//----------------------------------------------------------------------------//
// VMVX Codegen Pipelines
//----------------------------------------------------------------------------//
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/Passes.td b/compiler/src/iree/compiler/Codegen/VMVX/Passes.td
index 19b3c6a..db3f809 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/VMVX/Passes.td
@@ -21,7 +21,7 @@
def VMVXSelectLoweringStrategyPass :
Pass<"iree-vmvx-select-lowering-strategy", "ModuleOp"> {
let summary =
- "Select a IREE::HAL::DispatchLoweringPassPipeline for lowering the variant";
+ "Select a pipeline attribute for lowering the variant";
}
def VMVXLinkExecutablesPass :
@@ -32,7 +32,7 @@
def VMVXLowerExecutableTargetPass :
InterfacePass<"iree-vmvx-lower-executable-target", "mlir::FunctionOpInterface"> {
let summary =
- "Lower executable target using an IREE::HAL::DispatchLoweringPassPipeline";
+ "Lower executable target using a pipeline attribute";
}
def VMVXLowerLinalgMicrokernelsPass :
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/VMVXLowerExecutableTarget.cpp b/compiler/src/iree/compiler/Codegen/VMVX/VMVXLowerExecutableTarget.cpp
index b6267cd..ec45f4f 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/VMVXLowerExecutableTarget.cpp
+++ b/compiler/src/iree/compiler/Codegen/VMVX/VMVXLowerExecutableTarget.cpp
@@ -70,29 +70,27 @@
}
OpPassManager &pipeline = maybePipeline.value();
- // Check for a custom pipeline via PipelineAttrInterface.
Attribute pipelineAttr = translationInfo.getPassPipeline();
- if (auto customPipeline =
+ auto target = IREE::HAL::ExecutableTargetAttr::lookup(funcOp);
+ VMVXCodegenPipelineOptions vmvxOptions(
+ /*enableUKernels=*/target && hasUkernel(target.getConfiguration()));
+
+ // No pipeline specified, nothing to do.
+ if (isa<IREE::Codegen::NoPipelineAttr>(pipelineAttr)) {
+ return;
+ }
+
+ // Check for a pipeline via PipelineAttrInterface.
+ if (auto pipelineIface =
dyn_cast<IREE::Codegen::PipelineAttrInterface>(pipelineAttr)) {
- if (failed(customPipeline.buildPipeline(pipeline,
- /*options=*/nullptr))) {
- funcOp.emitOpError("failed to build custom pass pipeline");
+ if (failed(pipelineIface.buildPipeline(pipeline, &vmvxOptions))) {
+ funcOp.emitOpError("failed to build pass pipeline");
return signalPassFailure();
}
} else {
- auto target = IREE::HAL::ExecutableTargetAttr::lookup(funcOp);
- bool enableUKernels = target && hasUkernel(target.getConfiguration());
- switch (translationInfo.getDispatchLoweringPassPipeline()) {
- // No pipeline specified, nothing to do.
- case IREE::Codegen::DispatchLoweringPassPipeline::None:
- return;
- case IREE::Codegen::DispatchLoweringPassPipeline::VMVXDefault:
- addVMVXDefaultPassPipeline(pipeline, enableUKernels);
- break;
- default:
- funcOp.emitOpError("Unsupported pipeline on VMVX target.");
- return signalPassFailure();
- }
+ // Not an interface implementor -- reject any remaining legacy pipeline.
+ funcOp.emitOpError("Unsupported pipeline on VMVX target.");
+ return signalPassFailure();
}
LDBG() << "Using pass pipeline: ";
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/test/custom_pass_pipeline.mlir b/compiler/src/iree/compiler/Codegen/VMVX/test/custom_pass_pipeline.mlir
index a4e6579..18d2aba 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/test/custom_pass_pipeline.mlir
+++ b/compiler/src/iree/compiler/Codegen/VMVX/test/custom_pass_pipeline.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt --pass-pipeline='builtin.module(func.func(iree-vmvx-lower-executable-target))' %s | FileCheck %s
+// RUN: iree-opt --pass-pipeline='builtin.module(func.func(iree-vmvx-lower-executable-target))' --verify-diagnostics --split-input-file %s | FileCheck %s
// Verify that a custom pass pipeline specified via #iree_codegen.pass_pipeline
// attribute is executed by the VMVX lower executable target pass.
@@ -17,3 +17,16 @@
// CHECK-LABEL: func.func @test_custom_pipeline
// CHECK-SAME: (%[[ARG0:.+]]: index)
// CHECK-NEXT: return %[[ARG0]]
+
+// -----
+
+#executable_target = #hal.executable.target<"vmvx", "vmvx-bytecode-fb">
+
+// expected-error@unknown {{transform dialect codegen pipeline should be consumed by LowerExecutableUsingTransformDialect}}
+// expected-error@+1 {{'func.func' op failed to build pass pipeline}}
+func.func @transform_dialect_pipeline_leaks() attributes {
+ hal.executable.target = #executable_target,
+ translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen codegen_spec = @foo>
+} {
+ return
+}
diff --git a/compiler/src/iree/compiler/Codegen/VMVX/test/select_lowering_strategy.mlir b/compiler/src/iree/compiler/Codegen/VMVX/test/select_lowering_strategy.mlir
index abde135..f5fd73f 100644
--- a/compiler/src/iree/compiler/Codegen/VMVX/test/select_lowering_strategy.mlir
+++ b/compiler/src/iree/compiler/Codegen/VMVX/test/select_lowering_strategy.mlir
@@ -10,7 +10,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64, 64, 0]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @matmul_static
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.matmul
@@ -30,7 +30,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64, 64]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @copy_op_dynamic
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.generic
@@ -48,7 +48,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @static_1d_fft_stage2
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: iree_linalg_ext.fft
@@ -87,7 +87,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64, 64, 0]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @fusion_quant_matmul_generic
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.matmul
@@ -103,7 +103,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64, 64]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @unpack_outer_dynamic
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.unpack
@@ -132,7 +132,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [64, 64]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @elem_pack_ukernels
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.generic
@@ -159,7 +159,7 @@
return
}
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @copy_cst
// CHECK-SAME: translation_info = #[[TRANSLATION]]
@@ -200,7 +200,7 @@
}
// CHECK-DAG: #[[CONFIG:.+]] = #iree_cpu.lowering_config<distribution = [0]>
-// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = VMVXDefault>
+// CHECK-DAG: #[[TRANSLATION:.+]] = #iree_codegen.translation_info<pipeline = #iree_codegen.vmvx_pipeline>
// CHECK: func.func @already_distributed
// CHECK-SAME: translation_info = #[[TRANSLATION]]
// CHECK: linalg.generic
diff --git a/docs/website/docs/reference/tuning.md b/docs/website/docs/reference/tuning.md
index 6294cf0..238a981 100644
--- a/docs/website/docs/reference/tuning.md
+++ b/docs/website/docs/reference/tuning.md
@@ -284,7 +284,7 @@
transform.iree.match.dims_equal %m_dims, [32] : !transform.param<i64>
transform.iree.match.dims_equal %n_dims, [2048] : !transform.param<i64>
transform.iree.match.dims_equal %k_dims, [1024] : !transform.param<i64>
- %0 = transform.param.constant #iree_codegen.compilation_info<lowering_config = #iree_gpu.lowering_config<{mma_kind = #iree_gpu.mma_layout<MFMA_F32_32x32x8_F16>, promote_operands = [0, 1], reduction = [0, 0, 64], subgroup_basis = [[1, 2, 1], [0, 1, 2]], workgroup = [32, 256, 0]}>, translation_info = <pipeline = LLVMGPUVectorDistribute workgroup_size = [128, 1, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<prefetch_shared_memory = true>, llvm_func_attrs = {"amdgpu-waves-per-eu" = "2"}}>> -> !transform.any_param
+ %0 = transform.param.constant #iree_codegen.compilation_info<lowering_config = #iree_gpu.lowering_config<{mma_kind = #iree_gpu.mma_layout<MFMA_F32_32x32x8_F16>, promote_operands = [0, 1], reduction = [0, 0, 64], subgroup_basis = [[1, 2, 1], [0, 1, 2]], workgroup = [32, 256, 0]}>, translation_info = <pipeline = #iree_gpu.pipeline<VectorDistribute> workgroup_size = [128, 1, 1] subgroup_size = 64, {gpu_pipeline_options = #iree_gpu.pipeline_options<prefetch_shared_memory = true>, llvm_func_attrs = {"amdgpu-waves-per-eu" = "2"}}>> -> !transform.any_param
transform.yield %arg0, %0 : !transform.any_op, !transform.any_param
}
transform.named_sequence @__kernel_config(%arg0: !transform.any_op {transform.consumed}) -> !transform.any_op attributes {iree_codegen.tuning_spec_entrypoint} {
diff --git a/samples/transform_dialect/example_module.mlir b/samples/transform_dialect/example_module.mlir
index 2dd8161..a79ec1e 100644
--- a/samples/transform_dialect/example_module.mlir
+++ b/samples/transform_dialect/example_module.mlir
@@ -148,7 +148,7 @@
// RUN: FileCheck %s --check-prefixes=CODEGEN-PRINTER
// CODEGEN-PRINTER: IR printer: Setting matmul strategy to custom_transform_strategy
-// CODEGEN-PRINTER: translation_info = #iree_codegen.translation_info<pipeline = TransformDialectCodegen codegen_spec = @custom_transform_strategy>
+// CODEGEN-PRINTER: translation_info = #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen codegen_spec = @custom_transform_strategy>
// CODEGEN-PRINTER: IR printer: Setting reduce strategy to base vectorize top-level
// CODEGEN-PRINTER: translation_info = #iree_codegen.translation_info<pipeline = #iree_gpu.spirv_pipeline<BaseVectorize> workgroup_size = [16, 1, 1]>
diff --git a/samples/transform_dialect/transform_library.mlir b/samples/transform_dialect/transform_library.mlir
index f66a70b..fda287d 100644
--- a/samples/transform_dialect/transform_library.mlir
+++ b/samples/transform_dialect/transform_library.mlir
@@ -56,7 +56,7 @@
transform.named_sequence @custom_matmul(%matmul: !transform.any_op {transform.readonly}) {
%variant_op = transform.get_parent_op %matmul {op_name = "hal.executable.variant"} : (!transform.any_op) -> !transform.any_op
%funcs = transform.structured.match ops{["func.func"]} in %variant_op : (!transform.any_op) -> !transform.any_op
- %subgroup_reduce = transform.param.constant #iree_codegen.translation_info<pipeline = TransformDialectCodegen
+ %subgroup_reduce = transform.param.constant #iree_codegen.translation_info<pipeline = #iree_codegen.transform_dialect_codegen
codegen_spec = @custom_transform_strategy> -> !transform.any_param
transform.annotate %funcs "translation_info" = %subgroup_reduce : !transform.any_op, !transform.any_param
transform.print {name = "Setting matmul strategy to custom_transform_strategy"}
@@ -133,8 +133,9 @@
//
// To couple this with a transform dialect based codegen strategy, the target
// codegen strategy can be included inline with this library and relevant ops
- // can be annotated with `TransformDialectCodegen` as the lowering pipeline,
- // with a reference to the strategy to use (see an example above).
+ // can be annotated with `#iree_codegen.transform_dialect_codegen` as the
+ // lowering pipeline, with a reference to the strategy to use (see an example
+ // above).
transform.named_sequence @kernel_config(%variant_op: !transform.any_op {transform.consumed}) {
transform.foreach_match in %variant_op
@match_matmul -> @custom_matmul,
diff --git a/tests/transform_dialect/cpu/transform_library.mlir b/tests/transform_dialect/cpu/transform_library.mlir
index 92c8953..763919f 100644
--- a/tests/transform_dialect/cpu/transform_library.mlir
+++ b/tests/transform_dialect/cpu/transform_library.mlir
@@ -28,7 +28,7 @@
// CSE is needed on the workgroup_count region to pass this particular test.
transform.apply_cse to %memref_func : !transform.any_op
- %none_attr = transform.param.constant #iree_codegen.translation_info<pipeline = None> -> !transform.any_param
+ %none_attr = transform.param.constant #iree_codegen.translation_info<pipeline = #iree_codegen.no_pipeline> -> !transform.any_param
transform.annotate %memref_func "translation_info" = %none_attr : !transform.any_op, !transform.any_param
transform.yield
}
