diff --git a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h index 58dce89fdb578..8d81d8ec14ee7 100644 --- a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.h @@ -62,10 +62,6 @@ void populateExpandBFloat16Patterns(RewritePatternSet &patterns); /// Add patterns to expand Arith ops. void populateArithExpandOpsPatterns(RewritePatternSet &patterns); -/// Create a pass to replace signed ops with unsigned ones where they are proven -/// equivalent. -std::unique_ptr createArithUnsignedWhenEquivalentPass(); - /// Add patterns for int range based optimizations. void populateIntRangeOptimizationsPatterns(RewritePatternSet &patterns, DataFlowSolver &solver); diff --git a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.td b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.td index 1d37314885d93..d026d494cb50c 100644 --- a/mlir/include/mlir/Dialect/Arith/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/Arith/Transforms/Passes.td @@ -20,7 +20,7 @@ def ArithExpandOpsPass : Pass<"arith-expand"> { ]; } -def ArithUnsignedWhenEquivalent : Pass<"arith-unsigned-when-equivalent"> { +def ArithUnsignedWhenEquivalentPass : Pass<"arith-unsigned-when-equivalent"> { let summary = "Replace signed ops with unsigned ones where they are proven equivalent"; let description = [{ Replace signed ops with their unsigned equivalents when integer range analysis @@ -33,7 +33,6 @@ def ArithUnsignedWhenEquivalent : Pass<"arith-unsigned-when-equivalent"> { The affect ops include division, remainder, shifts, min, max, and integer comparisons. }]; - let constructor = "mlir::arith::createArithUnsignedWhenEquivalentPass()"; } def ArithIntRangeOpts : Pass<"int-range-optimizations"> { diff --git a/mlir/include/mlir/Dialect/Func/Transforms/Passes.h b/mlir/include/mlir/Dialect/Func/Transforms/Passes.h index 0248f068320c5..6fe9cc4bb2986 100644 --- a/mlir/include/mlir/Dialect/Func/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Func/Transforms/Passes.h @@ -22,12 +22,9 @@ class RewritePatternSet; namespace func { -#define GEN_PASS_DECL +#define GEN_PASS_DECL_DUPLICATEFUNCTIONELIMINATIONPASS #include "mlir/Dialect/Func/Transforms/Passes.h.inc" -/// Pass to deduplicate functions. -std::unique_ptr createDuplicateFunctionEliminationPass(); - //===----------------------------------------------------------------------===// // Registration //===----------------------------------------------------------------------===// diff --git a/mlir/include/mlir/Dialect/Func/Transforms/Passes.td b/mlir/include/mlir/Dialect/Func/Transforms/Passes.td index c3caf120d052e..4163997515bb0 100644 --- a/mlir/include/mlir/Dialect/Func/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/Func/Transforms/Passes.td @@ -19,7 +19,6 @@ def DuplicateFunctionEliminationPass : Pass<"duplicate-function-elimination", name. The pass chooses one representative per equivalence class, erases the remainder, and updates function calls accordingly. }]; - let constructor = "mlir::func::createDuplicateFunctionEliminationPass()"; } #endif // MLIR_DIALECT_FUNC_TRANSFORMS_PASSES_TD diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h index 09f62e1190f29..9b3754f582038 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/LegalizeForExport.h @@ -17,7 +17,7 @@ class Pass; namespace LLVM { -#define GEN_PASS_DECL_LLVMLEGALIZEFOREXPORT +#define GEN_PASS_DECL_LLVMLEGALIZEFOREXPORTPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" /// Make argument-taking successors of each block distinct. PHI nodes in LLVM @@ -27,10 +27,6 @@ namespace LLVM { /// a new dummy block for LLVM PHI nodes to tell the sources apart. void ensureDistinctSuccessors(Operation *op); -/// Creates a pass that legalizes the LLVM dialect operations so that they can -/// be translated to LLVM IR. -std::unique_ptr createLegalizeForExportPass(); - } // namespace LLVM } // namespace mlir diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h index 261d2a2389c16..3f625a4e484b3 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/OptimizeForNVVM.h @@ -13,16 +13,10 @@ namespace mlir { class Pass; - -namespace NVVM { - -#define GEN_PASS_DECL_NVVMOPTIMIZEFORTARGET +namespace LLVM { +#define GEN_PASS_DECL_NVVMOPTIMIZEFORTARGETPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" - -/// Creates a pass that optimizes LLVM IR for the NVVM target. -std::unique_ptr createOptimizeForTargetPass(); - -} // namespace NVVM +} // namespace LLVM } // namespace mlir #endif // MLIR_DIALECT_LLVMIR_TRANSFORMS_OPTIMIZENVVM_H diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td index b2a03ce61e7e4..961909d5c8d27 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/Passes.td @@ -24,13 +24,16 @@ def LLVMAddComdats : Pass<"llvm-add-comdats", "::mlir::ModuleOp"> { }]; } -def LLVMLegalizeForExport : Pass<"llvm-legalize-for-export"> { +def LLVMLegalizeForExportPass : Pass<"llvm-legalize-for-export"> { let summary = "Legalize LLVM dialect to be convertible to LLVM IR"; - let constructor = "::mlir::LLVM::createLegalizeForExportPass()"; + let description = [{ + Creates a pass that legalizes the LLVM dialect operations so that they can + be translated to LLVM IR. + }]; let dependentDialects = ["LLVM::LLVMDialect"]; } -def LLVMRequestCWrappers +def LLVMRequestCWrappersPass : Pass<"llvm-request-c-wrappers", "::mlir::func::FuncOp"> { let summary = "Request C wrapper emission for all functions"; let description = [{ @@ -40,12 +43,10 @@ def LLVMRequestCWrappers conversion of builtin functions to LLVM to avoid the attribute being dropped by other passes. }]; - let constructor = "::mlir::LLVM::createRequestCWrappersPass()"; } -def NVVMOptimizeForTarget : Pass<"llvm-optimize-for-nvvm-target"> { +def NVVMOptimizeForTargetPass : Pass<"llvm-optimize-for-nvvm-target"> { let summary = "Optimize NVVM IR"; - let constructor = "::mlir::NVVM::createOptimizeForTargetPass()"; } def DIScopeForLLVMFuncOpPass : Pass<"ensure-debug-info-scope-on-llvm-func", "::mlir::ModuleOp"> { diff --git a/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h index c6fd2642e307d..0becb4d6b785c 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h +++ b/mlir/include/mlir/Dialect/LLVMIR/Transforms/RequestCWrappers.h @@ -16,10 +16,9 @@ class Pass; namespace LLVM { -#define GEN_PASS_DECL_LLVMREQUESTCWRAPPERS +#define GEN_PASS_DECL_LLVMREQUESTCWRAPPERSPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" -std::unique_ptr createRequestCWrappersPass(); } // namespace LLVM } // namespace mlir diff --git a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h index 964c35b3f15b8..03cb30bc07357 100644 --- a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h +++ b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.h @@ -19,7 +19,8 @@ namespace tensor { //===----------------------------------------------------------------------===// /// Creates an instance of the `tensor` subset folding pass. -std::unique_ptr createFoldTensorSubsetOpsPass(); +#define GEN_PASS_DECL_FOLDTENSORSUBSETOPSPASS +#include "mlir/Dialect/Tensor/Transforms/Passes.h.inc" //===----------------------------------------------------------------------===// // Registration diff --git a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.td b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.td index be4c333836ec0..62f467f2513ac 100644 --- a/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.td +++ b/mlir/include/mlir/Dialect/Tensor/Transforms/Passes.td @@ -11,7 +11,7 @@ include "mlir/Pass/PassBase.td" -def FoldTensorSubsetOps : Pass<"fold-tensor-subset-ops"> { +def FoldTensorSubsetOpsPass : Pass<"fold-tensor-subset-ops"> { let summary = "Fold tensor subset ops into producer/consumer ops"; let description = [{ The pass folds tensor subset ops into producer/consumer ops. @@ -21,7 +21,6 @@ def FoldTensorSubsetOps : Pass<"fold-tensor-subset-ops"> { - vector.transfer_write into tensor.insert_slice }]; - let constructor = "mlir::tensor::createFoldTensorSubsetOpsPass()"; let dependentDialects = [ "affine::AffineDialect", "tensor::TensorDialect", "vector::VectorDialect" ]; diff --git a/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp b/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp index dabfffda390bb..bd1eac16070eb 100644 --- a/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp +++ b/mlir/lib/Dialect/Arith/Transforms/UnsignedWhenEquivalent.cpp @@ -18,7 +18,7 @@ namespace mlir { namespace arith { -#define GEN_PASS_DEF_ARITHUNSIGNEDWHENEQUIVALENT +#define GEN_PASS_DEF_ARITHUNSIGNEDWHENEQUIVALENTPASS #include "mlir/Dialect/Arith/Transforms/Passes.h.inc" } // namespace arith } // namespace mlir @@ -118,7 +118,7 @@ struct ConvertCmpIToUnsigned final : OpRewritePattern { }; struct ArithUnsignedWhenEquivalentPass - : public arith::impl::ArithUnsignedWhenEquivalentBase< + : public arith::impl::ArithUnsignedWhenEquivalentPassBase< ArithUnsignedWhenEquivalentPass> { void runOnOperation() override { @@ -151,7 +151,3 @@ void mlir::arith::populateUnsignedWhenEquivalentPatterns( ConvertOpToUnsigned, ConvertCmpIToUnsigned>( patterns.getContext(), solver); } - -std::unique_ptr mlir::arith::createArithUnsignedWhenEquivalentPass() { - return std::make_unique(); -} diff --git a/mlir/lib/Dialect/Func/Transforms/DuplicateFunctionElimination.cpp b/mlir/lib/Dialect/Func/Transforms/DuplicateFunctionElimination.cpp index fbb6abfd65b10..ea85c20cd1f8f 100644 --- a/mlir/lib/Dialect/Func/Transforms/DuplicateFunctionElimination.cpp +++ b/mlir/lib/Dialect/Func/Transforms/DuplicateFunctionElimination.cpp @@ -10,10 +10,12 @@ #include "mlir/Dialect/Func/Transforms/Passes.h" namespace mlir { -namespace { - +namespace func { #define GEN_PASS_DEF_DUPLICATEFUNCTIONELIMINATIONPASS #include "mlir/Dialect/Func/Transforms/Passes.h.inc" +} // namespace func + +namespace { // Define a notion of function equivalence that allows for reuse. Ignore the // symbol name for this purpose. @@ -80,7 +82,7 @@ struct DuplicateFuncOpEquivalenceInfo }; struct DuplicateFunctionEliminationPass - : public impl::DuplicateFunctionEliminationPassBase< + : public func::impl::DuplicateFunctionEliminationPassBase< DuplicateFunctionEliminationPass> { using DuplicateFunctionEliminationPassBase< @@ -115,9 +117,4 @@ struct DuplicateFunctionEliminationPass }; } // namespace - -std::unique_ptr mlir::func::createDuplicateFunctionEliminationPass() { - return std::make_unique(); -} - } // namespace mlir diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp index 1ac994fa5fb78..8dd0c28d98522 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/LegalizeForExport.cpp @@ -17,7 +17,7 @@ namespace mlir { namespace LLVM { -#define GEN_PASS_DEF_LLVMLEGALIZEFOREXPORT +#define GEN_PASS_DEF_LLVMLEGALIZEFOREXPORTPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" } // namespace LLVM } // namespace mlir @@ -77,14 +77,10 @@ void mlir::LLVM::ensureDistinctSuccessors(Operation *op) { namespace { struct LegalizeForExportPass - : public LLVM::impl::LLVMLegalizeForExportBase { + : public LLVM::impl::LLVMLegalizeForExportPassBase { void runOnOperation() override { LLVM::ensureDistinctSuccessors(getOperation()); LLVM::legalizeDIExpressionsRecursively(getOperation()); } }; } // namespace - -std::unique_ptr LLVM::createLegalizeForExportPass() { - return std::make_unique(); -} diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp index c1ec1df48e5b9..8db32ec1526c4 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/OptimizeForNVVM.cpp @@ -15,10 +15,10 @@ #include "mlir/Transforms/GreedyPatternRewriteDriver.h" namespace mlir { -namespace NVVM { -#define GEN_PASS_DEF_NVVMOPTIMIZEFORTARGET +namespace LLVM { +#define GEN_PASS_DEF_NVVMOPTIMIZEFORTARGETPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" -} // namespace NVVM +} // namespace LLVM } // namespace mlir using namespace mlir; @@ -40,7 +40,7 @@ struct ExpandDivF16 : public OpRewritePattern { }; struct NVVMOptimizeForTarget - : public NVVM::impl::NVVMOptimizeForTargetBase { + : public LLVM::impl::NVVMOptimizeForTargetPassBase { void runOnOperation() override; void getDependentDialects(DialectRegistry ®istry) const override { @@ -99,7 +99,3 @@ void NVVMOptimizeForTarget::runOnOperation() { if (failed(applyPatternsGreedily(getOperation(), std::move(patterns)))) return signalPassFailure(); } - -std::unique_ptr NVVM::createOptimizeForTargetPass() { - return std::make_unique(); -} diff --git a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp index 099d9c26e5580..a5fccbdc368fe 100644 --- a/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp +++ b/mlir/lib/Dialect/LLVMIR/Transforms/RequestCWrappers.cpp @@ -13,7 +13,7 @@ namespace mlir { namespace LLVM { -#define GEN_PASS_DEF_LLVMREQUESTCWRAPPERS +#define GEN_PASS_DEF_LLVMREQUESTCWRAPPERSPASS #include "mlir/Dialect/LLVMIR/Transforms/Passes.h.inc" } // namespace LLVM } // namespace mlir @@ -22,7 +22,7 @@ using namespace mlir; namespace { class RequestCWrappersPass - : public LLVM::impl::LLVMRequestCWrappersBase { + : public LLVM::impl::LLVMRequestCWrappersPassBase { public: void runOnOperation() override { getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(), @@ -30,7 +30,3 @@ class RequestCWrappersPass } }; } // namespace - -std::unique_ptr mlir::LLVM::createRequestCWrappersPass() { - return std::make_unique(); -} diff --git a/mlir/lib/Dialect/Tensor/Transforms/FoldTensorSubsetOps.cpp b/mlir/lib/Dialect/Tensor/Transforms/FoldTensorSubsetOps.cpp index 998b0fb6eb4b7..f7a58d8d71938 100644 --- a/mlir/lib/Dialect/Tensor/Transforms/FoldTensorSubsetOps.cpp +++ b/mlir/lib/Dialect/Tensor/Transforms/FoldTensorSubsetOps.cpp @@ -28,7 +28,7 @@ namespace mlir { namespace tensor { -#define GEN_PASS_DEF_FOLDTENSORSUBSETOPS +#define GEN_PASS_DEF_FOLDTENSORSUBSETOPSPASS #include "mlir/Dialect/Tensor/Transforms/Passes.h.inc" } // namespace tensor } // namespace mlir @@ -268,7 +268,8 @@ void tensor::populateFoldTensorSubsetIntoVectorTransferPatterns( namespace { struct FoldTensorSubsetOpsPass final - : public tensor::impl::FoldTensorSubsetOpsBase { + : public tensor::impl::FoldTensorSubsetOpsPassBase< + FoldTensorSubsetOpsPass> { void runOnOperation() override; }; @@ -279,7 +280,3 @@ void FoldTensorSubsetOpsPass::runOnOperation() { tensor::populateFoldTensorSubsetOpPatterns(patterns); (void)applyPatternsGreedily(getOperation(), std::move(patterns)); } - -std::unique_ptr tensor::createFoldTensorSubsetOpsPass() { - return std::make_unique(); -} diff --git a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp index 3b011d2b2b0ac..c3c871a539291 100644 --- a/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp +++ b/mlir/test/lib/Pass/TestVulkanRunnerPipeline.cpp @@ -65,7 +65,8 @@ void buildTestVulkanRunnerPipeline(OpPassManager &passManager, passManager.addPass(createGpuModuleToBinaryPass()); passManager.addPass(createFinalizeMemRefToLLVMConversionPass()); - passManager.nest().addPass(LLVM::createRequestCWrappersPass()); + passManager.nest().addPass( + LLVM::createLLVMRequestCWrappersPass()); // VulkanRuntimeWrappers.cpp requires these calling convention options. GpuToLLVMConversionPassOptions opt; opt.hostBarePtrCallConv = false;