Start porting the Conversion dir to tablegen pass registration. (#6011)
* Start porting the Conversion dir to tablegen pass registration.
* I had an immediate need for the HLOToLinalg directory to be converted.
* Just converted passes which were referenced in the same registration
function.
* Eventually, init_conversion.h goes away once the rest are done.
diff --git a/iree/compiler/Conversion/BUILD b/iree/compiler/Conversion/BUILD
index 2c33ea4..053cbd0 100644
--- a/iree/compiler/Conversion/BUILD
+++ b/iree/compiler/Conversion/BUILD
@@ -12,18 +12,64 @@
# See the License for the specific language governing permissions and
# limitations under the License.
+load("//build_tools/bazel:tblgen.bzl", "gentbl_cc_library")
+
package(
default_visibility = ["//visibility:public"],
features = ["layering_check"],
licenses = ["notice"], # Apache 2.0
)
+gentbl_cc_library(
+ name = "PassesIncGen",
+ tbl_outs = [
+ (
+ ["-gen-pass-decls"],
+ "Passes.h.inc",
+ ),
+ ],
+ tblgen = "@llvm-project//mlir:mlir-tblgen",
+ td_file = "Passes.td",
+ td_srcs = [
+ "@llvm-project//mlir:PassBaseTdFiles",
+ ],
+)
+
+cc_library(
+ name = "PassHeaders",
+ hdrs = [
+ "PassDetail.h",
+ "Passes.h",
+ "Passes.h.inc",
+ "Rewriters.h",
+ ],
+ deps = [
+ ":PassesIncGen",
+ "@llvm-project//mlir:Pass",
+ "@llvm-project//mlir:Transforms",
+ ],
+)
+
+cc_library(
+ name = "Conversion",
+ srcs = [
+ "Passes.cpp",
+ ],
+ deps = [
+ ":PassHeaders",
+ ":PassesIncGen",
+ "//iree/compiler/Conversion/Common",
+ "//iree/compiler/Conversion/HLOToLinalg",
+ ],
+)
+
cc_library(
name = "init_conversions",
hdrs = [
"init_conversions.h",
],
deps = [
+ ":PassHeaders",
"//iree/compiler/Conversion/LinalgToLLVM",
"//iree/compiler/Conversion/LinalgToLLVMGPU",
"//iree/compiler/Conversion/LinalgToLinalg",
diff --git a/iree/compiler/Conversion/CMakeLists.txt b/iree/compiler/Conversion/CMakeLists.txt
index 2de5a39..9d5132e 100644
--- a/iree/compiler/Conversion/CMakeLists.txt
+++ b/iree/compiler/Conversion/CMakeLists.txt
@@ -10,12 +10,48 @@
iree_add_all_subdirs()
+iree_tablegen_library(
+ NAME
+ PassesIncGen
+ TD_FILE
+ "Passes.td"
+ OUTS
+ -gen-pass-decls Passes.h.inc
+)
+
+iree_cc_library(
+ NAME
+ PassHeaders
+ HDRS
+ "PassDetail.h"
+ "Passes.h"
+ "Passes.h.inc"
+ "Rewriters.h"
+ DEPS
+ MLIRPass
+ MLIRTransforms
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
+ Conversion
+ SRCS
+ "Passes.cpp"
+ DEPS
+ ::PassHeaders
+ iree::compiler::Conversion::Common
+ iree::compiler::Conversion::HLOToLinalg
+ PUBLIC
+)
+
iree_cc_library(
NAME
init_conversions
HDRS
"init_conversions.h"
DEPS
+ ::PassHeaders
iree::compiler::Conversion::LinalgToLLVM
iree::compiler::Conversion::LinalgToLLVMGPU
iree::compiler::Conversion::LinalgToLinalg
@@ -26,3 +62,8 @@
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+# TODO: For some reason, these dependencies are not being added automatically.
+add_dependencies(
+ iree_compiler_Conversion_PassHeaders
+ iree_compiler_Conversion_PassesIncGen
+)
diff --git a/iree/compiler/Conversion/Common/BUILD b/iree/compiler/Conversion/Common/BUILD
index 3bd8795..a2c009b 100644
--- a/iree/compiler/Conversion/Common/BUILD
+++ b/iree/compiler/Conversion/Common/BUILD
@@ -58,6 +58,7 @@
"Transforms.h",
],
deps = [
+ "//iree/compiler/Conversion:PassHeaders",
"//iree/compiler/Conversion/CodegenUtils",
"//iree/compiler/Conversion/Common:FoldTensorExtractOpIncGen",
"//iree/compiler/Dialect/Flow/IR",
diff --git a/iree/compiler/Conversion/Common/CMakeLists.txt b/iree/compiler/Conversion/Common/CMakeLists.txt
index 7816f9c..df6fed8 100644
--- a/iree/compiler/Conversion/Common/CMakeLists.txt
+++ b/iree/compiler/Conversion/Common/CMakeLists.txt
@@ -58,6 +58,7 @@
MLIRVector
iree::compiler::Conversion::CodegenUtils
iree::compiler::Conversion::Common::FoldTensorExtractOpIncGen
+ iree::compiler::Conversion::PassHeaders
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::IREE::IR
diff --git a/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp b/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp
index 9a0243e..c59d102 100644
--- a/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp
+++ b/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp
@@ -15,7 +15,8 @@
#include <memory>
#include <utility>
-#include "iree/compiler/Conversion/Common/Passes.h"
+#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/IREE/IR/IREETypes.h"
#include "llvm/ADT/APFloat.h"
@@ -161,8 +162,7 @@
}
};
-struct ConvertF32ToF16Pass
- : public PassWrapper<ConvertF32ToF16Pass, OperationPass<ModuleOp>> {
+struct DemoteF32ToF16Pass : public DemoteF32ToF16Base<DemoteF32ToF16Pass> {
void runOnOperation() override {
MLIRContext *context = &getContext();
ModuleOp moduleOp = getOperation();
@@ -182,12 +182,8 @@
} // namespace
std::unique_ptr<OperationPass<ModuleOp>> createDemoteF32ToF16Pass() {
- return std::make_unique<ConvertF32ToF16Pass>();
+ return std::make_unique<DemoteF32ToF16Pass>();
}
-static PassRegistration<ConvertF32ToF16Pass> pass(
- "iree-convert-f32-to-f16",
- "Convert f32 operations and values into equivalent f16 ones.");
-
} // namespace iree_compiler
} // namespace mlir
diff --git a/iree/compiler/Conversion/Common/Passes.h b/iree/compiler/Conversion/Common/Passes.h
index bd29a23..459cbbb 100644
--- a/iree/compiler/Conversion/Common/Passes.h
+++ b/iree/compiler/Conversion/Common/Passes.h
@@ -62,10 +62,6 @@
/// backends that require linearized access.
std::unique_ptr<OperationPass<ModuleOp>> createFlattenMemRefSubspanPass();
-/// Create a pass to convert a model using f32 type to the equivalent one
-/// using 16.
-std::unique_ptr<OperationPass<ModuleOp>> createDemoteF32ToF16Pass();
-
/// Sets the number of workgroups to use for each entry point in the dispatch
/// region.
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
diff --git a/iree/compiler/Conversion/HLOToLinalg/BUILD b/iree/compiler/Conversion/HLOToLinalg/BUILD
index 0647911..e87f6f5 100644
--- a/iree/compiler/Conversion/HLOToLinalg/BUILD
+++ b/iree/compiler/Conversion/HLOToLinalg/BUILD
@@ -19,18 +19,18 @@
)
cc_library(
- name = "HLOToLinalgOnTensors",
+ name = "HLOToLinalg",
srcs = [
"FusionOfTensorOps.cpp",
"HLOToLinalgOnTensors.cpp",
- ],
- hdrs = [
- "HLOToLinalgOnTensorPasses.h",
+ "ResolveShapeOps.cpp",
],
deps = [
+ "//iree/compiler/Conversion:PassHeaders",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/IR",
"//iree/compiler/Dialect/HAL/IR:HALDialect",
+ "//iree/compiler/Dialect/IREE/IR",
"//iree/compiler/Dialect/Shape/IR",
"@llvm-project//mlir:Affine",
"@llvm-project//mlir:IR",
@@ -47,25 +47,3 @@
"@mlir-hlo//:legalize_to_linalg",
],
)
-
-cc_library(
- name = "HLOToLinalg",
- srcs = [
- "ResolveShapeOps.cpp",
- ],
- hdrs = [
- "Passes.h",
- ],
- deps = [
- "//iree/compiler/Dialect/IREE/IR",
- "//iree/compiler/Dialect/Shape/IR",
- "@llvm-project//llvm:Support",
- "@llvm-project//mlir:DialectUtils",
- "@llvm-project//mlir:IR",
- "@llvm-project//mlir:MemRefDialect",
- "@llvm-project//mlir:Pass",
- "@llvm-project//mlir:StandardOps",
- "@llvm-project//mlir:Support",
- "@llvm-project//mlir:Transforms",
- ],
-)
diff --git a/iree/compiler/Conversion/HLOToLinalg/CMakeLists.txt b/iree/compiler/Conversion/HLOToLinalg/CMakeLists.txt
index 1048910..4a82722 100644
--- a/iree/compiler/Conversion/HLOToLinalg/CMakeLists.txt
+++ b/iree/compiler/Conversion/HLOToLinalg/CMakeLists.txt
@@ -12,12 +12,11 @@
iree_cc_library(
NAME
- HLOToLinalgOnTensors
- HDRS
- "HLOToLinalgOnTensorPasses.h"
+ HLOToLinalg
SRCS
"FusionOfTensorOps.cpp"
"HLOToLinalgOnTensors.cpp"
+ "ResolveShapeOps.cpp"
DEPS
MLIRAffine
MLIRIR
@@ -30,31 +29,13 @@
MLIRSupport
MLIRTensor
MLIRTransforms
+ iree::compiler::Conversion::PassHeaders
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::HAL::IR::HALDialect
- iree::compiler::Dialect::Shape::IR
- tensorflow::mlir_hlo
- PUBLIC
-)
-
-iree_cc_library(
- NAME
- HLOToLinalg
- HDRS
- "Passes.h"
- SRCS
- "ResolveShapeOps.cpp"
- DEPS
- LLVMSupport
- MLIRIR
- MLIRMemRef
- MLIRPass
- MLIRStandard
- MLIRSupport
- MLIRTransforms
iree::compiler::Dialect::IREE::IR
iree::compiler::Dialect::Shape::IR
+ tensorflow::mlir_hlo
PUBLIC
)
diff --git a/iree/compiler/Conversion/HLOToLinalg/FusionOfTensorOps.cpp b/iree/compiler/Conversion/HLOToLinalg/FusionOfTensorOps.cpp
index b8755f3..8a40bd8 100644
--- a/iree/compiler/Conversion/HLOToLinalg/FusionOfTensorOps.cpp
+++ b/iree/compiler/Conversion/HLOToLinalg/FusionOfTensorOps.cpp
@@ -20,7 +20,8 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h"
+#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -41,7 +42,7 @@
/// Pass to fuse linalg on tensor operations as well as fusion of hal.interface*
/// operations with linalg.tensor_reshape operation.
struct FusionOfTensorOpsPass
- : public PassWrapper<FusionOfTensorOpsPass, OperationPass<>> {
+ : public FusionOfTensorOpsBase<FusionOfTensorOpsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry
.insert<AffineDialect, IREE::HAL::HALDialect, linalg::LinalgDialect>();
@@ -120,8 +121,5 @@
return std::make_unique<FusionOfTensorOpsPass>();
}
-static PassRegistration<FusionOfTensorOpsPass> pass(
- "iree-codegen-fusion-of-tensor-ops", "Fuse operations on tensors");
-
} // namespace iree_compiler
} // namespace mlir
diff --git a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h
deleted file mode 100644
index bd4826c..0000000
--- a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h
+++ /dev/null
@@ -1,47 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-//===- HLOToLinalgOnTensorsPasses.h - Passes to convert from HLO To Linalg ===//
-//
-// IREE specific passes used in the HLO to Linalg on tensors conversion as well
-// as fusion.
-//
-//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_HLOTOLINALGONTENSORS_PASSES_H_
-#define IREE_COMPILER_CONVERSION_HLOTOLINALGONTENSORS_PASSES_H_
-
-#include "mlir/IR/BuiltinDialect.h"
-#include "mlir/Pass/Pass.h"
-#include "mlir/Transforms/DialectConversion.h"
-
-namespace mlir {
-namespace iree_compiler {
-
-/// Creates a pass to fuse operations on tensors.
-std::unique_ptr<Pass> createFusionOfTensorOpsPass();
-
-/// Creates XLA-HLO to Linalg on tensors transformation pass.
-std::unique_ptr<OperationPass<FuncOp>> createHLOToLinalgOnTensorsPass(
- bool useLinalgOnTensorsPath = false);
-
-/// Populates the patterns that convert from MHLO to Linalg on tensors. Imports
-/// patterns from XLA, as well as some IREE specific modifications.
-void populateHLOToLinalgOnTensorsConversionPatterns(
- MLIRContext *context, TypeConverter &typeConverter,
- OwningRewritePatternList &patterns);
-
-} // namespace iree_compiler
-} // namespace mlir
-
-#endif // IREE_COMPILER_CONVERSION_HLOTOLINALGONTENSORS_PASSES_H_
diff --git a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensors.cpp b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensors.cpp
index bd01378..be48d53 100644
--- a/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensors.cpp
+++ b/iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensors.cpp
@@ -21,7 +21,9 @@
//===----------------------------------------------------------------------===//
#include <memory>
-#include "iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h"
+#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Conversion/Rewriters.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeDialect.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeOps.h"
@@ -291,26 +293,21 @@
} // namespace
struct ConvertHLOToLinalgOnTensorsPass
- : public PassWrapper<ConvertHLOToLinalgOnTensorsPass, FunctionPass> {
- ConvertHLOToLinalgOnTensorsPass(bool useLinalgOnTensorsPath = false)
- : useLinalgOnTensorsPath(useLinalgOnTensorsPath){};
-
+ : public ConvertHLOToLinalgOnTensorsBase<ConvertHLOToLinalgOnTensorsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::Flow::FlowDialect, linalg::LinalgDialect,
mhlo::MhloDialect, ShapeDialect, math::MathDialect,
memref::MemRefDialect>();
}
- void runOnFunction() override {
+ void runOnOperation() override {
OwningRewritePatternList patterns(&getContext());
MLIRContext *context = &getContext();
auto typeConverter = mhlo::createHloToLinalgSignedIntegerConverter();
populateHLOToLinalgOnTensorsConversionPatterns(context, *typeConverter,
patterns);
- if (useLinalgOnTensorsPath) {
- patterns.insert<PadTensorOpConversion>(context);
- }
+ patterns.insert<PadTensorOpConversion>(context);
ConversionTarget target(getContext());
target.addIllegalDialect<mhlo::MhloDialect>();
@@ -328,27 +325,14 @@
// Let the rest fall through.
target.markUnknownOpDynamicallyLegal([](Operation *) { return true; });
- if (useLinalgOnTensorsPath) {
- // Set linalg.pad_tensor illegal for now.
- target.addIllegalOp<linalg::PadTensorOp>();
- }
+ // Set linalg.pad_tensor illegal for now.
+ target.addIllegalOp<linalg::PadTensorOp>();
- if (failed(applyPartialConversion(getFunction(), target,
+ if (failed(applyPartialConversion(getOperation(), target,
std::move(patterns)))) {
return signalPassFailure();
}
}
-
- private:
- bool useLinalgOnTensorsPath;
-};
-
-/// This pass is just added for lit-testing when using the linalg on tensors
-/// path. Remove when the linalg on tensors path becomes default.
-struct ConvertHLOToLinalgOnTensorsPassExperimental
- : public ConvertHLOToLinalgOnTensorsPass {
- ConvertHLOToLinalgOnTensorsPassExperimental()
- : ConvertHLOToLinalgOnTensorsPass(true) {}
};
/// Convert mhlo.constant op into std.const.
@@ -383,24 +367,9 @@
typeConverter, context, PatternBenefit(1000));
}
-static llvm::cl::opt<bool> clUseLinalgOnTensorsPath(
- "iree-linalg-on-tensors-path",
- llvm::cl::desc("Convert from MHLO to Linalg on tensors for linalg on "
- "tensor codegen path"),
- llvm::cl::init(false));
-
-std::unique_ptr<OperationPass<FuncOp>> createHLOToLinalgOnTensorsPass(
- bool useLinalgOnTensorsPath) {
- return std::make_unique<ConvertHLOToLinalgOnTensorsPass>(
- useLinalgOnTensorsPath);
+std::unique_ptr<OperationPass<FuncOp>> createHLOToLinalgOnTensorsPass() {
+ return std::make_unique<ConvertHLOToLinalgOnTensorsPass>();
}
-static PassRegistration<ConvertHLOToLinalgOnTensorsPass> legalize_pass(
- "iree-codegen-hlo-to-linalg-on-tensors",
- "Convert from XLA-HLO ops to Linalg ops on tensors", []() {
- return std::make_unique<ConvertHLOToLinalgOnTensorsPass>(
- clUseLinalgOnTensorsPath);
- });
-
} // namespace iree_compiler
} // namespace mlir
diff --git a/iree/compiler/Conversion/HLOToLinalg/Passes.h b/iree/compiler/Conversion/HLOToLinalg/Passes.h
deleted file mode 100644
index 9584f4f..0000000
--- a/iree/compiler/Conversion/HLOToLinalg/Passes.h
+++ /dev/null
@@ -1,36 +0,0 @@
-// Copyright 2020 Google LLC
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// https://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-//===- Passes.h - Codegen pass to convert from XLA to Linalg on buffers----===//
-//
-// IREE specific passes used in the XLA to Linalg conversion
-//
-//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_HLOTOLINALG_PASSES_H_
-#define IREE_COMPILER_CONVERSION_HLOTOLINALG_PASSES_H_
-#include <memory>
-
-#include "mlir/IR/BuiltinOps.h"
-#include "mlir/Pass/Pass.h"
-
-namespace mlir {
-namespace iree_compiler {
-
-/// Resolves shape related ops (std.dim, shapex.tie_shape, etc.) by tracing
-/// them back to the original HAL interface bindings.
-std::unique_ptr<OperationPass<FuncOp>> createResolveShapeOpsPass();
-} // namespace iree_compiler
-} // namespace mlir
-
-#endif // IREE_COMPILER_CONVERSION_HLOTOLINALG_PASSES_H_
diff --git a/iree/compiler/Conversion/HLOToLinalg/ResolveShapeOps.cpp b/iree/compiler/Conversion/HLOToLinalg/ResolveShapeOps.cpp
index 4f9107d..eed39e4 100644
--- a/iree/compiler/Conversion/HLOToLinalg/ResolveShapeOps.cpp
+++ b/iree/compiler/Conversion/HLOToLinalg/ResolveShapeOps.cpp
@@ -20,6 +20,8 @@
//
//===----------------------------------------------------------------------===//
+#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeDialect.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeOps.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -85,17 +87,16 @@
}
};
-struct ResolveShapeOpsPass
- : public PassWrapper<ResolveShapeOpsPass, FunctionPass> {
+struct ResolveShapeOpsPass : public ResolveShapeOpsBase<ResolveShapeOpsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<ShapeDialect>();
}
- void runOnFunction() override;
+ void runOnOperation() override;
};
} // namespace
-void ResolveShapeOpsPass::runOnFunction() {
+void ResolveShapeOpsPass::runOnOperation() {
MLIRContext *context = &getContext();
OwningRewritePatternList dimPatterns(&getContext());
@@ -106,8 +107,8 @@
ConversionTarget target(*context);
target.addIllegalOp<memref::DimOp>();
target.markUnknownOpDynamicallyLegal([](Operation *) { return true; });
- if (failed(
- applyFullConversion(getFunction(), target, std::move(dimPatterns)))) {
+ if (failed(applyFullConversion(getOperation(), target,
+ std::move(dimPatterns)))) {
return signalPassFailure();
}
@@ -117,14 +118,12 @@
// Then elide all shapex.tie_shape ops and canonicalize shapex.ranked_dim
// given that we don't need the shape annotation anymore.
- (void)applyPatternsAndFoldGreedily(getFunction(), std::move(shapePatterns));
+ (void)applyPatternsAndFoldGreedily(getOperation(), std::move(shapePatterns));
}
std::unique_ptr<OperationPass<FuncOp>> createResolveShapeOpsPass() {
return std::make_unique<ResolveShapeOpsPass>();
}
-static PassRegistration<ResolveShapeOpsPass> pass("iree-codegen-resolve-shape",
- "resolve shape");
} // namespace iree_compiler
} // namespace mlir
diff --git a/iree/compiler/Conversion/HLOToLinalg/test/pad_tensor_to_tensor.mlir b/iree/compiler/Conversion/HLOToLinalg/test/pad_tensor_to_tensor.mlir
index f7a2bb9..e63eb68 100644
--- a/iree/compiler/Conversion/HLOToLinalg/test/pad_tensor_to_tensor.mlir
+++ b/iree/compiler/Conversion/HLOToLinalg/test/pad_tensor_to_tensor.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-codegen-hlo-to-linalg-on-tensors -iree-linalg-on-tensors-path -canonicalize %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-codegen-hlo-to-linalg-on-tensors -canonicalize %s | IreeFileCheck %s
module {
func @pad_tensor(%arg0 : tensor<?x?xf32>, %arg1 : tensor<f32>, %arg2 : index, %arg3 : index) -> tensor<?x?xf32> {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/BUILD b/iree/compiler/Conversion/LinalgToSPIRV/BUILD
index aad7268..5eeff1e 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/BUILD
+++ b/iree/compiler/Conversion/LinalgToSPIRV/BUILD
@@ -54,6 +54,7 @@
],
deps = [
":CodeGenOptionUtils",
+ "//iree/compiler/Conversion:PassHeaders",
"//iree/compiler/Conversion/CodegenUtils",
"//iree/compiler/Conversion/Common",
"//iree/compiler/Conversion/HLOToLinalg",
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt b/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt
index 93e63a2..64bb005 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt
+++ b/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt
@@ -75,6 +75,7 @@
iree::compiler::Conversion::Common
iree::compiler::Conversion::HLOToLinalg
iree::compiler::Conversion::LinalgToVector
+ iree::compiler::Conversion::PassHeaders
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::HAL::IR::HALDialect
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp b/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp
index 4f34aff..a24c5ab 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp
+++ b/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp
@@ -21,10 +21,10 @@
#include "iree/compiler/Conversion/LinalgToSPIRV/Passes.h"
#include "iree/compiler/Conversion/Common/Passes.h"
-#include "iree/compiler/Conversion/HLOToLinalg/Passes.h"
#include "iree/compiler/Conversion/LinalgToSPIRV/CodeGenOptionUtils.h"
#include "iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h"
#include "iree/compiler/Conversion/LinalgToVector/Passes.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "llvm/Support/CommandLine.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
diff --git a/iree/compiler/Conversion/PassDetail.h b/iree/compiler/Conversion/PassDetail.h
new file mode 100644
index 0000000..f4a2bb5
--- /dev/null
+++ b/iree/compiler/Conversion/PassDetail.h
@@ -0,0 +1,29 @@
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef IREE_COMPILER_CONVERSIONS_PASS_DETAIL_H_
+#define IREE_COMPILER_CONVERSIONS_PASS_DETAIL_H_
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+#define GEN_PASS_CLASSES
+#include "iree/compiler/Conversion/Passes.h.inc"
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_CONVERSIONS_PASS_DETAIL_H_
diff --git a/iree/compiler/Conversion/Passes.cpp b/iree/compiler/Conversion/Passes.cpp
new file mode 100644
index 0000000..376259e
--- /dev/null
+++ b/iree/compiler/Conversion/Passes.cpp
@@ -0,0 +1,31 @@
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "iree/compiler/Conversion/Passes.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+namespace {
+#define GEN_PASS_REGISTRATION
+#include "iree/compiler/Conversion/Passes.h.inc"
+} // namespace
+
+void registerConversionPasses() {
+ // Generated.
+ registerPasses();
+}
+
+} // namespace iree_compiler
+} // namespace mlir
diff --git a/iree/compiler/Conversion/Passes.h b/iree/compiler/Conversion/Passes.h
new file mode 100644
index 0000000..1b0898c
--- /dev/null
+++ b/iree/compiler/Conversion/Passes.h
@@ -0,0 +1,53 @@
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef IREE_COMPILER_CONVERSION_PASSES_H_
+#define IREE_COMPILER_CONVERSION_PASSES_H_
+
+#include <memory>
+
+#include "mlir/Pass/Pass.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+// Registers all conversion passes in this directory.
+void registerConversionPasses();
+
+//------------------------------------------------------------------------------
+// Conversions into Linalg
+//------------------------------------------------------------------------------
+
+/// Creates a pass to fuse Linalg operations on tensors.
+std::unique_ptr<Pass> createFusionOfTensorOpsPass();
+
+/// Creates XLA-HLO to Linalg on tensors transformation pass.
+std::unique_ptr<OperationPass<FuncOp>> createHLOToLinalgOnTensorsPass();
+
+/// Resolves shape related ops (std.dim, shapex.tie_shape, etc.) by tracing
+/// them back to the original HAL interface bindings.
+std::unique_ptr<OperationPass<FuncOp>> createResolveShapeOpsPass();
+
+//------------------------------------------------------------------------------
+// Misc/common conversions
+//------------------------------------------------------------------------------
+
+/// Create a pass to convert a model using f32 type to the equivalent one
+/// using f16.
+std::unique_ptr<OperationPass<ModuleOp>> createDemoteF32ToF16Pass();
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_CONVERSION_PASSES_H_
diff --git a/iree/compiler/Conversion/Passes.td b/iree/compiler/Conversion/Passes.td
new file mode 100644
index 0000000..26abb4e
--- /dev/null
+++ b/iree/compiler/Conversion/Passes.td
@@ -0,0 +1,45 @@
+
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef IREE_CONVERSION_PASSES
+#define IREE_CONVERSION_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+def ConvertHLOToLinalgOnTensors :
+ Pass<"iree-codegen-hlo-to-linalg-on-tensors", "FuncOp"> {
+ let summary = "Convert from XLA-HLO ops to Linalg ops on tensors";
+ let constructor = "mlir::iree_compiler::createHLOToLinalgOnTensorsPass()";
+}
+
+def DemoteF32ToF16 :
+ Pass<"iree-convert-f32-to-f16", "ModuleOp"> {
+ let summary = "Convert f32 operations and values into equivalent f16 ones.";
+ let constructor = "mlir::iree_compiler::createDemoteF32ToF16Pass()";
+}
+
+def FusionOfTensorOps :
+ Pass<"iree-codegen-fusion-of-tensor-ops", ""> {
+ let summary = "Fuse operations on tensors";
+ let constructor = "mlir::iree_compiler::createFusionOfTensorOpsPass()";
+}
+
+def ResolveShapeOps :
+ Pass<"iree-codegen-resolve-shape", "FuncOp"> {
+ let summary = "resolve shapes";
+ let constructor = "mlir::iree_compiler::createResolveShapeOpsPass()";
+}
+
+#endif // IREE_DIALECT_FLOW_PASSES
diff --git a/iree/compiler/Conversion/Rewriters.h b/iree/compiler/Conversion/Rewriters.h
new file mode 100644
index 0000000..9c4ceba
--- /dev/null
+++ b/iree/compiler/Conversion/Rewriters.h
@@ -0,0 +1,40 @@
+// Copyright 2021 Google LLC
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// https://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef IREE_COMPILER_CONVERSION_REWRITER_H_
+#define IREE_COMPILER_CONVERSION_REWRITER_H_
+
+#include "mlir/Transforms/DialectConversion.h"
+
+namespace mlir {
+namespace iree_compiler {
+
+/// Populates the patterns that convert from MHLO to Linalg on tensors. Imports
+/// patterns from XLA, as well as some IREE specific modifications.
+void populateHLOToLinalgOnTensorsConversionPatterns(
+ MLIRContext *context, TypeConverter &typeConverter,
+ OwningRewritePatternList &patterns);
+
+/// Populates IREE specific patterns to convert HLO broadcasting ops to Linalg.
+/// These are being maintained separately because they are a standalone unit
+/// that is both intricate and possible to upstream, should there be alignment
+/// to do so.
+void populateHLOBroadcastingToLinalgPatterns(
+ MLIRContext *context, TypeConverter &typeConverter,
+ OwningRewritePatternList &patterns);
+
+} // namespace iree_compiler
+} // namespace mlir
+
+#endif // IREE_COMPILER_CONVERSION_REWRITER_H_
diff --git a/iree/compiler/Conversion/init_conversions.h b/iree/compiler/Conversion/init_conversions.h
index 9740f55..c7143f3 100644
--- a/iree/compiler/Conversion/init_conversions.h
+++ b/iree/compiler/Conversion/init_conversions.h
@@ -16,11 +16,11 @@
#define IREE_COMPILER_CONVERSION_INIT_CONVERSIONS_H_
#include "iree/compiler/Conversion/Common/Passes.h"
-#include "iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h"
#include "iree/compiler/Conversion/LinalgToLLVM/Passes.h"
#include "iree/compiler/Conversion/LinalgToLinalg/Passes.h"
#include "iree/compiler/Conversion/LinalgToSPIRV/Passes.h"
#include "iree/compiler/Conversion/LinalgToVector/Passes.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Conversion/VectorToLLVM/Passes.h"
namespace mlir {
@@ -42,15 +42,6 @@
(void)init_once;
}
-inline void registerHLOToLinalgPasses() {
- static bool init_once = []() {
- createHLOToLinalgOnTensorsPass();
- createDemoteF32ToF16Pass();
- return true;
- }();
- (void)init_once;
-}
-
inline void registerLinalgToVectorPasses() {
static bool init_once = []() {
createVectorizeLinalgConvPass();
diff --git a/iree/compiler/Dialect/Flow/Transforms/BUILD b/iree/compiler/Dialect/Flow/Transforms/BUILD
index 1853f9a..1f9bdf3 100644
--- a/iree/compiler/Dialect/Flow/Transforms/BUILD
+++ b/iree/compiler/Dialect/Flow/Transforms/BUILD
@@ -22,7 +22,7 @@
)
gentbl_cc_library(
- name = "Passes_inc_gen",
+ name = "PassesIncGen",
tbl_outs = [
(
["-gen-pass-decls"],
@@ -63,8 +63,9 @@
"Passes.h.inc",
],
deps = [
- ":Passes_inc_gen",
- "//iree/compiler/Conversion/HLOToLinalg:HLOToLinalgOnTensors",
+ ":PassesIncGen",
+ "//iree/compiler/Conversion:PassHeaders",
+ "//iree/compiler/Conversion/HLOToLinalg",
"//iree/compiler/Conversion/LinalgToLinalg",
"//iree/compiler/Dialect/Flow/Conversion",
"//iree/compiler/Dialect/Flow/Conversion/HLOToFlow",
diff --git a/iree/compiler/Dialect/Flow/Transforms/CMakeLists.txt b/iree/compiler/Dialect/Flow/Transforms/CMakeLists.txt
index 4820ea1..9770e68 100644
--- a/iree/compiler/Dialect/Flow/Transforms/CMakeLists.txt
+++ b/iree/compiler/Dialect/Flow/Transforms/CMakeLists.txt
@@ -12,7 +12,7 @@
iree_tablegen_library(
NAME
- Passes_inc_gen
+ PassesIncGen
TD_FILE
"Passes.td"
OUTS
@@ -63,8 +63,9 @@
MLIRTosaToStandard
MLIRTransformUtils
MLIRTransforms
- iree::compiler::Conversion::HLOToLinalg::HLOToLinalgOnTensors
+ iree::compiler::Conversion::HLOToLinalg
iree::compiler::Conversion::LinalgToLinalg
+ iree::compiler::Conversion::PassHeaders
iree::compiler::Dialect::Flow::Conversion
iree::compiler::Dialect::Flow::Conversion::HLOToFlow
iree::compiler::Dialect::Flow::Conversion::StandardToFlow
@@ -85,3 +86,8 @@
PROPERTY COMPILE_FLAGS $<$<CXX_COMPILER_ID:GNU>:-fno-devirtualize>)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
+# TODO: For some reason, these dependencies are not being added automatically.
+add_dependencies(
+ iree_compiler_Dialect_Flow_Transforms_Transforms
+ iree_compiler_Dialect_Flow_Transforms_PassesIncGen
+)
diff --git a/iree/compiler/Dialect/Flow/Transforms/Passes.cpp b/iree/compiler/Dialect/Flow/Transforms/Passes.cpp
index 84cca68..21f12fc 100644
--- a/iree/compiler/Dialect/Flow/Transforms/Passes.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/Passes.cpp
@@ -16,8 +16,8 @@
#include <memory>
-#include "iree/compiler/Conversion/HLOToLinalg/HLOToLinalgOnTensorPasses.h"
#include "iree/compiler/Conversion/LinalgToLinalg/Passes.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/Shape/Conversion/Passes.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "mlir-hlo/Dialect/mhlo/transforms/passes.h"
@@ -200,7 +200,7 @@
// TODO(benvanik): move up to input; requires pre-partitioning conversion
// to be reworked first.
passManager.addNestedPass<FuncOp>(
- mlir::iree_compiler::createHLOToLinalgOnTensorsPass(true));
+ mlir::iree_compiler::createHLOToLinalgOnTensorsPass());
if (clEnable1x1ConvToMatmul) {
passManager.addNestedPass<FuncOp>(
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD b/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
index 93f6d87..1c14d49 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
@@ -28,6 +28,7 @@
"Passes.h",
],
deps = [
+ "//iree/compiler/Conversion:PassHeaders",
"//iree/compiler/Conversion/Common",
"//iree/compiler/Conversion/HLOToLinalg",
"//iree/compiler/Conversion/LinalgToLLVM",
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt b/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
index e302496..e95f514 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
@@ -43,6 +43,7 @@
iree::compiler::Conversion::Common
iree::compiler::Conversion::HLOToLinalg
iree::compiler::Conversion::LinalgToLLVM
+ iree::compiler::Conversion::PassHeaders
iree::compiler::Dialect::HAL::IR::HALDialect
iree::compiler::Dialect::HAL::Transforms
iree::compiler::Dialect::IREE::IR
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp b/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
index 01a8e15..9f96a7e 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
@@ -17,8 +17,8 @@
#include <memory>
#include "iree/compiler/Conversion/Common/Passes.h"
-#include "iree/compiler/Conversion/HLOToLinalg/Passes.h"
#include "iree/compiler/Conversion/LinalgToLLVM/Passes.h"
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Dialect/HAL/Transforms/Passes.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index 2ac6a86..9821862 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -150,6 +150,7 @@
":init_iree_passes_and_dialects",
":init_mlir_passes_and_dialects",
":init_xla_dialects",
+ "//iree/compiler/Conversion",
"//iree/compiler/Conversion:init_conversions",
"//iree/compiler/Dialect/HAL/Conversion:Passes",
],
@@ -306,6 +307,7 @@
":init_targets",
":init_translations",
":init_xla_dialects",
+ "//iree/compiler/Conversion",
"//iree/compiler/Conversion:init_conversions",
"//iree/compiler/Dialect/VM/Target:init_targets",
"//iree/compiler/Dialect/VM/Target/Bytecode",
diff --git a/iree/tools/CMakeLists.txt b/iree/tools/CMakeLists.txt
index cb31528..db596f3 100644
--- a/iree/tools/CMakeLists.txt
+++ b/iree/tools/CMakeLists.txt
@@ -222,6 +222,7 @@
::init_mlir_passes_and_dialects
::init_xla_dialects
iree::compiler::Conversion::init_conversions
+ iree::compiler::Conversion::Conversion
iree::compiler::Dialect::HAL::Conversion::Passes
INCLUDES
"${IREE_EMITC_INCLUDES}"
@@ -310,6 +311,7 @@
MLIRTargetLLVMIRExport
MLIRTranslation
iree::compiler::Conversion::init_conversions
+ iree::compiler::Conversion::Conversion
iree::compiler::Dialect::VM::Target::Bytecode
iree::compiler::Dialect::VM::Target::init_targets
iree::compiler::Translation::IREEVM
diff --git a/iree/tools/init_passes.h b/iree/tools/init_passes.h
index 8a102dd..d932d0b 100644
--- a/iree/tools/init_passes.h
+++ b/iree/tools/init_passes.h
@@ -19,6 +19,7 @@
#include <cstdlib>
+#include "iree/compiler/Conversion/Passes.h"
#include "iree/compiler/Conversion/init_conversions.h"
#include "iree/compiler/Dialect/HAL/Conversion/Passes.h"
#include "iree/tools/init_iree_passes.h"
@@ -31,10 +32,10 @@
inline void registerAllPasses() {
registerAllIreePasses();
registerCommonConversionPasses();
+ registerConversionPasses();
registerMlirPasses();
registerHALConversionPasses();
registerLinalgToSPIRVPasses();
- registerHLOToLinalgPasses();
registerLinalgToLLVMPasses();
registerLinalgToLinalgPasses();
registerVectorToLLVMPasses();
