llvm · chelini · May 13, 2025 · May 10, 2025
diff --git 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<Pass> 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
@@ -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
@@ -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<Pass> createDuplicateFunctionEliminationPass();
-
 //===----------------------------------------------------------------------===//
 // Registration
 //===----------------------------------------------------------------------===//

diff --git 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
@@ -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<Pass> 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
@@ -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<Pass> 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
@@ -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
@@ -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<Pass> 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
@@ -19,7 +19,8 @@ namespace tensor {
 //===----------------------------------------------------------------------===//
 
 /// Creates an instance of the `tensor` subset folding pass.
-std::unique_ptr<Pass> 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
@@ -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
@@ -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<CmpIOp> {
 };
 
 struct ArithUnsignedWhenEquivalentPass
-    : public arith::impl::ArithUnsignedWhenEquivalentBase<
+    : public arith::impl::ArithUnsignedWhenEquivalentPassBase<
           ArithUnsignedWhenEquivalentPass> {
 
   void runOnOperation() override {
@@ -151,7 +151,3 @@ void mlir::arith::populateUnsignedWhenEquivalentPatterns(
                ConvertOpToUnsigned<ExtSIOp, ExtUIOp>, ConvertCmpIToUnsigned>(
       patterns.getContext(), solver);
 }
-
-std::unique_ptr<Pass> mlir::arith::createArithUnsignedWhenEquivalentPass() {
-  return std::make_unique<ArithUnsignedWhenEquivalentPass>();
-}
diff --git 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<Pass> mlir::func::createDuplicateFunctionEliminationPass() {
-  return std::make_unique<DuplicateFunctionEliminationPass>();
-}
-
 } // namespace mlir
diff --git 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<LegalizeForExportPass> {
+    : public LLVM::impl::LLVMLegalizeForExportPassBase<LegalizeForExportPass> {
   void runOnOperation() override {
     LLVM::ensureDistinctSuccessors(getOperation());
     LLVM::legalizeDIExpressionsRecursively(getOperation());
   }
 };
 } // namespace
-
-std::unique_ptr<Pass> LLVM::createLegalizeForExportPass() {
-  return std::make_unique<LegalizeForExportPass>();
-}
diff --git 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<LLVM::FDivOp> {
 };
 
 struct NVVMOptimizeForTarget
-    : public NVVM::impl::NVVMOptimizeForTargetBase<NVVMOptimizeForTarget> {
+    : public LLVM::impl::NVVMOptimizeForTargetPassBase<NVVMOptimizeForTarget> {
   void runOnOperation() override;
 
   void getDependentDialects(DialectRegistry &registry) const override {
@@ -99,7 +99,3 @@ void NVVMOptimizeForTarget::runOnOperation() {
   if (failed(applyPatternsGreedily(getOperation(), std::move(patterns))))
     return signalPassFailure();
 }
-
-std::unique_ptr<Pass> NVVM::createOptimizeForTargetPass() {
-  return std::make_unique<NVVMOptimizeForTarget>();
-}
diff --git 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,15 +22,11 @@ using namespace mlir;
 
 namespace {
 class RequestCWrappersPass
-    : public LLVM::impl::LLVMRequestCWrappersBase<RequestCWrappersPass> {
+    : public LLVM::impl::LLVMRequestCWrappersPassBase<RequestCWrappersPass> {
 public:
   void runOnOperation() override {
     getOperation()->setAttr(LLVM::LLVMDialect::getEmitCWrapperAttrName(),
                             UnitAttr::get(&getContext()));
   }
 };
 } // namespace
-
-std::unique_ptr<Pass> mlir::LLVM::createRequestCWrappersPass() {
-  return std::make_unique<RequestCWrappersPass>();
-}
@@ -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<FoldTensorSubsetOpsPass> {
+    : 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<Pass> tensor::createFoldTensorSubsetOpsPass() {
-  return std::make_unique<FoldTensorSubsetOpsPass>();
-}
diff --git 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<func::FuncOp>().addPass(LLVM::createRequestCWrappersPass());
+  passManager.nest<func::FuncOp>().addPass(
+      LLVM::createLLVMRequestCWrappersPass());
   // VulkanRuntimeWrappers.cpp requires these calling convention options.
   GpuToLLVMConversionPassOptions opt;
   opt.hostBarePtrCallConv = false;