[mlir] Generalize Affine transforms to work on FunctionOpInterface.

[stellaraccident]

Summary

Most transforms have assumptions that they are processing children as part of a FuncOp currently, so migrating these to FunctionOpInterface is natural and does not change semantics. While some passes, upon inspection *could* operate on another container, I just uniformly changed them all to FunctionOpInterface for consistency and to avoid hidden gotchas.

I was somewhat surprised that on the CMake side none of this had a dependency on the Func dialect: I guess that was coming in transitively.

This is nearly an NFC, however, the change to InterfacePass means that the passes only work with PassManagers that have been explicitly nested properly (OperationPass has some legacy fallback paths for implicit nesting that do not exist in the new paths). I expect that this mainly impacts *-opt invocations, which must switch to a -pass-pipeline that is nested properly. However, it is possible that C++ code is also incorrect, in which case it will begin asserting and nesting must be done explicitly.

[stellaraccident]

Fix fir test.

[rriddle]

Nothing sticks out to me here as wrong/broken, but given this is affine stuff I'll defer to @bondhugula for approval.

[jpienaar]

Same as River (and thanks for generalizing)

flang/test/Fir/affine-promotion.fir

 // Test affine promotion pass
 
-// RUN: fir-opt --split-input-file --promote-to-affine --affine-loop-invariant-code-motion --cse %s | FileCheck %s
+// RUN: fir-opt --split-input-file -pass-pipeline="promote-to-affine,func.func(affine-loop-invariant-code-motion),cse" %s | FileCheck %s
 
 !arr_d1 = !fir.ref<!fir.array<?xf32>>
 #arr_len = affine_map<()[j1,k1] -> (k1 - j1 + 1)>
 
 func.func @loop_with_load_and_store(%a1: !arr_d1, %a2: !arr_d1, %a3: !arr_d1) {
   %c1 = arith.constant 1 : index
   %c0 = arith.constant 0 : index
   %len = arith.constant 100 : index

mlir/include/mlir/Dialect/Affine/LoopUtils.h

 
 #include "mlir/IR/Block.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Support/LogicalResult.h"
 
 namespace mlir {
 class AffineForOp;
 class AffineMap;
+class FunctionOpInterface;
 class LoopLikeOpInterface;
 struct MemRefRegion;
 class OpBuilder;
 class Value;
 class ValueRange;
 
-namespace func {
-class FuncOp;
-} // namespace func
-
 namespace scf {
 class ForOp;
 class ParallelOp;
 } // namespace scf
 
 /// Unrolls this for operation completely if the trip count is known to be
 /// constant. Returns failure otherwise.
 LogicalResult loopUnrollFull(AffineForOp forOp);
LogicalResult loopUnrollJamUpToFactor(AffineForOp forOp,
                                       uint64_t unrollJamFactor);
 
 /// Promotes the loop body of a AffineForOp to its containing block if the loop
 /// was known to have a single iteration.
 LogicalResult promoteIfSingleIteration(AffineForOp forOp);
 
 /// Promotes all single iteration AffineForOp's in the Function, i.e., moves
 /// their body into the containing Block.
-void promoteSingleIterationLoops(func::FuncOp f);
+void promoteSingleIterationLoops(FunctionOpInterface f);
 
 /// Skew the operations in an affine.for's body with the specified
 /// operation-wise shifts. The shifts are with respect to the original execution
 /// order, and are multiplied by the loop 'step' before being applied. If
 /// `unrollPrologueEpilogue` is set, fully unroll the prologue and epilogue
 /// loops when possible.
 LogicalResult affineForOpBodySkew(AffineForOp forOp, ArrayRef<uint64_t> shifts,
                                   bool unrollPrologueEpilogue = false);
 
 /// Identify valid and profitable bands of loops to tile. This is currently just
 /// a temporary placeholder to test the mechanics of tiled code generation.
 /// Returns all maximal outermost perfect loop nests to tile.
-void getTileableBands(func::FuncOp f,
+void getTileableBands(FunctionOpInterface f,
                       std::vector<SmallVector<AffineForOp, 6>> *bands);
 
 /// Tiles the specified band of perfectly nested loops creating tile-space loops
 /// and intra-tile loops. A band is a contiguous set of loops.
 LogicalResult
 tilePerfectlyNested(MutableArrayRef<AffineForOp> input,
                     ArrayRef<unsigned> tileSizes,
                     SmallVectorImpl<AffineForOp> *tiledNest = nullptr);

 ///       to %ub step %gridDim.x * blockDim.x * %step {
 ///      ...
 ///    }
 /// ```
 void mapLoopToProcessorIds(scf::ForOp forOp, ArrayRef<Value> processorId,
                            ArrayRef<Value> numProcessors);
 
 /// Gathers all AffineForOps in 'func.func' grouped by loop depth.
-void gatherLoops(func::FuncOp func,
+void gatherLoops(FunctionOpInterface func,
                  std::vector<SmallVector<AffineForOp, 2>> &depthToLoops);
 
 /// Creates an AffineForOp while ensuring that the lower and upper bounds are
 /// canonicalized, i.e., unused and duplicate operands are removed, any constant
 /// operands propagated/folded in, and duplicate bound maps dropped.
 AffineForOp createCanonicalizedAffineForOp(OpBuilder b, Location loc,
                                            ValueRange lbOperands,
                                            AffineMap lbMap,

mlir/include/mlir/Dialect/Affine/Passes.h

 
 #ifndef MLIR_DIALECT_AFFINE_PASSES_H
 #define MLIR_DIALECT_AFFINE_PASSES_H
 
 #include "mlir/Pass/Pass.h"
 #include <limits>
 
 namespace mlir {
-namespace func {
-class FuncOp;
-} // namespace func
-
 class AffineForOp;
+class FunctionOpInterface;
 
 /// Fusion mode to attempt. The default mode `Greedy` does both
 /// producer-consumer and sibling fusion.
 enum FusionMode { Greedy, ProducerConsumer, Sibling };
 
 /// Creates a simplification pass for affine structures (maps and sets). In
 /// addition, this pass also normalizes memrefs to have the trivial (identity)
 /// layout map.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createSimplifyAffineStructuresPass();
 
 /// Creates a loop invariant code motion pass that hoists loop invariant
 /// operations out of affine loops.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createAffineLoopInvariantCodeMotionPass();
 
 /// Creates a pass to convert all parallel affine.for's into 1-d affine.parallel
 /// ops.
-std::unique_ptr<OperationPass<func::FuncOp>> createAffineParallelizePass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+createAffineParallelizePass();
 
 /// Apply normalization transformations to affine loop-like ops.
-std::unique_ptr<OperationPass<func::FuncOp>> createAffineLoopNormalizePass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+createAffineLoopNormalizePass();
 
 /// Performs packing (or explicit copying) of accessed memref regions into
 /// buffers in the specified faster memory space through either pointwise copies
 /// or DMA operations.
-std::unique_ptr<OperationPass<func::FuncOp>> createAffineDataCopyGenerationPass(
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+createAffineDataCopyGenerationPass(
     unsigned slowMemorySpace, unsigned fastMemorySpace,
     unsigned tagMemorySpace = 0, int minDmaTransferSize = 1024,
     uint64_t fastMemCapacityBytes = std::numeric_limits<uint64_t>::max());
 /// Overload relying on pass options for initialization.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createAffineDataCopyGenerationPass();
 
 /// Creates a pass to replace affine memref accesses by scalars using store to
 /// load forwarding and redundant load elimination; consequently also eliminate
 /// dead allocs.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createAffineScalarReplacementPass();
 
 /// Creates a pass that transforms perfectly nested loops with independent
 /// bounds into a single loop.
-std::unique_ptr<OperationPass<func::FuncOp>> createLoopCoalescingPass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createLoopCoalescingPass();
 
 /// Creates a loop fusion pass which fuses loops according to type of fusion
 /// specified in `fusionMode`. Buffers of size less than or equal to
 /// `localBufSizeThreshold` are promoted to memory space `fastMemorySpace`.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createLoopFusionPass(unsigned fastMemorySpace = 0,
                      uint64_t localBufSizeThreshold = 0,
                      bool maximalFusion = false,
                      enum FusionMode fusionMode = FusionMode::Greedy);
 
 /// Creates a pass to perform tiling on loop nests.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createLoopTilingPass(uint64_t cacheSizeBytes);
 /// Overload relying on pass options for initialization.
-std::unique_ptr<OperationPass<func::FuncOp>> createLoopTilingPass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createLoopTilingPass();
 
 /// Creates a loop unrolling pass with the provided parameters.
 /// 'getUnrollFactor' is a function callback for clients to supply a function
 /// that computes an unroll factor - the callback takes precedence over unroll
 /// factors supplied through other means. If -1 is passed as the unrollFactor
 /// and no callback is provided, anything passed from the command-line (if at
 /// all) or the default unroll factor is used (LoopUnroll:kDefaultUnrollFactor).
-std::unique_ptr<OperationPass<func::FuncOp>> createLoopUnrollPass(
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createLoopUnrollPass(
     int unrollFactor = -1, bool unrollUpToFactor = false,
     bool unrollFull = false,
     const std::function<unsigned(AffineForOp)> &getUnrollFactor = nullptr);
 
 /// Creates a loop unroll jam pass to unroll jam by the specified factor. A
 /// factor of -1 lets the pass use the default factor or the one on the command
 /// line if provided.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createLoopUnrollAndJamPass(int unrollJamFactor = -1);
 
 /// Creates a pass to pipeline explicit movement of data across levels of the
 /// memory hierarchy.
-std::unique_ptr<OperationPass<func::FuncOp>> createPipelineDataTransferPass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+createPipelineDataTransferPass();
 
 /// Creates a pass to vectorize loops, operations and data types using a
 /// target-independent, n-D super-vector abstraction.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createSuperVectorizePass(ArrayRef<int64_t> virtualVectorSize);
 /// Overload relying on pass options for initialization.
-std::unique_ptr<OperationPass<func::FuncOp>> createSuperVectorizePass();
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createSuperVectorizePass();
 
 //===----------------------------------------------------------------------===//
 // Registration
 //===----------------------------------------------------------------------===//
 
 /// Generate the code for registering passes.
 #define GEN_PASS_REGISTRATION
 #include "mlir/Dialect/Affine/Passes.h.inc"
 
 } // namespace mlir
 
 #endif // MLIR_DIALECT_AFFINE_PASSES_H

mlir/include/mlir/Dialect/Affine/Passes.td

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_AFFINE_PASSES
 #define MLIR_DIALECT_AFFINE_PASSES
 
 include "mlir/Pass/PassBase.td"
 
-def AffineDataCopyGeneration : Pass<"affine-data-copy-generate", "func::FuncOp"> {
+def AffineDataCopyGeneration : InterfacePass<
+    "affine-data-copy-generate", "FunctionOpInterface"> {
   let summary = "Generate explicit copying for affine memory operations";
   let constructor = "mlir::createAffineDataCopyGenerationPass()";
   let dependentDialects = ["memref::MemRefDialect"];
   let options = [
     Option<"fastMemoryCapacity", "fast-mem-capacity", "uint64_t",
            /*default=*/"std::numeric_limits<uint64_t>::max()",
            "Set fast memory space capacity in KiB (default: unlimited)">,
     Option<"fastMemorySpace", "fast-mem-space", "unsigned",
Option<"skipNonUnitStrideLoops", "skip-non-unit-stride-loops", "bool",
            /*default=*/"false", "Testing purposes: avoid non-unit stride loop "
                                 "choice depths for copy placement">,
     Option<"tagMemorySpace", "tag-mem-space", "unsigned",
            /*default=*/"0",
            "Tag memory space identifier for copy generation (default: 0)">,
   ];
 }
 
-def AffineLoopFusion : Pass<"affine-loop-fusion", "func::FuncOp"> {
+def AffineLoopFusion : InterfacePass<
+    "affine-loop-fusion", "FunctionOpInterface"> {
   let summary = "Fuse affine loop nests";
   let description = [{
     This pass performs fusion of loop nests using a slicing-based approach. It
     combines two fusion strategies: producer-consumer fusion and sibling fusion.
     Producer-consumer fusion is aimed at fusing pairs of loops where the first
     one writes to a memref that the second reads. Sibling fusion targets pairs
     of loops that share no dependences between them but that load from the same
     memref. The fused loop nests, when possible, are rewritten to access
Option<"affineFusionMode", "mode", "enum FusionMode",
            "clEnumValN( mlir::FusionMode::ProducerConsumer, "
            "\"producer\", \"Perform only producer-consumer fusion\"), "
            "clEnumValN( mlir::FusionMode::Sibling, "
            "\"sibling\", \"Perform only sibling fusion\"))">,
     ];
   let dependentDialects = ["memref::MemRefDialect"];
 }
 
-def AffineLoopInvariantCodeMotion
-    : Pass<"affine-loop-invariant-code-motion", "func::FuncOp"> {
+def AffineLoopInvariantCodeMotion : InterfacePass<
+    "affine-loop-invariant-code-motion", "FunctionOpInterface"> {
   let summary = "Hoist loop invariant instructions outside of affine loops";
   let constructor = "mlir::createAffineLoopInvariantCodeMotionPass()";
 }
 
-def AffineLoopTiling : Pass<"affine-loop-tile", "func::FuncOp"> {
+def AffineLoopTiling : InterfacePass<
+    "affine-loop-tile", "FunctionOpInterface"> {
   let summary = "Tile affine loop nests";
   let constructor = "mlir::createLoopTilingPass()";
   let options = [
     Option<"cacheSizeInKiB", "cache-size", "uint64_t", /*default=*/"512",
            "Set size of cache to tile for in KiB (default: 512)">,
     Option<"separate", "separate", "bool", /*default=*/"false",
            "Separate full and partial tiles (default: false)">,
     Option<"tileSize", "tile-size", "unsigned", /*default=*/"",
            "Use this tile size for all loops">,
     ListOption<"tileSizes", "tile-sizes", "unsigned",
                "List of tile sizes for each perfect nest "
                "(overridden by -tile-size)",
                "llvm::cl::ZeroOrMore">,
   ];
 }
 
-def AffineLoopUnroll : Pass<"affine-loop-unroll", "func::FuncOp"> {
+def AffineLoopUnroll : InterfacePass<
+    "affine-loop-unroll", "FunctionOpInterface"> {
   let summary = "Unroll affine loops";
   let constructor = "mlir::createLoopUnrollPass()";
   let options = [
     Option<"unrollFactor", "unroll-factor", "unsigned", /*default=*/"4",
            "Use this unroll factor for all loops being unrolled">,
     Option<"unrollUpToFactor", "unroll-up-to-factor", "bool",
            /*default=*/"false", "Allow unrolling up to the factor specified">,
     Option<"unrollFull", "unroll-full", "bool", /*default=*/"false",
            "Fully unroll loops">,
     Option<"numRepetitions", "unroll-num-reps", "unsigned", /*default=*/"1",
            "Unroll innermost loops repeatedly this many times">,
     Option<"unrollFullThreshold", "unroll-full-threshold", "unsigned",
            /*default=*/"1",
            "Unroll all loops with trip count less than or equal to this">,
   ];
 }
 
-def AffineLoopUnrollAndJam : Pass<"affine-loop-unroll-jam", "func::FuncOp"> {
+def AffineLoopUnrollAndJam : InterfacePass<
+    "affine-loop-unroll-jam", "FunctionOpInterface"> {
   let summary = "Unroll and jam affine loops";
   let constructor = "mlir::createLoopUnrollAndJamPass()";
   let options = [
     Option<"unrollJamFactor", "unroll-jam-factor", "unsigned",
            /*default=*/"4",
            "Use this unroll jam factor for all loops (default 4)">,
   ];
 }
 
 def AffinePipelineDataTransfer
-    : Pass<"affine-pipeline-data-transfer", "func::FuncOp"> {
+    : InterfacePass<"affine-pipeline-data-transfer", "FunctionOpInterface"> {
   let summary = "Pipeline non-blocking data transfers between explicitly "
                 "managed levels of the memory hierarchy";
   let description = [{
     This pass performs a transformation to overlap non-blocking DMA operations
     in a loop with computations through double buffering. This is achieved by
     advancing dma_start operations with respect to other operations.
 
     Input
module {
         return
       }
     }
     ```
   }];
   let constructor = "mlir::createPipelineDataTransferPass()";
 }
 
-def AffineScalarReplacement : Pass<"affine-scalrep", "func::FuncOp"> {
+def AffineScalarReplacement : InterfacePass<
+    "affine-scalrep", "FunctionOpInterface"> {
   let summary = "Replace affine memref acceses by scalars by forwarding stores "
                 "to loads and eliminating redundant loads";
   let description = [{
     This pass performs store to load forwarding and redundant load elimination
     for affine memref accesses and potentially eliminates the entire memref
     if all its accesses are forwarded.
 
     Input
module {
         return %0 : memref<10x10xf32>
       }
     }
     ```
   }];
   let constructor = "mlir::createAffineScalarReplacementPass()";
 }
 
-def AffineVectorize : Pass<"affine-super-vectorize", "func::FuncOp"> {
+def AffineVectorize : InterfacePass<
+    "affine-super-vectorize", "FunctionOpInterface"> {
   let summary = "Vectorize to a target independent n-D vector abstraction";
   let constructor = "mlir::createSuperVectorizePass()";
   let dependentDialects = ["vector::VectorDialect"];
   let options = [
     ListOption<"vectorSizes", "virtual-vector-size", "int64_t",
                "Specify an n-D virtual vector size for vectorization",
                "llvm::cl::ZeroOrMore">,
     // Optionally, the fixed mapping from loop to fastest varying MemRef
ListOption<"fastestVaryingPattern", "test-fastest-varying", "int64_t",
                "llvm::cl::ZeroOrMore">,
     Option<"vectorizeReductions", "vectorize-reductions", "bool",
            /*default=*/"false",
            "Vectorize known reductions expressed via iter_args. "
            "Switched off by default.">
   ];
 }
 
-def AffineParallelize : Pass<"affine-parallelize", "func::FuncOp"> {
+def AffineParallelize : InterfacePass<
+    "affine-parallelize", "FunctionOpInterface"> {
   let summary = "Convert affine.for ops into 1-D affine.parallel";
   let constructor = "mlir::createAffineParallelizePass()";
   let options = [
     Option<"maxNested", "max-nested", "unsigned", /*default=*/"-1u",
            "Maximum number of nested parallel loops to produce. "
            "Defaults to unlimited (UINT_MAX).">,
     Option<"parallelReductions", "parallel-reductions", "bool",
            /*default=*/"false",
            "Whether to parallelize reduction loops. Defaults to false.">
   ];
 }
 
-def AffineLoopNormalize : Pass<"affine-loop-normalize", "func::FuncOp"> {
+def AffineLoopNormalize : InterfacePass<
+    "affine-loop-normalize", "FunctionOpInterface"> {
   let summary = "Apply normalization transformations to affine loop-like ops";
   let constructor = "mlir::createAffineLoopNormalizePass()";
 }
 
-def LoopCoalescing : Pass<"affine-loop-coalescing", "func::FuncOp"> {
+def LoopCoalescing : InterfacePass<
+    "affine-loop-coalescing", "FunctionOpInterface"> {
   let summary = "Coalesce nested loops with independent bounds into a single "
                 "loop";
   let constructor = "mlir::createLoopCoalescingPass()";
   let dependentDialects = ["arith::ArithmeticDialect"];
 }
 
-def SimplifyAffineStructures : Pass<"affine-simplify-structures", "func::FuncOp"> {
+def SimplifyAffineStructures : InterfacePass<
+    "affine-simplify-structures", "FunctionOpInterface"> {
   let summary = "Simplify affine expressions in maps/sets and normalize "
                 "memrefs";
   let constructor = "mlir::createSimplifyAffineStructuresPass()";
 }
 
 #endif // MLIR_DIALECT_AFFINE_PASSES

mlir/include/mlir/Dialect/Affine/Utils.h

 
 class AffineForOp;
 class AffineIfOp;
 class AffineParallelOp;
 class DominanceInfo;
 class Operation;
 class PostDominanceInfo;
 
-namespace func {
-class FuncOp;
-} // namespace func
-
 namespace memref {
 class AllocOp;
 } // namespace memref
 
 struct LogicalResult;
 
 using ReductionLoopMap = DenseMap<Operation *, SmallVector<LoopReduction, 2>>;
 
struct VectorizationStrategy {
   // Maps loops that implement vectorizable reductions to the corresponding
   // reduction descriptors.
   ReductionLoopMap reductionLoops;
 };
 
 /// Replace affine store and load accesses by scalars by forwarding stores to
 /// loads and eliminate invariant affine loads; consequently, eliminate dead
 /// allocs.
-void affineScalarReplace(func::FuncOp f, DominanceInfo &domInfo,
+void affineScalarReplace(Operation *parentOp, DominanceInfo &domInfo,
                          PostDominanceInfo &postDomInfo);
 
 /// Vectorizes affine loops in 'loops' using the n-D vectorization factors in
 /// 'vectorSizes'. By default, each vectorization factor is applied
 /// inner-to-outer to the loops of each loop nest. 'fastestVaryingPattern' can
 /// be optionally used to provide a different loop vectorization order.
 /// If `reductionLoops` is not empty, the given reduction loops may be
 /// vectorized along the reduction dimension.

mlir/lib/Dialect/Affine/Analysis/AffineAnalysis.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h"
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/IR/AffineValueMap.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/IR/AffineExprVisitor.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/Interfaces/ViewLikeInterface.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
 #define DEBUG_TYPE "affine-analysis"
 
while (currBlock && currBlock != endBlock &&
            !currBlock->getParentOp()->hasTrait<OpTrait::AffineScope>()) {
       ancestorBlocks.push_back(currBlock);
       currBlock = currBlock->getParentOp()->getBlock();
     }
   };
 
   if (numCommonLoops == 0) {
     Block *block = srcAccess.opInst->getBlock();
-    while (!llvm::isa<func::FuncOp>(block->getParentOp())) {
+    while (!llvm::isa<FunctionOpInterface>(block->getParentOp())) {
       block = block->getParentOp()->getBlock();
     }
     return block;
   }
   Value commonForIV = srcDomain.getValue(numCommonLoops - 1);
   AffineForOp forOp = getForInductionVarOwner(commonForIV);
   assert(forOp && "commonForValue was not an induction variable");
 

mlir/lib/Dialect/Affine/Transforms/AffineDataCopyGeneration.cpp

 };
 
 } // namespace
 
 /// Generates copies for memref's living in 'slowMemorySpace' into newly created
 /// buffers in 'fastMemorySpace', and replaces memory operations to the former
 /// by the latter. Only load op's handled for now.
 /// TODO: extend this to store op's.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineDataCopyGenerationPass(unsigned slowMemorySpace,
                                          unsigned fastMemorySpace,
                                          unsigned tagMemorySpace,
                                          int minDmaTransferSize,
                                          uint64_t fastMemCapacityBytes) {
   return std::make_unique<AffineDataCopyGeneration>(
       slowMemorySpace, fastMemorySpace, tagMemorySpace, minDmaTransferSize,
       fastMemCapacityBytes);
 }
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineDataCopyGenerationPass() {
   return std::make_unique<AffineDataCopyGeneration>();
 }
 
 /// Generate copies for this block. The block is partitioned into separate
 /// ranges: each range is either a sequence of one or more operations starting
 /// and ending with an affine load or store op, or just an affine.forop (which
 /// could have other affine for op's nested within).
if (curBegin != block->end()) {
     // Exclude the affine.yield - hence, the std::prev.
     (void)affineDataCopyGenerate(/*begin=*/curBegin,
                                  /*end=*/std::prev(block->end()), copyOptions,
                                  /*filterMemRef=*/llvm::None, copyNests);
   }
 }
 
 void AffineDataCopyGeneration::runOnOperation() {
-  func::FuncOp f = getOperation();
+  FunctionOpInterface f = getOperation();
   OpBuilder topBuilder(f.getBody());
   zeroIndex = topBuilder.create<arith::ConstantIndexOp>(f.getLoc(), 0);
 
   // Nests that are copy-in's or copy-out's; the root AffineForOps of those
   // nests are stored herein.
   DenseSet<Operation *> copyNests;
 
   // Clear recorded copy nests.

mlir/lib/Dialect/Affine/Transforms/AffineLoopInvariantCodeMotion.cpp

void LoopInvariantCodeMotion::runOnOperation() {
   // way, we first LICM from the inner loop, and place the ops in
   // the outer loop, which in turn can be further LICM'ed.
   getOperation().walk([&](AffineForOp op) {
     LLVM_DEBUG(op->print(llvm::dbgs() << "\nOriginal loop\n"));
     runOnAffineForOp(op);
   });
 }
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineLoopInvariantCodeMotionPass() {
   return std::make_unique<LoopInvariantCodeMotion>();
 }

mlir/lib/Dialect/Affine/Transforms/AffineLoopNormalize.cpp

getOperation().walk([](Operation *op) {
       else if (auto affineFor = dyn_cast<AffineForOp>(op))
         (void)normalizeAffineFor(affineFor);
     });
   }
 };
 
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineLoopNormalizePass() {
   return std::make_unique<AffineLoopNormalizePass>();
 }

mlir/lib/Dialect/Affine/Transforms/AffineParallelize.cpp

struct ParallelizationCandidate {
   /// The potentially parallelizable loop.
   AffineForOp loop;
   /// Desciprtors of reductions that can be parallelized in the loop.
   SmallVector<LoopReduction> reductions;
 };
 } // namespace
 
 void AffineParallelize::runOnOperation() {
-  func::FuncOp f = getOperation();
+  FunctionOpInterface f = getOperation();
 
   // The walker proceeds in pre-order to process the outer loops first
   // and control the number of outer parallel loops.
   std::vector<ParallelizationCandidate> parallelizableLoops;
   f.walk<WalkOrder::PreOrder>([&](AffineForOp loop) {
     SmallVector<LoopReduction> reductions;
     if (isLoopParallel(loop, parallelReductions ? &reductions : nullptr))
       parallelizableLoops.emplace_back(loop, std::move(reductions));
if (numParentParallelOps < maxNested) {
       }
     } else {
       LLVM_DEBUG(llvm::dbgs() << "[" DEBUG_TYPE "] too many nested loops\n"
                               << loop);
     }
   }
 }
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineParallelizePass() {
   return std::make_unique<AffineParallelize>();
 }

mlir/lib/Dialect/Affine/Transforms/AffineScalarReplacement.cpp

 namespace {
 struct AffineScalarReplacement
     : public AffineScalarReplacementBase<AffineScalarReplacement> {
   void runOnOperation() override;
 };
 
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createAffineScalarReplacementPass() {
   return std::make_unique<AffineScalarReplacement>();
 }
 
 void AffineScalarReplacement::runOnOperation() {
   affineScalarReplace(getOperation(), getAnalysis<DominanceInfo>(),
                       getAnalysis<PostDominanceInfo>());
 }

mlir/lib/Dialect/Affine/Transforms/LoopCoalescing.cpp

for (unsigned end = loops.size(); end > 0; --end) {
       // If a band was found and transformed, keep looking at the loops above
       // the outermost transformed loop.
       if (start != end - 1)
         end = start + 1;
     }
   }
 
   void runOnOperation() override {
-    func::FuncOp func = getOperation();
+    FunctionOpInterface func = getOperation();
     func.walk([&](Operation *op) {
       if (auto scfForOp = dyn_cast<scf::ForOp>(op))
         walkLoop(scfForOp);
       else if (auto affineForOp = dyn_cast<AffineForOp>(op))
         walkLoop(affineForOp);
     });
   }
 };
 
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>> mlir::createLoopCoalescingPass() {
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+mlir::createLoopCoalescingPass() {
   return std::make_unique<LoopCoalescingPass>();
 }

mlir/lib/Dialect/Affine/Transforms/LoopFusion.cpp

LoopFusion(unsigned fastMemorySpace, uint64_t localBufSizeThresholdBytes,
     this->affineFusionMode = affineFusionMode;
   }
 
   void runOnOperation() override;
 };
 
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createLoopFusionPass(unsigned fastMemorySpace,
                            uint64_t localBufSizeThreshold, bool maximalFusion,
                            enum FusionMode affineFusionMode) {
   return std::make_unique<LoopFusion>(fastMemorySpace, localBufSizeThreshold,
                                       maximalFusion, affineFusionMode);
 }
 
 namespace {
opToWalk->walk([&](Operation *op) {
         loadOpInsts.push_back(op);
       else if (isa<AffineWriteOpInterface>(op))
         storeOpInsts.push_back(op);
     });
   }
 };
 
 // MemRefDependenceGraph is a graph data structure where graph nodes are
-// top-level operations in a FuncOp which contain load/store ops, and edges
-// are memref dependences between the nodes.
+// top-level operations in a FuncOpInterface which contain load/store ops, and
+// edges are memref dependences between the nodes.
 // TODO: Add a more flexible dependence graph representation.
 // TODO: Add a depth parameter to dependence graph construction.
 struct MemRefDependenceGraph {
 public:
   // Node represents a node in the graph. A Node is either an entire loop nest
   // rooted at the top level which contains loads/stores, or a top level
   // load/store.
   struct Node {
public:
   DenseMap<Value, unsigned> memrefEdgeCount;
   // The next unique identifier to use for newly created graph nodes.
   unsigned nextNodeId = 0;
 
   MemRefDependenceGraph() = default;
 
   // Initializes the dependence graph based on operations in 'f'.
   // Returns true on success, false otherwise.
-  bool init(func::FuncOp f);
+  bool init(FunctionOpInterface f);
 
   // Returns the graph node for 'id'.
   Node *getNode(unsigned id) {
     auto it = nodes.find(id);
     assert(it != nodes.end());
     return &it->second;
   }
 

 }
 
 } // namespace
 
 // Initializes the data dependence graph by walking operations in 'f'.
 // Assigns each node in the graph a node id based on program order in 'f'.
 // TODO: Add support for taking a Block arg to construct the
 // dependence graph at a different depth.
-bool MemRefDependenceGraph::init(func::FuncOp f) {
+bool MemRefDependenceGraph::init(FunctionOpInterface f) {
   LLVM_DEBUG(llvm::dbgs() << "--- Initializing MDG ---\n");
   DenseMap<Value, SetVector<unsigned>> memrefAccesses;
 
   // TODO: support multi-block functions.
   if (!llvm::hasSingleElement(f))
     return false;
 
   DenseMap<Operation *, unsigned> forToNodeMap;
static Value createPrivateMemRef(AffineForOp forOp, Operation *srcStoreOpInst,
                                  unsigned dstLoopDepth,
                                  Optional<unsigned> fastMemorySpace,
                                  uint64_t localBufSizeThreshold) {
   auto *forInst = forOp.getOperation();
 
   // Create builder to insert alloc op just before 'forOp'.
   OpBuilder b(forInst);
   // Builder to create constants at the top level.
-  OpBuilder top(forInst->getParentOfType<func::FuncOp>().getBody());
+  OpBuilder top(forInst->getParentOfType<FunctionOpInterface>().getBody());
   // Create new memref type based on slice bounds.
   auto oldMemRef = cast<AffineWriteOpInterface>(srcStoreOpInst).getMemRef();
   auto oldMemRefType = oldMemRef.getType().cast<MemRefType>();
   unsigned rank = oldMemRefType.getRank();
 
   // Compute MemRefRegion for 'srcStoreOpInst' at depth 'dstLoopDepth'.
   MemRefRegion region(srcStoreOpInst->getLoc());
   bool validRegion = succeeded(region.compute(srcStoreOpInst, dstLoopDepth));
auto canFuseWithSibNode = [&](Node *sibNode, Value memref) {
       // access that lies between 'dstNode' and 'sibNode'.
       if (hasNonAffineUsersOnThePath(dstNode->id, sibNode->id, mdg) ||
           hasNonAffineUsersOnThePath(sibNode->id, dstNode->id, mdg))
         return false;
       return true;
     };
 
     // Search for siblings which load the same memref function argument.
-    auto fn = dstNode->op->getParentOfType<func::FuncOp>();
+    auto fn = dstNode->op->getParentOfType<FunctionOpInterface>();
     for (unsigned i = 0, e = fn.getNumArguments(); i != e; ++i) {
       for (auto *user : fn.getArgument(i).getUsers()) {
         if (auto loadOp = dyn_cast<AffineReadOpInterface>(user)) {
           // Gather loops surrounding 'use'.
           SmallVector<AffineForOp, 4> loops;
           getLoopIVs(*user, &loops);
           // Skip 'use' if it is not within a loop nest.
           if (loops.empty())

mlir/lib/Dialect/Affine/Transforms/LoopTiling.cpp

struct LoopTiling : public AffineLoopTilingBase<LoopTiling> {
   // If true, tile sizes are set to avoid max/min in bounds if possible.
   bool avoidMaxMinBounds = true;
 };
 
 } // namespace
 
 /// Creates a pass to perform loop tiling on all suitable loop nests of a
 /// Function.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createLoopTilingPass(uint64_t cacheSizeBytes) {
   return std::make_unique<LoopTiling>(cacheSizeBytes);
 }
-std::unique_ptr<OperationPass<func::FuncOp>> mlir::createLoopTilingPass() {
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
+mlir::createLoopTilingPass() {
   return std::make_unique<LoopTiling>();
 }
 
 /// Reduces each tile size to the largest divisor of the corresponding trip
 /// count (if the trip count is known).
 static void adjustToDivisorsOfTripCounts(ArrayRef<AffineForOp> band,
                                          SmallVectorImpl<unsigned> *tileSizes) {
   assert(band.size() == tileSizes->size() && "invalid tile size count");

mlir/lib/Dialect/Affine/Transforms/LoopUnroll.cpp

static bool isInnermostAffineForOp(AffineForOp op) {
   return !op.getBody()
               ->walk([&](AffineForOp nestedForOp) {
                 return WalkResult::interrupt();
               })
               .wasInterrupted();
 }
 
 /// Gathers loops that have no affine.for's nested within.
-static void gatherInnermostLoops(func::FuncOp f,
+static void gatherInnermostLoops(FunctionOpInterface f,
                                  SmallVectorImpl<AffineForOp> &loops) {
   f.walk([&](AffineForOp forOp) {
     if (isInnermostAffineForOp(forOp))
       loops.push_back(forOp);
   });
 }
 
 void LoopUnroll::runOnOperation() {
-  func::FuncOp func = getOperation();
+  FunctionOpInterface func = getOperation();
   if (func.isExternal())
     return;
 
   if (unrollFull && unrollFullThreshold.hasValue()) {
     // Store short loops as we walk.
     SmallVector<AffineForOp, 4> loops;
 
     // Gathers all loops with trip count <= minTripCount. Do a post order walk
LogicalResult LoopUnroll::runOnAffineForOp(AffineForOp forOp) {
   if (unrollFull)
     return loopUnrollFull(forOp);
   // Otherwise, unroll by the given unroll factor.
   if (unrollUpToFactor)
     return loopUnrollUpToFactor(forOp, unrollFactor);
   return loopUnrollByFactor(forOp, unrollFactor);
 }
 
-std::unique_ptr<OperationPass<func::FuncOp>> mlir::createLoopUnrollPass(
+std::unique_ptr<InterfacePass<FunctionOpInterface>> mlir::createLoopUnrollPass(
     int unrollFactor, bool unrollUpToFactor, bool unrollFull,
     const std::function<unsigned(AffineForOp)> &getUnrollFactor) {
   return std::make_unique<LoopUnroll>(
       unrollFactor == -1 ? None : Optional<unsigned>(unrollFactor),
       unrollUpToFactor, unrollFull, getUnrollFactor);
 }

mlir/lib/Dialect/Affine/Transforms/LoopUnrollAndJam.cpp

explicit LoopUnrollAndJam(Optional<unsigned> unrollJamFactor = None) {
     if (unrollJamFactor)
       this->unrollJamFactor = *unrollJamFactor;
   }
 
   void runOnOperation() override;
 };
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createLoopUnrollAndJamPass(int unrollJamFactor) {
   return std::make_unique<LoopUnrollAndJam>(
       unrollJamFactor == -1 ? None : Optional<unsigned>(unrollJamFactor));
 }
 
 void LoopUnrollAndJam::runOnOperation() {
   if (getOperation().isExternal())
     return;
 
   // Currently, just the outermost loop from the first loop nest is
   // unroll-and-jammed by this pass. However, runOnAffineForOp can be called on
   // any for operation.
   auto &entryBlock = getOperation().front();
   if (auto forOp = dyn_cast<AffineForOp>(entryBlock.front()))
     (void)loopUnrollJamByFactor(forOp, unrollJamFactor);
 }

mlir/lib/Dialect/Affine/Transforms/PassDetail.h

 //===- PassDetail.h - Affine Pass class details -----------------*- C++ -*-===//
 //
 // Part of the LLVM Project, 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 DIALECT_AFFINE_TRANSFORMS_PASSDETAIL_H_
 #define DIALECT_AFFINE_TRANSFORMS_PASSDETAIL_H_
 
 #include "mlir/Dialect/Affine/Passes.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/IR/DialectRegistry.h"
+#include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/Pass/Pass.h"
 
 namespace mlir {
 // Forward declaration from Dialect.h
 template <typename ConcreteDialect>
 void registerDialect(DialectRegistry &registry);
 
 namespace arith {

mlir/lib/Dialect/Affine/Transforms/PipelineDataTransfer.cpp

struct PipelineDataTransfer
 
   std::vector<AffineForOp> forOps;
 };
 
 } // namespace
 
 /// Creates a pass to pipeline explicit movement of data across levels of the
 /// memory hierarchy.
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createPipelineDataTransferPass() {
   return std::make_unique<PipelineDataTransfer>();
 }
 
 // Returns the position of the tag memref operand given a DMA operation.
 // Temporary utility: will be replaced when DmaStart/DmaFinish abstract op's are
 // added.  TODO
 static unsigned getTagMemRefPos(Operation &dmaOp) {

mlir/lib/Dialect/Affine/Transforms/SimplifyAffineStructures.cpp

AffineMap simplify(AffineMap map) {
     return mMap.getAffineMap();
   }
 
   DenseMap<Attribute, Attribute> simplifiedAttributes;
 };
 
 } // namespace
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 mlir::createSimplifyAffineStructuresPass() {
   return std::make_unique<SimplifyAffineStructures>();
 }
 
 void SimplifyAffineStructures::runOnOperation() {
-  auto func = getOperation();
+  FunctionOpInterface f = getOperation();
   simplifiedAttributes.clear();
-  RewritePatternSet patterns(func.getContext());
-  AffineApplyOp::getCanonicalizationPatterns(patterns, func.getContext());
-  AffineForOp::getCanonicalizationPatterns(patterns, func.getContext());
-  AffineIfOp::getCanonicalizationPatterns(patterns, func.getContext());
+  RewritePatternSet patterns(f.getContext());
+  AffineApplyOp::getCanonicalizationPatterns(patterns, f.getContext());
+  AffineForOp::getCanonicalizationPatterns(patterns, f.getContext());
+  AffineIfOp::getCanonicalizationPatterns(patterns, f.getContext());
   FrozenRewritePatternSet frozenPatterns(std::move(patterns));
 
   // The simplification of affine attributes will likely simplify the op. Try to
   // fold/apply canonicalization patterns when we have affine dialect ops.
   SmallVector<Operation *> opsToSimplify;
-  func.walk([&](Operation *op) {
+  f.walk([&](Operation *op) {
     for (auto attr : op->getAttrs()) {
       if (auto mapAttr = attr.getValue().dyn_cast<AffineMapAttr>())
         simplifyAndUpdateAttribute(op, attr.getName(), mapAttr);
       else if (auto setAttr = attr.getValue().dyn_cast<IntegerSetAttr>())
         simplifyAndUpdateAttribute(op, attr.getName(), setAttr);
     }
 
     if (isa<AffineForOp, AffineIfOp, AffineApplyOp>(op))
       opsToSimplify.push_back(op);
   });
   (void)applyOpPatternsAndFold(opsToSimplify, frozenPatterns, /*strict=*/true);
 }

mlir/lib/Dialect/Affine/Transforms/SuperVectorize.cpp

for (NestedMatch &match : intersectingMatches) {
       if (succeeded(vectorizeRootMatch(match, strategy)))
         break;
     }
   }
 
   LLVM_DEBUG(dbgs() << "\n");
 }
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createSuperVectorizePass(ArrayRef<int64_t> virtualVectorSize) {
   return std::make_unique<Vectorize>(virtualVectorSize);
 }
-std::unique_ptr<OperationPass<func::FuncOp>> createSuperVectorizePass() {
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createSuperVectorizePass() {
   return std::make_unique<Vectorize>();
 }
 
 /// Applies vectorization to the current function by searching over a bunch of
 /// predetermined patterns.
 void Vectorize::runOnOperation() {
-  func::FuncOp f = getOperation();
+  FunctionOpInterface f = getOperation();
   if (!fastestVaryingPattern.empty() &&
       fastestVaryingPattern.size() != vectorSizes.size()) {
     f.emitRemark("Fastest varying pattern specified with different size than "
                  "the vector size.");
     return signalPassFailure();
   }
 
   if (vectorizeReductions && vectorSizes.size() != 1) {
vectorizeAffineLoopNest(std::vector<SmallVector<AffineForOp, 2>> &loops,
                         const VectorizationStrategy &strategy) {
   // Thread-safe RAII local context, BumpPtrAllocator freed on exit.
   NestedPatternContext mlContext;
   if (failed(verifyLoopNesting(loops)))
     return failure();
   return vectorizeLoopNest(loops, strategy);
 }
 
-std::unique_ptr<OperationPass<func::FuncOp>>
+std::unique_ptr<InterfacePass<FunctionOpInterface>>
 createSuperVectorizePass(ArrayRef<int64_t> virtualVectorSize) {
   return std::make_unique<Vectorize>(virtualVectorSize);
 }
-std::unique_ptr<OperationPass<func::FuncOp>> createSuperVectorizePass() {
+std::unique_ptr<InterfacePass<FunctionOpInterface>> createSuperVectorizePass() {
   return std::make_unique<Vectorize>();
 }
 
 } // namespace mlir

mlir/lib/Dialect/Affine/Utils/LoopFusionUtils.cpp

 #include "mlir/Dialect/Affine/LoopFusionUtils.h"
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/AffineStructures.h"
 #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/LoopUtils.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/IR/Operation.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
 #define DEBUG_TYPE "loop-fusion-utils"
 

 
 /// Collect loop nest statistics (eg. loop trip count and operation count)
 /// in 'stats' for loop nest rooted at 'forOp'. Returns true on success,
 /// returns false otherwise.
 bool mlir::getLoopNestStats(AffineForOp forOpRoot, LoopNestStats *stats) {
   auto walkResult = forOpRoot.walk([&](AffineForOp forOp) {
     auto *childForOp = forOp.getOperation();
     auto *parentForOp = forOp->getParentOp();
-    if (!llvm::isa<func::FuncOp>(parentForOp)) {
+    if (!llvm::isa<FunctionOpInterface>(parentForOp)) {
       if (!isa<AffineForOp>(parentForOp)) {
         LLVM_DEBUG(llvm::dbgs() << "Expected parent AffineForOp\n");
         return WalkResult::interrupt();
       }
       // Add mapping to 'forOp' from its parent AffineForOp.
       stats->loopMap[parentForOp].push_back(forOp);
     }
 

mlir/lib/Dialect/Affine/Utils/LoopUtils.cpp

 #include "mlir/Dialect/Affine/LoopUtils.h"
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/IR/AffineValueMap.h"
 #include "mlir/Dialect/Affine/Utils.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
LogicalResult mlir::promoteIfSingleIteration(AffineForOp forOp) {
   if (forOp.getLowerBoundMap().getNumResults() != 1)
     return failure();
 
   // Replaces all IV uses to its single iteration value.
   auto iv = forOp.getInductionVar();
   auto *parentBlock = forOp->getBlock();
   if (!iv.use_empty()) {
     if (forOp.hasConstantLowerBound()) {
-      OpBuilder topBuilder(forOp->getParentOfType<func::FuncOp>().getBody());
+      OpBuilder topBuilder(
+          forOp->getParentOfType<FunctionOpInterface>().getBody());
       auto constOp = topBuilder.create<arith::ConstantIndexOp>(
           forOp.getLoc(), forOp.getConstantLowerBound());
       iv.replaceAllUsesWith(constOp);
     } else {
       auto lbOperands = forOp.getLowerBoundOperands();
       auto lbMap = forOp.getLowerBoundMap();
       OpBuilder builder(forOp);
       if (lbMap == builder.getDimIdentityMap()) {
for (unsigned i = 0; i < std::numeric_limits<unsigned>::max(); ++i) {
     if (!root)
       return;
   }
 }
 
 /// Identify valid and profitable bands of loops to tile. This is currently just
 /// a temporary placeholder to test the mechanics of tiled code generation.
 /// Returns all maximal outermost perfect loop nests to tile.
-void mlir::getTileableBands(func::FuncOp f,
+void mlir::getTileableBands(FunctionOpInterface f,
                             std::vector<SmallVector<AffineForOp, 6>> *bands) {
   // Get maximal perfect nest of 'affine.for' insts starting from root
   // (inclusive).
-  for (AffineForOp forOp : f.getOps<AffineForOp>()) {
+  for (AffineForOp forOp : f.getBody().getOps<AffineForOp>()) {
     SmallVector<AffineForOp, 6> band;
     getPerfectlyNestedLoops(band, forOp);
     bands->push_back(band);
   }
 }
 
 /// Unrolls this loop completely.
 LogicalResult mlir::loopUnrollFull(AffineForOp forOp) {
static LogicalResult generateCopy(
     Block::iterator end, Block *copyPlacementBlock,
     Block::iterator copyInPlacementStart, Block::iterator copyOutPlacementStart,
     AffineCopyOptions copyOptions, DenseMap<Value, Value> &fastBufferMap,
     DenseSet<Operation *> &copyNests, uint64_t *sizeInBytes,
     Block::iterator *nBegin, Block::iterator *nEnd) {
   *nBegin = begin;
   *nEnd = end;
 
-  func::FuncOp f = begin->getParentOfType<func::FuncOp>();
+  auto f = begin->getParentOfType<FunctionOpInterface>();
   OpBuilder topBuilder(f.getBody());
   Value zeroIndex = topBuilder.create<arith::ConstantIndexOp>(f.getLoc(), 0);
 
   if (begin == end)
     return success();
 
   // Is the copy out point at the end of the block where we are doing
   // explicit copying.
   bool isCopyOutAtEndOfBlock = (end == copyOutPlacementStart);
 
   // Copies for read regions are going to be inserted at 'begin'.
   OpBuilder prologue(copyPlacementBlock, copyInPlacementStart);
   // Copies for write regions are going to be inserted at 'end'.
   OpBuilder epilogue(copyPlacementBlock, copyOutPlacementStart);
   OpBuilder &b = region.isWrite() ? epilogue : prologue;
 
   // Builder to create constants at the top level.
-  auto func = copyPlacementBlock->getParent()->getParentOfType<func::FuncOp>();
+  auto func =
+      copyPlacementBlock->getParent()->getParentOfType<FunctionOpInterface>();
   OpBuilder top(func.getBody());
 
   auto loc = region.loc;
   auto memref = region.memref;
   auto memRefType = memref.getType().cast<MemRefType>();
 
   if (!memRefType.getLayout().isIdentity()) {
     LLVM_DEBUG(llvm::dbgs() << "Non-identity layout map not yet supported\n");
for (auto &op : *block) {
     if (auto forOp = dyn_cast<AffineForOp>(op)) {
       depthToLoops[currLoopDepth].push_back(forOp);
       gatherLoopsInBlock(forOp.getBody(), currLoopDepth + 1, depthToLoops);
     }
   }
 }
 
 /// Gathers all AffineForOps in 'func.func' grouped by loop depth.
-void mlir::gatherLoops(func::FuncOp func,
+void mlir::gatherLoops(FunctionOpInterface func,
                        std::vector<SmallVector<AffineForOp, 2>> &depthToLoops) {
   for (auto &block : func)
     gatherLoopsInBlock(&block, /*currLoopDepth=*/0, depthToLoops);
 
   // Remove last loop level from output since it's empty.
   if (!depthToLoops.empty()) {
     assert(depthToLoops.back().empty() && "Last loop level is not empty?");
     depthToLoops.pop_back();

mlir/lib/Dialect/Affine/Utils/Utils.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Affine/Utils.h"
 
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/IR/AffineValueMap.h"
 #include "mlir/Dialect/Affine/LoopUtils.h"
-#include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/IR/AffineExprVisitor.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Dominance.h"
+#include "mlir/IR/FunctionInterfaces.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 
 #define DEBUG_TYPE "affine-utils"
 
 using namespace mlir;
 using namespace presburger;
 

 /// Returns the outermost affine.for/parallel op that the `ifOp` is invariant
 /// on. The `ifOp` could be hoisted and placed right before such an operation.
 /// This method assumes that the ifOp has been canonicalized (to be correct and
 /// effective).
 static Operation *getOutermostInvariantForOp(AffineIfOp ifOp) {
   // Walk up the parents past all for op that this conditional is invariant on.
   auto ifOperands = ifOp.getOperands();
   auto *res = ifOp.getOperation();
-  while (!isa<func::FuncOp>(res->getParentOp())) {
+  while (!isa<FunctionOpInterface>(res->getParentOp())) {
     auto *parentOp = res->getParentOp();
     if (auto forOp = dyn_cast<AffineForOp>(parentOp)) {
       if (llvm::is_contained(ifOperands, forOp.getInductionVar()))
         break;
     } else if (auto parallelOp = dyn_cast<AffineParallelOp>(parentOp)) {
       for (auto iv : parallelOp.getIVs())
         if (llvm::is_contained(ifOperands, iv))
           break;

 // multiple sources to forward from.
 //
 // TODO: more forwarding can be done when support for
 // loop/conditional live-out SSA values is available.
 // TODO: do general dead store elimination for memref's. This pass
 // currently only eliminates the stores only if no other loads/uses (other
 // than dealloc) remain.
 //
-void mlir::affineScalarReplace(func::FuncOp f, DominanceInfo &domInfo,
+void mlir::affineScalarReplace(Operation *parentOp, DominanceInfo &domInfo,
                                PostDominanceInfo &postDomInfo) {
   // Load op's whose results were replaced by those forwarded from stores.
   SmallVector<Operation *, 8> opsToErase;
 
   // A list of memref's that are potentially dead / could be eliminated.
   SmallPtrSet<Value, 4> memrefsToErase;
 
   // Walk all load's and perform store to load forwarding.
-  f.walk([&](AffineReadOpInterface loadOp) {
+  parentOp->walk([&](AffineReadOpInterface loadOp) {
     if (failed(
             forwardStoreToLoad(loadOp, opsToErase, memrefsToErase, domInfo))) {
       loadCSE(loadOp, opsToErase, domInfo);
     }
   });
 
   // Erase all load op's whose results were replaced with store fwd'ed ones.
   for (auto *op : opsToErase)
     op->erase();
   opsToErase.clear();
 
   // Walk all store's and perform unused store elimination
-  f.walk([&](AffineWriteOpInterface storeOp) {
+  parentOp->walk([&](AffineWriteOpInterface storeOp) {
     findUnusedStore(storeOp, opsToErase, postDomInfo);
   });
   // Erase all store op's which don't impact the program
   for (auto *op : opsToErase)
     op->erase();
 
   // Check if the store fwd'ed memrefs are now left with only stores and can
   // thus be completely deleted. Note: the canonicalize pass should be able
LogicalResult mlir::replaceAllMemRefUsesWith(
   // Assert same elemental type.
   assert(oldMemRef.getType().cast<MemRefType>().getElementType() ==
          newMemRef.getType().cast<MemRefType>().getElementType());
 
   std::unique_ptr<DominanceInfo> domInfo;
   std::unique_ptr<PostDominanceInfo> postDomInfo;
   if (domOpFilter)
     domInfo = std::make_unique<DominanceInfo>(
-        domOpFilter->getParentOfType<func::FuncOp>());
+        domOpFilter->getParentOfType<FunctionOpInterface>());
 
   if (postDomOpFilter)
     postDomInfo = std::make_unique<PostDominanceInfo>(
-        postDomOpFilter->getParentOfType<func::FuncOp>());
+        postDomOpFilter->getParentOfType<FunctionOpInterface>());
 
   // Walk all uses of old memref; collect ops to perform replacement. We use a
   // DenseSet since an operation could potentially have multiple uses of a
   // memref (although rare), and the replacement later is going to erase ops.
   DenseSet<Operation *> opsToReplace;
   for (auto *op : oldMemRef.getUsers()) {
     // Skip this use if it's not dominated by domOpFilter.
     if (domOpFilter && !domInfo->dominates(domOpFilter, op))

mlir/test/Dialect/Affine/SuperVectorize/uniform_divergent.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=128" -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=128})" \
+// RUN:   -split-input-file | FileCheck %s
 
 // Specific tests to check vectorization of uniform/divergent values.
 
 // CHECK-LABEL: @uniform_arg
 // CHECK-SAME:  %[[in:.*]]: memref<512xf32>,
 // CHECK-SAME:  %[[uniform:.*]]: f32
 func.func @uniform_arg(%in : memref<512xf32>, %uniform : f32) {
   affine.for %i = 0 to 512 {

mlir/test/Dialect/Affine/SuperVectorize/vectorize_1d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=128 test-fastest-varying=0" -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=128 test-fastest-varying=0})" \
+// RUN:   -split-input-file | FileCheck %s
 
 // CHECK-DAG: #[[$map_id1:map[0-9]+]] = affine_map<(d0) -> (d0)>
 // CHECK-DAG: #[[$map_proj_d0d1_0:map[0-9]+]] = affine_map<(d0, d1) -> (0)>
 
 // CHECK-LABEL: func @vec1d_1
 func.func @vec1d_1(%A : memref<?x?xf32>, %B : memref<?x?x?xf32>) {
 // CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
 // CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index

mlir/test/Dialect/Affine/SuperVectorize/vectorize_2d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=4,8" | FileCheck %s -check-prefix=VECT
-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,256 test-fastest-varying=1,0" | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=4,8})" | FileCheck %s -check-prefix=VECT
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,256 test-fastest-varying=1,0})" | FileCheck %s
 
 // Permutation maps used in vectorization.
 // CHECK-DAG: #[[$map_id1:map[0-9]+]] = affine_map<(d0) -> (d0)>
 // CHECK-DAG: #[[$map_proj_d0d1_zerod1:map[0-9]+]] = affine_map<(d0, d1) -> (0, d1)>
 // CHECK-DAG: #[[$map_proj_d0d1_d0zero:map[0-9]+]] = affine_map<(d0, d1) -> (d0, 0)>
 // VECT-DAG: #[[$map_id1:map[0-9]+]] = affine_map<(d0) -> (d0)>
 // VECT-DAG: #[[$map_proj_d0d1_zerod1:map[0-9]+]] = affine_map<(d0, d1) -> (0, d1)>
 // VECT-DAG: #[[$map_proj_d0d1_d0zero:map[0-9]+]] = affine_map<(d0, d1) -> (d0, 0)>

mlir/test/Dialect/Affine/SuperVectorize/vectorize_3d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,64,256 test-fastest-varying=2,1,0" | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,64,256 test-fastest-varying=2,1,0})" | FileCheck %s
 
 func.func @vec3d(%A : memref<?x?x?xf32>) {
   %c0 = arith.constant 0 : index
   %c1 = arith.constant 1 : index
   %c2 = arith.constant 2 : index
   %0 = memref.dim %A, %c0 : memref<?x?x?xf32>
   %1 = memref.dim %A, %c1 : memref<?x?x?xf32>
   %2 = memref.dim %A, %c2 : memref<?x?x?xf32>

mlir/test/Dialect/Affine/SuperVectorize/vectorize_outer_loop_2d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,256 test-fastest-varying=2,0" | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,256 test-fastest-varying=2,0})" | FileCheck %s
 
 // Permutation maps used in vectorization.
 // CHECK: #[[map_proj_d0d1d2_d0d2:map[0-9]*]] = affine_map<(d0, d1, d2) -> (d0, d2)>
 
 func.func @vec2d(%A : memref<?x?x?xf32>) {
   %c0 = arith.constant 0 : index
   %c1 = arith.constant 1 : index
   %c2 = arith.constant 2 : index

mlir/test/Dialect/Affine/SuperVectorize/vectorize_outer_loop_transpose_2d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,256 test-fastest-varying=0,2" | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,256 test-fastest-varying=0,2})" | FileCheck %s
 
 // Permutation maps used in vectorization.
 // CHECK: #[[map_proj_d0d1d2_d2d0:map[0-9]*]] = affine_map<(d0, d1, d2) -> (d2, d0)>
 
 func.func @vec2d(%A : memref<?x?x?xf32>) {
   %c0 = arith.constant 0 : index
   %c1 = arith.constant 1 : index
   %c2 = arith.constant 2 : index

mlir/test/Dialect/Affine/SuperVectorize/vectorize_reduction.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=128 test-fastest-varying=0 vectorize-reductions=true" -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=128 test-fastest-varying=0 vectorize-reductions=true})" \
+// RUN:   -split-input-file | FileCheck %s
 
 // The inner reduction loop '%j' is vectorized.
 
 func.func @vecdim_reduction(%in: memref<256x512xf32>, %out: memref<256xf32>) {
  %cst = arith.constant 0.000000e+00 : f32
  affine.for %i = 0 to 256 {
    %final_red = affine.for %j = 0 to 512 iter_args(%red_iter = %cst) -> (f32) {
      %ld = affine.load %in[%i, %j] : memref<256x512xf32>

mlir/test/Dialect/Affine/SuperVectorize/vectorize_reduction_2d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,256 test-fastest-varying=1,0 vectorize-reductions=true" -verify-diagnostics
+// RUN: mlir-opt %s -verify-diagnostics -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,256 test-fastest-varying=1,0 vectorize-reductions=true})"
 
 // TODO: Vectorization of reduction loops along the reduction dimension is not
 //       supported for higher-rank vectors yet, so we are just checking that an
 //       error message is produced.
 
 // expected-error@+1 {{Vectorizing reductions is supported only for 1-D vectors}}
 func.func @vecdim_reduction_2d(%in: memref<256x512x1024xf32>, %out: memref<256xf32>) {
  %cst = arith.constant 0.000000e+00 : f32

mlir/test/Dialect/Affine/SuperVectorize/vectorize_transpose_2d.mlir

-// RUN: mlir-opt %s -affine-super-vectorize="virtual-vector-size=32,256 test-fastest-varying=0,1" | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-super-vectorize{virtual-vector-size=32,256 test-fastest-varying=0,1})" | FileCheck %s
 
 // Permutation maps used in vectorization.
 // CHECK-DAG: #[[map_proj_d0d1d2_d2d1:map[0-9]*]] = affine_map<(d0, d1, d2) -> (d2, d1)>
 
 func.func @vec2d(%A : memref<?x?x?xf32>) {
   %c0 = arith.constant 0 : index
   %c1 = arith.constant 1 : index
   %c2 = arith.constant 2 : index

mlir/test/Dialect/Affine/affine-data-copy.mlir

-// RUN: mlir-opt %s -split-input-file -affine-data-copy-generate="generate-dma=false fast-mem-space=0 skip-non-unit-stride-loops" | FileCheck %s
+// RUN: mlir-opt %s -split-input-file -pass-pipeline="func.func(affine-data-copy-generate{generate-dma=false fast-mem-space=0 skip-non-unit-stride-loops})" | FileCheck %s
 // Small buffer size to trigger fine copies.
-// RUN: mlir-opt %s -split-input-file -affine-data-copy-generate="generate-dma=false fast-mem-space=0 fast-mem-capacity=1" | FileCheck --check-prefix=CHECK-SMALL %s
+// RUN: mlir-opt %s -split-input-file -pass-pipeline="func.func(affine-data-copy-generate{generate-dma=false fast-mem-space=0 fast-mem-capacity=1})" | FileCheck --check-prefix=CHECK-SMALL %s
 
 // Test affine data copy with a memref filter. We use a test pass that invokes
 // affine data copy utility on the input loop nest.
 // '-test-affine-data-copy-memref-filter' passes the first memref found in an
 // affine.load op in the innermost loop as a filter.
 // RUN: mlir-opt %s -split-input-file -test-affine-data-copy='memref-filter' | FileCheck %s --check-prefix=FILTER
 // RUN: mlir-opt %s -split-input-file -test-affine-data-copy='for-memref-region' | FileCheck %s --check-prefix=MEMREF_REGION
 

mlir/test/Dialect/Affine/affine-loop-invariant-code-motion.mlir

-// RUN: mlir-opt %s -affine-loop-invariant-code-motion -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-loop-invariant-code-motion)" -split-input-file | FileCheck %s
 
 func.func @nested_loops_both_having_invariant_code() {
   %m = memref.alloc() : memref<10xf32>
   %cf7 = arith.constant 7.0 : f32
   %cf8 = arith.constant 8.0 : f32
 
   affine.for %arg0 = 0 to 10 {
     %v0 = arith.addf %cf7, %cf8 : f32

mlir/test/Dialect/Affine/affine-loop-normalize.mlir

-// RUN: mlir-opt %s -affine-loop-normalize -split-input-file | FileCheck %s
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-loop-normalize)" -split-input-file | FileCheck %s
 
 // Normalize steps to 1 and lower bounds to 0.
 
 // CHECK-DAG: [[$MAP0:#map[0-9]+]] = affine_map<(d0) -> (d0 * 3)>
 // CHECK-DAG: [[$MAP1:#map[0-9]+]] = affine_map<(d0) -> (d0 * 2 + 1)>
 // CHECK-DAG: [[$MAP2:#map[0-9]+]] = affine_map<(d0, d1) -> (d0 + d1)>
 
 // CHECK-LABEL: func @normalize_parallel()

mlir/test/Dialect/Affine/dma-generate.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -affine-data-copy-generate="generate-dma fast-mem-space=2 skip-non-unit-stride-loops" -verify-diagnostics | FileCheck %s
-// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -affine-data-copy-generate="generate-dma fast-mem-capacity=16 fast-mem-space=2" | FileCheck %s --check-prefix FAST-MEM-16KB
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -verify-diagnostics \
+// RUN:   -pass-pipeline="func.func(affine-data-copy-generate{generate-dma fast-mem-space=2 skip-non-unit-stride-loops})" | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file \
+// RUN:   -pass-pipeline="func.func(affine-data-copy-generate{generate-dma fast-mem-capacity=16 fast-mem-space=2})" | FileCheck %s --check-prefix FAST-MEM-16KB
 
 // We run most test cases with -copy-skip-non-unit-stride-loops to allow testing
 // DMA generation at inner levels easily - since the DMA generation would
 // otherwise always generate DMAs at the outermost level (default for fast mem
 // capacity is infinite). Using a specific capacity makes it harder to write
 // a test case as one would have to calculate total footprints. With
 // -copy-skip-non-unit-stride-loops, non-unit strides will always be skipped and
 // its inner loops will be traversed till a unit stride loop is found (or the

mlir/test/Dialect/Affine/loop-coalescing.mlir

-// RUN: mlir-opt -split-input-file -allow-unregistered-dialect -affine-loop-coalescing %s | FileCheck %s
+// RUN: mlir-opt -split-input-file -allow-unregistered-dialect \
+// RUN:   -pass-pipeline="func.func(affine-loop-coalescing)" %s | FileCheck %s
 
 // CHECK-LABEL: @one_3d_nest
 func.func @one_3d_nest() {
   // Capture original bounds.  Note that for zero-based step-one loops, the
   // upper bound is also the number of iterations.
   // CHECK: %[[orig_lb:.*]] = arith.constant 0
   // CHECK: %[[orig_step:.*]] = arith.constant 1
   // CHECK: %[[orig_ub_k:.*]] = arith.constant 3

mlir/test/Dialect/Affine/loop-tiling-validity.mlir

-// RUN: mlir-opt %s  -split-input-file -affine-loop-tile="tile-size=32" -verify-diagnostics | FileCheck %s
+// RUN: mlir-opt %s  -split-input-file -pass-pipeline="func.func(affine-loop-tile{tile-size=32})" \
+// RUN:   -verify-diagnostics | FileCheck %s
 
 // -----
 
 // There is no dependence violated in this case. No error should be raised.
 
 // CHECK-DAG: [[$LB:#map[0-9]+]] = affine_map<(d0) -> (d0)>
 // CHECK-DAG: [[$UB:#map[0-9]+]] = affine_map<(d0) -> (d0 + 32)>
 

mlir/test/Dialect/Affine/loop-tiling.mlir

-// RUN: mlir-opt %s -split-input-file -affine-loop-tile="tile-size=32" | FileCheck %s
-// RUN: mlir-opt %s -split-input-file -affine-loop-tile="cache-size=512" | FileCheck %s --check-prefix=MODEL
-// RUN: mlir-opt %s -split-input-file -affine-loop-tile="tile-size=32 separate" | FileCheck %s --check-prefix=SEPARATE
+// RUN: mlir-opt %s -split-input-file -pass-pipeline="func.func(affine-loop-tile{tile-size=32})" | FileCheck %s
+// RUN: mlir-opt %s -split-input-file -pass-pipeline="func.func(affine-loop-tile{cache-size=512})" | FileCheck %s --check-prefix=MODEL
+// RUN: mlir-opt %s -split-input-file -pass-pipeline="func.func(affine-loop-tile{tile-size=32 separate})" | FileCheck %s --check-prefix=SEPARATE
 
 // -----
 
 // CHECK-DAG: [[$UB:#map[0-9]+]] = affine_map<(d0) -> (d0 + 32)>
 // CHECK-DAG: [[$UB_MIN:#map[0-9]+]] = affine_map<(d0) -> (d0 + 32, 50)>
 // CHECK-DAG: [[$ID:#map[0-9]+]] = affine_map<(d0) -> (d0)>
 // CHECK-DAG: [[$ID_PLUS_21:#map[0-9]+]] = affine_map<(d0) -> (d0 + 21)>
 

mlir/test/Dialect/Affine/parallelize.mlir

-// RUN: mlir-opt %s -allow-unregistered-dialect -affine-parallelize | FileCheck %s
-// RUN: mlir-opt %s -allow-unregistered-dialect -affine-parallelize='max-nested=1' | FileCheck --check-prefix=MAX-NESTED %s
-// RUN: mlir-opt %s -allow-unregistered-dialect -affine-parallelize='parallel-reductions=1' | FileCheck --check-prefix=REDUCE %s
+// RUN: mlir-opt %s -allow-unregistered-dialect -pass-pipeline="func.func(affine-parallelize)" | FileCheck %s
+// RUN: mlir-opt %s -allow-unregistered-dialect -pass-pipeline="func.func(affine-parallelize{max-nested=1})" | FileCheck --check-prefix=MAX-NESTED %s
+// RUN: mlir-opt %s -allow-unregistered-dialect -pass-pipeline="func.func(affine-parallelize{parallel-reductions=1})" | FileCheck --check-prefix=REDUCE %s
 
 // CHECK-LABEL:    func @reduce_window_max() {
 func.func @reduce_window_max() {
   %cst = arith.constant 0.000000e+00 : f32
   %0 = memref.alloc() : memref<1x8x8x64xf32>
   %1 = memref.alloc() : memref<1x18x18x64xf32>
   affine.for %arg0 = 0 to 1 {
     affine.for %arg1 = 0 to 8 {

mlir/test/Dialect/Affine/scalrep.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-scalrep | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-scalrep)" | FileCheck %s
 
 // CHECK-DAG: [[$MAP0:#map[0-9]+]] = affine_map<(d0, d1) -> (d1 + 1)>
 // CHECK-DAG: [[$MAP1:#map[0-9]+]] = affine_map<(d0, d1) -> (d0)>
 // CHECK-DAG: [[$MAP2:#map[0-9]+]] = affine_map<(d0, d1) -> (d1)>
 // CHECK-DAG: [[$MAP3:#map[0-9]+]] = affine_map<(d0, d1) -> (d0 - 1)>
 // CHECK-DAG: [[$MAP4:#map[0-9]+]] = affine_map<(d0) -> (d0 + 1)>
 
 // CHECK-LABEL: func @simple_store_load() {

mlir/test/Dialect/Affine/simplify-structures.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -affine-simplify-structures | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -pass-pipeline="func.func(affine-simplify-structures)" | FileCheck %s
 
 // CHECK-DAG: #[[$SET_2D:.*]] = affine_set<(d0, d1) : (d0 - 100 == 0, d1 - 10 == 0, -d0 + 100 >= 0, d1 >= 0)>
 // CHECK-DAG: #[[$SET_7_11:.*]] = affine_set<(d0, d1) : (d0 * 7 + d1 * 5 + 88 == 0, d0 * 5 - d1 * 11 + 60 == 0, d0 * 11 + d1 * 7 - 24 == 0, d0 * 7 + d1 * 5 + 88 == 0)>
 
 // An external function that we will use in bodies to avoid DCE.
 func.func private @external() -> ()
 
 // CHECK-LABEL: func @test_gaussian_elimination_empty_set0() {

mlir/test/Dialect/Affine/unroll-jam.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll-jam="unroll-jam-factor=2" | FileCheck %s
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll-jam="unroll-jam-factor=4" | FileCheck --check-prefix=UJAM-FOUR %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll-jam{unroll-jam-factor=2})" | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll-jam{unroll-jam-factor=4})" | FileCheck --check-prefix=UJAM-FOUR %s
 
 // CHECK-DAG: [[$MAP_PLUS_1:#map[0-9]+]] = affine_map<(d0) -> (d0 + 1)>
 // CHECK-DAG: [[$MAP_DIV_OFFSET:#map[0-9]+]] = affine_map<()[s0] -> (((s0 - 1) floordiv 2) * 2 + 1)>
 // CHECK-DAG: [[$MAP_SYM_UB:#map[0-9]+]] = affine_map<()[s0, s1] -> (s0, s1, 1024)>
 
 // UJAM-FOUR-DAG: [[$UBMAP:#map[0-9]+]] = affine_map<()[s0] -> (s0 + 8)>
 // UJAM-FOUR-DAG: [[$MAP_PLUS_1:#map[0-9]+]] = affine_map<(d0) -> (d0 + 1)>
 // UJAM-FOUR-DAG: [[$MAP_PLUS_2:#map[0-9]+]] = affine_map<(d0) -> (d0 + 2)>

mlir/test/Dialect/Affine/unroll.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll="unroll-full" | FileCheck %s --check-prefix UNROLL-FULL
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll="unroll-full unroll-full-threshold=2" | FileCheck %s --check-prefix SHORT
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll="unroll-factor=4" | FileCheck %s --check-prefix UNROLL-BY-4
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-unroll="unroll-factor=1" | FileCheck %s --check-prefix UNROLL-BY-1
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll{unroll-full})" | FileCheck %s --check-prefix UNROLL-FULL
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll{unroll-full unroll-full-threshold=2})" | FileCheck %s --check-prefix SHORT
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll{unroll-factor=4})" | FileCheck %s --check-prefix UNROLL-BY-4
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-unroll{unroll-factor=1})" | FileCheck %s --check-prefix UNROLL-BY-1
 
 // UNROLL-FULL-DAG: [[$MAP0:#map[0-9]+]] = affine_map<(d0) -> (d0 + 1)>
 // UNROLL-FULL-DAG: [[$MAP1:#map[0-9]+]] = affine_map<(d0) -> (d0 + 2)>
 // UNROLL-FULL-DAG: [[$MAP2:#map[0-9]+]] = affine_map<(d0) -> (d0 + 3)>
 // UNROLL-FULL-DAG: [[$MAP3:#map[0-9]+]] = affine_map<(d0) -> (d0 + 4)>
 // UNROLL-FULL-DAG: [[$MAP4:#map[0-9]+]] = affine_map<(d0, d1) -> (d0 + 1)>
 // UNROLL-FULL-DAG: [[$MAP5:#map[0-9]+]] = affine_map<(d0, d1) -> (d0 + 3)>
 // UNROLL-FULL-DAG: [[$MAP6:#map[0-9]+]] = affine_map<(d0)[s0] -> (d0 + s0 + 1)>

mlir/test/Dialect/SCF/loop-unroll.mlir

 // RUN: mlir-opt %s -test-loop-unrolling='unroll-factor=2' | FileCheck %s --check-prefix UNROLL-BY-2
 // RUN: mlir-opt %s -test-loop-unrolling='unroll-factor=3' | FileCheck %s --check-prefix UNROLL-BY-3
 // RUN: mlir-opt %s -test-loop-unrolling='unroll-factor=2 loop-depth=0' | FileCheck %s --check-prefix UNROLL-OUTER-BY-2
 // RUN: mlir-opt %s -test-loop-unrolling='unroll-factor=2 loop-depth=1' | FileCheck %s --check-prefix UNROLL-INNER-BY-2
 // RUN: mlir-opt %s -test-loop-unrolling='unroll-factor=2 annotate=true' | FileCheck %s --check-prefix UNROLL-BY-2-ANNOTATE
-// RUN: mlir-opt %s --affine-loop-unroll='unroll-factor=6 unroll-up-to-factor=true' | FileCheck %s --check-prefix UNROLL-UP-TO
+// RUN: mlir-opt %s -pass-pipeline="func.func(affine-loop-unroll{unroll-factor=6 unroll-up-to-factor=true})" | FileCheck %s --check-prefix UNROLL-UP-TO
 
 func.func @dynamic_loop_unroll(%arg0 : index, %arg1 : index, %arg2 : index,
                           %arg3: memref<?xf32>) {
   %0 = arith.constant 7.0 : f32
   scf.for %i0 = %arg0 to %arg1 step %arg2 {
     memref.store %0, %arg3[%i0] : memref<?xf32>
   }
   return

mlir/test/Transforms/loop-fusion-2.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion -split-input-file | FileCheck %s
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion="fusion-maximal" -split-input-file | FileCheck %s --check-prefix=MAXIMAL
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion)" -split-input-file | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion{fusion-maximal})" -split-input-file | FileCheck %s --check-prefix=MAXIMAL
 
 // Part I of fusion tests in  mlir/test/Transforms/loop-fusion.mlir.
 // Part III of fusion tests in mlir/test/Transforms/loop-fusion-3.mlir
 // Part IV of fusion tests in mlir/test/Transforms/loop-fusion-4.mlir
 
 // -----
 
 // CHECK-LABEL: func @should_fuse_at_depth_above_loop_carried_dependence(%{{.*}}: memref<64x4xf32>, %{{.*}}: memref<64x4xf32>) {

mlir/test/Transforms/loop-fusion-3.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion -split-input-file | FileCheck %s
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion="fusion-maximal" -split-input-file | FileCheck %s --check-prefix=MAXIMAL
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion)" -split-input-file | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion{fusion-maximal})" -split-input-file | FileCheck %s --check-prefix=MAXIMAL
 
-// Part I of fusion tests in  mlir/test/Transforms/loop-fusion.mlir. 
+// Part I of fusion tests in  mlir/test/Transforms/loop-fusion.mlir.
 // Part II of fusion tests in mlir/test/Transforms/loop-fusion-2.mlir
 // Part IV of fusion tests in mlir/test/Transforms/loop-fusion-4.mlir
 
 // -----
 
 // Test case from github bug 777.
 // CHECK-LABEL: func @mul_add_0
 func.func @mul_add_0(%arg0: memref<3x4xf32>, %arg1: memref<4x3xf32>, %arg2: memref<3x3xf32>, %arg3: memref<3x3xf32>) {

mlir/test/Transforms/loop-fusion-4.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion="mode=producer" -split-input-file | FileCheck %s --check-prefix=PRODUCER-CONSUMER
-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion="fusion-maximal mode=sibling" -split-input-file | FileCheck %s --check-prefix=SIBLING-MAXIMAL
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion{mode=producer})" -split-input-file | FileCheck %s --check-prefix=PRODUCER-CONSUMER
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion{fusion-maximal mode=sibling})" -split-input-file | FileCheck %s --check-prefix=SIBLING-MAXIMAL
 
 // Part I of fusion tests in  mlir/test/Transforms/loop-fusion.mlir.
 // Part II of fusion tests in mlir/test/Transforms/loop-fusion-2.mlir
 // Part III of fusion tests in mlir/test/Transforms/loop-fusion-3.mlir
 
 // Expects fusion of producer into consumer at depth 4 and subsequent removal of
 // source loop.
 // PRODUCER-CONSUMER-LABEL: func @unflatten4d

mlir/test/Transforms/loop-fusion.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -affine-loop-fusion -split-input-file | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -pass-pipeline="func.func(affine-loop-fusion)" -split-input-file | FileCheck %s
 
 // Part II of fusion tests in  mlir/test/Transforms/loop-fusion=2.mlir.
 // Part III of fusion tests in mlir/test/Transforms/loop-fusion-3.mlir
 // Part IV of fusion tests in mlir/test/Transforms/loop-fusion-4.mlir
 
 // TODO: Add more tests:
 // *) Add nested fusion test cases when non-constant loop bound support is
 //    added to iteration domain in dependence check.

mlir/test/Transforms/pipeline-data-transfer.mlir

-// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -affine-pipeline-data-transfer | FileCheck %s
+// RUN: mlir-opt -allow-unregistered-dialect %s -split-input-file -pass-pipeline="func.func(affine-pipeline-data-transfer)" | FileCheck %s
 
 // -----
 
 // CHECK-DAG: [[$MOD_2:#map[0-9]+]] = affine_map<(d0) -> (d0 mod 2)>
 // CHECK-DAG: [[$MAP_MINUS_1:#map[0-9]+]] = affine_map<(d0) -> (d0 - 1)>
 
 // CHECK-LABEL: func @loop_nest_dma() {
 func.func @loop_nest_dma() {

utils/bazel/llvm-project-overlay/mlir/BUILD.bazel

srcs = glob([
         "lib/Dialect/Affine/Analysis/*.h",
     ]),
     hdrs = glob(["include/mlir/Dialect/Affine/Analysis/*.h"]),
     includes = ["include"],
     deps = [
         ":AffineDialect",
         ":Analysis",
         ":ArithmeticDialect",
-        ":FuncDialect",
         ":IR",
         ":Support",
         ":ViewLikeInterface",
         "//llvm:Support",
     ],
 )
 
 cc_library(
hdrs = [
         "include/mlir/Dialect/Affine/LoopUtils.h",
         "include/mlir/Dialect/Affine/Utils.h",
     ],
     includes = ["include"],
     deps = [
         ":AffineAnalysis",
         ":AffineDialect",
         ":Analysis",
-        ":FuncDialect",
         ":IR",
         ":MemRefDialect",
         ":SCFDialect",
         ":Support",
         ":TransformUtils",
         "//llvm:Support",
     ],
 )
cc_library(
     deps = [
         ":AffineAnalysis",
         ":AffineDialect",
         ":AffinePassIncGen",
         ":AffineUtils",
         ":Analysis",
         ":ArithmeticDialect",
         ":ArithmeticUtils",
-        ":FuncDialect",
         ":IR",
         ":MemRefDialect",
         ":Pass",
         ":SCFDialect",
         ":SCFUtils",
         ":Support",
         ":Transforms",
         ":VectorDialect",
cc_library(
     ]) + [":ConversionPassDetail"],
     hdrs = glob(["include/mlir/Conversion/AffineToStandard/*.h"]),
     includes = ["include"],
     deps = [
         ":AffineDialect",
         ":AffineUtils",
         ":ArithmeticDialect",
         ":ConversionPassIncGen",
-        ":FuncDialect",
         ":IR",
         ":MemRefDialect",
         ":Pass",
         ":SCFDialect",
         ":Support",
         ":Transforms",
         ":VectorDialect",
     ],