[CodeGen][Common] Switch to new pass generation tablegen definitions. (#18166)
This is not a NFC change because there are non-trivial changes in three
passes. The rest are just switching to the new tablegen definition
style.
The revision applies few cleanups during refactoring, which includes:
- Make all the passes have `iree-codegen-*` prefix.
- Make filename match the pass name.
- Switch namespaces to the single-line syntax.
The non-trivial changes happen in `IREEComprehensiveBufferizePass`,
`TileAndDistributeToWorkgroupsPass`, and
`TransformDialectInterpreterPass`.
In the `IREEComprehensiveBufferizePass`, the default allocationFn and
memcpyFn values are no longer `std::nullopt`. They are the default
functions that implemented in the file. It matches the behavior. The
only difference is that it does not rely on users to pass `std::nullopt`
to trigger the "default" behavior. Instead, it initialize the default
value of the class members be them. This saves one level of nesting
logic.
In the `TileAndDistributeToWorkgroupsPass`, it switches to use
`clEnumValN`. Without the change, the values in `.td` are magic numbers.
Now we explicitly tie them to linalg::DistributionMethod.
In `TransformDialectInterpreterPass`, it no longer uses
`--iree-codegen-transform-dialect-library` to get the library path.
Instead, users should use the predefined option (i.e.,
`library-file-name`) to specify the path. The old behavior was a hack
because the pass itself should not access external variable if possible.
The `--iree-codegen-transform-dialect-library` is still valid, but the
lit test does not use it anymore. Because it is very tricky to teach
tablegen to generate such behavior.
---------
Signed-off-by: hanhanW <hanhan0912@gmail.com>
diff --git a/compiler/src/iree/compiler/Codegen/Common/AddFastMathFlags.cpp b/compiler/src/iree/compiler/Codegen/Common/AddFastMathFlags.cpp
index b66169e..79f5629 100644
--- a/compiler/src/iree/compiler/Codegen/Common/AddFastMathFlags.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/AddFastMathFlags.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/IR/Matchers.h"
@@ -13,8 +12,10 @@
#define DEBUG_TYPE "iree-codegen-add-fast-math-flags"
-using namespace mlir;
-using namespace mlir::iree_compiler;
+namespace mlir::iree_compiler {
+
+#define GEN_PASS_DEF_ADDFASTMATHFLAGSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
/// Add `contract` FMF to operations that support it.
static void addContractFMF(Operation *op) {
@@ -29,10 +30,11 @@
/// Add the corresponding fast-math flags to operations given a floating-point
/// optimization mode.
// TODO: For now we only allow default flags, such as arithmetic reassociation.
-struct AddFastMathFlagsPass
- : public AddFastMathFlagsBase<AddFastMathFlagsPass> {
+struct AddFastMathFlagsPass final
+ : impl::AddFastMathFlagsPassBase<AddFastMathFlagsPass> {
public:
- using AddFastMathFlagsBase::AddFastMathFlagsBase;
+ using impl::AddFastMathFlagsPassBase<
+ AddFastMathFlagsPass>::AddFastMathFlagsPassBase;
void runOnOperation() override {
getOperation()->walk([](Operation *op) { addContractFMF(op); });
@@ -40,8 +42,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<LLVM::LLVMFuncOp>>
-mlir::iree_compiler::createAddFastMathFlagsPass() {
- return std::make_unique<AddFastMathFlagsPass>();
-}
+} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/BUILD.bazel b/compiler/src/iree/compiler/Codegen/Common/BUILD.bazel
index ec3eaec..08e7e9d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/BUILD.bazel
+++ b/compiler/src/iree/compiler/Codegen/Common/BUILD.bazel
@@ -28,7 +28,6 @@
iree_compiler_cc_library(
name = "PassHeaders",
hdrs = [
- "PassDetail.h",
"Passes.h",
"Passes.h.inc",
],
@@ -133,9 +132,10 @@
"PropagateReshapesByExpansion.cpp",
"ReconcileTranslationInfo.cpp",
"RematerializeParallelOps.cpp",
- "RemoveTrivialLoops.cpp",
+ "RemoveSingleIterationLoop.cpp",
"ReplaceSlowMinMaxOps.cpp",
"SplitFullPartialTransferPass.cpp",
+ "TensorToVectorVectorizePad.cpp",
"TestExecutablePreprocessing.cpp",
"TestPartitionableLoopsInterface.cpp",
"TileAndDistributeToWorkgroupsPass.cpp",
@@ -144,7 +144,6 @@
"TypePropagationPass.cpp",
"UserConfig.cpp",
"VectorizeMemrefCopy.cpp",
- "VectorizePad.cpp",
],
hdrs = [
"BufferizationAnalysis.h",
diff --git a/compiler/src/iree/compiler/Codegen/Common/BubbleUpOrdinalOps.cpp b/compiler/src/iree/compiler/Codegen/Common/BubbleUpOrdinalOps.cpp
index 5fdf49b..bdc3067 100644
--- a/compiler/src/iree/compiler/Codegen/Common/BubbleUpOrdinalOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/BubbleUpOrdinalOps.cpp
@@ -11,7 +11,6 @@
//
//===---------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -19,6 +18,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_BUBBLEUPORDINALOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Replace the following sequence
@@ -62,8 +64,8 @@
}
};
-struct BubbleUpOrdinalOpsPass
- : public BubbleUpOrdinalOpsBase<BubbleUpOrdinalOpsPass> {
+struct BubbleUpOrdinalOpsPass final
+ : impl::BubbleUpOrdinalOpsPassBase<BubbleUpOrdinalOpsPass> {
void runOnOperation() override;
};
} // namespace
@@ -77,9 +79,4 @@
return signalPassFailure();
}
}
-
-std::unique_ptr<Pass> createBubbleUpOrdinalOpsPass() {
- return std::make_unique<BubbleUpOrdinalOpsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/BufferizeCopyOnlyDispatchesPass.cpp b/compiler/src/iree/compiler/Codegen/Common/BufferizeCopyOnlyDispatchesPass.cpp
index f7dc6b5..9271cdf 100644
--- a/compiler/src/iree/compiler/Codegen/Common/BufferizeCopyOnlyDispatchesPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/BufferizeCopyOnlyDispatchesPass.cpp
@@ -11,7 +11,6 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/PassUtils.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
@@ -30,16 +29,18 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_BUFFERIZECOPYONLYDISPATCHESPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Pass to bufferize early copy-only dispatches. This allows backends
/// to use the `linalg.generic` operation generated for lowering the dispatch.
-struct BufferizeCopyOnlyDispatchesPass
- : public BufferizeCopyOnlyDispatchesBase<BufferizeCopyOnlyDispatchesPass> {
- BufferizeCopyOnlyDispatchesPass() = default;
- BufferizeCopyOnlyDispatchesPass(const BufferizeCopyOnlyDispatchesPass &pass) {
- }
-
+struct BufferizeCopyOnlyDispatchesPass final
+ : impl::BufferizeCopyOnlyDispatchesPassBase<
+ BufferizeCopyOnlyDispatchesPass> {
+ using impl::BufferizeCopyOnlyDispatchesPassBase<
+ BufferizeCopyOnlyDispatchesPass>::BufferizeCopyOnlyDispatchesPassBase;
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, bufferization::BufferizationDialect,
IREE::Flow::FlowDialect, linalg::LinalgDialect,
@@ -109,9 +110,4 @@
}
}
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createBufferizeCopyOnlyDispatchesPass() {
- return std::make_unique<BufferizeCopyOnlyDispatchesPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/CMakeLists.txt b/compiler/src/iree/compiler/Codegen/Common/CMakeLists.txt
index 44c25aa..13c9da2 100644
--- a/compiler/src/iree/compiler/Codegen/Common/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Codegen/Common/CMakeLists.txt
@@ -23,7 +23,6 @@
NAME
PassHeaders
HDRS
- "PassDetail.h"
"Passes.h"
"Passes.h.inc"
DEPS
@@ -124,9 +123,10 @@
"PropagateReshapesByExpansion.cpp"
"ReconcileTranslationInfo.cpp"
"RematerializeParallelOps.cpp"
- "RemoveTrivialLoops.cpp"
+ "RemoveSingleIterationLoop.cpp"
"ReplaceSlowMinMaxOps.cpp"
"SplitFullPartialTransferPass.cpp"
+ "TensorToVectorVectorizePad.cpp"
"TestExecutablePreprocessing.cpp"
"TestPartitionableLoopsInterface.cpp"
"TileAndDistributeToWorkgroupsPass.cpp"
@@ -135,7 +135,6 @@
"TypePropagationPass.cpp"
"UserConfig.cpp"
"VectorizeMemrefCopy.cpp"
- "VectorizePad.cpp"
DEPS
::PassHeaders
::PassesIncGen
diff --git a/compiler/src/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp b/compiler/src/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
index 4881d14..345e86b 100644
--- a/compiler/src/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
@@ -12,7 +12,6 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
@@ -26,11 +25,14 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CLEANUPBUFFERALLOCVIEWPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Runs canonicalization patterns on interface load/store ops.
-struct CleanupBufferAllocViewPass
- : public CleanupBufferAllocViewBase<CleanupBufferAllocViewPass> {
+struct CleanupBufferAllocViewPass final
+ : impl::CleanupBufferAllocViewPassBase<CleanupBufferAllocViewPass> {
void runOnOperation() override {
RewritePatternSet patterns(&getContext());
populateReshapeToInterfaceTensorPatterns(patterns);
@@ -43,10 +45,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createCleanupBufferAllocViewPass() {
- return std::make_unique<CleanupBufferAllocViewPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/CommonDialectRegistration.cpp b/compiler/src/iree/compiler/Codegen/Common/CommonDialectRegistration.cpp
index 6a96ecb..260428a 100644
--- a/compiler/src/iree/compiler/Codegen/Common/CommonDialectRegistration.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/CommonDialectRegistration.cpp
@@ -5,7 +5,6 @@
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
#include "iree-dialects/Dialect/LinalgTransform/StructuredTransformOpsExt.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/TransformExtensions/CommonExtensions.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h"
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
index dda4917..22f08cf 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConcretizePadResultShape.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "llvm/Support/Debug.h"
@@ -24,6 +23,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CONCRETIZEPADRESULTSHAPEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Gets the given `attrOrValue` as an index value by creating constant ops
/// for attributes.
static Value getAsIndexValue(OpFoldResult attrOrValue, OpBuilder &builder,
@@ -126,12 +128,9 @@
};
class ConcretizePadResultShapePass final
- : public ConcretizePadResultShapeBase<ConcretizePadResultShapePass> {
+ : public impl::ConcretizePadResultShapePassBase<
+ ConcretizePadResultShapePass> {
public:
- ConcretizePadResultShapePass() = default;
- ConcretizePadResultShapePass(const ConcretizePadResultShapePass &pass) =
- default;
-
void runOnOperation() override {
MLIRContext *context = &getContext();
auto funcOp = getOperation();
@@ -155,11 +154,6 @@
} // namespace
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createConcretizePadResultShapePass() {
- return std::make_unique<ConcretizePadResultShapePass>();
-}
-
void populateConcretizePadResultShapePatterns(RewritePatternSet &patterns,
ArrayRef<int64_t> numWorkgroups) {
MLIRContext *context = patterns.getContext();
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp
index 3c52c16..572b719 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ArithToF32.cpp
@@ -14,7 +14,6 @@
#include <memory>
#include <utility>
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
@@ -39,6 +38,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CONVERTBF16ARITHTOF32PASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
Value convertRankedFloat(OpBuilder &builder, Type type, ValueRange inputs,
@@ -235,9 +237,10 @@
}
};
-struct ConvertBf16ArithToF32Pass
- : public ConvertBf16ArithToF32Base<ConvertBf16ArithToF32Pass> {
- using ConvertBf16ArithToF32Base::ConvertBf16ArithToF32Base;
+struct ConvertBf16ArithToF32Pass final
+ : impl::ConvertBf16ArithToF32PassBase<ConvertBf16ArithToF32Pass> {
+ using impl::ConvertBf16ArithToF32PassBase<
+ ConvertBf16ArithToF32Pass>::ConvertBf16ArithToF32PassBase;
void runOnOperation() override {
MLIRContext *context = &this->getContext();
RewritePatternSet patterns(context);
@@ -312,9 +315,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<>> createConvertBf16ArithToF32Pass() {
- return std::make_unique<ConvertBf16ArithToF32Pass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ToUInt16Buffers.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ToUInt16Buffers.cpp
index 73d1d03..5f25384 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ToUInt16Buffers.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvertBf16ToUInt16Buffers.cpp
@@ -11,7 +11,6 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
@@ -24,6 +23,7 @@
#include "mlir/Dialect/Arith/Transforms/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/Func/Transforms/FuncConversions.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/MemRef/Transforms/Transforms.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
@@ -34,10 +34,13 @@
#include "mlir/Transforms/DialectConversion.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#define DEBUG_TYPE "iree-spirv-emulate-bf16"
+#define DEBUG_TYPE "iree-codegen-convert-bf16-to-uint16-buffers"
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CONVERTBF16TOUINT16BUFFERSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
class Bf16EmulationConverter : public TypeConverter {
@@ -249,7 +252,7 @@
//===----------------------------------------------------------------------===//
struct ConvertBf16ToUInt16BuffersPass final
- : public ConvertBf16ToUInt16BuffersBase<ConvertBf16ToUInt16BuffersPass> {
+ : impl::ConvertBf16ToUInt16BuffersPassBase<ConvertBf16ToUInt16BuffersPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<vector::VectorDialect>();
}
@@ -312,13 +315,4 @@
};
} // namespace
-
-//===----------------------------------------------------------------------===//
-// Public interface
-//===----------------------------------------------------------------------===//
-
-std::unique_ptr<OperationPass<>> createConvertBf16ToUInt16BuffersPass() {
- return std::make_unique<ConvertBf16ToUInt16BuffersPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
index a5b61ab..69f5510 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvertToDestinationPassingStylePass.cpp
@@ -13,7 +13,6 @@
//===----------------------------------------------------------------------===//
#include "iree/compiler/Codegen/Common/BufferizationAnalysis.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
@@ -48,19 +47,21 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CONVERTTODESTINATIONPASSINGSTYLEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class ConvertToDestinationPassingStylePass
- : public ConvertToDestinationPassingStyleBase<
+class ConvertToDestinationPassingStylePass final
+ : public impl::ConvertToDestinationPassingStylePassBase<
ConvertToDestinationPassingStylePass> {
public:
- ConvertToDestinationPassingStylePass() = default;
- ConvertToDestinationPassingStylePass(bool useWARForCooperativeMatrixCodegen) {
+ using impl::ConvertToDestinationPassingStylePassBase<
+ ConvertToDestinationPassingStylePass>::
+ ConvertToDestinationPassingStylePassBase;
+ explicit ConvertToDestinationPassingStylePass(
+ bool useWARForCooperativeMatrixCodegen) {
this->useWARForCooperativeMatrixCodegen = useWARForCooperativeMatrixCodegen;
}
- ConvertToDestinationPassingStylePass(
- const ConvertToDestinationPassingStylePass &pass) {
- useWARForCooperativeMatrixCodegen = pass.useWARForCooperativeMatrixCodegen;
- }
void getDependentDialects(DialectRegistry ®istry) const override {
registry
diff --git a/compiler/src/iree/compiler/Codegen/Common/ConvolutionToIGEMM.cpp b/compiler/src/iree/compiler/Codegen/Common/ConvolutionToIGEMM.cpp
index 7eb6003..76b03ef 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ConvolutionToIGEMM.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ConvolutionToIGEMM.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h"
@@ -18,12 +17,15 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_CONVOLUTIONTOIGEMMPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
using iree_compiler::IREE::LinalgExt::IREELinalgExtDialect;
-class ConvolutionToIGEMMPass
- : public ConvolutionToIGEMMBase<ConvolutionToIGEMMPass> {
+class ConvolutionToIGEMMPass final
+ : public impl::ConvolutionToIGEMMPassBase<ConvolutionToIGEMMPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<tensor::TensorDialect, IREELinalgExtDialect>();
}
@@ -95,10 +97,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createConvolutionToIGEMMPass() {
- return std::make_unique<ConvolutionToIGEMMPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/DecomposeAffineOpsPass.cpp b/compiler/src/iree/compiler/Codegen/Common/DecomposeAffineOpsPass.cpp
index 60282c9..d83dc39 100644
--- a/compiler/src/iree/compiler/Codegen/Common/DecomposeAffineOpsPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/DecomposeAffineOpsPass.cpp
@@ -4,16 +4,19 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
+#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Affine/Transforms/Transforms.h"
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_DECOMPOSEAFFINEOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-struct DecomposeAffineOpsPass
- : public DecomposeAffineOpsBase<DecomposeAffineOpsPass> {
+struct DecomposeAffineOpsPass final
+ : impl::DecomposeAffineOpsPassBase<DecomposeAffineOpsPass> {
void runOnOperation() override;
};
@@ -27,8 +30,4 @@
(void)decompose(rewriter, op);
});
}
-
-std::unique_ptr<Pass> createDecomposeAffineOpsPass() {
- return std::make_unique<DecomposeAffineOpsPass>();
-}
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/DecomposeConvolutionToLowerDimOps.cpp b/compiler/src/iree/compiler/Codegen/Common/DecomposeConvolutionToLowerDimOps.cpp
index e4a6c8c..70c3383 100644
--- a/compiler/src/iree/compiler/Codegen/Common/DecomposeConvolutionToLowerDimOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/DecomposeConvolutionToLowerDimOps.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "llvm/ADT/STLExtras.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -16,6 +15,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_DECOMPOSECONVOLUTIONTOLOWERDIMOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
static bool foldHDim(linalg::DepthwiseConv2DNhwcHwcOp convOp) {
@@ -120,8 +122,8 @@
loweringConfigAttr.value().getNativeVectorSize());
}
-class DecomposeConvolutionToLowerDimOpsPass
- : public DecomposeConvolutionToLowerDimOpsBase<
+class DecomposeConvolutionToLowerDimOpsPass final
+ : public impl::DecomposeConvolutionToLowerDimOpsPassBase<
DecomposeConvolutionToLowerDimOpsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, linalg::LinalgDialect>();
@@ -169,9 +171,4 @@
};
} // namespace
-
-std::unique_ptr<Pass> createDecomposeConvolutionToLowerDimOpsPass() {
- return std::make_unique<DecomposeConvolutionToLowerDimOpsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/DecomposeLinalgGeneric.cpp b/compiler/src/iree/compiler/Codegen/Common/DecomposeLinalgGeneric.cpp
index 0239333..de3005d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/DecomposeLinalgGeneric.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/DecomposeLinalgGeneric.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
@@ -14,10 +13,13 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_DECOMPOSELINALGGENERICPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class DecomposeLinalgGenericPass
- : public DecomposeLinalgGenericBase<DecomposeLinalgGenericPass> {
+class DecomposeLinalgGenericPass final
+ : public impl::DecomposeLinalgGenericPassBase<DecomposeLinalgGenericPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, linalg::LinalgDialect>();
}
@@ -34,9 +36,4 @@
};
} // namespace
-
-std::unique_ptr<Pass> createDecomposeLinalgGenericPass() {
- return std::make_unique<DecomposeLinalgGenericPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/DecomposePackUnPackOps.cpp b/compiler/src/iree/compiler/Codegen/Common/DecomposePackUnPackOps.cpp
index 80ca58e..e0c56a0 100644
--- a/compiler/src/iree/compiler/Codegen/Common/DecomposePackUnPackOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/DecomposePackUnPackOps.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -26,6 +25,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_DECOMPOSEPACKUNPACKOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// A wrapper pattern that calls linalg::lowerPack on tensor::PackOp. It lowers
@@ -62,9 +64,11 @@
}
};
-struct DecomposePackUnPackOpsPass
- : public DecomposePackUnPackOpsBase<DecomposePackUnPackOpsPass> {
- DecomposePackUnPackOpsPass(bool tileOuterToOne) {
+struct DecomposePackUnPackOpsPass final
+ : impl::DecomposePackUnPackOpsPassBase<DecomposePackUnPackOpsPass> {
+ using impl::DecomposePackUnPackOpsPassBase<
+ DecomposePackUnPackOpsPass>::DecomposePackUnPackOpsPassBase;
+ explicit DecomposePackUnPackOpsPass(bool tileOuterToOne) {
this->tileOuterToOne = tileOuterToOne;
}
void getDependentDialects(DialectRegistry ®istry) const override {
diff --git a/compiler/src/iree/compiler/Codegen/Common/DecomposeSoftmax.cpp b/compiler/src/iree/compiler/Codegen/Common/DecomposeSoftmax.cpp
index 79fb38c..25bc161 100644
--- a/compiler/src/iree/compiler/Codegen/Common/DecomposeSoftmax.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/DecomposeSoftmax.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
@@ -15,6 +14,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_DECOMPOSESOFTMAXPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Given an N-dimensional tensor x, this op converts
/// softmax(x) to the following sequence of operations:
@@ -96,9 +98,11 @@
return success();
}
-struct DecomposeSoftmaxPass : DecomposeSoftmaxBase<DecomposeSoftmaxPass> {
- DecomposeSoftmaxPass() = default;
- DecomposeSoftmaxPass(bool useFusion) { this->useFusion = useFusion; }
+struct DecomposeSoftmaxPass
+ : impl::DecomposeSoftmaxPassBase<DecomposeSoftmaxPass> {
+ using impl::DecomposeSoftmaxPassBase<
+ DecomposeSoftmaxPass>::DecomposeSoftmaxPassBase;
+ explicit DecomposeSoftmaxPass(bool useFusion) { this->useFusion = useFusion; }
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect>();
diff --git a/compiler/src/iree/compiler/Codegen/Common/EmulateNarrowType.cpp b/compiler/src/iree/compiler/Codegen/Common/EmulateNarrowType.cpp
index b0b3545..1824fb0 100644
--- a/compiler/src/iree/compiler/Codegen/Common/EmulateNarrowType.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/EmulateNarrowType.cpp
@@ -3,7 +3,7 @@
// Licensed under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
+
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/Transforms.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
@@ -25,6 +25,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_EMULATENARROWTYPEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
//===----------------------------------------------------------------------===//
@@ -96,8 +99,8 @@
// Pass Definition
//===----------------------------------------------------------------------===//
-struct EmulateNarrowTypePass
- : public EmulateNarrowTypeBase<EmulateNarrowTypePass> {
+struct EmulateNarrowTypePass final
+ : impl::EmulateNarrowTypePassBase<EmulateNarrowTypePass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<arith::ArithDialect, func::FuncDialect,
memref::MemRefDialect, vector::VectorDialect,
@@ -160,13 +163,4 @@
}
};
} // namespace
-
-//===----------------------------------------------------------------------===//
-// Public interface
-//===----------------------------------------------------------------------===//
-
-std::unique_ptr<OperationPass<>> createEmulateNarrowTypePass() {
- return std::make_unique<EmulateNarrowTypePass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/EraseDeadAllocAndStores.cpp b/compiler/src/iree/compiler/Codegen/Common/EraseDeadAllocAndStores.cpp
index 4a22610..cdab370 100644
--- a/compiler/src/iree/compiler/Codegen/Common/EraseDeadAllocAndStores.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/EraseDeadAllocAndStores.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -14,12 +13,17 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_ERASEDEADALLOCANDSTORESPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class EraseDeadAllocAndStoresPass
- : public EraseDeadAllocAndStoresBase<EraseDeadAllocAndStoresPass> {
+class EraseDeadAllocAndStoresPass final
+ : public impl::EraseDeadAllocAndStoresPassBase<
+ EraseDeadAllocAndStoresPass> {
public:
- using EraseDeadAllocAndStoresBase::EraseDeadAllocAndStoresBase;
+ using impl::EraseDeadAllocAndStoresPassBase<
+ EraseDeadAllocAndStoresPass>::EraseDeadAllocAndStoresPassBase;
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<scf::SCFDialect, vector::VectorDialect>();
@@ -34,10 +38,4 @@
}
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createEraseDeadAllocAndStoresPass() {
- return std::make_unique<EraseDeadAllocAndStoresPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/EraseHALDescriptorTypeFromMemRef.cpp b/compiler/src/iree/compiler/Codegen/Common/EraseHALDescriptorTypeFromMemRef.cpp
index 38ecf16..2f38171 100644
--- a/compiler/src/iree/compiler/Codegen/Common/EraseHALDescriptorTypeFromMemRef.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/EraseHALDescriptorTypeFromMemRef.cpp
@@ -13,7 +13,6 @@
#include <memory>
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALTypes.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
@@ -25,6 +24,10 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_ERASEHALDESCRIPTORTYPEFROMMEMREFPASS
+#define GEN_PASS_DEF_CONVERTHALDESCRIPTORTYPETOGPUADDRESSSPACEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Returns true if the given `type` is considered as legal.
@@ -37,7 +40,7 @@
}
struct EraseHALDescriptorTypeFromMemRefPass final
- : public EraseHALDescriptorTypeFromMemRefBase<
+ : impl::EraseHALDescriptorTypeFromMemRefPassBase<
EraseHALDescriptorTypeFromMemRefPass> {
void runOnOperation() override {
AttrTypeReplacer replacer;
@@ -65,7 +68,7 @@
};
struct ConvertHALDescriptorTypeToGPUAddressSpacePass final
- : public ConvertHALDescriptorTypeToGPUAddressSpaceBase<
+ : impl::ConvertHALDescriptorTypeToGPUAddressSpacePassBase<
ConvertHALDescriptorTypeToGPUAddressSpacePass> {
void runOnOperation() override {
AttrTypeReplacer replacer;
@@ -96,13 +99,4 @@
};
} // namespace
-
-std::unique_ptr<Pass> createEraseHALDescriptorTypeFromMemRefPass() {
- return std::make_unique<EraseHALDescriptorTypeFromMemRefPass>();
-}
-
-std::unique_ptr<Pass> createConvertHALDescriptorTypeToGPUAddressSpacePass() {
- return std::make_unique<ConvertHALDescriptorTypeToGPUAddressSpacePass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ExtractAddressComputation.cpp b/compiler/src/iree/compiler/Codegen/Common/ExtractAddressComputation.cpp
index 10ca17d..7c34f03 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ExtractAddressComputation.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ExtractAddressComputation.cpp
@@ -6,7 +6,6 @@
#include "iree/compiler/Codegen/Common/ExtractAddressComputation.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
@@ -14,12 +13,13 @@
#include "mlir/Dialect/Utils/StaticValueUtils.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#define DEBUG_TYPE "extract-address-computation"
-
-using namespace mlir;
+#define DEBUG_TYPE "iree-codegen-extract-address-computation"
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_EXTRACTADDRESSCOMPUTATIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
//===----------------------------------------------------------------------===//
// Helper functions for the `load base[off0...]`
// => `load (subview base[off0...])[0...]` pattern.
@@ -96,8 +96,8 @@
//===----------------------------------------------------------------------===//
namespace {
-struct ExtractAddressComputationPass
- : public ExtractAddressComputationBase<ExtractAddressComputationPass> {
+struct ExtractAddressComputationPass final
+ : impl::ExtractAddressComputationPassBase<ExtractAddressComputationPass> {
void runOnOperation() override;
};
} // namespace
@@ -110,8 +110,4 @@
return signalPassFailure();
}
}
-
-std::unique_ptr<Pass> createExtractAddressComputationPass() {
- return std::make_unique<ExtractAddressComputationPass>();
-}
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp b/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
index 6b97aa5..2f7a39e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
@@ -33,7 +33,6 @@
#include <memory>
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/UKernelOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
@@ -62,6 +61,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_FLATTENMEMREFSUBSPANPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
//===----------------------------------------------------------------------===//
@@ -740,11 +742,8 @@
// Pass
//===----------------------------------------------------------------------===//
-struct FlattenMemRefSubspanPass
- : public FlattenMemRefSubspanBase<FlattenMemRefSubspanPass> {
- FlattenMemRefSubspanPass() {}
- FlattenMemRefSubspanPass(const FlattenMemRefSubspanPass &pass) {}
-
+struct FlattenMemRefSubspanPass final
+ : impl::FlattenMemRefSubspanPassBase<FlattenMemRefSubspanPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, memref::MemRefDialect>();
}
@@ -903,9 +902,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<ModuleOp>> createFlattenMemRefSubspanPass() {
- return std::make_unique<FlattenMemRefSubspanPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/FoldAffineMinInDistributedLoops.cpp b/compiler/src/iree/compiler/Codegen/Common/FoldAffineMinInDistributedLoops.cpp
index dd5d986..9cfaec8 100644
--- a/compiler/src/iree/compiler/Codegen/Common/FoldAffineMinInDistributedLoops.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/FoldAffineMinInDistributedLoops.cpp
@@ -14,7 +14,6 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
@@ -35,6 +34,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_FOLDAFFINEMININDISTRIBUTEDLOOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
#ifndef NDEBUG
inline raw_ostream &operator<<(raw_ostream &os,
const LoopTilingAndDistributionInfo &info) {
@@ -154,7 +156,7 @@
};
struct FoldAffineMinInDistributedLoopsPass final
- : public FoldAffineMinInDistributedLoopsBase<
+ : impl::FoldAffineMinInDistributedLoopsPassBase<
FoldAffineMinInDistributedLoopsPass> {
void runOnOperation() override {
RewritePatternSet patterns(&getContext());
@@ -181,10 +183,4 @@
numWorkgroups);
}
}
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createFoldAffineMinInDistributedLoopsPass() {
- return std::make_unique<FoldAffineMinInDistributedLoopsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp b/compiler/src/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
index c4e516f..9e59996 100644
--- a/compiler/src/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
@@ -3,7 +3,7 @@
// Licensed under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
+
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Bufferization/IR/Bufferization.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -15,6 +15,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_FOLDTENSOREXTRACTOPPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
#include "iree/compiler/Codegen/Common/FoldTensorExtractOp.cpp.inc"
} // namespace
@@ -48,8 +51,8 @@
/// In theory this could live upstream, but given that there is disagreement
/// about the validity of `tensor_to_memref` usage/canonicalizations, keeping
/// this pattern here.
-class FoldTensorExtractOpPass
- : public FoldTensorExtractOpBase<FoldTensorExtractOpPass> {
+class FoldTensorExtractOpPass final
+ : public impl::FoldTensorExtractOpPassBase<FoldTensorExtractOpPass> {
void runOnOperation() override;
};
} // namespace
@@ -60,9 +63,4 @@
if (failed(applyPatternsAndFoldGreedily(getOperation(), std::move(patterns))))
signalPassFailure();
}
-
-std::unique_ptr<OperationPass<>> createFoldTensorExtractOpPass() {
- return std::make_unique<FoldTensorExtractOpPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
index 3bd452d..a452a55 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "llvm/Support/MathExtras.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
@@ -17,6 +16,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_FOROPCANONICALIZATIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Pattern to combine instructions across ForOp boundary. It is common when
@@ -287,8 +289,8 @@
}
};
-struct ForOpCanonicalizationPass
- : public ForOpCanonicalizationBase<ForOpCanonicalizationPass> {
+struct ForOpCanonicalizationPass final
+ : impl::ForOpCanonicalizationPassBase<ForOpCanonicalizationPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<scf::SCFDialect, vector::VectorDialect>();
}
@@ -312,10 +314,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createForOpCanonicalizationPass() {
- return std::make_unique<ForOpCanonicalizationPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/FuseTensorPadWithConsumer.cpp b/compiler/src/iree/compiler/Codegen/Common/FuseTensorPadWithConsumer.cpp
index dd44b08..d14b515 100644
--- a/compiler/src/iree/compiler/Codegen/Common/FuseTensorPadWithConsumer.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/FuseTensorPadWithConsumer.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
#include "mlir/Pass/Pass.h"
@@ -12,10 +11,13 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_FUSETENSORPADWITHCONSUMERPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct FuseTensorPadWithConsumerPass final
- : public FuseTensorPadWithConsumerBase<FuseTensorPadWithConsumerPass> {
+ : impl::FuseTensorPadWithConsumerPassBase<FuseTensorPadWithConsumerPass> {
void runOnOperation() override {
MLIRContext *context = &getContext();
auto funcOp = getOperation();
@@ -30,10 +32,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createFuseTensorPadWithConsumerPass() {
- return std::make_unique<FuseTensorPadWithConsumerPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/GenericVectorization.cpp b/compiler/src/iree/compiler/Codegen/Common/GenericVectorization.cpp
index 6c24885..ebe4ef2 100644
--- a/compiler/src/iree/compiler/Codegen/Common/GenericVectorization.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/GenericVectorization.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/TileSizeSelection.h"
#include "mlir/Dialect/Affine/LoopUtils.h"
@@ -24,6 +23,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_GENERICVECTORIZATIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct VectorizationTileSizes {
@@ -310,20 +312,11 @@
return success(maxFlatVecSize < maxVectorSize);
}
-class GenericVectorizationPass
- : public GenericVectorizationBase<GenericVectorizationPass> {
+class GenericVectorizationPass final
+ : public impl::GenericVectorizationPassBase<GenericVectorizationPass> {
public:
- using GenericVectorizationBase::GenericVectorizationBase;
- GenericVectorizationPass(const GenericVectorizationPassOptions &options) {
- this->enableVectorMasking.setValue(options.enableVectorMasking);
- this->useConfiguredVectorSizes.setValue(options.useConfiguredVectorSizes);
- this->vectorizePadding.setValue(options.vectorizePadding);
- this->vectorizeGatherAccesses.setValue(options.vectorizeGatherAccesses);
- this->enableCleanup.setValue(options.enableCleanup);
- this->generateContract.setValue(options.generateContract);
- this->foldCastIntoContract.setValue(options.foldCastIntoContract);
- this->maxVectorSize.setValue(options.maxVectorSize);
- }
+ using impl::GenericVectorizationPassBase<
+ GenericVectorizationPass>::GenericVectorizationPassBase;
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<tensor::TensorDialect, linalg::LinalgDialect,
@@ -438,14 +431,4 @@
}
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createGenericVectorizationPass() {
- return std::make_unique<GenericVectorizationPass>();
-}
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createGenericVectorizationPass(const GenericVectorizationPassOptions &options) {
- return std::make_unique<GenericVectorizationPass>(options);
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/HoistStaticallyBoundAllocations.cpp b/compiler/src/iree/compiler/Codegen/Common/HoistStaticallyBoundAllocations.cpp
index 1d3b43c..1208ea6 100644
--- a/compiler/src/iree/compiler/Codegen/Common/HoistStaticallyBoundAllocations.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/HoistStaticallyBoundAllocations.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/Transforms.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "llvm/Support/Debug.h"
+#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Linalg/Utils/Utils.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/IR/PatternMatch.h"
@@ -16,10 +16,17 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_HOISTSTATICALLYBOUNDALLOCATIONSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct HoistStaticallyBoundAllocationsPass
- : HoistStaticallyBoundAllocationsBase<HoistStaticallyBoundAllocationsPass> {
+ : impl::HoistStaticallyBoundAllocationsPassBase<
+ HoistStaticallyBoundAllocationsPass> {
+ using impl::HoistStaticallyBoundAllocationsPassBase<
+ HoistStaticallyBoundAllocationsPass>::
+ HoistStaticallyBoundAllocationsPassBase;
void runOnOperation() override;
};
@@ -39,9 +46,4 @@
vscaleRange);
}
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createHoistStaticallyBoundAllocationsPass() {
- return std::make_unique<HoistStaticallyBoundAllocationsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp b/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
index af9fb4c..7a52de0 100644
--- a/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/HoistUnrolledVectorExtractInsertSlice.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/SCF/IR/SCF.h"
#include "mlir/Dialect/SCF/Transforms/Patterns.h"
@@ -22,6 +21,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_HOISTUNROLLEDVECTOREXTRACTINSERTSLICEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Returns all the users of `srcTensor` if they are artifacts from vector
/// unrolling. It is true only if
/// 1. All the users are vector.extract_strided_slice ops.
@@ -199,13 +201,10 @@
}
namespace {
-class HoistUnrolledVectorExtractInsertSlicePass
- : public HoistUnrolledVectorExtractInsertSliceBase<
+class HoistUnrolledVectorExtractInsertSlicePass final
+ : public impl::HoistUnrolledVectorExtractInsertSlicePassBase<
HoistUnrolledVectorExtractInsertSlicePass> {
public:
- using HoistUnrolledVectorExtractInsertSliceBase::
- HoistUnrolledVectorExtractInsertSliceBase;
-
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<scf::SCFDialect, vector::VectorDialect>();
}
@@ -231,10 +230,4 @@
}
}
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createHoistUnrolledVectorExtractInsertSlicePass() {
- return std::make_unique<HoistUnrolledVectorExtractInsertSlicePass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/IREEComprehensiveBufferizePass.cpp b/compiler/src/iree/compiler/Codegen/Common/IREEComprehensiveBufferizePass.cpp
index 9b46b84..6893cc9 100644
--- a/compiler/src/iree/compiler/Codegen/Common/IREEComprehensiveBufferizePass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/IREEComprehensiveBufferizePass.cpp
@@ -10,7 +10,6 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Interfaces/BufferizationInterfaces.h"
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
@@ -45,9 +44,36 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_ELIMINATEEMPTYTENSORSPASS
+#define GEN_PASS_DEF_IREECOMPREHENSIVEBUFFERIZEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
+// Default allocation functions.
+static FailureOr<Value> defaultAllocationFn(OpBuilder &builder, Location loc,
+ MemRefType allocationType,
+ ValueRange dynamicSizes,
+ unsigned int alignment) {
+ MemRefType type = allocationType;
+ if (auto storage = type.getMemorySpace()) {
+ // We cannot allocate to generate a resultant MemRef type with descriptor
+ // type memory space; that's runtime allocations. So erase and fallback to
+ // the default 0 memory space. It is fine given this is just the default
+ // allocator; backends are expected to control by themselves.
+ if (llvm::isa<IREE::HAL::DescriptorTypeAttr>(storage))
+ type = MemRefType::get(type.getShape(), type.getElementType(),
+ type.getLayout());
+ }
+ return builder.create<memref::AllocOp>(loc, type, dynamicSizes).getResult();
+}
+static LogicalResult defaultMemCpyFn(OpBuilder &builder, Location loc,
+ Value from, Value to) {
+ Operation *copyOp = createLinalgCopyOp(builder, loc, from, to);
+ return success(static_cast<bool>(copyOp));
+}
+
namespace {
-class EliminateEmptyTensorsPass
- : public EliminateEmptyTensorsBase<EliminateEmptyTensorsPass> {
+class EliminateEmptyTensorsPass final
+ : public impl::EliminateEmptyTensorsPassBase<EliminateEmptyTensorsPass> {
public:
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::Flow::FlowDialect, tensor::TensorDialect>();
@@ -57,13 +83,15 @@
};
/// Pass to convert from tensor based ops to memref based ops.
-class IREEComprehensiveBufferizePass
- : public IREEComprehensiveBufferizeBase<IREEComprehensiveBufferizePass> {
+class IREEComprehensiveBufferizePass final
+ : public impl::IREEComprehensiveBufferizePassBase<
+ IREEComprehensiveBufferizePass> {
public:
+ using impl::IREEComprehensiveBufferizePassBase<
+ IREEComprehensiveBufferizePass>::IREEComprehensiveBufferizePassBase;
explicit IREEComprehensiveBufferizePass(
- std::optional<BufferizationOptions::AllocationFn> allocationFn =
- std::nullopt,
- std::optional<BufferizationOptions::MemCpyFn> memCpyFn = std::nullopt)
+ BufferizationOptions::AllocationFn allocationFn,
+ BufferizationOptions::MemCpyFn memCpyFn)
: allocationFn(allocationFn), memCpyFn(memCpyFn) {}
void getDependentDialects(DialectRegistry ®istry) const override {
@@ -88,34 +116,11 @@
void runOnOperation() override;
private:
- const std::optional<BufferizationOptions::AllocationFn> allocationFn;
- const std::optional<BufferizationOptions::MemCpyFn> memCpyFn;
+ const BufferizationOptions::AllocationFn allocationFn = defaultAllocationFn;
+ const BufferizationOptions::MemCpyFn memCpyFn = defaultMemCpyFn;
};
} // namespace
-// Default allocation functions.
-static FailureOr<Value> defaultAllocationFn(OpBuilder &builder, Location loc,
- MemRefType allocationType,
- ValueRange dynamicSizes,
- unsigned int alignment) {
- MemRefType type = allocationType;
- if (auto storage = type.getMemorySpace()) {
- // We cannot allocate to generate a resultant MemRef type with descriptor
- // type memory space; that's runtime allocations. So erase and fallback to
- // the default 0 memory space. It is fine given this is just the default
- // allocator; backends are expected to control by themselves.
- if (llvm::isa<IREE::HAL::DescriptorTypeAttr>(storage))
- type = MemRefType::get(type.getShape(), type.getElementType(),
- type.getLayout());
- }
- return builder.create<memref::AllocOp>(loc, type, dynamicSizes).getResult();
-}
-static LogicalResult defaultMemCpyFn(OpBuilder &builder, Location loc,
- Value from, Value to) {
- Operation *copyOp = createLinalgCopyOp(builder, loc, from, to);
- return success(static_cast<bool>(copyOp));
-}
-
static IREEOneShotBufferizationOptions getBufferizationOptions() {
IREEOneShotBufferizationOptions options;
@@ -227,11 +232,6 @@
}
std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createEliminateEmptyTensorsPass() {
- return std::make_unique<EliminateEmptyTensorsPass>();
-}
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
createIREEComprehensiveBufferizePass(
std::optional<BufferizationOptions::AllocationFn> allocationFn,
std::optional<BufferizationOptions::MemCpyFn> memCpyFn) {
@@ -239,8 +239,8 @@
allocationFn = defaultAllocationFn;
if (!memCpyFn)
memCpyFn = defaultMemCpyFn;
- return std::make_unique<IREEComprehensiveBufferizePass>(allocationFn,
- memCpyFn);
+ return std::make_unique<IREEComprehensiveBufferizePass>(allocationFn.value(),
+ memCpyFn.value());
}
void addIREEPostBufferizationPasses(OpPassManager &funcPassManager) {
diff --git a/compiler/src/iree/compiler/Codegen/Common/IREEExpandStridedMetadata.cpp b/compiler/src/iree/compiler/Codegen/Common/IREEExpandStridedMetadata.cpp
index d4b66fb..7ae89bd 100644
--- a/compiler/src/iree/compiler/Codegen/Common/IREEExpandStridedMetadata.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/IREEExpandStridedMetadata.cpp
@@ -4,9 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#define DEBUG_TYPE "iree-codegen-expand-strided-metadata"
-
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenOps.h"
@@ -18,8 +15,13 @@
#include "mlir/Dialect/MemRef/Transforms/Transforms.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
+#define DEBUG_TYPE "iree-codegen-expand-strided-metadata"
+
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_IREEEXPANDSTRIDEDMETADATAPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Helper struct to return the offset, sizes and strides
/// of a `source` of a `memref.extract_strided_metadata` op.
@@ -242,8 +244,11 @@
}
};
-struct IREEExpandStridedMetadataPass
- : public IREEExpandStridedMetadataBase<IREEExpandStridedMetadataPass> {
+struct IREEExpandStridedMetadataPass final
+ : impl::IREEExpandStridedMetadataPassBase<IREEExpandStridedMetadataPass> {
+ using impl::IREEExpandStridedMetadataPassBase<
+ IREEExpandStridedMetadataPass>::IREEExpandStridedMetadataPassBase;
+
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, arith::ArithDialect,
IREE::Codegen::IREECodegenDialect, memref::MemRefDialect>();
@@ -289,9 +294,4 @@
}
}
}
-
-std::unique_ptr<Pass> createIREEExpandStridedMetadataPass() {
- return std::make_unique<IREEExpandStridedMetadataPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/InstrumentMemoryAccesses.cpp b/compiler/src/iree/compiler/Codegen/Common/InstrumentMemoryAccesses.cpp
index cd5a25f..d399182 100644
--- a/compiler/src/iree/compiler/Codegen/Common/InstrumentMemoryAccesses.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/InstrumentMemoryAccesses.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
@@ -16,10 +15,13 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_INSTRUMENTMEMORYACCESSESPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct InstrumentMemoryAccessesPass
- : InstrumentMemoryAccessesBase<InstrumentMemoryAccessesPass> {
+ : impl::InstrumentMemoryAccessesPassBase<InstrumentMemoryAccessesPass> {
void runOnOperation() override {
// Lookup the root instrumentation op. If not present it means the dispatch
// is not instrumented and we can skip it.
@@ -83,10 +85,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createInstrumentMemoryAccessesPass() {
- return std::make_unique<InstrumentMemoryAccessesPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp b/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
index 592ddb8..dd38ef2 100644
--- a/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/LowerExecutableUsingTransformDialect.cpp
@@ -5,14 +5,16 @@
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
#include "iree-dialects/Dialect/LinalgTransform/Passes.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_LOWEREXECUTABLEUSINGTRANSFORMDIALECTPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class LowerExecutableUsingTransformDialectPass
- : public LowerExecutableUsingTransformDialectBase<
+class LowerExecutableUsingTransformDialectPass final
+ : public impl::LowerExecutableUsingTransformDialectPassBase<
LowerExecutableUsingTransformDialectPass> {
public:
void runOnOperation() override;
@@ -62,10 +64,4 @@
return signalPassFailure();
}
}
-
-std::unique_ptr<OperationPass<ModuleOp>>
-createLowerExecutableUsingTransformDialectPass() {
- return std::make_unique<LowerExecutableUsingTransformDialectPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/LowerUKernelsToCalls.cpp b/compiler/src/iree/compiler/Codegen/Common/LowerUKernelsToCalls.cpp
index 76aca28..ca54052 100644
--- a/compiler/src/iree/compiler/Codegen/Common/LowerUKernelsToCalls.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/LowerUKernelsToCalls.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Interfaces/UKernelOpInterface.h"
#include "llvm/ADT/MapVector.h"
@@ -14,9 +13,12 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_LOWERUKERNELOPSTOCALLSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct LowerUKernelOpsToCallsPass
- : LowerUKernelOpsToCallsBase<LowerUKernelOpsToCallsPass> {
+ : impl::LowerUKernelOpsToCallsPassBase<LowerUKernelOpsToCallsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<memref::MemRefDialect, func::FuncDialect>();
}
@@ -53,9 +55,4 @@
rewriter.replaceOp(r.first, r.second->getResults());
}
}
-
-std::unique_ptr<OperationPass<ModuleOp>> createLowerUKernelOpsToCallsPass() {
- return std::make_unique<LowerUKernelOpsToCallsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoNop.cpp b/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoNop.cpp
index 08d1bb1..c3df306 100644
--- a/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoNop.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoNop.cpp
@@ -5,7 +5,6 @@
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
#include "iree/compiler/Codegen/Common/EncodingUtils.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/PassUtils.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/Encoding/IR/EncodingOps.h"
@@ -18,11 +17,14 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_MATERIALIZEENCODINGINTONOPPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
using namespace IREE::Encoding;
namespace {
-struct MaterializeEncodingIntoNopPass
- : public MaterializeEncodingIntoNopBase<MaterializeEncodingIntoNopPass> {
+struct MaterializeEncodingIntoNopPass final
+ : impl::MaterializeEncodingIntoNopPassBase<MaterializeEncodingIntoNopPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect, tensor::TensorDialect>();
}
@@ -70,11 +72,6 @@
};
} // namespace
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createMaterializeEncodingIntoNopPass() {
- return std::make_unique<MaterializeEncodingIntoNopPass>();
-}
-
void addEncodingToNopPasses(FunctionLikeNest &passManager) {
passManager.addPass(createMaterializeEncodingIntoNopPass)
.addPass(createBufferizeCopyOnlyDispatchesPass)
diff --git a/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoPackUnPack.cpp b/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoPackUnPack.cpp
index 99fb0d3..cf81495 100644
--- a/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoPackUnPack.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/MaterializeEncodingIntoPackUnPack.cpp
@@ -9,7 +9,6 @@
//===---------------------------------------------------------------------===//
#include "iree/compiler/Codegen/Common/EncodingUtils.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Encoding/IR/EncodingOps.h"
diff --git a/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp b/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
index e75b8f1..c3cc131 100644
--- a/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/MaterializeUserConfigs.cpp
@@ -5,7 +5,7 @@
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
#include <iterator>
-#include "iree/compiler/Codegen/Common/PassDetail.h"
+
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/UserConfig.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
@@ -35,6 +35,9 @@
"this will default to `__kernel_config`."),
llvm::cl::init(""));
+#define GEN_PASS_DEF_MATERIALIZEUSERCONFIGSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
static const char kTranslationInfoAttrName[] = "translation_info";
@@ -66,8 +69,8 @@
return StrategyRunResult::Success;
}
-struct MaterializeUserConfigsPass
- : public MaterializeUserConfigsBase<MaterializeUserConfigsPass> {
+struct MaterializeUserConfigsPass final
+ : impl::MaterializeUserConfigsPassBase<MaterializeUserConfigsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registerTransformDialectTranslationDependentDialects(registry);
}
@@ -204,9 +207,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<ModuleOp>> createMaterializeUserConfigsPass() {
- return std::make_unique<MaterializeUserConfigsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/MemrefCopyToLinalg.cpp b/compiler/src/iree/compiler/Codegen/Common/MemrefCopyToLinalg.cpp
index b06f1af..4e6c07b 100644
--- a/compiler/src/iree/compiler/Codegen/Common/MemrefCopyToLinalg.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/MemrefCopyToLinalg.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
@@ -14,6 +13,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_MEMREFCOPYTOLINALGPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
struct MemrefCopyOpToLinalg : public OpRewritePattern<memref::CopyOp> {
using OpRewritePattern<memref::CopyOp>::OpRewritePattern;
@@ -30,8 +32,8 @@
}
};
-struct MemrefCopyToLinalgPass
- : public MemrefCopyToLinalgPassBase<MemrefCopyToLinalgPass> {
+struct MemrefCopyToLinalgPass final
+ : impl::MemrefCopyToLinalgPassBase<MemrefCopyToLinalgPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect>();
}
@@ -48,10 +50,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createMemrefCopyToLinalgPass() {
- return std::make_unique<MemrefCopyToLinalgPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/NormalizeLoopBounds.cpp b/compiler/src/iree/compiler/Codegen/Common/NormalizeLoopBounds.cpp
index 61f02e6..af6d561 100644
--- a/compiler/src/iree/compiler/Codegen/Common/NormalizeLoopBounds.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/NormalizeLoopBounds.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arith/Utils/Utils.h"
@@ -18,6 +17,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_NORMALIZELOOPBOUNDSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
static OpFoldResult emitNormalizedUpperBound(RewriterBase &rewriter,
Location loc, OpFoldResult lb,
OpFoldResult ub,
@@ -165,10 +167,11 @@
}
namespace {
-struct NormalizeLoopBoundsPass
- : public NormalizeLoopBoundsPassBase<NormalizeLoopBoundsPass> {
- NormalizeLoopBoundsPass(bool nFor, bool nForall)
- : normalizeFor(nFor), normalizeForall(nForall) {}
+struct NormalizeLoopBoundsPass final
+ : impl::NormalizeLoopBoundsPassBase<NormalizeLoopBoundsPass> {
+ using impl::NormalizeLoopBoundsPassBase<
+ NormalizeLoopBoundsPass>::NormalizeLoopBoundsPassBase;
+
void runOnOperation() override {
Operation *op = getOperation();
IRRewriter rewriter(op);
@@ -183,17 +186,6 @@
});
}
}
-
-private:
- bool normalizeFor;
- bool normalizeForall;
};
} // namespace
-
-std::unique_ptr<Pass> createNormalizeLoopBoundsPass(bool normalizeFor,
- bool normalizeForall) {
- return std::make_unique<NormalizeLoopBoundsPass>(normalizeFor,
- normalizeForall);
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/OptimizeTensorInsertExtractSlices.cpp b/compiler/src/iree/compiler/Codegen/Common/OptimizeTensorInsertExtractSlices.cpp
index 2b608be..82e074e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/OptimizeTensorInsertExtractSlices.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/OptimizeTensorInsertExtractSlices.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
#include "mlir/Dialect/Linalg/Transforms/Hoisting.h"
@@ -22,15 +21,15 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_OPTIMIZETENSORINSERTEXTRACTSLICESPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class OptimizeTensorInsertExtractSlicesPass
- : public OptimizeTensorInsertExtractSlicesBase<
+class OptimizeTensorInsertExtractSlicesPass final
+ : public impl::OptimizeTensorInsertExtractSlicesPassBase<
OptimizeTensorInsertExtractSlicesPass> {
public:
- using OptimizeTensorInsertExtractSlicesBase::
- OptimizeTensorInsertExtractSlicesBase;
-
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<scf::SCFDialect, vector::VectorDialect>();
}
@@ -233,10 +232,4 @@
}
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createOptimizeTensorInsertExtractSlicesPass() {
- return std::make_unique<OptimizeTensorInsertExtractSlicesPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp b/compiler/src/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
index c35d70d..28db15d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Dialect/Affine/LoopUtils.h"
@@ -25,6 +24,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_OPTIMIZEVECTORTRANSFERPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
// Pattern to canonialize tranpose where only one dimension is not unit
@@ -59,9 +61,11 @@
[&](LoopLikeOpInterface loopLike) { moveLoopInvariantCode(loopLike); });
}
-struct OptimizeVectorTransferPass
- : public OptimizeVectorTransferBase<OptimizeVectorTransferPass> {
- OptimizeVectorTransferPass(bool flatten) : flatten(flatten) {}
+struct OptimizeVectorTransferPass final
+ : impl::OptimizeVectorTransferPassBase<OptimizeVectorTransferPass> {
+ using impl::OptimizeVectorTransferPassBase<
+ OptimizeVectorTransferPass>::OptimizeVectorTransferPassBase;
+
void runOnOperation() override {
auto funcOp = getOperation();
LDBG("before optimize vector transfer\n" << funcOp);
@@ -119,29 +123,7 @@
memref::eraseDeadAllocAndStores(rewriter, funcOp);
LDBG("after erasing unused allocs and stores\n" << funcOp);
}
-
- LogicalResult initializeOptions(
- StringRef options,
- function_ref<LogicalResult(const Twine &)> errorHandler) override {
- if (failed(Pass::initializeOptions(options, errorHandler))) {
- return failure();
- }
- // `flatten` may have been set to `true` in the constructor already.
- // The |= is so we preserve that rather than overwrite it with the default
- // value `false` of `optionFlatten`.
- flatten |= optionFlatten;
- return success();
- }
-
-private:
- bool flatten;
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createOptimizeVectorTransferPass(bool flatten) {
- return std::make_unique<OptimizeVectorTransferPass>(flatten);
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/PadDynamicAlloc.cpp b/compiler/src/iree/compiler/Codegen/Common/PadDynamicAlloc.cpp
index 17b286e..e5f10d3 100644
--- a/compiler/src/iree/compiler/Codegen/Common/PadDynamicAlloc.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/PadDynamicAlloc.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -14,6 +13,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_PADDYNAMICALLOCPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// If a value is defined by `%dim = affine_max(0, %src)` kind of op return
/// `%src` otherwise return `%dim`.
/// This is useful to handle common pattern generated by bufferization to
@@ -76,7 +78,8 @@
namespace {
-struct PadDynamicAllocPass : public PadDynamicAllocBase<PadDynamicAllocPass> {
+struct PadDynamicAllocPass final
+ : impl::PadDynamicAllocPassBase<PadDynamicAllocPass> {
void runOnOperation() override {
auto funcOp = getOperation();
MLIRContext *context = &getContext();
@@ -91,10 +94,4 @@
}
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createPadDynamicAlloc() {
- return std::make_unique<PadDynamicAllocPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/PassDetail.h b/compiler/src/iree/compiler/Codegen/Common/PassDetail.h
deleted file mode 100644
index 09919ae..0000000
--- a/compiler/src/iree/compiler/Codegen/Common/PassDetail.h
+++ /dev/null
@@ -1,25 +0,0 @@
-// Copyright 2023 The IREE Authors
-//
-// Licensed under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-
-#ifndef IREE_COMPILER_CODEGEN_COMMON_PASS_DETAIL_H_
-#define IREE_COMPILER_CODEGEN_COMMON_PASS_DETAIL_H_
-
-#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
-#include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
-#include "mlir/Dialect/MemRef/IR/MemRef.h"
-#include "mlir/Dialect/Transform/IR/TransformOps.h"
-#include "mlir/Interfaces/FunctionInterfaces.h"
-#include "mlir/Pass/Pass.h"
-
-namespace mlir::iree_compiler {
-
-#define GEN_PASS_CLASSES
-#include "iree/compiler/Codegen/Common/Passes.h.inc"
-
-} // namespace mlir::iree_compiler
-
-#endif // IREE_COMPILER_CODEGEN_COMMON_PASS_DETAIL_H_
diff --git a/compiler/src/iree/compiler/Codegen/Common/Passes.h b/compiler/src/iree/compiler/Codegen/Common/Passes.h
index ade1550..b6db41b 100644
--- a/compiler/src/iree/compiler/Codegen/Common/Passes.h
+++ b/compiler/src/iree/compiler/Codegen/Common/Passes.h
@@ -19,6 +19,7 @@
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/Transform/IR/TransformDialect.h"
#include "mlir/Pass/Pass.h"
namespace mlir::iree_compiler {
@@ -47,152 +48,21 @@
void addConstantBufferizePasses(OpPassManager &funcPassManager);
-std::unique_ptr<OperationPass<LLVM::LLVMFuncOp>> createAddFastMathFlagsPass();
+/// Populate Encoding to Nop pass and canonicalizer pass to the pipeline
+void addEncodingToNopPasses(FunctionLikeNest &passManager);
-/// Pass to bubble up ordinal operations to allow workgroup count computation
-/// based on slices to correlate back to workload computation.
-std::unique_ptr<Pass> createBubbleUpOrdinalOpsPass();
+//------------------------------------------------------------------------------
+// Wrappers that not use tablegen options. See Passes.td for details.
+//------------------------------------------------------------------------------
-/// Pass to perform canonicalizations/cleanups related to HAL interface/buffer
-/// allocations and view operations.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createCleanupBufferAllocViewPass();
-/// Pass to bufferize dispatches that are copying from one interface to
-/// another. This will create a `linalg.generic` op which is a copy that can
-/// then be used by backends to handle appropriately.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createBufferizeCopyOnlyDispatchesPass();
-
-/// Pass to perform canonicalizations/cleanups related to HAL interface/buffer
-/// allocations and view operations.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createCleanupBufferAllocViewPass();
-
-/// Concretizes tensor.pad op's result shape if its source op implements
-/// OffsetSizeAndStrideOpInterface. For example, pad(extract_slice).
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createConcretizePadResultShapePass();
-
-/// Convert BF16 buffer ops and conversions to simulated behavior with uint16.
-std::unique_ptr<OperationPass<>> createConvertBf16ArithToF32Pass();
-
-/// Convert BF16 buffer ops and conversions to simulated behavior with uint16.
-std::unique_ptr<OperationPass<>> createConvertBf16ToUInt16BuffersPass();
-
-/// Converts entry point function within dispatch regions to use
-/// destination-passing style, which is better suited for the upstream
-/// comprehensive bufferization pass.
std::unique_ptr<InterfacePass<FunctionOpInterface>>
createConvertToDestinationPassingStylePass(
- bool useWARForCooperativeMatrixCodegen = false);
-
-/// Converts convolution operations to a GEMM with an im2col op on the image.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createConvolutionToIGEMMPass();
-
-// Decompose affine.apply operations into sub affine.apply that can be
-// hoisted in different loops.
-std::unique_ptr<Pass> createDecomposeAffineOpsPass();
-
-// Decomposes high-D convolution ops into low-D ones.
-std::unique_ptr<Pass> createDecomposeConvolutionToLowerDimOpsPass();
-
-// Decomposes linalg generics on tensors into generics containing no more than
-// one op in the body.
-std::unique_ptr<Pass> createDecomposeLinalgGenericPass();
-
-/// Creates a pass to decompose tensor.pack and tensor.unpack ops. The pass does
-/// tiling and generalization. See implementation for more details.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createDecomposePackUnPackOpsPass(bool tileOuterToOne = false);
-
-/// Creates a pass to convert the softmax op into a sequence of linalg generic
-/// ops.
-std::unique_ptr<Pass> createDecomposeSoftmaxPass(bool useFusion = true);
-
-/// A pass to eliminate tensor.empty ops that could turn into allocations
-/// during bufferization.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createEliminateEmptyTensorsPass();
-
-/// A pass to emulate memref load operations that use narrow integer types
-/// with equivalent operations on supported wide integer types.
-std::unique_ptr<OperationPass<>> createEmulateNarrowTypePass();
-
-/// Creates a pass to erase dead alloc ops where all uses are just store ops.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createEraseDeadAllocAndStoresPass();
-
-std::unique_ptr<Pass> createEraseHALDescriptorTypeFromMemRefPass();
-
-std::unique_ptr<Pass> createConvertHALDescriptorTypeToGPUAddressSpacePass();
-
-// Extract address computations into their own separate instructions.
-std::unique_ptr<Pass> createExtractAddressComputationPass();
-
-// Extract address computations (including the ones with GPU instructions) into
-// their own separate instructions.
-std::unique_ptr<Pass> createExtractAddressComputationGPUPass();
-
-/// Flattens n-D MemRef subspan ops to 1-D MemRef and folds the byte offsets
-/// on subspan ops to the consumer load/store ops, in preparation for lowering
-/// to backends that require linearized access.
-std::unique_ptr<OperationPass<ModuleOp>> createFlattenMemRefSubspanPass();
-
-/// Creates a pass to fold `affine.min` ops in tiled and distributed loops.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createFoldAffineMinInDistributedLoopsPass();
-
-/// After running the upstream TensorConstantBufferize pass, remove
-/// tensor_loads introduced for use only in tensor_extract. These can be
-/// folded to use a load of the created memref object that holds the constant
-/// values.
-std::unique_ptr<OperationPass<>> createFoldTensorExtractOpPass();
-
-/// An ad-hoc pass to canonicalize selected loop carried dependencies on
-/// scf.for.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createForOpCanonicalizationPass();
+ bool useWARForCooperativeMatrixCodegen);
std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createMaterializeEncodingIntoNopPass();
+createDecomposePackUnPackOpsPass(bool tileOuterToOne);
-/// Fuses tensor.pad ops into their consumer ops' tiled loop nests.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createFuseTensorPadWithConsumerPass();
-
-struct GenericVectorizationPassOptions {
- bool enableVectorMasking = false;
- // Controls whether the op lowering configuration (if present) should be used
- // to specify the masked vector sizes.
- bool useConfiguredVectorSizes = true;
- bool vectorizePadding = false;
- bool vectorizeGatherAccesses = false;
- // The flag controls whether it touches the structure generated from tiling,
- // which affects later steps like bufferization and vector hoisting.
- bool enableCleanup = true;
- // Enable conversion for reduction ops to contraction ops.
- bool generateContract = true;
- // Enable folding casting ops into contraction ops. Note that the resulting
- // mixed-type contraction ops are only handled by certain backends.
- bool foldCastIntoContract = false;
- // Max vector size allowed to avoid creating large vectors.
- int64_t maxVectorSize = std::numeric_limits<int64_t>::max();
-};
-/// Creates a pass to perform vectorization on LinAlg and tensor ops.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createGenericVectorizationPass();
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createGenericVectorizationPass(const GenericVectorizationPassOptions &options);
-
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createOptimizeTensorInsertExtractSlicesPass();
-
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createHoistStaticallyBoundAllocationsPass();
-
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createHoistUnrolledVectorExtractInsertSlicePass();
+std::unique_ptr<Pass> createDecomposeSoftmaxPass(bool useFusion);
/// Pass to perform linalg on tensor bufferization. The function passed into
/// the pass through the `allocationFn` argument is invoked whenever a new
@@ -204,104 +74,28 @@
/// striding) and default memory space.
std::unique_ptr<InterfacePass<FunctionOpInterface>>
createIREEComprehensiveBufferizePass(
- std::optional<BufferizationOptions::AllocationFn> allocationFn =
- std::nullopt,
- std::optional<BufferizationOptions::MemCpyFn> memCpyFn = std::nullopt);
-
-/// Pass to resolve `memref.expand_strided_metadata` operations.
-std::unique_ptr<Pass> createIREEExpandStridedMetadataPass();
-
-/// Instruments memory reads and writes for address tracking.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createInstrumentMemoryAccessesPass();
-
-/// Pass to lower an executable using transform dialect sequence provided in the
-/// module
-std::unique_ptr<OperationPass<ModuleOp>>
-createLowerExecutableUsingTransformDialectPass();
-
-/// Pass to lower ukernel operations into their defined function calls.
-std::unique_ptr<OperationPass<ModuleOp>> createLowerUKernelOpsToCallsPass();
-
-/// Creates a pass to convert memref.copy to linalg op.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createMemrefCopyToLinalgPass();
-
-/// Extracts lowering configs and translation info from user configs.
-std::unique_ptr<OperationPass<ModuleOp>> createMaterializeUserConfigsPass();
-
-/// Normalizes the iteration range of `scf.for` and `scf.forall` loops to
-/// [0, ub) += 1.
-std::unique_ptr<Pass>
-createNormalizeLoopBoundsPass(bool normalizeFor = true,
- bool normalizeForall = true);
-
-/// Pass to optimize vector transfer_read and transfer_write.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createOptimizeVectorTransferPass(bool flatten = false);
-
-/// Pad dynamic alloc op to convert them into static one.
-std::unique_ptr<InterfacePass<FunctionOpInterface>> createPadDynamicAlloc();
-
-/// Pass to convert math operations to their polynomial approximation.
-std::unique_ptr<OperationPass<>> createPolynomialApproximationPass();
-
-/// Pass to propagate reshapes by expansion through all ops without explicit
-/// lowering configurations.
-std::unique_ptr<OperationPass<>> createPropagateReshapesByExpansionPass();
-
-/// Pass to reconcile TranslationInfo across multiple functions in a dispatch
-/// and set the appropriate values on the surrounding HAL ops.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableVariantOp>>
-createReconcileTranslationInfoPass();
-
-/// Pass to fuse parallel linalg operations.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createRematerializeParallelOpsPass();
-
-/// Creates a pass to remove single iteration distributed loops.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createRemoveSingleIterationLoopPass();
-
-/// Create a pass that replaces maximumf/minimumf with minumf/maxnumf ops.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createReplaceSlowMinMaxOpsPass();
-
-/// Pass to optimize vector transfer_read and transfer_write. See Passes.td for
-/// `option` details.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createSplitFullPartialTransferPass();
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createSplitFullPartialTransferPass(StringRef option);
-
-/// Tests iree-hal-preprocess-executables-with behavior.
-std::unique_ptr<OperationPass<void>> createTestExecutablePreprocessingPass();
-
-/// Pass to test Partitionable loop interface
-std::unique_ptr<OperationPass<void>>
-createTestPartitionableLoopsInterfacePass();
-
-/// Pass to tile and distribute to workgroups.
-std::unique_ptr<InterfacePass<FunctionOpInterface>>
-createTileAndDistributeToWorkgroupsPass(
- int32_t maxWorkgroupParallelDims = kNumMaxParallelDims,
- linalg::DistributionMethod distributionMethod =
- linalg::DistributionMethod::Cyclic);
+ std::optional<BufferizationOptions::AllocationFn> allocationFn,
+ std::optional<BufferizationOptions::MemCpyFn> memCpyFn);
/// Create an IREE-specific Transform dialect interpreter pass with all
/// registrations necessary for IREE.
std::unique_ptr<Pass>
-createTransformDialectInterpreterPass(StringRef transformSequenceName = "");
+createTransformDialectInterpreterPass(StringRef transformSequenceName);
-/// Pass to propagate type to avoid generating load/stores of illegal types.
-std::unique_ptr<InterfacePass<FunctionOpInterface>> createTypePropagationPass();
+/// Pass to tile and distribute to workgroups.
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+createTileAndDistributeToWorkgroupsPass(
+ int32_t maxWorkgroupParallelDims,
+ linalg::DistributionMethod distributionMethod);
-/// Pass to vectorize memref.copy ops.
-std::unique_ptr<OperationPass<void>> createVectorizeMemrefCopyPass();
+// TODO(hanchung): Move it where it is defined (i.e., Codegen/LLVMGPU).
+// Extract address computations (including the ones with GPU instructions) into
+// their own separate instructions.
+std::unique_ptr<Pass> createExtractAddressComputationGPUPass();
-/// Creates a pass to vectorize a very specific form of tensor.pad ops with
-/// control flows.
-std::unique_ptr<InterfacePass<FunctionOpInterface>> createVectorizePadPass();
+//----------------------------------------------------------------------------//
+// CodeGen Common Patterns
+//----------------------------------------------------------------------------//
/// Populates patterns with patterns to concretize tensor.pad op's result
/// shape. `numWorkgroups`, if not empty, will be used as bounds for simplifying
@@ -338,12 +132,16 @@
void populateVectorTransferTensorSliceTransforms(RewritePatternSet &patterns,
PatternBenefit benefit = 1);
+//----------------------------------------------------------------------------//
+// Register CodeGen Common Passes
+//----------------------------------------------------------------------------//
+
+#define GEN_PASS_DECL
+#include "iree/compiler/Codegen/Common/Passes.h.inc" // IWYU pragma: keep
+
/// Method to register all passes.
void registerCodegenCommonPasses();
-/// Populate Encoding to Nop pass and canonicalizer pass to the pipeline
-void addEncodingToNopPasses(FunctionLikeNest &passManager);
-
} // namespace mlir::iree_compiler
#endif // IREE_COMPILER_CODEGEN_COMMON_PASSES_H_
diff --git a/compiler/src/iree/compiler/Codegen/Common/Passes.td b/compiler/src/iree/compiler/Codegen/Common/Passes.td
index 5bda8d1..a10351e 100644
--- a/compiler/src/iree/compiler/Codegen/Common/Passes.td
+++ b/compiler/src/iree/compiler/Codegen/Common/Passes.td
@@ -13,56 +13,62 @@
// Common passes for all backends (keep alphabetical)
//===---------------------------------------------------------------------===//
-def AddFastMathFlags
+def AddFastMathFlagsPass
: Pass<"iree-codegen-add-fast-math-flags", "LLVM::LLVMFuncOp"> {
let summary = "Add fast math flags to all the operations supporting them, "
"given a floating-point mode.";
- let constructor = "mlir::iree_compiler::createAddFastMathFlagsPass()";
}
-def BubbleUpOrdinalOps : Pass<"iree-codegen-bubble-up-ordinal-ops", ""> {
+def BubbleUpOrdinalOpsPass : Pass<"iree-codegen-bubble-up-ordinal-ops", ""> {
let summary = "Bubbles op ordinal ops to allow for workgroup count computation";
- let constructor = "mlir::iree_compiler::createBubbleUpOrdinalOpsPass()";
+ let description = [{
+ Pass to bubble up ordinal operations to allow workgroup count computation
+ based on slices to correlate back to workload computation.
+ }];
}
-def BufferizeCopyOnlyDispatches :
+def BufferizeCopyOnlyDispatchesPass :
InterfacePass<"iree-codegen-bufferize-copy-only-dispatches", "mlir::FunctionOpInterface"> {
let summary =
"Bufferize dispatches that copy to/from interfaces to convert to a linalg.copy op";
- let constructor = "mlir::iree_compiler::createBufferizeCopyOnlyDispatchesPass()";
+ let description = [{
+ Pass to bufferize dispatches that are copying from one interface to
+ another. This will create a `linalg.generic` op which is a copy that can
+ then be used by backends to handle appropriately.
+ }];
}
-def CleanupBufferAllocView :
+def CleanupBufferAllocViewPass :
InterfacePass<"iree-codegen-cleanup-buffer-alloc-view", "mlir::FunctionOpInterface"> {
let summary =
"Performs cleanups over HAL interface/buffer allocation/view operations";
- let constructor = "mlir::iree_compiler::createCleanupBufferAllocViewPass()";
}
-def ConcretizePadResultShape :
+def ConcretizePadResultShapePass :
InterfacePass<"iree-codegen-concretize-pad-result-shape", "mlir::FunctionOpInterface"> {
let summary =
"Concretizes tensor.pad op's result shape if its source op"
"implements OffsetSizeAndStrideOpInterface.";
- let constructor = "mlir::iree_compiler::createConcretizePadResultShapePass()";
}
-def ConvertBf16ArithToF32 : Pass<"iree-convert-bf16-arith-to-f32", ""> {
+def ConvertBf16ArithToF32Pass : Pass<"iree-convert-bf16-arith-to-f32", ""> {
let summary = "Convert bf16 arithmetic operations to f32";
- let constructor = "mlir::iree_compiler::createConvertBf16ArithToF32Pass()";
}
-def ConvertBf16ToUInt16Buffers :
- Pass<"iree-convert-bf16-to-uint16-buffers", ""> {
- let summary = "Convert bf16 buffers to uint16 equivalents";
- let constructor = "mlir::iree_compiler::createConvertBf16ToUInt16BuffersPass()";
+def ConvertBf16ToUInt16BuffersPass :
+ Pass<"iree-codegen-convert-bf16-to-uint16-buffers", ""> {
+ let summary = "Convert BF16 buffer ops and conversions to simulated behavior with uint16.";
}
-def ConvertToDestinationPassingStyle :
+def ConvertToDestinationPassingStylePass :
InterfacePass<"iree-codegen-convert-to-destination-passing-style", "mlir::FunctionOpInterface"> {
let summary =
"Transforms the code to make the dispatch use destination-passing style";
- let constructor = "mlir::iree_compiler::createConvertToDestinationPassingStylePass()";
+ let description = [{
+ Converts entry point function within dispatch regions to use
+ destination-passing style, which is better suited for the upstream
+ comprehensive bufferization pass.
+ }];
let options = [
Option<"useWARForCooperativeMatrixCodegen", "use-war-for-cooperative-matrix-codegen",
"bool", /*default=*/"false",
@@ -70,25 +76,16 @@
];
}
-def ConvolutionToIGEMM :
- InterfacePass<"iree-codegen-convolution-to-igemm", "mlir::FunctionOpInterface"> {
- let summary =
- "Transforms convolution operations into an implicit GEMM format.";
- let constructor = "mlir::iree_compiler::createConvolutionToIGEMMPass()";
-}
-
-def DecomposeAffineOps: Pass<"decompose-affine-ops"> {
+def DecomposeAffineOpsPass: Pass<"iree-codegen-decompose-affine-ops"> {
let summary = "Decompose `affine.apply` operations into sub `affine.apply`";
let description = [{
Decompose `affine.apply` operations into sub `affine.apply` where each
sub expression references values that are defined in the same loop scope.
The sub expression are then stitched back together following the loop
nest order.
-
The goal of this pass is to break down `affine.apply` expressions such
that the resulting sub expressions can be hoisted out in their respective
loop.
-
E.g., Let's say we have
```mlir
%res = affine.apply
@@ -113,27 +110,28 @@
```
Now the sequence of instructions leading to and including
`%inv1x32_plus_inv2` can be hoisted out of the loop.
-
This pass requires `scf.for` structures to still be around otherwise
the break down will be meaningless.
-
Note: The decomposition performed by this pass will be undone by
canonicalization. Make sure to lower the resulting ops before that.
}];
- let constructor = "mlir::iree_compiler::createDecomposeAffineOpsPass()";
let dependentDialects = [
- "affine::AffineDialect"
+ "mlir::affine::AffineDialect"
];
}
-def DecomposeConvolutionToLowerDimOps :
- Pass<"iree-codegen-decompose-convolution-to-lower-dim-ops", ""> {
- let summary = "Decomposes linalg convolution ops to lower dim ops";
- let constructor =
- "mlir::iree_compiler::createDecomposeConvolutionToLowerDimOpsPass()";
+def ConvolutionToIGEMMPass :
+ InterfacePass<"iree-codegen-convolution-to-igemm", "mlir::FunctionOpInterface"> {
+ let summary =
+ "Transforms convolution operations into an implicit GEMM format.";
}
-def DecomposeLinalgGeneric :
+def DecomposeConvolutionToLowerDimOpsPass :
+ Pass<"iree-codegen-decompose-convolution-to-lower-dim-ops", ""> {
+ let summary = "Decomposes linalg convolution ops to lower dim ops";
+}
+
+def DecomposeLinalgGenericPass :
Pass<"iree-codegen-decompose-linalg-generic", ""> {
let summary = "Decomposes linalg generic ops into individual ops";
let description = [{
@@ -145,24 +143,21 @@
Operates on tensor based linalg ops.
}];
- let constructor = "mlir::iree_compiler::createDecomposeLinalgGenericPass()";
}
-def DecomposePackUnPackOps :
+def DecomposePackUnPackOpsPass :
InterfacePass<"iree-codegen-decompose-pack-unpack-ops", "mlir::FunctionOpInterface"> {
let summary = "Decompose pack/unpack ops into vectorizable ops";
- let constructor = "mlir::iree_compiler::createDecomposePackUnPackOpsPass()";
let options = [
Option<"tileOuterToOne", "tile-outer-to-one", "bool", "false",
"Always apply tiling to make outer dimension be ones">
];
}
-def DecomposeSoftmax :
+def DecomposeSoftmaxPass :
InterfacePass<"iree-codegen-decompose-softmax", "mlir::FunctionOpInterface"> {
let summary =
"Decomposes softmax op into a sequence of linalg ops";
- let constructor = "mlir::iree_compiler::createDecomposeSoftmaxPass()";
let options = [
Option<"useFusion", "use-fusion",
"bool", /*default=*/"true",
@@ -170,57 +165,51 @@
];
}
-def ReconcileTranslationInfo
+def ReconcileTranslationInfoPass
: Pass<"iree-codegen-reconcile-translation-info", "IREE::HAL::ExecutableVariantOp"> {
let summary =
"Reconcile information (like workgroup_size, subgroup_size) across "
"`TranslationInfo` set on each function in the dispatch and merge them"
"and set them at the appropriate places in the surrounding HAL ops";
- let constructor = "mlir::iree_compiler::createReconcileTranslationInfoPass()";
}
-def ReplaceSlowMinMaxOps
+def ReplaceSlowMinMaxOpsPass
: InterfacePass<"iree-codegen-replace-slow-min-max-ops", "mlir::FunctionOpInterface"> {
let summary =
"Replace slow min/max operations that propagate NaNs and distinguish "
"between +/-0.0 with faster min/max operations that ignore them.";
- let constructor = "mlir::iree_compiler::createReplaceSlowMinMaxOpsPass()";
}
-def EliminateEmptyTensors :
+def EliminateEmptyTensorsPass :
InterfacePass<"iree-eliminate-empty-tensors", "mlir::FunctionOpInterface"> {
let summary = "Eliminate tensor.empty ops to avoid buffer allocations";
- let constructor = "mlir::iree_compiler::createEliminateEmptyTensorsPass()";
}
-def EmulateNarrowType :
+def EmulateNarrowTypePass :
Pass<"iree-codegen-emulate-narrow-type", ""> {
let summary = "Emulate narrow integer operations using wide integer operations";
- let constructor = "mlir::iree_compiler::createEmulateNarrowTypePass()";
+ let description = [{
+ A pass to emulate memref load operations that use narrow integer types
+ with equivalent operations on supported wide integer types.
+ }];
}
-def EraseDeadAllocAndStores :
+def EraseDeadAllocAndStoresPass :
InterfacePass<"iree-codegen-erase-dead-alloc-and-stores", "mlir::FunctionOpInterface"> {
let summary = "Erase alloc ops if all the uses are just stores";
- let constructor =
- "mlir::iree_compiler::createEraseDeadAllocAndStoresPass()";
}
-def EraseHALDescriptorTypeFromMemRef
+def EraseHALDescriptorTypeFromMemRefPass
: Pass<"iree-codegen-erase-hal-descriptor-type-from-memref"> {
let summary = "Erase #hal.descriptor_type from MemRef memory space";
- let constructor =
- "mlir::iree_compiler::createEraseHALDescriptorTypeFromMemRefPass()";
}
-def ConvertHALDescriptorTypeToGPUAddressSpace
+def ConvertHALDescriptorTypeToGPUAddressSpacePass
: Pass<"iree-codegen-convert-hal-descriptor-type-to-gpu-address-space"> {
let summary = "Convert #hal.descriptor_type to #gpu.address_space<global>";
- let constructor =
- "mlir::iree_compiler::createConvertHALDescriptorTypeToGPUAddressSpacePass()";
}
-def ExtractAddressComputation: Pass<"extract-address-computation"> {
+def ExtractAddressComputationPass : Pass<"iree-codegen-extract-address-computation"> {
let summary = "Extract address computations from memory accesses";
let description = [{
Extract the address computation from the instructions with memory
@@ -237,44 +226,49 @@
memref.load %new_base[%c0,...]
```
}];
- let constructor = "mlir::iree_compiler::createExtractAddressComputationPass()";
let dependentDialects = [
"memref::MemRefDialect"
];
}
-def FlattenMemRefSubspan : Pass<"iree-codegen-flatten-memref-subspan", "ModuleOp"> {
+def FlattenMemRefSubspanPass : Pass<"iree-codegen-flatten-memref-subspan", "ModuleOp"> {
let summary =
"Flatten n-D MemRef subspan ops to 1-D ones and fold byte offsets";
- let constructor = "mlir::iree_compiler::createFlattenMemRefSubspanPass()";
+ let description = [{
+ Flattens n-D MemRef subspan ops to 1-D MemRef and folds the byte offsets
+ on subspan ops to the consumer load/store ops, in preparation for lowering
+ to backends that require linearized access.
+ }];
}
-def FoldAffineMinInDistributedLoops :
+def FoldAffineMinInDistributedLoopsPass :
InterfacePass<"iree-codegen-fold-affinemin-in-distributed-loops", "mlir::FunctionOpInterface"> {
let summary = "Fold `affine.min` ops in distributed loops";
- let constructor = "mlir::iree_compiler::createFoldAffineMinInDistributedLoopsPass()";
}
-def FoldTensorExtractOp :
+def FoldTensorExtractOpPass :
Pass<"iree-codegen-fold-tensor-extract-op", ""> {
let summary = "Fold `tensor.extract` operations prior to lowering to LLVM";
- let constructor = "mlir::iree_compiler::createFoldTensorExtractOpPass()";
+ let description = [{
+ After running the upstream TensorConstantBufferize pass, remove
+ tensor_loads introduced for use only in tensor_extract. These can be
+ folded to use a load of the created memref object that holds the constant
+ values.
+ }];
}
-def ForOpCanonicalization :
+def ForOpCanonicalizationPass :
InterfacePass<"iree-codegen-canonicalize-scf-for", "mlir::FunctionOpInterface"> {
let summary =
"Adhoc canonicalization of selected loop-carried values/dependencies for scf.for ops";
- let constructor = "mlir::iree_compiler::createForOpCanonicalizationPass()";
}
-def FuseTensorPadWithConsumer :
+def FuseTensorPadWithConsumerPass :
InterfacePass<"iree-codegen-fuse-tensor-pad-with-consumer", "mlir::FunctionOpInterface"> {
let summary = "Fuse tensor.pad op into its consumer op's tiled loop nest";
- let constructor = "mlir::iree_compiler::createFuseTensorPadWithConsumerPass()";
}
-def GenericVectorization :
+def GenericVectorizationPass :
InterfacePass<"iree-codegen-generic-vectorization", "mlir::FunctionOpInterface"> {
let summary = "Pass to perform vectorization on tensor/linalg ops.";
let options = [
@@ -297,29 +291,23 @@
/*default=*/"2147483647",
"Max vector size allowed to avoid creating large vectors.">
];
- let constructor =
- "mlir::iree_compiler::createGenericVectorizationPass()";
}
-def OptimizeTensorInsertExtractSlices
+def OptimizeTensorInsertExtractSlicesPass
: InterfacePass<"iree-codegen-optimize-tensor-insert-extract-slices",
"mlir::FunctionOpInterface"> {
let summary = "Optimize tensor.insert_slice/tensor.extract_slice operations "
"(e.g. hoist and fold)";
- let constructor =
- "mlir::iree_compiler::createOptimizeTensorInsertExtractSlicesPass()";
}
-def HoistUnrolledVectorExtractInsertSlice :
+def HoistUnrolledVectorExtractInsertSlicePass :
InterfacePass<"iree-codegen-hoist-vector-extract-insert-slice", "mlir::FunctionOpInterface"> {
let summary = "Hoist unrolled vector (extract, insert) pairs out of scf.for op";
- let constructor = "mlir::iree_compiler::createHoistUnrolledVectorExtractInsertSlicePass()";
}
-def HoistStaticallyBoundAllocations :
- InterfacePass<"iree-hoist-statically-bound-allocations", "mlir::FunctionOpInterface"> {
+def HoistStaticallyBoundAllocationsPass :
+ InterfacePass<"iree-codegen-hoist-statically-bound-allocations", "mlir::FunctionOpInterface"> {
let summary = "Hoist statically bound alloca ops to the entry block of functions";
- let constructor = "mlir::iree_compiler::createHoistStaticallyBoundAllocationsPass()";
// Note: These options only exist to help with testing, real world uses should look at the target.
// There also should be no observable change if the input IR is not scalable.
let options = [
@@ -335,10 +323,9 @@
];
}
-def IREEComprehensiveBufferize :
+def IREEComprehensiveBufferizePass :
InterfacePass<"iree-codegen-iree-comprehensive-bufferize", "mlir::FunctionOpInterface"> {
let summary = "Convert from to Linalg ops on tensors to buffers";
- let constructor = "mlir::iree_compiler::createIREEComprehensiveBufferizePass()";
let options = [
Option<"testAnalysisOnly", "test-analysis-only", "bool",
/*default=*/"false",
@@ -349,43 +336,37 @@
];
}
-def IREEExpandStridedMetadata :
+def IREEExpandStridedMetadataPass :
Pass<"iree-codegen-expand-strided-metadata", ""> {
let summary = "Resolve memref.extract_strided_metadata operations";
- let constructor = "mlir::iree_compiler::createIREEExpandStridedMetadataPass()";
let options = [
Option<"allowUnresolved", "allow-unresolved", "bool", /*default=*/"false",
"Allow unresolved strided metadata op (for testing)">,
];
}
-def InstrumentMemoryAccesses :
+def InstrumentMemoryAccessesPass :
InterfacePass<"iree-codegen-instrument-memory-accesses", "mlir::FunctionOpInterface"> {
let summary = "Instruments memory reads and writes for address tracking when dispatch instrumentation is enabled.";
- let constructor = "mlir::iree_compiler::createInstrumentMemoryAccessesPass()";
}
-def LowerExecutableUsingTransformDialect :
+def LowerExecutableUsingTransformDialectPass :
Pass<"iree-codegen-lower-executable-using-transform-dialect", "ModuleOp"> {
let summary = "Lower executables using the transform dialect recipe provided in the module.";
- let constructor = "mlir::iree_compiler::createLowerExecutableUsingTransformDialectPass()";
}
-def LowerUKernelOpsToCalls :
+def LowerUKernelOpsToCallsPass :
Pass<"iree-codegen-lower-ukernel-ops-to-calls", "ModuleOp"> {
let summary = "Lower micro-kernel wrapper ops into function calls";
- let constructor = "mlir::iree_compiler::createLowerUKernelOpsToCallsPass()";
}
-def MaterializeEncodingIntoNop :
+def MaterializeEncodingIntoNopPass :
InterfacePass<"iree-codegen-materialize-encoding-into-nop", "mlir::FunctionOpInterface"> {
- let summary = "";
- let constructor = "mlir::iree_compiler::createMaterializeEncodingIntoNopPass()";
+ let summary = "Drop the encodings from tensor types with encodings.";
}
-def MaterializeUserConfigs : Pass<"iree-codegen-materialize-user-configs", "ModuleOp"> {
+def MaterializeUserConfigsPass : Pass<"iree-codegen-materialize-user-configs", "ModuleOp"> {
let summary = "Sets the lowering configs and translation info from user configs";
- let constructor = "mlir::iree_compiler::createMaterializeUserConfigsPass()";
let dependentDialects = [
"transform::TransformDialect"
];
@@ -394,14 +375,15 @@
def MemrefCopyToLinalgPass :
InterfacePass<"iree-codegen-memrefcopy-to-linalg", "mlir::FunctionOpInterface"> {
let summary = "Convert memref.copy to linalg op";
- let constructor =
- "mlir::iree_compiler::createMemrefCopyToLinalgPass()";
}
def NormalizeLoopBoundsPass :
Pass<"iree-codegen-normalize-loop-bounds", ""> {
let summary = "Normalize the loop bounds of `scf.for` and `scf.forall`";
- let constructor = "mlir::iree_compiler::createNormalizeLoopBoundsPass()";
+ let description = [{
+ Normalizes the iteration range of `scf.for` and `scf.forall` loops to
+ [0, ub) += 1.
+ }];
let options = [
Option<"normalizeFor", "normalize-for", "bool", "true",
"Enable normalization for `scf.for` loops">,
@@ -414,13 +396,12 @@
];
}
-def OptimizeVectorTransfer :
+def OptimizeVectorTransferPass :
InterfacePass<"iree-codegen-optimize-vector-transfer", "mlir::FunctionOpInterface"> {
let summary =
"Run optimization transformations on vector transfer operations";
- let constructor = "mlir::iree_compiler::createOptimizeVectorTransferPass()";
let options = [
- Option<"optionFlatten", "flatten", "bool", "false",
+ Option<"flatten", "flatten", "bool", "false",
"Flatten the vector type of vector transfers where possible (contiguous row-major data).">,
];
let dependentDialects = [
@@ -428,43 +409,40 @@
];
}
-def PadDynamicAlloc :
+def PadDynamicAllocPass :
InterfacePass<"iree-codegen-pad-dynamic-alloc", "mlir::FunctionOpInterface"> {
let summary = "Pass to pad dynamic alloc into static one.";
- let constructor = "mlir::iree_compiler::createPadDynamicAlloc()";
}
def PolynomialApproximationPass :
Pass<"iree-codegen-polynomial-approximation", ""> {
let summary = "Convert math operations to their polynomial approximation";
- let constructor =
- "mlir::iree_compiler::createPolynomialApproximationPass()";
}
def PropagateReshapesByExpansionPass :
Pass<"iree-codegen-propagate-reshapes-by-expansion", ""> {
let summary = "Propagates reshaping operations by expansion.";
- let constructor = "mlir::iree_compiler::createPropagateReshapesByExpansionPass()";
+ let description = [{
+ Pass to propagate reshapes by expansion through all ops without explicit
+ lowering configurations.
+ }];
}
-def RematerializeParallelOps :
+def RematerializeParallelOpsPass :
InterfacePass<"iree-codegen-rematerialize-parallel-ops", "mlir::FunctionOpInterface"> {
let summary = "Pass to rematerialize and merge parallel ops into consumers.";
- let constructor = "mlir::iree_compiler::createRematerializeParallelOpsPass()";
}
-def RemoveSingleIterationLoop :
+def RemoveSingleIterationLoopPass :
InterfacePass<"iree-codegen-remove-single-iteration-loop", "mlir::FunctionOpInterface"> {
let summary = "Remove distributed loop with single iteration.";
- let constructor = "mlir::iree_compiler::createRemoveSingleIterationLoopPass()";
}
-def SplitFullPartialTransfer :
+def SplitFullPartialTransferPass :
InterfacePass<"iree-codegen-split-full-partial-transfer", "mlir::FunctionOpInterface"> {
let summary =
"Split a vector.transfer operation into an in-bounds (i.e., no "
"out-of-bounds masking) fastpath and a slowpath.";
- let constructor = "mlir::iree_compiler::createSplitFullPartialTransferPass()";
let options = [
Option<"splitVectorTransfersTo", "split-transfers", "std::string",
/*default=*/"",
@@ -477,45 +455,48 @@
];
}
-def TensorToVectorVectorizePad :
+def TensorToVectorVectorizePadPass :
InterfacePass<"iree-codegen-vectorize-tensor-pad", "mlir::FunctionOpInterface"> {
let summary = "Vectorize a very specific form of tensor.pad with "
"control flows";
- let constructor =
- "mlir::iree_compiler::createVectorizePadPass()";
}
-def TestExecutablePreprocessing :
+def TestExecutablePreprocessingPass :
Pass<"iree-codegen-test-executable-preprocessing", ""> {
let summary = "Tests iree-hal-preprocess-executables-with behavior.";
- let constructor = "mlir::iree_compiler::createTestExecutablePreprocessingPass()";
}
-def TestPartitionableLoopsInterface :
+def TestPartitionableLoopsInterfacePass :
Pass<"iree-codegen-test-partitionable-loops-interface", ""> {
let summary = "Test the PartitionableLoopsInterface";
- let constructor = "mlir::iree_compiler::createTestPartitionableLoopsInterfacePass()";
}
-def TileAndDistributeToWorkgroups :
+def TileAndDistributeToWorkgroupsPass :
InterfacePass<"iree-codegen-tile-and-distribute-to-workgroups", "mlir::FunctionOpInterface"> {
let summary = "Tile and distribute operations to workgroups";
- let constructor = "mlir::iree_compiler::createTileAndDistributeToWorkgroupsPass()";
let options = [
Option<"maxWorkgroupParallelDims", "max-workgroup-parallel-dims", "int32_t",
- /*default=*/ "",
+ /*default=*/ "kNumMaxParallelDims",
"Maximum number of dims to distribute workgroups across.">,
- Option<"distributionMethod", "distribution-method", "int32_t",
- /*default=*/ "0",
- "Pick the distribution method">
+ Option<"distributionMethod", "distribution-method", "linalg::DistributionMethod",
+ /*default=*/ "linalg::DistributionMethod::Cyclic",
+ "Pick the distribution method. See linalg::DistributionMethod for details",
+ [{::llvm::cl::values(
+ clEnumValN(linalg::DistributionMethod::Cyclic,
+ "0", "Use Cyclic strategy"),
+ clEnumValN(linalg::DistributionMethod::CyclicNumProcsGeNumIters,
+ "1", "Use CyclicNumProcGeNumIter strategy"),
+ clEnumValN(linalg::DistributionMethod::CyclicNumProcsEqNumIters,
+ "2", "Use CyclicNumProcEqNumIter strategy"),
+ clEnumValN(linalg::DistributionMethod::None,
+ "3", "Use None strategy")
+ )}]>,
];
}
-def TransformDialectInterpreter :
+def TransformDialectInterpreterPass :
Pass<"iree-transform-dialect-interpreter"> {
let summary = "Pass to apply transform dialect operations.";
- let constructor =
- "mlir::iree_compiler::createTransformDialectInterpreterPass()";
let description = [{
This pass runs the transform dialect interpreter and applies the named
sequence transformation specified by the provided name (defaults to
@@ -523,7 +504,7 @@
}];
let options = [
Option<"entryPoint", "entry-point", "std::string",
- /*default=*/[{"transform::TransformDialect::kTransformEntryPointSymbolName.str()"}],
+ /*default=*/"::mlir::transform::TransformDialect::kTransformEntryPointSymbolName.str()",
"Entry point of the pass pipeline.">,
Option<"libraryFileName", "library-file-name", "std::string",
/*default=*/[{""}],
@@ -531,16 +512,14 @@
];
}
-def TypePropagation :
+def TypePropagationPass :
InterfacePass<"iree-codegen-type-propagation", "mlir::FunctionOpInterface"> {
let summary = "Propogate the type of tensor to avoid load/stores of illegal bit widths";
- let constructor = "mlir::iree_compiler::createTypePropagationPass()";
}
def VectorizeMemrefCopyPass :
Pass<"iree-codegen-vectorize-memref-copy", ""> {
let summary = "Vectorizes memref copy operations.";
- let constructor = "mlir::iree_compiler::createVectorizeMemrefCopyPass()";
}
#endif // IREE_CODEGEN_COMMON_PASSES
diff --git a/compiler/src/iree/compiler/Codegen/Common/PolynomialApproximationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/PolynomialApproximationPass.cpp
index b401b46..605ec5d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/PolynomialApproximationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/PolynomialApproximationPass.cpp
@@ -3,7 +3,7 @@
// Licensed under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
+
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Math/Transforms/Approximation.h"
#include "mlir/Dialect/Math/Transforms/Passes.h"
@@ -20,11 +20,15 @@
"Skip polynomial lowering for math op natively available on GPU"),
llvm::cl::init(false));
+#define GEN_PASS_DEF_POLYNOMIALAPPROXIMATIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// math dialect elementry functions -> polynomial form.
-class PolynomialApproximationPass
- : public PolynomialApproximationPassBase<PolynomialApproximationPass> {
+class PolynomialApproximationPass final
+ : public impl::PolynomialApproximationPassBase<
+ PolynomialApproximationPass> {
void runOnOperation() override {
RewritePatternSet mathPatterns(&getContext());
populateExpandTanPattern(mathPatterns);
@@ -51,9 +55,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<>> createPolynomialApproximationPass() {
- return std::make_unique<PolynomialApproximationPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/PropagateReshapesByExpansion.cpp b/compiler/src/iree/compiler/Codegen/Common/PropagateReshapesByExpansion.cpp
index 860f30f..67d2358 100644
--- a/compiler/src/iree/compiler/Codegen/Common/PropagateReshapesByExpansion.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/PropagateReshapesByExpansion.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
@@ -13,10 +12,13 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_PROPAGATERESHAPESBYEXPANSIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-struct PropagateReshapesByExpansionPass
- : public PropagateReshapesByExpansionPassBase<
+struct PropagateReshapesByExpansionPass final
+ : impl::PropagateReshapesByExpansionPassBase<
PropagateReshapesByExpansionPass> {
void runOnOperation() override;
};
@@ -71,8 +73,4 @@
}
}
-std::unique_ptr<OperationPass<>> createPropagateReshapesByExpansionPass() {
- return std::make_unique<PropagateReshapesByExpansionPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ReconcileTranslationInfo.cpp b/compiler/src/iree/compiler/Codegen/Common/ReconcileTranslationInfo.cpp
index a5d15cf..e8e18f5 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ReconcileTranslationInfo.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ReconcileTranslationInfo.cpp
@@ -14,16 +14,19 @@
// up. In case of inconsistencies, this pass will throw an error.
//===---------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_RECONCILETRANSLATIONINFOPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-class ReconcileTranslationInfoPass
- : public ReconcileTranslationInfoBase<ReconcileTranslationInfoPass> {
+class ReconcileTranslationInfoPass final
+ : public impl::ReconcileTranslationInfoPassBase<
+ ReconcileTranslationInfoPass> {
public:
void runOnOperation() override;
};
@@ -146,9 +149,4 @@
});
}
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableVariantOp>>
-createReconcileTranslationInfoPass() {
- return std::make_unique<ReconcileTranslationInfoPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/RematerializeParallelOps.cpp b/compiler/src/iree/compiler/Codegen/Common/RematerializeParallelOps.cpp
index b71ae9b..e168f6c 100644
--- a/compiler/src/iree/compiler/Codegen/Common/RematerializeParallelOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/RematerializeParallelOps.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
@@ -14,6 +13,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_REMATERIALIZEPARALLELOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
static bool isScalarOrTensorOfSizeOne(Type t) {
@@ -61,8 +63,8 @@
}
};
-struct RematerializeParallelOpsPass
- : public RematerializeParallelOpsBase<RematerializeParallelOpsPass> {
+struct RematerializeParallelOpsPass final
+ : impl::RematerializeParallelOpsPassBase<RematerializeParallelOpsPass> {
void runOnOperation() override {
auto funcOp = getOperation();
RewritePatternSet fusionPatterns(funcOp.getContext());
@@ -76,10 +78,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createRematerializeParallelOpsPass() {
- return std::make_unique<RematerializeParallelOpsPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp b/compiler/src/iree/compiler/Codegen/Common/RemoveSingleIterationLoop.cpp
similarity index 93%
rename from compiler/src/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp
rename to compiler/src/iree/compiler/Codegen/Common/RemoveSingleIterationLoop.cpp
index 9eed5b3..8e7b0b7 100644
--- a/compiler/src/iree/compiler/Codegen/Common/RemoveTrivialLoops.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/RemoveSingleIterationLoop.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
#include "iree/compiler/Codegen/Transforms/Transforms.h"
@@ -22,6 +21,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_REMOVESINGLEITERATIONLOOPPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Converts a symbolic GPU processor dimension to its numeric one.
static unsigned dimToIndex(gpu::Dimension dim) {
switch (dim) {
@@ -111,7 +113,8 @@
namespace {
class RemoveSingleIterationLoopPass final
- : public RemoveSingleIterationLoopBase<RemoveSingleIterationLoopPass> {
+ : public impl::RemoveSingleIterationLoopPassBase<
+ RemoveSingleIterationLoopPass> {
void runOnOperation() override {
auto funcOp = getOperation();
@@ -129,10 +132,4 @@
}
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createRemoveSingleIterationLoopPass() {
- return std::make_unique<RemoveSingleIterationLoopPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp b/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
index d9cbd50..fc27a46 100644
--- a/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/ReplaceSlowMinMaxOps.cpp
@@ -4,15 +4,16 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/Transforms.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-using namespace mlir;
-using namespace mlir::iree_compiler;
+namespace mlir::iree_compiler {
+
+#define GEN_PASS_DEF_REPLACESLOWMINMAXOPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
namespace {
@@ -56,17 +57,17 @@
}
};
-struct ReplaceSlowMinMaxOpsPass
- : public ReplaceSlowMinMaxOpsBase<ReplaceSlowMinMaxOpsPass> {
+struct ReplaceSlowMinMaxOpsPass final
+ : impl::ReplaceSlowMinMaxOpsPassBase<ReplaceSlowMinMaxOpsPass> {
public:
- using ReplaceSlowMinMaxOpsBase::ReplaceSlowMinMaxOpsBase;
+ using impl::ReplaceSlowMinMaxOpsPassBase<
+ ReplaceSlowMinMaxOpsPass>::ReplaceSlowMinMaxOpsPassBase;
void runOnOperation() override;
};
} // namespace
-void mlir::iree_compiler::populateReplaceSlowMinMaxOpsPatterns(
- RewritePatternSet &patterns) {
+void populateReplaceSlowMinMaxOpsPatterns(RewritePatternSet &patterns) {
patterns.add<
ReplaceSlowWithFastMinMaxOpPattern<arith::MinimumFOp, arith::MinNumFOp>,
ReplaceSlowWithFastMinMaxOpPattern<arith::MaximumFOp, arith::MaxNumFOp>,
@@ -84,7 +85,4 @@
}
}
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-mlir::iree_compiler::createReplaceSlowMinMaxOpsPass() {
- return std::make_unique<ReplaceSlowMinMaxOpsPass>();
-}
+} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/SplitFullPartialTransferPass.cpp b/compiler/src/iree/compiler/Codegen/Common/SplitFullPartialTransferPass.cpp
index 88b6705..1b597e1 100644
--- a/compiler/src/iree/compiler/Codegen/Common/SplitFullPartialTransferPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/SplitFullPartialTransferPass.cpp
@@ -6,7 +6,6 @@
#include <string>
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Vector/Transforms/LoweringPatterns.h"
#include "mlir/Dialect/Vector/Transforms/VectorRewritePatterns.h"
@@ -16,14 +15,15 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_SPLITFULLPARTIALTRANSFERPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-struct SplitFullPartialTransferPass
- : public SplitFullPartialTransferBase<SplitFullPartialTransferPass> {
- SplitFullPartialTransferPass() = default;
- SplitFullPartialTransferPass(StringRef option) {
- this->splitVectorTransfersTo = std::string(option);
- }
+struct SplitFullPartialTransferPass final
+ : impl::SplitFullPartialTransferPassBase<SplitFullPartialTransferPass> {
+ using impl::SplitFullPartialTransferPassBase<
+ SplitFullPartialTransferPass>::SplitFullPartialTransferPassBase;
void runOnOperation() override {
MLIRContext *ctx = &getContext();
@@ -46,14 +46,4 @@
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createSplitFullPartialTransferPass() {
- return std::make_unique<SplitFullPartialTransferPass>();
-}
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createSplitFullPartialTransferPass(StringRef option) {
- return std::make_unique<SplitFullPartialTransferPass>(option);
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/VectorizePad.cpp b/compiler/src/iree/compiler/Codegen/Common/TensorToVectorVectorizePad.cpp
similarity index 96%
rename from compiler/src/iree/compiler/Codegen/Common/VectorizePad.cpp
rename to compiler/src/iree/compiler/Codegen/Common/TensorToVectorVectorizePad.cpp
index 45372b5..3023aa4 100644
--- a/compiler/src/iree/compiler/Codegen/Common/VectorizePad.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TensorToVectorVectorizePad.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
@@ -26,6 +25,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_TENSORTOVECTORVECTORIZEPADPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Gets the given `attrOrValue` as an index value by creating constant ops
/// for attributes.
static Value getAsIndexValue(OpFoldResult attrOrValue, OpBuilder &builder,
@@ -229,8 +231,11 @@
}
};
-struct TensorToVectorVectorizePadPass
- : public TensorToVectorVectorizePadBase<TensorToVectorVectorizePadPass> {
+struct TensorToVectorVectorizePadPass final
+ : impl::TensorToVectorVectorizePadPassBase<TensorToVectorVectorizePadPass> {
+ using impl::TensorToVectorVectorizePadPassBase<
+ TensorToVectorVectorizePadPass>::TensorToVectorVectorizePadPassBase;
+
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, arith::ArithDialect,
linalg::LinalgDialect, scf::SCFDialect,
@@ -254,10 +259,4 @@
PatternBenefit baseBenefit) {
patterns.add<VectorizePadWithConditions>(patterns.getContext(), baseBenefit);
}
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createVectorizePadPass() {
- return std::make_unique<TensorToVectorVectorizePadPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/TestExecutablePreprocessing.cpp b/compiler/src/iree/compiler/Codegen/Common/TestExecutablePreprocessing.cpp
index 77e6be3..217d20b 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TestExecutablePreprocessing.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TestExecutablePreprocessing.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
@@ -12,10 +11,14 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_TESTEXECUTABLEPREPROCESSINGPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-struct TestExecutablePreprocessingPass
- : public TestExecutablePreprocessingBase<TestExecutablePreprocessingPass> {
+struct TestExecutablePreprocessingPass final
+ : impl::TestExecutablePreprocessingPassBase<
+ TestExecutablePreprocessingPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::HAL::HALDialect>();
}
@@ -45,9 +48,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<void>> createTestExecutablePreprocessingPass() {
- return std::make_unique<TestExecutablePreprocessingPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp b/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
index 36d7362..3b89434 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TestPartitionableLoopsInterface.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Interfaces/PartitionableLoopsInterface.h"
#include "iree/compiler/Dialect/Util/IR/UtilDialect.h"
@@ -15,6 +14,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_TESTPARTITIONABLELOOPSINTERFACEPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// For ops that implement the `PartitionableLoopsInterface` that have the
@@ -45,8 +47,8 @@
}
};
-struct TestPartitionableLoopsInterfacePass
- : public TestPartitionableLoopsInterfaceBase<
+struct TestPartitionableLoopsInterfacePass final
+ : impl::TestPartitionableLoopsInterfacePassBase<
TestPartitionableLoopsInterfacePass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::Util::UtilDialect>();
@@ -63,10 +65,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<void>>
-createTestPartitionableLoopsInterfacePass() {
- return std::make_unique<TestPartitionableLoopsInterfacePass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/TileAndDistributeToWorkgroupsPass.cpp b/compiler/src/iree/compiler/Codegen/Common/TileAndDistributeToWorkgroupsPass.cpp
index 2ce82fa..9e69888 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TileAndDistributeToWorkgroupsPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TileAndDistributeToWorkgroupsPass.cpp
@@ -15,7 +15,6 @@
//===---------------------------------------------------------------------===//
#include "iree/compiler/Codegen/Common/EncodingUtils.h"
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Common/Transforms.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenAttrs.h"
@@ -45,6 +44,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_TILEANDDISTRIBUTETOWORKGROUPSPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Method to return the configuration to use for first-level tile and
/// distribute. Returns the
/// - Root op of the dispatch. If no root op was found returns `nullptr`.
@@ -260,14 +262,17 @@
namespace {
-struct TileAndDistributeToWorkgroupsPass
- : public TileAndDistributeToWorkgroupsBase<
+struct TileAndDistributeToWorkgroupsPass final
+ : impl::TileAndDistributeToWorkgroupsPassBase<
TileAndDistributeToWorkgroupsPass> {
+ using impl::TileAndDistributeToWorkgroupsPassBase<
+ TileAndDistributeToWorkgroupsPass>::TileAndDistributeToWorkgroupsPassBase;
+
TileAndDistributeToWorkgroupsPass(
int32_t maxWorkgroupParallelDims,
linalg::DistributionMethod distributionMethod) {
this->maxWorkgroupParallelDims = maxWorkgroupParallelDims;
- this->distributionMethod = (int32_t)distributionMethod;
+ this->distributionMethod = distributionMethod;
}
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<affine::AffineDialect, IREE::Flow::FlowDialect,
diff --git a/compiler/src/iree/compiler/Codegen/Common/TransformDialectInterpreterPass.cpp b/compiler/src/iree/compiler/Codegen/Common/TransformDialectInterpreterPass.cpp
index a28c8d6..c528b6d 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TransformDialectInterpreterPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TransformDialectInterpreterPass.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Codegen/Dialect/Codegen/IR/IREECodegenDialect.h"
#include "mlir/Dialect/Transform/IR/TransformDialect.h"
@@ -13,18 +12,26 @@
using namespace mlir;
+namespace mlir::iree_compiler {
+
+#define GEN_PASS_DEF_TRANSFORMDIALECTINTERPRETERPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
/// Pass declaration.
/// Interpreter pass that applies transform dialect ops for codegen.
/// This needs to be its own pass because the registration mechanism and ops
/// available are different than for other interpreters.
-class TransformDialectInterpreterPass
- : public iree_compiler::TransformDialectInterpreterBase<
+class TransformDialectInterpreterPass final
+ : public impl::TransformDialectInterpreterPassBase<
TransformDialectInterpreterPass> {
public:
- TransformDialectInterpreterPass(StringRef libraryFileName = StringRef(),
- StringRef entryPoint = StringRef()) {
+ using impl::TransformDialectInterpreterPassBase<
+ TransformDialectInterpreterPass>::TransformDialectInterpreterPassBase;
+
+ TransformDialectInterpreterPass(StringRef libraryFileName,
+ StringRef entryPoint) {
this->libraryFileName = libraryFileName.str();
this->entryPoint = entryPoint.str();
}
@@ -68,6 +75,7 @@
}
};
} // namespace
+} // namespace mlir::iree_compiler
namespace mlir::iree_compiler {
diff --git a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
index 9edb2d9..f3fba90 100644
--- a/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
@@ -24,7 +24,6 @@
//
//===---------------------------------------------------------------------===//
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtDialect.h"
#include "iree/compiler/Dialect/LinalgExt/IR/LinalgExtOps.h"
@@ -40,6 +39,9 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_TYPEPROPAGATIONPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
/// Insert instructions to convert from one element type to another.
static Value convertElementType(OpBuilder &b, Location loc, Type targetType,
Value source) {
@@ -568,8 +570,8 @@
}
};
-struct TypePropagationPass : public TypePropagationBase<TypePropagationPass> {
- TypePropagationPass() = default;
+struct TypePropagationPass final
+ : impl::TypePropagationPassBase<TypePropagationPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<arith::ArithDialect>();
}
@@ -627,10 +629,4 @@
}
};
} // namespace
-
-std::unique_ptr<InterfacePass<mlir::FunctionOpInterface>>
-createTypePropagationPass() {
- return std::make_unique<TypePropagationPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/VectorizeMemrefCopy.cpp b/compiler/src/iree/compiler/Codegen/Common/VectorizeMemrefCopy.cpp
index c6517cd..4265919 100644
--- a/compiler/src/iree/compiler/Codegen/Common/VectorizeMemrefCopy.cpp
+++ b/compiler/src/iree/compiler/Codegen/Common/VectorizeMemrefCopy.cpp
@@ -4,7 +4,6 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Codegen/Common/PassDetail.h"
#include "iree/compiler/Codegen/Common/Passes.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
@@ -13,10 +12,13 @@
namespace mlir::iree_compiler {
+#define GEN_PASS_DEF_VECTORIZEMEMREFCOPYPASS
+#include "iree/compiler/Codegen/Common/Passes.h.inc"
+
namespace {
-struct VectorizeMemrefCopyPass
- : public VectorizeMemrefCopyPassBase<VectorizeMemrefCopyPass> {
+struct VectorizeMemrefCopyPass final
+ : impl::VectorizeMemrefCopyPassBase<VectorizeMemrefCopyPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<arith::ArithDialect, vector::VectorDialect>();
}
@@ -31,9 +33,4 @@
};
} // namespace
-
-std::unique_ptr<OperationPass<void>> createVectorizeMemrefCopyPass() {
- return std::make_unique<VectorizeMemrefCopyPass>();
-}
-
} // namespace mlir::iree_compiler
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/batch_matmuls.mlir b/compiler/src/iree/compiler/Codegen/Common/test/batch_matmuls.mlir
index 2d91350..52a9c45 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/batch_matmuls.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/batch_matmuls.mlir
@@ -1,6 +1,5 @@
// RUN: iree-opt %s \
-// RUN: --iree-codegen-transform-dialect-library=%p/batch_matmul_match_spec.mlir \
-// RUN: --iree-transform-dialect-interpreter \
+// RUN: --iree-transform-dialect-interpreter="library-file-name=%p/batch_matmul_match_spec.mlir" \
// RUN: --split-input-file --verify-diagnostics
!lhs = tensor<128x80x32xf32>
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/convert_bf16_to_uint16_buffers.mlir b/compiler/src/iree/compiler/Codegen/Common/test/convert_bf16_to_uint16_buffers.mlir
index 399d509..1f99752 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/convert_bf16_to_uint16_buffers.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/convert_bf16_to_uint16_buffers.mlir
@@ -1,5 +1,5 @@
// RUN: iree-opt --split-input-file \
-// RUN: --iree-convert-bf16-to-uint16-buffers %s | FileCheck %s
+// RUN: --iree-codegen-convert-bf16-to-uint16-buffers %s | FileCheck %s
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
#hal.descriptor_set.layout<0, bindings = [
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/convolutions.mlir b/compiler/src/iree/compiler/Codegen/Common/test/convolutions.mlir
index be4bb2f..86630c3 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/convolutions.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/convolutions.mlir
@@ -1,6 +1,5 @@
// RUN: iree-opt %s \
-// RUN: --iree-codegen-transform-dialect-library=%p/convolution_match_spec.mlir \
-// RUN: --iree-transform-dialect-interpreter \
+// RUN: --iree-transform-dialect-interpreter="library-file-name=%p/convolution_match_spec.mlir" \
// RUN: --split-input-file --verify-diagnostics
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/decompose_affine_ops.mlir b/compiler/src/iree/compiler/Codegen/Common/test/decompose_affine_ops.mlir
index f0bb6c9..92567c2 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/decompose_affine_ops.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/decompose_affine_ops.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -decompose-affine-ops %s | FileCheck %s
+// RUN: iree-opt -iree-codegen-decompose-affine-ops %s | FileCheck %s
// Check that we have one affine.apply per loop dependence:
// IV0 with STRIDE0, IV1 with STRID1, etc.
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/extract_address_computation.mlir b/compiler/src/iree/compiler/Codegen/Common/test/extract_address_computation.mlir
index 399a067..8bef6c3 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/extract_address_computation.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/extract_address_computation.mlir
@@ -1,6 +1,6 @@
-// RUN: iree-opt -extract-address-computation %s --split-input-file | FileCheck %s
-// RUN: iree-opt -extract-address-computation -expand-strided-metadata \
-// RUN: -loop-invariant-code-motion -decompose-affine-ops -loop-invariant-code-motion \
+// RUN: iree-opt -iree-codegen-extract-address-computation %s --split-input-file | FileCheck %s
+// RUN: iree-opt -iree-codegen-extract-address-computation -expand-strided-metadata \
+// RUN: -loop-invariant-code-motion -iree-codegen-decompose-affine-ops -loop-invariant-code-motion \
// RUN: %s --split-input-file | FileCheck --check-prefix=INTEGRATION %s
// Simple test: check that we extract the address computation of a load into
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/hoist_statically_bound_allocations.mlir b/compiler/src/iree/compiler/Codegen/Common/test/hoist_statically_bound_allocations.mlir
index 1bcc5c0..56ad104 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/hoist_statically_bound_allocations.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/hoist_statically_bound_allocations.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt --split-input-file --pass-pipeline="builtin.module(func.func(iree-hoist-statically-bound-allocations{vscale-min=1 vscale-max=16}))" %s | FileCheck %s
-// RUN: iree-opt --split-input-file --pass-pipeline="builtin.module(func.func(iree-hoist-statically-bound-allocations{vscale-max=0}))" %s | FileCheck %s --check-prefix=CHECK-UNBOUNDED-VSCALE
+// RUN: iree-opt --split-input-file --pass-pipeline="builtin.module(func.func(iree-codegen-hoist-statically-bound-allocations{vscale-min=1 vscale-max=16}))" %s | FileCheck %s
+// RUN: iree-opt --split-input-file --pass-pipeline="builtin.module(func.func(iree-codegen-hoist-statically-bound-allocations{vscale-max=0}))" %s | FileCheck %s --check-prefix=CHECK-UNBOUNDED-VSCALE
// Note: Scalable allocations are not hoisted if vscale is unbounded.
diff --git a/compiler/src/iree/compiler/Codegen/Common/test/reductions.mlir b/compiler/src/iree/compiler/Codegen/Common/test/reductions.mlir
index aeefd04..c4d2d7a 100644
--- a/compiler/src/iree/compiler/Codegen/Common/test/reductions.mlir
+++ b/compiler/src/iree/compiler/Codegen/Common/test/reductions.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt %s --iree-codegen-transform-dialect-library=%p/reductions_codegen_spec.mlir --iree-transform-dialect-interpreter --split-input-file | FileCheck %s
-// RUN: iree-opt %s --iree-codegen-transform-dialect-library=%p/reductions_match_spec.mlir --iree-transform-dialect-interpreter --split-input-file --verify-diagnostics
+// RUN: iree-opt %s --iree-transform-dialect-interpreter="library-file-name=%p/reductions_codegen_spec.mlir" --split-input-file | FileCheck %s
+// RUN: iree-opt %s --iree-transform-dialect-interpreter="library-file-name=%p/reductions_match_spec.mlir" --split-input-file --verify-diagnostics
// Check that the same transform script applies to reductions with optional
// leading and trailing elementwise operations, potentially reordered
diff --git a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
index 31accfc..aeb9d7f 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMCPU/Passes.cpp
@@ -401,7 +401,7 @@
}
{
- funcPassManager.addPass(createVectorizePadPass());
+ funcPassManager.addPass(createTensorToVectorVectorizePadPass());
if (pipelineOpt.decomposePackUnPackOps) {
funcPassManager.addPass(createDecomposePackUnPackOpsPass());
funcPassManager.addPass(createCanonicalizerPass());
@@ -465,7 +465,7 @@
}
{
- funcPassManager.addPass(createVectorizePadPass());
+ funcPassManager.addPass(createTensorToVectorVectorizePadPass());
GenericVectorizationPassOptions options;
options.useConfiguredVectorSizes = pipelineOpt.useConfiguredVectorSizes;
options.enableVectorMasking = pipelineOpt.enableVectorMasking;
@@ -481,7 +481,8 @@
}
// Eliminate redundant transfer_read/write to avoid stack allocations.
- funcPassManager.addPass(createOptimizeVectorTransferPass(/*flatten=*/true));
+ funcPassManager.addPass(createOptimizeVectorTransferPass(
+ OptimizeVectorTransferPassOptions{/*flatten=*/true}));
addCPUBufferizePasses(funcPassManager);
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
index 4517eb2..8b74d1b 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/Passes.cpp
@@ -341,7 +341,8 @@
// Normalize loop bounds for later lowerings.
funcPassManager.addPass(iree_compiler::createNormalizeLoopBoundsPass(
- /*normalizeFor=*/false, /*normalizeForall=*/true));
+ NormalizeLoopBoundsPassOptions{/*normalizeFor=*/false,
+ /*normalizeForall=*/true}));
funcPassManager.addPass(createCanonicalizerPass());
funcPassManager.addPass(createCSEPass());
funcPassManager.addPass(createLoopInvariantCodeMotionPass());
@@ -973,7 +974,7 @@
// Pad allocations with dynamic dimension after linalg lowering but before
// lowering SCF and affine ops.
- .addPass(createPadDynamicAlloc)
+ .addPass(createPadDynamicAllocPass)
.addPass(createLowerAffinePass)
.addPass(createCanonicalizerPass)
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention.mlir
index 353ec9e..e029174 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt %s --pass-pipeline='builtin.module(iree-transform-dialect-interpreter)' \
-// RUN: --iree-gpu-test-target=sm_60 --iree-codegen-transform-dialect-library=%p/attention_transform_spec.mlir| \
+// RUN: iree-opt %s --pass-pipeline='builtin.module(iree-transform-dialect-interpreter{library-file-name=%p/attention_transform_spec.mlir})' \
+// RUN: --iree-gpu-test-target=sm_60 | \
// RUN: FileCheck --check-prefix=CHECK %s
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
diff --git a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention_mfma.mlir b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention_mfma.mlir
index db40536..7e69786 100644
--- a/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention_mfma.mlir
+++ b/compiler/src/iree/compiler/Codegen/LLVMGPU/test/attention_mfma.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt %s --pass-pipeline='builtin.module(iree-transform-dialect-interpreter)' \
-// RUN: --iree-gpu-test-target=gfx908 --iree-codegen-transform-dialect-library=%p/attention_mfma_transform_spec.mlir | \
+// RUN: iree-opt %s --pass-pipeline='builtin.module(iree-transform-dialect-interpreter{library-file-name=%p/attention_mfma_transform_spec.mlir})' \
+// RUN: --iree-gpu-test-target=gfx908 | \
// RUN: FileCheck --check-prefix=CHECK %s
#pipeline_layout = #hal.pipeline.layout<push_constants = 0, sets = [
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 91536b7..da19ad2 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
@@ -1,10 +1,8 @@
-// RUN: iree-opt %s --pass-pipeline="builtin.module(iree-transform-dialect-interpreter)" \
-// RUN: --iree-codegen-transform-dialect-library=%p/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir \
+// RUN: iree-opt %s --pass-pipeline="builtin.module(iree-transform-dialect-interpreter{library-file-name=%p/transform_dialect_codegen_vector_warp_execute_on_lane_0_spec.mlir})" \
// RUN: --allow-unregistered-dialect | \
// RUN: FileCheck %s --check-prefix=WARP-EXECUTE
-// RUN: iree-opt %s --pass-pipeline="builtin.module(iree-transform-dialect-interpreter)" \
-// RUN: --iree-codegen-transform-dialect-library=%p/transform_dialect_codegen_vector_distribution_spec.mlir \
+// RUN: iree-opt %s --pass-pipeline="builtin.module(iree-transform-dialect-interpreter{library-file-name=%p/transform_dialect_codegen_vector_distribution_spec.mlir})" \
// RUN: --allow-unregistered-dialect | \
// RUN: FileCheck %s
diff --git a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp
index f8093dd..a9d5086 100644
--- a/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp
+++ b/compiler/src/iree/compiler/Codegen/SPIRV/Passes.cpp
@@ -181,7 +181,7 @@
// Math dialect elementry functions -> polynomial form.
.addPass(createPolynomialApproximationPass)
- .addPass(createPadDynamicAlloc);
+ .addPass(createPadDynamicAllocPass);
// TODO: query this from the target.
auto getIndexBitwidth = [](mlir::FunctionOpInterface) { return 32; };
