[mlir] Add operations to BlockAndValueMapping and rename it to IRMapping

[Mogball]

Summary

The patch adds operations to BlockAndValueMapping and renames it to IRMapping. When operations are cloned, old operations are mapped to the cloned operations. This allows mapping from an operation to a cloned operation. Example:

Operation *opWithRegion = ...
Operation *opInsideRegion = &opWithRegion->front().front();

IRMapping map
Operation *newOpWithRegion = opWithRegion->clone(map);
Operation *newOpInsideRegion = map.lookupOrNull(opInsideRegion);

Migration instructions:
All includes to mlir/IR/BlockAndValueMapping.h should be replaced with mlir/IR/IRMapping.h. All uses of BlockAndValueMapping need to be renamed to IRMapping.

[rriddle]

Something I'd considered in the past for this was just to evolve BlockAndValueMapping into something that could also map operations (i.e. something like an IRMapping). That feels like it would be nice here, to avoid duplicating the underlying code for cloning.

[Mogball]

IRMapping

[jpienaar]

For existing users they just need to include the new header and sed s/BlockAndValueMapping/IRMapping/g ?

[Mogball]

Yup

[mehdi_amini]

In D139665#4034784, @jpienaar wrote:

For existing users they just need to include the new header and sed s/BlockAndValueMapping/IRMapping/g ?

Can you add the instructions in the commit message?

mlir/include/mlir/IR/IRMapping.h

21

IR entities?

94

[Mogball]

review comments and update description with migration instructions

[rriddle]

Nice, LG!

[Mogball]

rebase

flang/include/flang/Optimizer/Dialect/FIRDialect.h

 //===----------------------------------------------------------------------===//
 
 #ifndef FORTRAN_OPTIMIZER_DIALECT_FIRDIALECT_H
 #define FORTRAN_OPTIMIZER_DIALECT_FIRDIALECT_H
 
 #include "mlir/IR/Dialect.h"
 
 namespace mlir {
-class BlockAndValueMapping;
+class IRMapping;
 } // namespace mlir
 
 namespace fir {
 
 /// FIR dialect
 class FIROpsDialect final : public mlir::Dialect {
 public:
   explicit FIROpsDialect(mlir::MLIRContext *ctx);
public:
   explicit FIRCodeGenDialect(mlir::MLIRContext *ctx);
   virtual ~FIRCodeGenDialect();
 
   static llvm::StringRef getDialectNamespace() { return "fircg"; }
 };
 
 /// Support for inlining on FIR.
 bool canLegallyInline(mlir::Operation *op, mlir::Region *reg, bool,
-                      mlir::BlockAndValueMapping &map);
+                      mlir::IRMapping &map);
 bool canLegallyInline(mlir::Operation *, mlir::Operation *, bool);
 
 } // namespace fir
 
 #endif // FORTRAN_OPTIMIZER_DIALECT_FIRDIALECT_H

flang/include/flang/Optimizer/Transforms/Passes.h

 #define FORTRAN_OPTIMIZER_TRANSFORMS_PASSES_H
 
 #include "flang/Optimizer/Dialect/FIROps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Pass/PassRegistry.h"
 #include <memory>
 
 namespace mlir {
-class BlockAndValueMapping;
+class IRMapping;
 class GreedyRewriteConfig;
 class Operation;
 class Pass;
 class Region;
 } // namespace mlir
 
 namespace fir {
 

flang/lib/Optimizer/Builder/HLFIRTools.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "flang/Optimizer/Builder/HLFIRTools.h"
 #include "flang/Optimizer/Builder/FIRBuilder.h"
 #include "flang/Optimizer/Builder/MutableBox.h"
 #include "flang/Optimizer/Builder/Todo.h"
 #include "flang/Optimizer/HLFIR/HLFIROps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include <optional>
 
 // Return explicit extents. If the base is a fir.box, this won't read it to
 // return the extents and will instead return an empty vector.
 static llvm::SmallVector<mlir::Value>
 getExplicitExtents(fir::FortranVariableOpInterface var) {
   llvm::SmallVector<mlir::Value> result;
   if (mlir::Value shape = var.getShape()) {

 
 hlfir::YieldElementOp
 hlfir::inlineElementalOp(mlir::Location loc, fir::FirOpBuilder &builder,
                          hlfir::ElementalOp elemental,
                          mlir::ValueRange oneBasedIndices) {
   // hlfir.elemental region is a SizedRegion<1>.
   assert(elemental.getRegion().hasOneBlock() &&
          "expect elemental region to have one block");
-  mlir::BlockAndValueMapping mapper;
+  mlir::IRMapping mapper;
   mapper.map(elemental.getIndices(), oneBasedIndices);
   mlir::Operation *newOp;
   for (auto &op : elemental.getRegion().back().getOperations())
     newOp = builder.clone(op, mapper);
   auto yield = mlir::dyn_cast_or_null<hlfir::YieldElementOp>(newOp);
   assert(yield && "last ElementalOp operation must be am hlfir.yield_element");
   return yield;
 }

flang/lib/Optimizer/Dialect/FIRDialect.cpp

struct FIRInlinerInterface : public mlir::DialectInlinerInterface {
   bool isLegalToInline(mlir::Operation *call, mlir::Operation *callable,
                        bool wouldBeCloned) const final {
     return fir::canLegallyInline(call, callable, wouldBeCloned);
   }
 
   /// This hook checks to see if the operation `op` is legal to inline into the
   /// given region `reg`.
   bool isLegalToInline(mlir::Operation *op, mlir::Region *reg,
-                       bool wouldBeCloned,
-                       mlir::BlockAndValueMapping &map) const final {
+                       bool wouldBeCloned, mlir::IRMapping &map) const final {
     return fir::canLegallyInline(op, reg, wouldBeCloned, map);
   }
 
   /// This hook is called when a terminator operation has been inlined.
   /// We handle the return (a Fortran FUNCTION) by replacing the values
   /// previously returned by the call operation with the operands of the
   /// return.
   void handleTerminator(mlir::Operation *op,

flang/lib/Optimizer/Dialect/Inliner.cpp

 
 static llvm::cl::opt<bool>
     aggressivelyInline("inline-all",
                        llvm::cl::desc("aggressively inline everything"),
                        llvm::cl::init(false));
 
 /// Should we inline the callable `op` into region `reg`?
 bool fir::canLegallyInline(mlir::Operation *, mlir::Region *, bool,
-                           mlir::BlockAndValueMapping &) {
+                           mlir::IRMapping &) {
   return aggressivelyInline;
 }
 
 bool fir::canLegallyInline(mlir::Operation *, mlir::Operation *, bool) {
   return aggressivelyInline;
 }

mlir/docs/Interfaces.md

class DialectInlinerInterface :
     public DialectInterface::Base<DialectInlinerInterface> {
 public:
   /// Returns true if the given region 'src' can be inlined into the region
   /// 'dest' that is attached to an operation registered to the current dialect.
   /// 'valueMapping' contains any remapped values from within the 'src' region.
   /// This can be used to examine what values will replace entry arguments into
   /// the 'src' region, for example.
   virtual bool isLegalToInline(Region *dest, Region *src,
-                               BlockAndValueMapping &valueMapping) const {
+                               IRMapping &valueMapping) const {
     return false;
   }
 };
 
 /// Override the inliner interface to add support for the AffineDialect to
 /// enable inlining affine operations.
 struct AffineInlinerInterface : public DialectInlinerInterface {
   /// Affine structures have specific inlining constraints.
   bool isLegalToInline(Region *dest, Region *src,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     ...
   }
 };
 
 /// Register the interface with the dialect.
 AffineDialect::AffineDialect(MLIRContext *context) ... {
   addInterfaces<AffineInlinerInterface>();
 }

 
 ```c++
 class InlinerInterface : public
     DialectInterfaceCollection<DialectInlinerInterface> {
   /// The hooks for this class mirror the hooks for the DialectInlinerInterface,
   /// with default implementations that call the hook on the interface for a
   /// given dialect.
   virtual bool isLegalToInline(Region *dest, Region *src,
-                               BlockAndValueMapping &valueMapping) const {
+                               IRMapping &valueMapping) const {
     auto *handler = getInterfaceFor(dest->getContainingOp());
     return handler ? handler->isLegalToInline(dest, src, valueMapping) : false;
   }
 };
 
 MLIRContext *ctx = ...;
 InlinerInterface interface(ctx);
 if(!interface.isLegalToInline(...))

mlir/docs/Rationale/UsageOfConst.md

 though matching and binding values themselves makes perfect sense for both const
 and non-const values. Actually fixing this would cause massive code bloat and
 complexity.
 
 Other parts of the code are just outright incorrect. For example, the operation
 cloning methods are defined on `Operation` like this:
 
 ```C++
-Operation *clone(BlockAndValueMapping &mapper, MLIRContext *context) const;
+Operation *clone(IRMapping &mapper, MLIRContext *context) const;
 
 Operation *clone(MLIRContext *context) const;
 ```
 
 While it makes sense for a clone method to be `const` conceptually (the original
 operation isn't modified) this is a violation of the model, since the returned
 operation must be mutable, and provides access to the full graph of operands as
 the original operation, violating the graph based const model we were shooting
1.  All of the const-duplicated accessors would be eliminated, e.g.
     approximately ~130 lines of code from just Operation.h alone.
 1.  Const-only predicates would be changed to be non-const, e.g.
     `Operation::isTerminator() const` would have the const removed.
 1.  Iterators and other types and functions that are templated to support
     `const` can have those template arguments removed.
 1.  Types like `OpPointer` and `ConstOpPointer` that exist solely to propagate
     const can be entirely removed from the codebase.
 1.  We can close bugs complaining about const incorrectness in the IR.
-

mlir/docs/Tutorials/Toy/Ch-4.md

bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// This hook checks to see if the given operation is legal to inline into the
   /// given region. For Toy this hook can simply return true, as all Toy
   /// operations are inlinable.
   bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     return true;
   }
 
   /// This hook cheks if the given 'src' region can be inlined into the 'dest'
   /// region. The regions here are the bodies of the callable functions. For
   /// Toy, any function can be inlined, so we simply return true.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     return true;
   }
 
   /// This hook is called when a terminator operation has been inlined. The only
   /// terminator that we have in the Toy dialect is the return
   /// operation(toy.return). We handle the return by replacing the values
   /// previously returned by the call operation with the operands of the
   /// return.

mlir/examples/toy/Ch4/mlir/Dialect.cpp

struct ToyInlinerInterface : public DialectInlinerInterface {
 
   /// All call operations within toy can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// All operations within toy can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   // All functions within toy can be inlined.
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   /// Handle the given inlined terminator(toy.return) by replacing it with a new

mlir/examples/toy/Ch5/mlir/Dialect.cpp

struct ToyInlinerInterface : public DialectInlinerInterface {
 
   /// All call operations within toy can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// All operations within toy can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   // All functions within toy can be inlined.
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   /// Handle the given inlined terminator(toy.return) by replacing it with a new

mlir/examples/toy/Ch6/mlir/Dialect.cpp

struct ToyInlinerInterface : public DialectInlinerInterface {
 
   /// All call operations within toy can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// All operations within toy can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   // All functions within toy can be inlined.
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   /// Handle the given inlined terminator(toy.return) by replacing it with a new

mlir/examples/toy/Ch7/mlir/Dialect.cpp

struct ToyInlinerInterface : public DialectInlinerInterface {
 
   /// All call operations within toy can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// All operations within toy can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   // All functions within toy can be inlined.
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   /// Handle the given inlined terminator(toy.return) by replacing it with a new

mlir/include/mlir/Dialect/Func/IR/FuncOps.td

let extraClassDeclaration = [{
 
     /// Create a deep copy of this function and all of its blocks, remapping any
     /// operands that use values outside of the function using the map that is
     /// provided (leaving them alone if no entry is present). If the mapper
     /// contains entries for function arguments, these arguments are not
     /// included in the new function. Replaces references to cloned sub-values
     /// with the corresponding value that is copied, and adds those mappings to
     /// the mapper.
-    FuncOp clone(BlockAndValueMapping &mapper);
+    FuncOp clone(IRMapping &mapper);
     FuncOp clone();
 
     /// Clone the internal blocks and attributes from this function into dest.
     /// Any cloned blocks are appended to the back of dest. This function
     /// asserts that the attributes of the current function and dest are
     /// compatible.
-    void cloneInto(FuncOp dest, BlockAndValueMapping &mapper);
+    void cloneInto(FuncOp dest, IRMapping &mapper);
 
     //===------------------------------------------------------------------===//
     // CallableOpInterface
     //===------------------------------------------------------------------===//
 
     /// Returns the region on the current operation that is callable. This may
     /// return null in the case of an external callable object, e.g. an external
     /// function.

mlir/include/mlir/Dialect/Linalg/IR/LinalgInterfaces.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef MLIR_DIALECT_LINALG_IR_LINALGINTERFACES_H_
 #define MLIR_DIALECT_LINALG_IR_LINALGINTERFACES_H_
 
 #include "mlir/Dialect/Utils/StructuredOpsUtils.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/Interfaces/DestinationStyleOpInterface.h"
 #include "mlir/Interfaces/InferTypeOpInterface.h"
 #include "mlir/Interfaces/ViewLikeInterface.h"
 
 namespace mlir {
 namespace linalg {

mlir/include/mlir/IR/BlockAndValueMapping.h

-//===- BlockAndValueMapping.h -----------------------------------*- 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
-//
-//===----------------------------------------------------------------------===//
-//
-// This file defines a utility class for maintaining a mapping for multiple
-// value types.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef MLIR_IR_BLOCKANDVALUEMAPPING_H
-#define MLIR_IR_BLOCKANDVALUEMAPPING_H
-
-#include "mlir/IR/Block.h"
-
-namespace mlir {
-// This is a utility class for mapping one set of values to another. New
-// mappings can be inserted via 'map'. Existing mappings can be
-// found via the 'lookup*' functions. There are two variants that differ only in
-// return value when an existing is not found for the provided key.
-// 'lookupOrNull' returns nullptr where as 'lookupOrDefault' will return the
-// lookup key.
-class BlockAndValueMapping {
-public:
-  /// Inserts a new mapping for 'from' to 'to'. If there is an existing mapping,
-  /// it is overwritten.
-  void map(Block *from, Block *to) { blockMap[from] = to; }
-  void map(Value from, Value to) { valueMap[from] = to; }
-
-  template <
-      typename S, typename T,
-      std::enable_if_t<!std::is_assignable<Value, S>::value &&
-                       !std::is_assignable<Block *, S>::value> * = nullptr>
-  void map(S &&from, T &&to) {
-    for (auto pair : llvm::zip(from, to))
-      map(std::get<0>(pair), std::get<1>(pair));
-  }
-
-  /// Erases a mapping for 'from'.
-  void erase(Block *from) { blockMap.erase(from); }
-  void erase(Value from) { valueMap.erase(from); }
-
-  /// Checks to see if a mapping for 'from' exists.
-  bool contains(Block *from) const { return blockMap.count(from); }
-  bool contains(Value from) const { return valueMap.count(from); }
-
-  /// Lookup a mapped value within the map. If a mapping for the provided value
-  /// does not exist then return nullptr.
-  Block *lookupOrNull(Block *from) const {
-    return lookupOrValue(from, (Block *)nullptr);
-  }
-  Value lookupOrNull(Value from) const { return lookupOrValue(from, Value()); }
-
-  /// Lookup a mapped value within the map. If a mapping for the provided value
-  /// does not exist then return the provided value.
-  Block *lookupOrDefault(Block *from) const {
-    return lookupOrValue(from, from);
-  }
-  Value lookupOrDefault(Value from) const { return lookupOrValue(from, from); }
-
-  /// Lookup a mapped value within the map. This asserts the provided value
-  /// exists within the map.
-  Block *lookup(Block *from) const { return lookupImpl(from); }
-  Value lookup(Value from) const { return lookupImpl(from); }
-
-  /// Clears all mappings held by the mapper.
-  void clear() { valueMap.clear(); }
-
-  /// Return the held value mapping.
-  const DenseMap<Value, Value> &getValueMap() const { return valueMap; }
-
-  /// Return the held block mapping.
-  const DenseMap<Block *, Block *> &getBlockMap() const { return blockMap; }
-
-private:
-  template<typename T>
-  T lookupImpl(T from) const {
-    T result = lookupOrNull(from);
-    assert(result && "expected 'from' to be contained within the map");
-    return result;
-  }
-
-  /// Utility lookupOrValue that looks up an existing key or returns the
-  /// provided value.
-  Block *lookupOrValue(Block *from, Block *value) const {
-    auto it = blockMap.find(from);
-    return it != blockMap.end() ? it->second : value;
-  }
-  Value lookupOrValue(Value from, Value value) const {
-    auto it = valueMap.find(from);
-    return it != valueMap.end() ? it->second : value;
-  }
-
-  DenseMap<Value, Value> valueMap;
-  DenseMap<Block *, Block *> blockMap;
-};
-
-} // namespace mlir
-
-#endif // MLIR_IR_BLOCKANDVALUEMAPPING_H

mlir/include/mlir/IR/Builders.h

 #define MLIR_IR_BUILDERS_H
 
 #include "mlir/IR/OpDefinition.h"
 #include "llvm/Support/Compiler.h"
 
 namespace mlir {
 
 class AffineExpr;
-class BlockAndValueMapping;
+class IRMapping;
 class UnknownLoc;
 class FileLineColLoc;
 class Type;
 class PrimitiveType;
 class IntegerType;
 class FloatType;
 class FunctionType;
 class IndexType;
public:
   /// Note: This function does not erase the operation on a successful fold.
   LogicalResult tryFold(Operation *op, SmallVectorImpl<Value> &results);
 
   /// Creates a deep copy of the specified operation, remapping any operands
   /// that use values outside of the operation using the map that is provided
   /// ( leaving them alone if no entry is present).  Replaces references to
   /// cloned sub-operations to the corresponding operation that is copied,
   /// and adds those mappings to the map.
-  Operation *clone(Operation &op, BlockAndValueMapping &mapper);
+  Operation *clone(Operation &op, IRMapping &mapper);
   Operation *clone(Operation &op);
 
   /// Creates a deep copy of this operation but keep the operation regions
   /// empty. Operands are remapped using `mapper` (if present), and `mapper` is
   /// updated to contain the results.
-  Operation *cloneWithoutRegions(Operation &op, BlockAndValueMapping &mapper) {
+  Operation *cloneWithoutRegions(Operation &op, IRMapping &mapper) {
     return insert(op.cloneWithoutRegions(mapper));
   }
   Operation *cloneWithoutRegions(Operation &op) {
     return insert(op.cloneWithoutRegions());
   }
   template <typename OpT>
   OpT cloneWithoutRegions(OpT op) {
     return cast<OpT>(cloneWithoutRegions(*op.getOperation()));

mlir/include/mlir/IR/IRMapping.h

+//===- IRMapping.h ----------------------------------------------*- 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
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines a utility class for maintaining a mapping of SSA values,
+// blocks, and operations.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_IR_IRMAPPING_H
+#define MLIR_IR_IRMAPPING_H
+
+#include "mlir/IR/Block.h"
+
+namespace mlir {
+/// This is a utility class for mapping one set of IR entities to another. New
+/// mappings can be inserted via 'map'. Existing mappings can be
+/// found via the 'lookup*' functions. There are three variants that differ only
+/// in return value when an existing is not found for the provided key: SSA
+/// values, blocks, and operations. 'lookupOrNull' returns nullptr where as
+/// 'lookupOrDefault' will return the lookup key.
+class IRMapping {
+public:
+  /// Inserts a new mapping for 'from' to 'to'. If there is an existing mapping,
+  /// it is overwritten.
+  void map(Value from, Value to) { valueMap[from] = to; }
+  void map(Block *from, Block *to) { blockMap[from] = to; }
+  void map(Operation *from, Operation *to) { operationMap[from] = to; }
+
+  template <typename S, typename T,
+            std::enable_if_t<!std::is_assignable_v<Value, S> &&
+                             !std::is_assignable_v<Block *, S> &&
+                             !std::is_assignable_v<Operation *, S>> * = nullptr>
+  void map(S &&from, T &&to) {
+    for (auto [fromValue, toValue] : llvm::zip(from, to))
+      map(fromValue, toValue);
+  }
+
+  /// Erases a mapping for 'from'.
+  template <typename T>
+  void erase(T from) {
+    getMap<T>().erase(from);
+  }
+
+  /// Checks to see if a mapping for 'from' exists.
+  template <typename T>
+  bool contains(T from) const {
+    return getMap<T>().count(from);
+  }
+
+  /// Lookup a mapped value within the map. If a mapping for the provided value
+  /// does not exist then return nullptr.
+  template <typename T>
+  auto lookupOrNull(T from) const {
+    return lookupOrValue(from, T(nullptr));
+  }
+
+  /// Lookup a mapped value within the map. If a mapping for the provided value
+  /// does not exist then return the provided value.
+  template <typename T>
+  auto lookupOrDefault(T from) const {
+    return lookupOrValue(from, from);
+  }
+
+  /// Lookup a mapped value within the map. This asserts the provided value
+  /// exists within the map.
+  template <typename T>
+  auto lookup(T from) const {
+    auto result = lookupOrNull(from);
+    assert(result && "expected 'from' to be contained within the map");
+    return result;
+  }
+
+  /// Clears all mappings held by the mapper.
+  void clear() { valueMap.clear(); }
+
+  /// Return the held value mapping.
+  const DenseMap<Value, Value> &getValueMap() const { return valueMap; }
+
+  /// Return the held block mapping.
+  const DenseMap<Block *, Block *> &getBlockMap() const { return blockMap; }
+
+  /// Return the held operation mapping.
+  const DenseMap<Operation *, Operation *> &getOperationMap() const {
+    return operationMap;
+  }
+
+private:
+  /// Return the map for the given value type.
+  template <typename T>
+  auto &getMap() const {
+    if constexpr (std::is_convertible_v<T, Value>)
+      return const_cast<DenseMap<Value, Value> &>(valueMap);
+    else if constexpr (std::is_convertible_v<T, Block *>)
+      return const_cast<DenseMap<Block *, Block *> &>(blockMap);
+    else
+      return const_cast<DenseMap<Operation *, Operation *> &>(operationMap);
+  }
+
+  /// Utility lookupOrValue that looks up an existing key or returns the
+  /// provided value.
+  template <typename T>
+  auto lookupOrValue(T from, T value) const {
+    auto &map = getMap<T>();
+    auto it = map.find(from);
+    return it != map.end() ? it->second : value;
+  }
+
+  DenseMap<Value, Value> valueMap;
+  DenseMap<Block *, Block *> blockMap;
+  DenseMap<Operation *, Operation *> operationMap;
+};
+
+} // namespace mlir
+
+#endif // MLIR_IR_IRMAPPING_H

mlir/include/mlir/IR/Operation.h

public:
   /// the options parameter.
   ///
   /// Calling this method from multiple threads is generally safe if through the
   /// process of cloning no new uses of 'Value's from outside the operation are
   /// created. Cloning an isolated-from-above operation with no operands, such
   /// as top level function operations, is therefore always safe. Using the
   /// mapper, it is possible to avoid adding uses to outside operands by
   /// remapping them to 'Value's owned by the caller thread.
-  Operation *clone(BlockAndValueMapping &mapper,
+  Operation *clone(IRMapping &mapper,
                    CloneOptions options = CloneOptions::all());
   Operation *clone(CloneOptions options = CloneOptions::all());
 
   /// Create a partial copy of this operation without traversing into attached
   /// regions. The new operation will have the same number of regions as the
   /// original one, but they will be left empty.
   /// Operands are remapped using `mapper` (if present), and `mapper` is updated
   /// to contain the results.
-  Operation *cloneWithoutRegions(BlockAndValueMapping &mapper);
+  Operation *cloneWithoutRegions(IRMapping &mapper);
 
   /// Create a partial copy of this operation without traversing into attached
   /// regions. The new operation will have the same number of regions as the
   /// original one, but they will be left empty.
   Operation *cloneWithoutRegions();
 
   /// Returns the operation block that contains this operation.
   Block *getBlock() { return block; }

mlir/include/mlir/IR/PatternMatch.h

virtual void inlineRegionBefore(Region &region, Region &parent,
                                   Region::iterator before);
   void inlineRegionBefore(Region &region, Block *before);
 
   /// Clone the blocks that belong to "region" before the given position in
   /// another region "parent". The two regions must be different. The caller is
   /// responsible for creating or updating the operation transferring flow of
   /// control to the region and passing it the correct block arguments.
   virtual void cloneRegionBefore(Region &region, Region &parent,
-                                 Region::iterator before,
-                                 BlockAndValueMapping &mapping);
+                                 Region::iterator before, IRMapping &mapping);
   void cloneRegionBefore(Region &region, Region &parent,
                          Region::iterator before);
   void cloneRegionBefore(Region &region, Block *before);
 
   /// This method replaces the uses of the results of `op` with the values in
   /// `newValues` when the provided `functor` returns true for a specific use.
   /// The number of values in `newValues` is required to match the number of
   /// results of `op`. `allUsesReplaced`, if non-null, is set to true if all of

mlir/include/mlir/IR/Region.h

 #define MLIR_IR_REGION_H
 
 #include "mlir/IR/Block.h"
 
 namespace mlir {
 class TypeRange;
 template <typename ValueRangeT>
 class ValueTypeRange;
-class BlockAndValueMapping;
+class IRMapping;
 
 /// This class contains a list of basic blocks and a link to the parent
 /// operation it is attached to.
 class Region {
 public:
   Region() = default;
   explicit Region(Operation *container);
   ~Region();
public:
   /// cloned blocks are appended to the back of dest. If the mapper
   /// contains entries for block arguments, these arguments are not included
   /// in the respective cloned block.
   ///
   /// Calling this method from multiple threads is generally safe if through the
   /// process of cloning, no new uses of 'Value's from outside the region are
   /// created. Using the mapper, it is possible to avoid adding uses to outside
   /// operands by remapping them to 'Value's owned by the caller thread.
-  void cloneInto(Region *dest, BlockAndValueMapping &mapper);
+  void cloneInto(Region *dest, IRMapping &mapper);
   /// Clone this region into 'dest' before the given position in 'dest'.
-  void cloneInto(Region *dest, Region::iterator destPos,
-                 BlockAndValueMapping &mapper);
+  void cloneInto(Region *dest, Region::iterator destPos, IRMapping &mapper);
 
   /// Takes body of another region (that region will have no body after this
   /// operation completes).  The current body of this region is cleared.
   void takeBody(Region &other) {
     dropAllReferences();
     blocks.clear();
     blocks.splice(blocks.end(), other.getBlocks());
   }

mlir/include/mlir/Interfaces/DestinationStyleOpInterface.h

 //===- DestinationStyleOpInterface.h ----------------------------*- 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 MLIR_INTERFACES_DESTINATIONSTYLEOPINTERFACE_H_
 #define MLIR_INTERFACES_DESTINATIONSTYLEOPINTERFACE_H_
 
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/IR/Value.h"
 #include "llvm/ADT/SmallVector.h"
 
 namespace mlir {
 /// OpOperand vector that implicitly converts to a Value vector.
 struct OpOperandVector : public llvm::SmallVector<OpOperand *> {
   operator SmallVector<Value>();

mlir/include/mlir/Transforms/DialectConversion.h

void inlineRegionBefore(Region &region, Region &parent,
                           Region::iterator before) override;
   using PatternRewriter::inlineRegionBefore;
 
   /// PatternRewriter hook for cloning blocks of one region into another. The
   /// given region to clone *must* not have been modified as part of conversion
   /// yet, i.e. it must be within an operation that is either in the process of
   /// conversion, or has not yet been converted.
   void cloneRegionBefore(Region &region, Region &parent,
-                         Region::iterator before,
-                         BlockAndValueMapping &mapping) override;
+                         Region::iterator before, IRMapping &mapping) override;
   using PatternRewriter::cloneRegionBefore;
 
   /// PatternRewriter hook for inserting a new operation.
   void notifyOperationInserted(Operation *op) override;
 
   /// PatternRewriter hook for updating the root operation in-place.
   /// Note: These methods only track updates to the top-level operation itself,
   /// and not nested regions. Updates to regions will still require notification

mlir/include/mlir/Transforms/InliningUtils.h

 
 #include "mlir/IR/DialectInterface.h"
 #include "mlir/IR/Location.h"
 #include "mlir/IR/Region.h"
 
 namespace mlir {
 
 class Block;
-class BlockAndValueMapping;
+class IRMapping;
 class CallableOpInterface;
 class CallOpInterface;
 class OpBuilder;
 class Operation;
 class Region;
 class TypeRange;
 class Value;
 class ValueRange;
public:
   /// 'dest' that is attached to an operation registered to the current dialect.
   /// 'wouldBeCloned' is set to true if the given 'src' region is set to be
   /// cloned during the inlining process, or false if the region is set to be
   /// moved in-place(i.e. no duplicates would be created). 'valueMapping'
   /// contains any remapped values from within the 'src' region. This can be
   /// used to examine what values will replace entry arguments into the 'src'
   /// region for example.
   virtual bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                               BlockAndValueMapping &valueMapping) const {
+                               IRMapping &valueMapping) const {
     return false;
   }
 
   /// Returns true if the given operation 'op', that is registered to this
   /// dialect, can be inlined into the given region, false otherwise.
   /// 'wouldBeCloned' is set to true if the given 'op' is set to be cloned
   /// during the inlining process, or false if the operation is set to be moved
   /// in-place(i.e. no duplicates would be created). 'valueMapping' contains any
   /// remapped values from within the 'src' region. This can be used to examine
   /// what values may potentially replace the operands to 'op'.
   virtual bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
-                               BlockAndValueMapping &valueMapping) const {
+                               IRMapping &valueMapping) const {
     return false;
   }
 
   /// This hook is invoked on an operation that contains regions. It should
   /// return true if the analyzer should recurse within the regions of this
   /// operation when computing legality and cost, false otherwise. The default
   /// implementation returns true.
   virtual bool shouldAnalyzeRecursively(Operation *op) const { return true; }
public:
 
   //===--------------------------------------------------------------------===//
   // Analysis Hooks
   //===--------------------------------------------------------------------===//
 
   virtual bool isLegalToInline(Operation *call, Operation *callable,
                                bool wouldBeCloned) const;
   virtual bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                               BlockAndValueMapping &valueMapping) const;
+                               IRMapping &valueMapping) const;
   virtual bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
-                               BlockAndValueMapping &valueMapping) const;
+                               IRMapping &valueMapping) const;
   virtual bool shouldAnalyzeRecursively(Operation *op) const;
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   virtual void handleTerminator(Operation *op, Block *newDest) const;
   virtual void handleTerminator(Operation *op,

 /// remappings for the entry arguments to the region. 'resultsToReplace'
 /// corresponds to any results that should be replaced by terminators within the
 /// inlined region. 'regionResultTypes' specifies the expected return types of
 /// the terminators in the region. 'inlineLoc' is an optional Location that, if
 /// provided, will be used to update the inlined operations' location
 /// information. 'shouldCloneInlinedRegion' corresponds to whether the source
 /// region should be cloned into the 'inlinePoint' or spliced directly.
 LogicalResult inlineRegion(InlinerInterface &interface, Region *src,
-                           Operation *inlinePoint, BlockAndValueMapping &mapper,
+                           Operation *inlinePoint, IRMapping &mapper,
                            ValueRange resultsToReplace,
                            TypeRange regionResultTypes,
                            Optional<Location> inlineLoc = std::nullopt,
                            bool shouldCloneInlinedRegion = true);
 LogicalResult inlineRegion(InlinerInterface &interface, Region *src,
                            Block *inlineBlock, Block::iterator inlinePoint,
-                           BlockAndValueMapping &mapper,
-                           ValueRange resultsToReplace,
+                           IRMapping &mapper, ValueRange resultsToReplace,
                            TypeRange regionResultTypes,
                            Optional<Location> inlineLoc = std::nullopt,
                            bool shouldCloneInlinedRegion = true);
 
 /// This function is an overload of the above 'inlineRegion' that allows for
 /// providing the set of operands ('inlinedOperands') that should be used
 /// in-favor of the region arguments when inlining.
 LogicalResult inlineRegion(InlinerInterface &interface, Region *src,

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

 
 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/Utils.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_CONVERTAFFINETOSTANDARD
 #include "mlir/Conversion/Passes.h.inc"

mlir/lib/Conversion/FuncToLLVM/FuncToLLVM.cpp

 #include "mlir/Conversion/LLVMCommon/Pattern.h"
 #include "mlir/Conversion/LLVMCommon/VectorPattern.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/LLVMIR/FunctionCallUtils.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
 #include "mlir/IR/Attributes.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributeInterfaces.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Support/LogicalResult.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/TypeSwitch.h"

mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp

 #include "mlir/Conversion/LLVMCommon/TypeConverter.h"
 #include "mlir/Conversion/MemRefToLLVM/AllocLikeConversion.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/LLVMIR/FunctionCallUtils.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Pass/Pass.h"
 #include "llvm/ADT/SmallBitVector.h"
 #include <optional>
 
 namespace mlir {
 #define GEN_PASS_DEF_MEMREFTOLLVMCONVERSIONPASS
 #include "mlir/Conversion/Passes.h.inc"
 } // namespace mlir
matchAndRewrite(memref::GenericAtomicRMWOp atomicOp, OpAdaptor adaptor,
     Value init = rewriter.create<LLVM::LoadOp>(loc, dataPtr);
     rewriter.create<LLVM::BrOp>(loc, init, loopBlock);
 
     // Prepare the body of the loop block.
     rewriter.setInsertionPointToStart(loopBlock);
 
     // Clone the GenericAtomicRMWOp region and extract the result.
     auto loopArgument = loopBlock->getArgument(0);
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     mapping.map(atomicOp.getCurrentValue(), loopArgument);
     Block &entryBlock = atomicOp.body().front();
     for (auto &nestedOp : entryBlock.without_terminator()) {
       Operation *clone = rewriter.clone(nestedOp, mapping);
       mapping.map(nestedOp.getResults(), clone->getResults());
     }
     Value result = mapping.lookup(entryBlock.getTerminator()->getOperand(0));
 
matchAndRewrite(memref::GenericAtomicRMWOp atomicOp, OpAdaptor adaptor,
     return success();
   }
 
 private:
   // Clones a segment of ops [start, end) and erases the original.
   void moveOpsRange(ValueRange oldResult, ValueRange newResult,
                     Block::iterator start, Block::iterator end,
                     ConversionPatternRewriter &rewriter) const {
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     mapping.map(oldResult, newResult);
     SmallVector<Operation *, 2> opsToErase;
     for (auto it = start; it != end; ++it) {
       rewriter.clone(*it, mapping);
       opsToErase.push_back(&*it);
     }
     for (auto *it : opsToErase)
       rewriter.eraseOp(it);

mlir/lib/Conversion/SCFToControlFlow/SCFToControlFlow.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h"
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/MLIRContext.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_SCFTOCONTROLFLOW
 #include "mlir/Conversion/Passes.h.inc"

mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp

 #include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/IR/AffineExpr.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/Passes.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/Sequence.h"
 #include "llvm/Support/Debug.h"
 

 ///   the new for loop.
 /// In either case,
 /// - append the instructions from the loops body to worklist, in reverse order.
 /// To note the end of the current scope in case a loop or conditional was
 /// inserted, a sentinel (the `gpu.launch` operation) is inserted into the
 /// worklist. This signals the processor of the worklist to pop the rewriter
 /// one scope-level up.
 static LogicalResult processParallelLoop(
-    ParallelOp parallelOp, gpu::LaunchOp launchOp,
-    BlockAndValueMapping &cloningMap, SmallVectorImpl<Operation *> &worklist,
+    ParallelOp parallelOp, gpu::LaunchOp launchOp, IRMapping &cloningMap,
+    SmallVectorImpl<Operation *> &worklist,
     DenseMap<gpu::Processor, Value> &bounds, PatternRewriter &rewriter) {
   // TODO: Verify that this is a valid GPU mapping.
   // processor ids: 0-2 block [x/y/z], 3-5 -> thread [x/y/z], 6-> sequential
   ArrayAttr mapping =
       parallelOp->getAttrOfType<ArrayAttr>(gpu::getMappingAttrName());
 
   // TODO: Support reductions.
   if (!mapping || parallelOp.getNumResults() != 0)
Value constantOne =
       rewriter.create<arith::ConstantIndexOp>(parallelOp.getLoc(), 1);
   gpu::LaunchOp launchOp = rewriter.create<gpu::LaunchOp>(
       parallelOp.getLoc(), constantOne, constantOne, constantOne, constantOne,
       constantOne, constantOne);
   rewriter.setInsertionPointToEnd(&launchOp.getBody().front());
   rewriter.create<gpu::TerminatorOp>(loc);
   rewriter.setInsertionPointToStart(&launchOp.getBody().front());
 
-  BlockAndValueMapping cloningMap;
+  IRMapping cloningMap;
   llvm::DenseMap<gpu::Processor, Value> launchBounds;
   SmallVector<Operation *, 16> worklist;
   if (failed(processParallelLoop(parallelOp, launchOp, cloningMap, worklist,
                                  launchBounds, rewriter)))
     return failure();
 
   // Whether we have seen any side-effects. Reset when leaving an inner scope.
   bool seenSideeffects = false;

mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp

 //===- ShapeToStandard.cpp - conversion from Shape to Standard dialect ----===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/ShapeToStandard/ShapeToStandard.h"
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Shape/IR/Shape.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "llvm/ADT/STLExtras.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_CONVERTSHAPETOSTANDARD
 #include "mlir/Conversion/Passes.h.inc"
ReduceOpConverter::matchAndRewrite(shape::ReduceOp op, OpAdaptor adaptor,
   auto loop = rewriter.create<scf::ForOp>(
       loc, zero, rank, one, op.getInitVals(),
       [&](OpBuilder &b, Location loc, Value iv, ValueRange args) {
         Value extent = b.create<tensor::ExtractOp>(loc, adaptor.getShape(), iv);
 
         SmallVector<Value, 2> mappedValues{iv, extent};
         mappedValues.append(args.begin(), args.end());
 
-        BlockAndValueMapping mapping;
+        IRMapping mapping;
         Block *reduceBody = op.getBody();
         mapping.map(reduceBody->getArguments(), mappedValues);
         for (auto &nested : reduceBody->without_terminator())
           b.clone(nested, mapping);
 
         SmallVector<Value, 2> mappedResults;
         for (auto result : reduceBody->getTerminator()->getOperands())
           mappedResults.push_back(mapping.lookup(result));

mlir/lib/Conversion/TosaToSCF/TosaToSCF.cpp

 //===- TosaToSCF.cpp - Lowering Tosa to SCF Dialect -----------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // These rewriters lower from the Tosa to the SCF dialect.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Conversion/TosaToSCF/TosaToSCF.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Tosa/IR/TosaOps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 
 using namespace mlir;
 using namespace tosa;
 
 static void inlineIfCase(Region &srcRegion, Region &dstRegion,
                          OperandRange operands, PatternRewriter &rewriter) {

mlir/lib/Dialect/Affine/IR/AffineOps.cpp

 //===- AffineOps.cpp - MLIR Affine Operations -----------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/IR/AffineValueMap.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/IR/AffineExprVisitor.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Interfaces/ShapedOpInterfaces.h"
 #include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/ScopeExit.h"
 #include "llvm/ADT/SmallBitVector.h"

 
 /// Checks if `value` known to be a legal affine dimension or symbol in `src`
 /// region remains legal if the operation that uses it is inlined into `dest`
 /// with the given value mapping. `legalityCheck` is either `isValidDim` or
 /// `isValidSymbol`, depending on the value being required to remain a valid
 /// dimension or symbol.
 static bool
 remainsLegalAfterInline(Value value, Region *src, Region *dest,
-                        const BlockAndValueMapping &mapping,
+                        const IRMapping &mapping,
                         function_ref<bool(Value, Region *)> legalityCheck) {
   // If the value is a valid dimension for any other reason than being
   // a top-level value, it will remain valid: constants get inlined
   // with the function, transitive affine applies also get inlined and
   // will be checked themselves, etc.
   if (!isTopLevelValue(value, src))
     return true;
 
remainsLegalAfterInline(Value value, Region *src, Region *dest,
   return matchPattern(value.getDefiningOp(), m_Constant(&operandCst)) ||
          isDimLikeOp;
 }
 
 /// Checks if all values known to be legal affine dimensions or symbols in `src`
 /// remain so if their respective users are inlined into `dest`.
 static bool
 remainsLegalAfterInline(ValueRange values, Region *src, Region *dest,
-                        const BlockAndValueMapping &mapping,
+                        const IRMapping &mapping,
                         function_ref<bool(Value, Region *)> legalityCheck) {
   return llvm::all_of(values, [&](Value v) {
     return remainsLegalAfterInline(v, src, dest, mapping, legalityCheck);
   });
 }
 
 /// Checks if an affine read or write operation remains legal after inlining
 /// from `src` to `dest`.
 template <typename OpTy>
 static bool remainsLegalAfterInline(OpTy op, Region *src, Region *dest,
-                                    const BlockAndValueMapping &mapping) {
+                                    const IRMapping &mapping) {
   static_assert(llvm::is_one_of<OpTy, AffineReadOpInterface,
                                 AffineWriteOpInterface>::value,
                 "only ops with affine read/write interface are supported");
 
   AffineMap map = op.getAffineMap();
   ValueRange dimOperands = op.getMapOperands().take_front(map.getNumDims());
   ValueRange symbolOperands =
       op.getMapOperands().take_back(map.getNumSymbols());
   if (!remainsLegalAfterInline(
           dimOperands, src, dest, mapping,
           static_cast<bool (*)(Value, Region *)>(isValidDim)))
     return false;
   if (!remainsLegalAfterInline(
           symbolOperands, src, dest, mapping,
           static_cast<bool (*)(Value, Region *)>(isValidSymbol)))
     return false;
   return true;
 }
 
 /// Checks if an affine apply operation remains legal after inlining from `src`
 /// to `dest`.
 //  Use "unused attribute" marker to silence clang-tidy warning stemming from
 //  the inability to see through "llvm::TypeSwitch".
 template <>
-bool LLVM_ATTRIBUTE_UNUSED
-remainsLegalAfterInline(AffineApplyOp op, Region *src, Region *dest,
-                        const BlockAndValueMapping &mapping) {
+bool LLVM_ATTRIBUTE_UNUSED remainsLegalAfterInline(AffineApplyOp op,
+                                                   Region *src, Region *dest,
+                                                   const IRMapping &mapping) {
   // If it's a valid dimension, we need to check that it remains so.
   if (isValidDim(op.getResult(), src))
     return remainsLegalAfterInline(
         op.getMapOperands(), src, dest, mapping,
         static_cast<bool (*)(Value, Region *)>(isValidDim));
 
   // Otherwise it must be a valid symbol, check that it remains so.
   return remainsLegalAfterInline(
struct AffineInlinerInterface : public DialectInlinerInterface {
   // Analysis Hooks
   //===--------------------------------------------------------------------===//
 
   /// Returns true if the given region 'src' can be inlined into the region
   /// 'dest' that is attached to an operation registered to the current dialect.
   /// 'wouldBeCloned' is set if the region is cloned into its new location
   /// rather than moved, indicating there may be other users.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     // We can inline into affine loops and conditionals if this doesn't break
     // affine value categorization rules.
     Operation *destOp = dest->getParentOp();
     if (!isa<AffineParallelOp, AffineForOp, AffineIfOp>(destOp))
       return false;
 
     // Multi-block regions cannot be inlined into affine constructs, all of
     // which require single-block regions.
bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
     }
 
     return true;
   }
 
   /// Returns true if the given operation 'op', that is registered to this
   /// dialect, can be inlined into the given region, false otherwise.
   bool isLegalToInline(Operation *op, Region *region, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     // Always allow inlining affine operations into a region that is marked as
     // affine scope, or into affine loops and conditionals. There are some edge
     // cases when inlining *into* affine structures, but that is handled in the
     // other 'isLegalToInline' hook above.
     Operation *parentOp = region->getParentOp();
     return parentOp->hasTrait<OpTrait::AffineScope>() ||
            isa<AffineForOp, AffineParallelOp, AffineIfOp>(parentOp);
   }

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

 #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/IR/AffineValueMap.h"
 #include "mlir/Dialect/Affine/LoopUtils.h"
 #include "mlir/Dialect/Affine/Utils.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_AFFINELOOPTILING
 #include "mlir/Dialect/Affine/Passes.h.inc"
 } // namespace mlir
 

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

 
 #include "mlir/Dialect/Affine/Analysis/AffineAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.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/IRMapping.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/Support/CommandLine.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_AFFINELOOPUNROLLANDJAM
 #include "mlir/Dialect/Affine/Passes.h.inc"
 } // namespace mlir
 

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

 #include "mlir/Dialect/Affine/Analysis/LoopAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/NestedMatcher.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/Utils.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/Vector/IR/VectorOps.h"
 #include "mlir/Dialect/Vector/Utils/VectorUtils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Support/LLVM.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/Support/Debug.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_AFFINEVECTORIZE
 #include "mlir/Dialect/Affine/Passes.h.inc"
struct VectorizationState {
 
   // Used to build and insert all the new operations created. The insertion
   // point is preserved and updated along the vectorization process.
   OpBuilder builder;
 
   // Maps input scalar operations to their vector counterparts.
   DenseMap<Operation *, Operation *> opVectorReplacement;
   // Maps input scalar values to their vector counterparts.
-  BlockAndValueMapping valueVectorReplacement;
+  IRMapping valueVectorReplacement;
   // Maps input scalar values to their new scalar counterparts in the vector
   // loop nest.
-  BlockAndValueMapping valueScalarReplacement;
+  IRMapping valueScalarReplacement;
   // Maps results of reduction loops to their new scalar counterparts.
   DenseMap<Value, Value> loopResultScalarReplacement;
 
   // Maps the newly created vector loops to their vector dimension.
   DenseMap<Operation *, unsigned> vecLoopToVecDim;
 
   // Maps the new vectorized loops to the corresponding vector masks if it is
   // required.

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/LoopAnalysis.h"
 #include "mlir/Dialect/Affine/Analysis/Utils.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Affine/LoopUtils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
 #define DEBUG_TYPE "loop-fusion-utils"
 
 using namespace mlir;
 

 
 /// Fuses 'srcForOp' into 'dstForOp' with destination loop block insertion point
 /// and source slice loop bounds specified in 'srcSlice'.
 void mlir::fuseLoops(AffineForOp srcForOp, AffineForOp dstForOp,
                      const ComputationSliceState &srcSlice,
                      bool isInnermostSiblingInsertion) {
   // Clone 'srcForOp' into 'dstForOp' at 'srcSlice->insertPoint'.
   OpBuilder b(srcSlice.insertPoint->getBlock(), srcSlice.insertPoint);
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   b.clone(*srcForOp, mapper);
 
   // Update 'sliceLoopNest' upper and lower bounds from computed 'srcSlice'.
   SmallVector<AffineForOp, 4> sliceLoops;
   for (unsigned i = 0, e = srcSlice.ivs.size(); i < e; ++i) {
     auto loopIV = mapper.lookupOrNull(srcSlice.ivs[i]);
     if (!loopIV)
       continue;

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

 #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/IRMapping.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"
static AffineForOp generateShiftedLoop(
   assert(lbMap.getNumInputs() == lbOperands.size());
   assert(ubMap.getNumInputs() == ubOperands.size());
 
   auto loopChunk = b.create<AffineForOp>(srcForOp.getLoc(), lbOperands, lbMap,
                                          ubOperands, ubMap, srcForOp.getStep());
   auto loopChunkIV = loopChunk.getInductionVar();
   auto srcIV = srcForOp.getInductionVar();
 
-  BlockAndValueMapping operandMap;
+  IRMapping operandMap;
 
   auto bodyBuilder = OpBuilder::atBlockTerminator(loopChunk.getBody());
   for (const auto &it : llvm::drop_begin(opGroupQueue, offset)) {
     uint64_t shift = it.first;
     auto ops = it.second;
     // All 'same shift' operations get added with their operands being
     // remapped to results of cloned operations, and their IV used remapped.
     // Generate the remapping if the shift is not zero: remappedIV = newIV -
static void generateUnrolledLoop(
   // Keep a pointer to the last non-terminator operation in the original block
   // so that we know what to clone (since we are doing this in-place).
   Block::iterator srcBlockEnd = std::prev(loopBodyBlock->end(), 2);
 
   // Unroll the contents of 'forOp' (append unrollFactor - 1 additional copies).
   SmallVector<Value, 4> lastYielded(yieldedValues);
 
   for (unsigned i = 1; i < unrollFactor; i++) {
-    BlockAndValueMapping operandMap;
+    IRMapping operandMap;
 
     // Prepare operand map.
     operandMap.map(iterArgs, lastYielded);
 
     // If the induction variable is used, create a remapping to the value for
     // this unrolled instance.
     if (!forOpIV.use_empty()) {
       Value ivUnroll = ivRemapFn(i, forOpIV, builder);
if (forOp.getLowerBoundMap().getNumResults() != 1 ||
       return failure();
     if (failed(generateCleanupLoopForUnroll(forOp, unrollJamFactor)))
       assert(false && "cleanup loop lower bound map for single result lower "
                       "and upper bound maps can always be determined");
   }
 
   // `operandMaps[i - 1]` carries old->new operand mapping for the ith unrolled
   // iteration. There are (`unrollJamFactor` - 1) iterations.
-  SmallVector<BlockAndValueMapping, 4> operandMaps(unrollJamFactor - 1);
+  SmallVector<IRMapping, 4> operandMaps(unrollJamFactor - 1);
 
   // For any loop with iter_args, replace it with a new loop that has
   // `unrollJamFactor` copies of its iterOperands, iter_args and yield
   // operands.
   SmallVector<AffineForOp, 4> newLoopsWithIterArgs;
   OpBuilder builder(forOp.getContext());
   for (AffineForOp oldForOp : loopsWithIterArgs) {
     SmallVector<Value, 4> dupIterOperands, dupIterArgs, dupYieldOperands;
for (auto loop : inputNest) {
     AffineForOp fullTileLoop = createCanonicalizedAffineForOp(
         b, loop.getLoc(), lbVmap.getOperands(), lbVmap.getAffineMap(),
         ubVmap.getOperands(), ubVmap.getAffineMap());
     b = OpBuilder::atBlockTerminator(fullTileLoop.getBody());
     fullTileLoops.push_back(fullTileLoop);
   }
 
   // Add the body for the full tile loop nest.
-  BlockAndValueMapping operandMap;
+  IRMapping operandMap;
   for (const auto &loopEn : llvm::enumerate(inputNest))
     operandMap.map(loopEn.value().getInductionVar(),
                    fullTileLoops[loopEn.index()].getInductionVar());
   b = OpBuilder::atBlockTerminator(fullTileLoops.back().getBody());
   for (auto &op : inputNest.back().getBody()->without_terminator())
     b.clone(op, operandMap);
   return success();
 }

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

 #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/Arith/Utils/Utils.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/IRMapping.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 
 #define DEBUG_TYPE "affine-utils"
 
 using namespace mlir;
 using namespace presburger;
 
static AffineIfOp hoistAffineIfOp(AffineIfOp ifOp, Operation *hoistOverOp) {
   if (hoistOverOp == ifOp)
     return ifOp;
 
   // Create the hoisted 'if' first. Then, clone the op we are hoisting over for
   // the else block. Then drop the else block of the original 'if' in the 'then'
   // branch while promoting its then block, and analogously drop the 'then'
   // block of the original 'if' from the 'else' branch while promoting its else
   // block.
-  BlockAndValueMapping operandMap;
+  IRMapping operandMap;
   OpBuilder b(hoistOverOp);
   auto hoistedIfOp = b.create<AffineIfOp>(ifOp.getLoc(), ifOp.getIntegerSet(),
                                           ifOp.getOperands(),
                                           /*elseBlock=*/true);
 
   // Create a clone of hoistOverOp to use for the else branch of the hoisted
   // conditional. The else block may get optimized away if empty.
   Operation *hoistOverOpClone = nullptr;

mlir/lib/Dialect/Arith/IR/ArithDialect.cpp

 
 namespace {
 /// This class defines the interface for handling inlining for arithmetic
 /// dialect operations.
 struct ArithInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// All arithmetic dialect ops can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 void arith::ArithDialect::initialize() {
   addOperations<
 #define GET_OP_LIST

mlir/lib/Dialect/Async/IR/Async.cpp

 //===- Async.cpp - MLIR Async Operations ----------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Async/IR/Async.h"
 
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/DialectImplementation.h"
 #include "mlir/IR/FunctionImplementation.h"
+#include "mlir/IR/IRMapping.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/ADT/TypeSwitch.h"
 
 using namespace mlir;
 using namespace mlir::async;
 
 #include "mlir/Dialect/Async/IR/AsyncOpsDialect.cpp.inc"
 

mlir/lib/Dialect/Async/Transforms/AsyncParallelFor.cpp

 #include "mlir/Dialect/Async/Passes.h"
 
 #include "PassDetail.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Async/IR/Async.h"
 #include "mlir/Dialect/Async/Transforms.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include <utility>
 
return [&, loopIdx](OpBuilder &nestedBuilder, Location loc, Value iv,
                                workLoopBuilder(loopIdx + 1));
         }
 
         b.create<scf::YieldOp>(loc);
         return;
       }
 
       // Copy the body of the parallel op into the inner-most loop.
-      BlockAndValueMapping mapping;
+      IRMapping mapping;
       mapping.map(op.getInductionVars(), computeBlockInductionVars);
       mapping.map(computeFuncType.captures, captures);
 
       for (auto &bodyOp : op.getLoopBody().getOps())
         b.clone(bodyOp, mapping);
     };
   };
 

mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp

 
 #include "PassDetail.h"
 #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Async/IR/Async.h"
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/SetVector.h"
 #include "llvm/Support/Debug.h"
 #include <optional>
 
// Prepare for coroutine conversion by creating the body of the function.
     SmallVector<Value, 4> unwrappedOperands(numOperands);
     for (size_t i = 0; i < numOperands; ++i) {
       Value operand = func.getArgument(numDependencies + i);
       unwrappedOperands[i] = builder.create<AwaitOp>(loc, operand).getResult();
     }
 
     // Map from function inputs defined above the execute op to the function
     // arguments.
-    BlockAndValueMapping valueMapping;
+    IRMapping valueMapping;
     valueMapping.map(functionInputs, func.getArguments());
     valueMapping.map(execute.getBodyRegion().getArguments(), unwrappedOperands);
 
     // Clone all operations from the execute operation body into the outlined
     // function body.
     for (Operation &op : execute.getBodyRegion().getOps())
       builder.clone(op, valueMapping);
   }

mlir/lib/Dialect/Bufferization/IR/BufferizableOpInterface.cpp

 //===- BufferizableOpInterface.cpp - Bufferizable Ops  ---=----------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Bufferization/IR/BufferizableOpInterface.h"
 #include "mlir/Dialect/Bufferization/IR/Bufferization.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/IR/AsmState.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Interfaces/ControlFlowInterfaces.h"
 #include "llvm/Support/Debug.h"
 
 //===----------------------------------------------------------------------===//
 // BufferizableOpInterface
LogicalResult BufferizationOptions::createMemCpy(OpBuilder &b, Location loc,
   if (memCpyFn)
     return (*memCpyFn)(b, loc, from, to);
 
   b.create<memref::CopyOp>(loc, from, to);
   return success();
 }
 
 //===----------------------------------------------------------------------===//
-// Bufferization-specific BlockAndValueMapping support with debugging.
+// Bufferization-specific IRMapping support with debugging.
 //===----------------------------------------------------------------------===//
 
 bool bufferization::isFunctionArgument(Value value) {
   auto bbArg = value.dyn_cast<BlockArgument>();
   if (!bbArg)
     return false;
   return isa<func::FuncOp>(bbArg.getOwner()->getParentOp());
 }

mlir/lib/Dialect/Bufferization/IR/BufferizationDialect.cpp

 // Bufferization Dialect Interfaces
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct BufferizationInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// Operations in Bufferization dialect are always legal to inline.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // Bufferization Dialect
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/ControlFlow/IR/ControlFlowOps.cpp

 //===- ControlFlowOps.cpp - ControlFlow Operations ------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/CommonFolders.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/BuiltinTypes.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/APFloat.h"
struct ControlFlowInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
   ~ControlFlowInlinerInterface() override = default;
 
   /// All control flow operations can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   /// ControlFlow terminator operations don't really need any special handing.
   void handleTerminator(Operation *op, Block *newDest) const final {}
 };
 } // namespace
 

mlir/lib/Dialect/Func/IR/FuncOps.cpp

 //===- FuncOps.cpp - Func Dialect Operations ------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/FunctionImplementation.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/APFloat.h"
struct FuncInlinerInterface : public DialectInlinerInterface {
 
   /// All call operations can be inlined.
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// All operations can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   /// All functions can be inlined.
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   //===--------------------------------------------------------------------===//
   // Transformation Hooks
   //===--------------------------------------------------------------------===//
 
   /// Handle the given inlined terminator by replacing it with a new operation

 void FuncOp::print(OpAsmPrinter &p) {
   function_interface_impl::printFunctionOp(
       p, *this, /*isVariadic=*/false, getFunctionTypeAttrName(),
       getArgAttrsAttrName(), getResAttrsAttrName());
 }
 
 /// Clone the internal blocks from this function into dest and all attributes
 /// from this function to dest.
-void FuncOp::cloneInto(FuncOp dest, BlockAndValueMapping &mapper) {
+void FuncOp::cloneInto(FuncOp dest, IRMapping &mapper) {
   // Add the attributes of this function to dest.
   llvm::MapVector<StringAttr, Attribute> newAttrMap;
   for (const auto &attr : dest->getAttrs())
     newAttrMap.insert({attr.getName(), attr.getValue()});
   for (const auto &attr : (*this)->getAttrs())
     newAttrMap.insert({attr.getName(), attr.getValue()});
 
   auto newAttrs = llvm::to_vector(llvm::map_range(
       newAttrMap, [](std::pair<StringAttr, Attribute> attrPair) {
         return NamedAttribute(attrPair.first, attrPair.second);
       }));
   dest->setAttrs(DictionaryAttr::get(getContext(), newAttrs));
 
   // Clone the body.
   getBody().cloneInto(&dest.getBody(), mapper);
 }
 
 /// Create a deep copy of this function and all of its blocks, remapping
 /// any operands that use values outside of the function using the map that is
 /// provided (leaving them alone if no entry is present). Replaces references
 /// to cloned sub-values with the corresponding value that is copied, and adds
 /// those mappings to the mapper.
-FuncOp FuncOp::clone(BlockAndValueMapping &mapper) {
+FuncOp FuncOp::clone(IRMapping &mapper) {
   // Create the new function.
   FuncOp newFunc = cast<FuncOp>(getOperation()->cloneWithoutRegions());
 
   // If the function has a body, then the user might be deleting arguments to
   // the function by specifying them in the mapper. If so, we don't add the
   // argument to the input type vector.
   if (!isExternal()) {
     FunctionType oldType = getFunctionType();
if (!isExternal()) {
     }
   }
 
   /// Clone the current function into the new one and return it.
   cloneInto(newFunc, mapper);
   return newFunc;
 }
 FuncOp FuncOp::clone() {
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   return clone(mapper);
 }
 
 //===----------------------------------------------------------------------===//
 // ReturnOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult ReturnOp::verify() {

mlir/lib/Dialect/GPU/IR/GPUDialect.cpp

 
 namespace {
 /// This class defines the interface for handling inlining with gpu
 /// operations.
 struct GPUInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// All gpu dialect ops can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 void GPUDialect::initialize() {
   addTypes<AsyncTokenType>();
   addTypes<MMAMatrixType>();

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/GPU/TransformOps/GPUTransformOps.h"
 #include "mlir/Dialect/PDL/IR/PDL.h"
 #include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Transform/IR/TransformDialect.h"
 #include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
 #include "mlir/Dialect/Transform/IR/TransformUtils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 
 using namespace mlir;
 using namespace mlir::gpu;
 using namespace mlir::transform;
 
 /// Check if given mapping attributes are one of the desired attributes
 static DiagnosedSilenceableFailure
 checkAttributeType(ArrayRef<DeviceMappingAttrInterface> threadMappingAttributes,
SmallVector<Value> gridDimValues = scf::ForeachThreadOp::getValuesSortedByKey(
       blockMapping, numBlocks, comparator);
   for (Value v : gridDimValues)
     gridDims.push_back(v.getDefiningOp<arith::ConstantIndexOp>().value());
 
   // Step 3. Generate the blockIds using the provided generator and map the
   // induction variables to the newly created ops.
   SmallVector<Value> blockOps;
   blockIdGenerator(rewriter, foreachThreadOp, blockOps);
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   for (auto [blockIdx, blockDim] :
        llvm::zip(foreachThreadOp.getThreadIndices(), blockMapping)) {
     bvm.map(blockIdx,
             blockOps[static_cast<int64_t>(
                 blockDim.cast<DeviceMappingAttrInterface>().getMappingId())]);
   }
 
   // Step 4. Move the body of foreachThreadOp.
static DiagnosedSilenceableFailure rewriteOneForeachThreadToGpuThreads(
 
   // Step 3. Create the gpu.thread ops and map the induction variables to the
   // newly created ops.
   IndexType indexType = rewriter.getIndexType();
   SmallVector<Value> threadOps{
       rewriter.create<ThreadIdOp>(loc, indexType, Dimension::x),
       rewriter.create<ThreadIdOp>(loc, indexType, Dimension::y),
       rewriter.create<ThreadIdOp>(loc, indexType, Dimension::z)};
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   for (auto [blockIdx, blockDim] :
        llvm::zip(foreachThreadOp.getThreadIndices(), threadMapping)) {
     bvm.map(
         blockIdx,
         threadOps[blockDim.cast<DeviceMappingAttrInterface>().getMappingId()]);
   }
 
   // Step 4. Maybe create conditionals to predicate the region.

mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/GPU/Transforms/Passes.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Pass/Pass.h"
 
 using namespace mlir;
 
 namespace {
 
 struct GpuAllReduceRewriter {
private:
   /// block is expected to have 2 arguments. The gpu.yield return the
   /// accumulated value of the same type.
   AccumulatorFactory getFactory(Region &body) {
     return AccumulatorFactory([&](Value lhs, Value rhs) {
       Block *block = rewriter.getInsertionBlock();
       Block *split = rewriter.splitBlock(block, rewriter.getInsertionPoint());
 
       // Insert accumulator body between split block.
-      BlockAndValueMapping mapping;
+      IRMapping mapping;
       mapping.map(body.getArgument(0), lhs);
       mapping.map(body.getArgument(1), rhs);
       rewriter.cloneRegionBefore(body, *split->getParent(),
                                  split->getIterator(), mapping);
 
       // Add branch before inserted body, into body.
       block = block->getNextNode();
       create<cf::BranchOp>(block, ValueRange());

mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/GPU/Transforms/Passes.h"
 
 #include "mlir/Dialect/Async/IR/Async.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/GPU/Transforms/Utils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/TypeSwitch.h"
 
 namespace mlir {
LogicalResult rewriteAsyncOp(gpu::AsyncOpInterface asyncOp) {
     resultTypes.reserve(1 + op->getNumResults());
     copy(op->getResultTypes(), std::back_inserter(resultTypes));
     resultTypes.push_back(tokenType);
     auto *newOp = Operation::create(op->getLoc(), op->getName(), resultTypes,
                                     op->getOperands(), op->getAttrDictionary(),
                                     op->getSuccessors(), op->getNumRegions());
 
     // Clone regions into new op.
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     for (auto pair : llvm::zip_first(op->getRegions(), newOp->getRegions()))
       std::get<0>(pair).cloneInto(&std::get<1>(pair), mapping);
 
     // Replace the op with the async clone.
     auto results = newOp->getResults();
     currentToken = results.back();
     builder.insert(newOp);
     op->replaceAllUsesWith(results.drop_back());
async::ExecuteOp addExecuteResults(async::ExecuteOp executeOp,
   transform(results, std::back_inserter(resultTypes),
             [](Value value) { return value.getType(); });
 
   // Clone executeOp with the extra results.
   OpBuilder builder(executeOp);
   auto newOp = builder.create<async::ExecuteOp>(
       executeOp.getLoc(), TypeRange{resultTypes}.drop_front() /*drop token*/,
       executeOp.getDependencies(), executeOp.getBodyOperands());
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   newOp.getRegion().getBlocks().clear();
   executeOp.getRegion().cloneInto(&newOp.getRegion(), mapper);
 
   // Replace executeOp with cloned one.
   executeOp.getOperation()->replaceAllUsesWith(
       newOp.getResults().drop_back(results.size()));
   executeOp.erase();
 

mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp

 #include "mlir/AsmParser/AsmParser.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/DLTI/DLTI.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/GPU/Transforms/Utils.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/SymbolTable.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include <limits>
 
 namespace mlir {
 #define GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS
for (auto dim : {gpu::Dimension::x, gpu::Dimension::y, gpu::Dimension::z})
     values.push_back(builder.create<OpTy>(loc, builder.getIndexType(), dim));
 }
 
 /// Adds operations generating block/thread ids and grid/block dimensions at the
 /// beginning of the `launchFuncOpBody` region. Add mapping from argument in
 /// entry block of `launchOpBody`, to the corresponding result value of the
 /// added operations.
 static void injectGpuIndexOperations(Location loc, Region &launchFuncOpBody,
-                                     Region &launchOpBody,
-                                     BlockAndValueMapping &map) {
+                                     Region &launchOpBody, IRMapping &map) {
   OpBuilder builder(loc->getContext());
   Block &firstBlock = launchOpBody.front();
   builder.setInsertionPointToStart(&launchFuncOpBody.front());
   SmallVector<Value, 12> indexOps;
   createForAllDimensions<gpu::BlockIdOp>(builder, loc, indexOps);
   createForAllDimensions<gpu::ThreadIdOp>(builder, loc, indexOps);
   createForAllDimensions<gpu::GridDimOp>(builder, loc, indexOps);
   createForAllDimensions<gpu::BlockDimOp>(builder, loc, indexOps);
for (Value operand : sinkCandidates) {
     Operation *operandOp = operand.getDefiningOp();
     if (!operandOp)
       continue;
     extractBeneficiaryOps(operandOp, sinkCandidates, toBeSunk, availableValues,
                           isSinkingBeneficiary);
   }
 
   // Insert operations so that the defs get cloned before uses.
-  BlockAndValueMapping map;
+  IRMapping map;
   OpBuilder builder(launchOpBody);
   for (Operation *op : toBeSunk) {
     Operation *clonedOp = builder.clone(*op, map);
     // Only replace uses within the launch op.
     for (auto pair : llvm::zip(op->getResults(), clonedOp->getResults()))
       replaceAllUsesInRegionWith(std::get<0>(pair), std::get<1>(pair),
                                  launchOp.getBody());
   }
if (auto blockBounds =
           maybeConstantDimsAttr(launchOp.getBlockSizeOperandValues()))
     outlinedFunc->setAttr(gpu::GPUFuncOp::getKnownBlockSizeAttrName(),
                           blockBounds);
   if (auto gridBounds =
           maybeConstantDimsAttr(launchOp.getGridSizeOperandValues()))
     outlinedFunc->setAttr(gpu::GPUFuncOp::getKnownGridSizeAttrName(),
                           gridBounds);
 
-  BlockAndValueMapping map;
+  IRMapping map;
 
   // Map the arguments corresponding to the launch parameters like blockIdx,
   // threadIdx, etc.
   Region &outlinedFuncBody = outlinedFunc.getBody();
   injectGpuIndexOperations(loc, outlinedFuncBody, launchOpBody, map);
 
   // Map arguments from gpu.launch region to the arguments of the gpu.func
   // operation.

mlir/lib/Dialect/LLVMIR/IR/LLVMDialect.cpp

 // DialectInlinerInterface
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct LLVMInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// Conservatively only allow inlining of pure ops.
-  bool isLegalToInline(Operation *op, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *op, Region *, bool, IRMapping &) const final {
     return isPure(op);
   }
 };
 } // end anonymous namespace
 
 //===----------------------------------------------------------------------===//
 // LLVMDialect initialization, type parsing, and registration.
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/Linalg/IR/LinalgDialect.cpp

 namespace {
 
 struct LinalgInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   // We don't have any special restrictions on what can be inlined into
   // destination regions (e.g. while/conditional bodies). Always allow it.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     return true;
   }
   // Operations in Linalg dialect are always legal to inline.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
   // Handle the given inlined terminator by replacing it with a new operation
   // as necessary. Required when the region has only one block.
   void handleTerminator(Operation *op,
                         ArrayRef<Value> valuesToRepl) const final {}
 };
 

mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp

static Operation *tileAndFuseFirstExtractUseThroughContainingOpBlockArgument(
   if (failed(tensor::getOrCreateDestinations(
           rewriter, tileableProducer->getLoc(), tileableProducer,
           destinationTensors))) {
     diag.attachNote(tileableProducer->getLoc())
         << "failed to get destination tensors for: " << *tileableProducer;
     return nullptr;
   }
 
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   bvm.map(destinationTensors[resultNumber], bbArg);
   auto tileableProducerClone =
       cast<TilingInterface>(rewriter.clone(*tileableProducer, bvm));
   auto scopeGuard =
       llvm::make_scope_exit([&]() { rewriter.eraseOp(tileableProducerClone); });
 
   // Tile the producer.
   FailureOr<Value> tiledProducer =

mlir/lib/Dialect/Linalg/Transforms/DropUnitDims.cpp

LogicalResult matchAndRewrite(GenericOp genericOp,
     auto newOp = rewriter.create<GenericOp>(
         loc, genericOp.getResultTypes(), newInputOperands, newOutputOperands,
         newIndexingMaps, genericOp.getIteratorTypesArray(),
         /*bodyBuild=*/nullptr, linalg::getPrunedAttributeList(genericOp));
 
     Region &region = newOp.getRegion();
     Block *block = new Block();
     region.push_back(block);
-    BlockAndValueMapping mapper;
+    IRMapping mapper;
     OpBuilder::InsertionGuard guard(rewriter);
     rewriter.setInsertionPointToStart(block);
     for (auto bbarg : genericOp.getRegionInputArgs())
       mapper.map(bbarg, block->addArgument(bbarg.getType(), loc));
 
     for (OpOperand *op : candidates) {
       BlockArgument bbarg = genericOp.getMatchingBlockArgument(op);
       mapper.map(bbarg, block->addArgument(bbarg.getType(), loc));

mlir/lib/Dialect/Linalg/Transforms/ElementwiseOpFusion.cpp

static void generateFusedElementwiseOpRegion(
     unsigned nloops, llvm::SmallDenseSet<int> &preservedProducerResults) {
   auto producer = cast<GenericOp>(fusedOperand->get().getDefiningOp());
   auto consumer = cast<GenericOp>(fusedOperand->getOwner());
   // Build the region of the fused op.
   Block &producerBlock = producer->getRegion(0).front();
   Block &consumerBlock = consumer->getRegion(0).front();
   Block *fusedBlock = new Block();
   fusedOp.getRegion().push_back(fusedBlock);
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   OpBuilder::InsertionGuard guard(rewriter);
   rewriter.setInsertionPointToStart(fusedBlock);
 
   // 2. Add an index operation for every fused loop dimension and use the
   // `consumerToProducerLoopsMap` to map the producer indices.
   if (producer.hasIndexSemantics()) {
     // Add an index operation for every fused loop dimension.
     unsigned numFusedOpLoops =
for (OpOperand *opOperand : genericOp.getDpsInputOperands()) {
           genericOp.getIteratorTypes(),
           /*doc=*/nullptr,
           /*library_call=*/nullptr);
 
       // Map the block argument corresponding to the replaced argument with the
       // scalar constant.
       Region &region = genericOp->getRegion(0);
       Block &entryBlock = *region.begin();
-      BlockAndValueMapping mapping;
+      IRMapping mapping;
       mapping.map(entryBlock.getArgument(opOperand->getOperandNumber()),
                   scalarConstant);
       Region &fusedRegion = fusedOp->getRegion(0);
       rewriter.cloneRegionBefore(region, fusedRegion, fusedRegion.begin(),
                                  mapping);
       rewriter.replaceOp(genericOp, fusedOp->getResults());
       return success();
     }

mlir/lib/Dialect/Linalg/Transforms/HoistPadding.cpp

FailureOr<Value> mlir::linalg::hoistPaddingOnTensors(
   //      deeper in the stack.
   //   2. Create a GenericOp if `transposeVector` is non-empty.
   //   3. Create a InsertSliceOp at the top of the stack.
   //   4. Iteratively pop and yield the result of the InsertSliceOp across
   //      the cloned loops.
   SmallVector<Value> clonedLoopIvs, leadingPackedTensorIndexings;
   clonedLoopIvs.reserve(nPackedLoops);
   leadingPackedTensorIndexings.reserve(nPackedLoops);
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   // Stack step 1. iteratively clone loops and push `packedTensor`.
   for (Operation *op : analysis.backwardSlice) {
     // Specifically sit out in the extract_slice(packedTensor) case: this is the
     // piece we seek to replace.
     if (auto sliceOp = dyn_cast<tensor::ExtractSliceOp>(op))
       if (bvm.lookupOrDefault(sliceOp.getSource()) == packedTensor)
         continue;
     // Clone all operations except it is a loop.

mlir/lib/Dialect/Linalg/Transforms/Loops.cpp

 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/Linalg/IR/Linalg.h"
 #include "mlir/Dialect/Linalg/Transforms/Transforms.h"
 #include "mlir/Dialect/Linalg/Utils/Utils.h"
 #include "mlir/Dialect/SCF/Transforms/Transforms.h"
 #include "mlir/Dialect/SCF/Utils/AffineCanonicalizationUtils.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/FoldUtils.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/ADT/TypeSwitch.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_LINALGLOWERTOAFFINELOOPS

 }
 
 template <typename LoadOpTy, typename StoreOpTy, typename OpType>
 static void inlineRegionAndEmitStore(OpBuilder &b, Location loc, OpType op,
                                      ArrayRef<Value> indexedValues,
                                      ArrayRef<SmallVector<Value>> indexing,
                                      ArrayRef<Value> outputBuffers) {
   auto &block = op->getRegion(0).front();
-  BlockAndValueMapping map;
+  IRMapping map;
   map.map(block.getArguments(), indexedValues);
   for (auto &op : block.without_terminator()) {
     auto *newOp = b.clone(op, map);
     map.map(op.getResults(), newOp->getResults());
   }
 
   Operation *terminator = block.getTerminator();
   for (OpOperand &operand : terminator->getOpOperands()) {

mlir/lib/Dialect/Linalg/Transforms/Tiling.cpp

for (OpOperand *initOperand : destinationStyleOp.getDpsInitOperands()) {
       outOffsets[reductionDim] = foreachThreadOp.getThreadIndices().front();
       // TODO: use SubsetExtractOpInterface once it is available.
       tiledDpsInitOperands.push_back(b.create<tensor::ExtractSliceOp>(
           loc, initOperand->get().getType().cast<RankedTensorType>(),
           destBbArgs[destNum], outOffsets, sizes, strides));
     }
 
     // 4.b. Clone the op and update init operands.
-    // We cannot use a BlockAndValueMapping here because it can replace
+    // We cannot use a IRMapping here because it can replace
     // different OpOperands with the same value.
     Operation *clonedOp = b.clone(*op.getOperation());
     b.updateRootInPlace(clonedOp, [&]() {
       for (auto [initOperandPtr, tiledInitValue] : llvm::zip_equal(
                cast<DestinationStyleOpInterface>(clonedOp).getDpsInitOperands(),
                tiledDpsInitOperands)) {
         initOperandPtr->set(tiledInitValue);
       }

mlir/lib/Dialect/Linalg/Transforms/TilingInterfaceImpl.cpp

static SmallVector<Value> getIndicesForAccess(OpBuilder &b, Location loc,
   return indices;
 }
 
 /// Method to inline the payload of a `linalgOp` given the iteration space
 /// point and values for the arguments of the payload.
 static LogicalResult inlinePayload(OpBuilder &b, LinalgOp linalgOp,
                                    ValueRange ivs, ValueRange argValues) {
   Block *body = linalgOp.getBlock();
-  BlockAndValueMapping map;
+  IRMapping map;
   map.map(body->getArguments(), argValues);
   for (auto &op : body->without_terminator()) {
     if (auto indexOp = dyn_cast<IndexOp>(&op)) {
       map.map(indexOp.getResult(), ivs[indexOp.getDim()]);
       continue;
     }
     b.clone(op, map);
   }
SmallVector<utils::IteratorType> newIteratorTypes =
         linalgOp.getIteratorTypesArray();
     newIteratorTypes[reductionDims[0]] = utils::IteratorType::parallel;
     SmallVector<AffineMap> newMaps = linalgOp.getIndexingMapsArray();
     newMaps.back() = AffineMap::get(newMaps.back().getNumDims(), 0, outputExpr,
                                     linalgOp.getContext());
     auto genericOp =
         b.create<GenericOp>(loc, TypeRange({out.getType()}), tiledOperands,
                             ValueRange({out}), newMaps, newIteratorTypes);
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     op->getRegion(0).cloneInto(&genericOp.getRegion(),
                                genericOp.getRegion().begin(), mapping);
     return genericOp.getOperation();
   }
 
   Operation *mergeReductions(Operation *op, OpBuilder &b, Location loc,
                              ValueRange partialReduce,
                              ArrayRef<int> reductionDims) const {

mlir/lib/Dialect/Linalg/Transforms/Transforms.cpp

Value GeneralizePadOpPattern::createFillOrGenerateOp(
   auto padValue = padOp.getConstantPaddingValue();
   if (padValue)
     return rewriter.create<FillOp>(padOp.getLoc(), padValue, dest).result();
 
   // Fill could not be optimized: Lower to tensor::GenerateOp with region.
   auto generateOp = rewriter.create<tensor::GenerateOp>(
       padOp.getLoc(), padOp.getResultType(), dynSizes);
   // Copy region to new op.
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   padOp.getRegion().cloneInto(&generateOp.getRegion(), bvm);
   return generateOp;
 }
 
 LogicalResult
 GeneralizePadOpPattern::matchAndRewrite(tensor::PadOp padOp,
                                         PatternRewriter &rewriter) const {
   // Given an OpFoldResult, return an index-typed value.

mlir/lib/Dialect/Linalg/Transforms/Vectorization.cpp

 
 // Custom vectorization precondition function type. This is intented to be used
 // with CustomVectorizationHook. Returns success if the corresponding custom
 // hook can vectorize the op.
 using CustomVectorizationPrecondition =
     std::function<LogicalResult(Operation *, bool)>;
 
 // Custom vectorization function type. Produce a vector form of Operation*
-// assuming all its vectorized operands are already in the BlockAndValueMapping.
+// assuming all its vectorized operands are already in the IRMapping.
 // Return nullptr if the Operation cannot be vectorized.
-using CustomVectorizationHook = std::function<VectorizationResult(
-    Operation *, const BlockAndValueMapping &)>;
+using CustomVectorizationHook =
+    std::function<VectorizationResult(Operation *, const IRMapping &)>;
 
 /// Helper function to vectorize the terminator of a `linalgOp`. New result
 /// vector values are appended to `newResults`. Return
 /// VectorizationStatus::NoReplace to signal the vectorization algorithm that it
 /// should not try to map produced operations and instead return the results
 /// using the `newResults` vector making them available to the vectorization
 /// algorithm for RAUW. This function is meant to be used as a
 /// CustomVectorizationHook.
 static VectorizationResult
 vectorizeLinalgYield(RewriterBase &rewriter, Operation *op,
-                     const BlockAndValueMapping &bvm, VectorizationState &state,
+                     const IRMapping &bvm, VectorizationState &state,
                      LinalgOp linalgOp, SmallVectorImpl<Value> &newResults) {
   auto yieldOp = dyn_cast<linalg::YieldOp>(op);
   if (!yieldOp)
     return VectorizationResult{VectorizationStatus::Failure, nullptr};
   for (const auto &output : llvm::enumerate(yieldOp.getValues())) {
     // TODO: Scan for an opportunity for reuse.
     // TODO: use a map.
     Value vectorValue = bvm.lookup(output.value());

 ///    offset = extractOp.indices[0]
 ///    for (i = 1; i < numIndices; i++)
 ///      offset = extractOp.dimSize[i] * offset + extractOp.indices[i];
 ///
 /// For tensor<45 x 80 x 15 x f32> and index [1, 2, 3], this leads to:
 ///  offset = ( ( 1 ) * 80 +  2 ) * 15  + 3
 static Value
 calculateGatherOffset(OpBuilder &b, tensor::ExtractOp extractOp,
-                      const BlockAndValueMapping &bvm,
+                      const IRMapping &bvm,
                       const SmallVectorImpl<int64_t> &targetShape) {
   // The vector of indices for GatherOp should be shaped as the output vector
   auto indexVecType = VectorType::get(targetShape, b.getIndexType());
   auto loc = extractOp.getLoc();
 
   Value offset = b.create<vector::BroadcastOp>(
       loc, indexVecType, bvm.lookup(extractOp.getIndices()[0]));
 
calculateGatherOffset(OpBuilder &b, tensor::ExtractOp extractOp,
 
   return offset;
 }
 
 /// Helper function to vectorize the tensor.extract operations. Returns
 /// VectorizationStatus::NewOp to signal the vectorization algorithm that it
 /// should map the produced operations. This function is meant to be used as a
 /// CustomVectorizationHook.
-static VectorizationResult
-vectorizeTensorExtract(RewriterBase &rewriter, Operation *op, LinalgOp linalgOp,
-                       const BlockAndValueMapping &bvm) {
+static VectorizationResult vectorizeTensorExtract(RewriterBase &rewriter,
+                                                  Operation *op,
+                                                  LinalgOp linalgOp,
+                                                  const IRMapping &bvm) {
   tensor::ExtractOp extractOp = dyn_cast<tensor::ExtractOp>(op);
   if (!extractOp)
     return VectorizationResult{VectorizationStatus::Failure, nullptr};
   auto loc = extractOp.getLoc();
 
   // Compute the static loop sizes of the extract op.
   auto targetShape = linalgOp.computeStaticLoopSizes();
 
static VectorizationResult vectorizeTensorExtract(RewriterBase &rewriter,
 
   return VectorizationResult{VectorizationStatus::NewOp, gatherOp};
 }
 
 /// Emit reduction operations if the shapes of the value to reduce is different
 /// that the result shape.
 // Note: this is a true builder that notifies the OpBuilder listener.
 // TODO: Consider moving as a static helper on the ReduceOp.
-static Operation *reduceIfNeeded(OpBuilder &b, LinalgOp linalgOp,
-                                 Operation *op, Value reduceValue,
-                                 Value initialValue,
-                                 const BlockAndValueMapping &bvm) {
+static Operation *reduceIfNeeded(OpBuilder &b, LinalgOp linalgOp, Operation *op,
+                                 Value reduceValue, Value initialValue,
+                                 const IRMapping &bvm) {
   Value reduceVec = bvm.lookup(reduceValue);
   Value outputVec = bvm.lookup(initialValue);
   auto reduceType = reduceVec.getType().dyn_cast<VectorType>();
   auto outputType = outputVec.getType().dyn_cast<VectorType>();
   // Reduce only if needed as the value may already have been reduce for
   // contraction vectorization.
   if (!reduceType ||
       (outputType && reduceType.getShape() == outputType.getShape()))

 ///
 /// This function assumes all operands of `op` have been vectorized and are in
 /// the `bvm` mapping. As a consequence, this function is meant to be called  on
 /// a topologically-sorted list of ops.
 /// This function does not update `bvm` but returns a VectorizationStatus that
 /// instructs the caller what `bvm` update needs to occur.
 static VectorizationResult
 vectorizeOneOp(RewriterBase &rewriter, LinalgOp linalgOp, Operation *op,
-               const BlockAndValueMapping &bvm,
+               const IRMapping &bvm,
                ArrayRef<CustomVectorizationHook> customVectorizationHooks) {
   LDBG("vectorize op " << *op << "\n");
 
   // 1. Try to apply any CustomVectorizationHook.
   if (!customVectorizationHooks.empty()) {
     for (auto &customFunc : customVectorizationHooks) {
       VectorizationResult result = customFunc(op, bvm);
       if (result.status == VectorizationStatus::Failure)

 vectorizeAsLinalgGeneric(RewriterBase &rewriter, VectorizationState &state,
                          LinalgOp linalgOp,
                          SmallVectorImpl<Value> &newResults) {
   LDBG("Vectorizing operation as linalg generic\n");
   Block *block = linalgOp.getBlock();
 
   // 2. Values defined above the region can only be broadcast for now. Make them
   // map to themselves.
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   SetVector<Value> valuesSet;
   mlir::getUsedValuesDefinedAbove(linalgOp->getRegion(0), valuesSet);
   bvm.map(valuesSet.getArrayRef(), valuesSet.getArrayRef());
 
   if (linalgOp.getNumDpsInits() == 0)
     return failure();
 
   // 3. Turn all BBArgs into vector.transfer_read / load.
LDBG("New vectorized bbarg(" << bbarg.getArgNumber() << "): " << readValue
                                  << "\n");
     bvm.map(bbarg, readValue);
     bvm.map(opOperand->get(), readValue);
   }
 
   SmallVector<CustomVectorizationHook> hooks;
   // 4a. Register CustomVectorizationHook for yieldOp.
   CustomVectorizationHook vectorizeYield =
-      [&](Operation *op,
-          const BlockAndValueMapping &bvm) -> VectorizationResult {
+      [&](Operation *op, const IRMapping &bvm) -> VectorizationResult {
     return vectorizeLinalgYield(rewriter, op, bvm, state, linalgOp, newResults);
   };
   hooks.push_back(vectorizeYield);
 
   // 4b. Register CustomVectorizationHook for indexOp.
   CustomVectorizationHook vectorizeIndex =
-      [&](Operation *op,
-          const BlockAndValueMapping &bvm) -> VectorizationResult {
+      [&](Operation *op, const IRMapping &bvm) -> VectorizationResult {
     return vectorizeLinalgIndex(rewriter, op, linalgOp);
   };
   hooks.push_back(vectorizeIndex);
 
   // 4c. Register CustomVectorizationHook for extractOp.
   CustomVectorizationHook vectorizeExtract =
-      [&](Operation *op,
-          const BlockAndValueMapping &bvm) -> VectorizationResult {
+      [&](Operation *op, const IRMapping &bvm) -> VectorizationResult {
     return vectorizeTensorExtract(rewriter, op, linalgOp, bvm);
   };
   hooks.push_back(vectorizeExtract);
 
   // 5. Iteratively call `vectorizeOneOp` to each op in the slice.
   for (Operation &op : block->getOperations()) {
     VectorizationResult result =
         vectorizeOneOp(rewriter, linalgOp, &op, bvm, hooks);

mlir/lib/Dialect/Math/IR/MathDialect.cpp

 
 namespace {
 /// This class defines the interface for handling inlining with math
 /// operations.
 struct MathInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   /// All operations within math ops can be inlined.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 void mlir::math::MathDialect::initialize() {
   addOperations<
 #define GET_OP_LIST
 #include "mlir/Dialect/Math/IR/MathOps.cpp.inc"
       >();
   addInterfaces<MathInlinerInterface>();
 }

mlir/lib/Dialect/MemRef/IR/MemRefDialect.cpp

 //===----------------------------------------------------------------------===//
 // MemRefDialect Dialect Interfaces
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct MemRefInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     return true;
   }
   bool isLegalToInline(Operation *, Region *, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 void mlir::memref::MemRefDialect::initialize() {
   addOperations<
 #define GET_OP_LIST

mlir/lib/Dialect/SCF/IR/SCF.cpp

 //===- SCF.cpp - Structured Control Flow Operations -----------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Arith/Utils/Utils.h"
 #include "mlir/Dialect/Bufferization/IR/Bufferization.h"
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/DeviceMappingInterface.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/FunctionInterfaces.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/TypeSwitch.h"
 
 using namespace mlir;
 using namespace mlir::scf;
 
 #include "mlir/Dialect/SCF/IR/SCFOpsDialect.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // SCFDialect Dialect Interfaces
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct SCFInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
   // We don't have any special restrictions on what can be inlined into
   // destination regions (e.g. while/conditional bodies). Always allow it.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     return true;
   }
   // Operations in scf dialect are always legal to inline since they are
   // pure.
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
   // Handle the given inlined terminator by replacing it with a new operation
   // as necessary. Required when the region has only one block.
   void handleTerminator(Operation *op,
                         ArrayRef<Value> valuesToRepl) const final {
     auto retValOp = dyn_cast<scf::YieldOp>(op);
     if (!retValOp)
LogicalResult matchAndRewrite(scf::ForOp forOp,
                                 PatternRewriter &rewriter) const final {
     bool canonicalize = false;
     Block &block = forOp.getRegion().front();
     auto yieldOp = cast<scf::YieldOp>(block.getTerminator());
 
     // An internal flat vector of block transfer
     // arguments `newBlockTransferArgs` keeps the 1-1 mapping of original to
     // transformed block argument mappings. This plays the role of a
-    // BlockAndValueMapping for the particular use case of calling into
+    // IRMapping for the particular use case of calling into
     // `mergeBlockBefore`.
     SmallVector<bool, 4> keepMask;
     keepMask.reserve(yieldOp.getNumOperands());
     SmallVector<Value, 4> newBlockTransferArgs, newIterArgs, newYieldValues,
         newResultValues;
     newBlockTransferArgs.reserve(1 + forOp.getNumIterOperands());
     newBlockTransferArgs.push_back(Value()); // iv placeholder with null value
     newIterArgs.reserve(forOp.getNumIterOperands());

 
 namespace {
 // Collapse loop dimensions that perform a single iteration.
 struct CollapseSingleIterationLoops : public OpRewritePattern<ParallelOp> {
   using OpRewritePattern<ParallelOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(ParallelOp op,
                                 PatternRewriter &rewriter) const override {
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     // Compute new loop bounds that omit all single-iteration loop dimensions.
     SmallVector<Value, 2> newLowerBounds;
     SmallVector<Value, 2> newUpperBounds;
     SmallVector<Value, 2> newSteps;
     newLowerBounds.reserve(op.getLowerBound().size());
     newUpperBounds.reserve(op.getUpperBound().size());
     newSteps.reserve(op.getStep().size());
     for (auto [lowerBound, upperBound, step, iv] :
LogicalResult matchAndRewrite(ParallelOp op,
     if (!op.getInitVals().empty() || !innerOp.getInitVals().empty())
       return failure();
 
     auto bodyBuilder = [&](OpBuilder &builder, Location /*loc*/,
                            ValueRange iterVals, ValueRange) {
       Block &innerBody = innerOp.getLoopBody().front();
       assert(iterVals.size() ==
              (outerBody.getNumArguments() + innerBody.getNumArguments()));
-      BlockAndValueMapping mapping;
+      IRMapping mapping;
       mapping.map(outerBody.getArguments(),
                   iterVals.take_front(outerBody.getNumArguments()));
       mapping.map(innerBody.getArguments(),
                   iterVals.take_back(innerBody.getNumArguments()));
       for (Operation &op : innerBody.without_terminator())
         builder.clone(op, mapping);
     };
 

mlir/lib/Dialect/SCF/Transforms/LoopPipelining.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/SCF/Transforms/Patterns.h"
 #include "mlir/Dialect/SCF/Transforms/Transforms.h"
 #include "mlir/Dialect/SCF/Utils/Utils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/MapVector.h"
 
 using namespace mlir;
 using namespace mlir::scf;
 
Value offset = rewriter.create<arith::ConstantIndexOp>(
       loc, (maxStage - useStage) * step);
   Value iv = rewriter.create<arith::AddIOp>(loc, newForIv, offset);
   replaceInOp(newOp, newForIv, iv);
   rewriter.setInsertionPointAfter(newOp);
 }
 
 /// If the value is a loop carried value coming from stage N + 1 remap, it will
 /// become a direct use.
-static void updateIterArgUses(RewriterBase &rewriter, BlockAndValueMapping &bvm,
+static void updateIterArgUses(RewriterBase &rewriter, IRMapping &bvm,
                               Operation *newOp, ForOp oldForOp, ForOp newForOp,
                               unsigned useStage,
                               const DenseMap<Operation *, unsigned> &stages) {
 
   for (unsigned i = 0; i < oldForOp.getNumRegionIterArgs(); i++) {
     Value yieldedVal = oldForOp.getBody()->getTerminator()->getOperand(i);
     Operation *dep = yieldedVal.getDefiningOp();
     if (!dep)
       continue;
     auto stageDep = stages.find(dep);
     if (stageDep == stages.end() || stageDep->second == useStage)
       continue;
     if (stageDep->second != useStage + 1)
       continue;
     Value replacement = bvm.lookup(yieldedVal);
     replaceInOp(newOp, newForOp.getRegionIterArg(i), replacement);
   }
 }
 
 /// For operands defined in a previous stage we need to remap it to use the
 /// correct region argument. We look for the right version of the Value based
 /// on the stage where it is used.
 static void updateCrossStageUses(
-    RewriterBase &rewriter, Operation *newOp, BlockAndValueMapping &bvm,
-    ForOp newForOp, unsigned useStage,
-    const DenseMap<Operation *, unsigned> &stages,
+    RewriterBase &rewriter, Operation *newOp, IRMapping &bvm, ForOp newForOp,
+    unsigned useStage, const DenseMap<Operation *, unsigned> &stages,
     const llvm::DenseMap<std::pair<Value, unsigned>, unsigned> &loopArgMap) {
   // Because we automatically cloned the sub-regions, there's no simple way
   // to walk the nested regions in pairs of (oldOps, newOps), so we just
   // traverse the set of remapped loop arguments, filter which ones are
   // relevant, and replace any uses.
   for (auto [remapPair, newIterIdx] : loopArgMap) {
     auto [crossArgValue, stageIdx] = remapPair;
     Operation *def = crossArgValue.getDefiningOp();
const llvm::MapVector<Value, LoopPipelinerInternal::LiverangeInfo>
         &crossStageValues,
     const llvm::DenseMap<std::pair<Value, unsigned>, unsigned> &loopArgMap,
     PatternRewriter &rewriter) {
   valueMapping.clear();
 
   // Create the kernel, we clone instruction based on the order given by
   // user and remap operands coming from a previous stages.
   rewriter.setInsertionPoint(newForOp.getBody(), newForOp.getBody()->begin());
-  BlockAndValueMapping mapping;
+  IRMapping mapping;
   mapping.map(forOp.getInductionVar(), newForOp.getInductionVar());
   for (const auto &arg : llvm::enumerate(forOp.getRegionIterArgs())) {
     mapping.map(arg.value(), newForOp.getRegionIterArgs()[arg.index()]);
   }
   SmallVector<Value> predicates(maxStage + 1, nullptr);
   if (!peelEpilogue) {
     // Create a predicate for each stage except the last stage.
     for (unsigned i = 0; i < maxStage; i++) {

mlir/lib/Dialect/SCF/Transforms/LoopRangeFolding.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SCF/Transforms/Passes.h"
 
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/SCF/Transforms/Transforms.h"
 #include "mlir/Dialect/SCF/Utils/Utils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_SCFFORLOOPRANGEFOLDING
 #include "mlir/Dialect/SCF/Transforms/Passes.h.inc"
 } // namespace mlir
 
 using namespace mlir;
 using namespace mlir::scf;
while (true) {
       Operation *user = *indVar.getUsers().begin();
       if (!isa<arith::AddIOp, arith::MulIOp>(user))
         break;
 
       if (!llvm::all_of(user->getOperands(), canBeFolded))
         break;
 
       OpBuilder b(op);
-      BlockAndValueMapping lbMap;
+      IRMapping lbMap;
       lbMap.map(indVar, op.getLowerBound());
-      BlockAndValueMapping ubMap;
+      IRMapping ubMap;
       ubMap.map(indVar, op.getUpperBound());
-      BlockAndValueMapping stepMap;
+      IRMapping stepMap;
       stepMap.map(indVar, op.getStep());
 
       if (isa<arith::AddIOp>(user)) {
         Operation *lbFold = b.clone(*user, lbMap);
         Operation *ubFold = b.clone(*user, ubMap);
 
         op.setLowerBound(lbFold->getResult(0));
         op.setUpperBound(ubFold->getResult(0));

mlir/lib/Dialect/SCF/Transforms/LoopSpecialization.cpp

 #include "mlir/Dialect/Affine/Analysis/AffineStructures.h"
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/SCF/Transforms/Transforms.h"
 #include "mlir/Dialect/SCF/Utils/AffineCanonicalizationUtils.h"
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
 #include "mlir/IR/AffineExpr.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/ADT/DenseMap.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_SCFFORLOOPPEELING
 #define GEN_PASS_DEF_SCFFORLOOPSPECIALIZATION
 #define GEN_PASS_DEF_SCFPARALLELLOOPSPECIALIZATION
for (AffineExpr expr : minOp.getMap().getResults()) {
         minConstant = std::min(minConstant, constantIndex.getValue());
     }
     if (minConstant == std::numeric_limits<int64_t>::max())
       return;
     constantIndices.push_back(minConstant);
   }
 
   OpBuilder b(op);
-  BlockAndValueMapping map;
+  IRMapping map;
   Value cond;
   for (auto bound : llvm::zip(op.getUpperBound(), constantIndices)) {
     Value constant =
         b.create<arith::ConstantIndexOp>(op.getLoc(), std::get<1>(bound));
     Value cmp = b.create<arith::CmpIOp>(op.getLoc(), arith::CmpIPredicate::eq,
                                         std::get<0>(bound), constant);
     cond = cond ? b.create<arith::AndIOp>(op.getLoc(), cond, cmp) : cmp;
     map.map(std::get<0>(bound), constant);
static void specializeForLoopForUnrolling(ForOp op) {
   for (AffineExpr expr : minOp.getMap().getResults()) {
     if (auto constantIndex = expr.dyn_cast<AffineConstantExpr>())
       minConstant = std::min(minConstant, constantIndex.getValue());
   }
   if (minConstant == std::numeric_limits<int64_t>::max())
     return;
 
   OpBuilder b(op);
-  BlockAndValueMapping map;
+  IRMapping map;
   Value constant = b.create<arith::ConstantIndexOp>(op.getLoc(), minConstant);
   Value cond = b.create<arith::CmpIOp>(op.getLoc(), arith::CmpIPredicate::eq,
                                        bound, constant);
   map.map(bound, constant);
   auto ifOp = b.create<scf::IfOp>(op.getLoc(), cond, /*withElseRegion=*/true);
   ifOp.getThenBodyBuilder().clone(*op.getOperation(), map);
   ifOp.getElseBodyBuilder().clone(*op.getOperation());
   op.erase();

mlir/lib/Dialect/SCF/Transforms/ParallelLoopFusion.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SCF/Transforms/Passes.h"
 
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
 #include "mlir/Dialect/SCF/Transforms/Transforms.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 
 namespace mlir {
 #define GEN_PASS_DEF_SCFPARALLELLOOPFUSION
 #include "mlir/Dialect/SCF/Transforms/Passes.h.inc"
 } // namespace mlir
 
return matchOperands(firstPloop.getLowerBound(),
          matchOperands(firstPloop.getStep(), secondPloop.getStep());
 }
 
 /// Checks if the parallel loops have mixed access to the same buffers. Returns
 /// `true` if the first parallel loop writes to the same indices that the second
 /// loop reads.
 static bool haveNoReadsAfterWriteExceptSameIndex(
     ParallelOp firstPloop, ParallelOp secondPloop,
-    const BlockAndValueMapping &firstToSecondPloopIndices) {
+    const IRMapping &firstToSecondPloopIndices) {
   DenseMap<Value, SmallVector<ValueRange, 1>> bufferStores;
   firstPloop.getBody()->walk([&](memref::StoreOp store) {
     bufferStores[store.getMemRef()].push_back(store.getIndices());
   });
   auto walkResult = secondPloop.getBody()->walk([&](memref::LoadOp load) {
     // Stop if the memref is defined in secondPloop body. Careful alias analysis
     // is needed.
     auto *memrefDef = load.getMemRef().getDefiningOp();
static bool haveNoReadsAfterWriteExceptSameIndex(
   });
   return !walkResult.wasInterrupted();
 }
 
 /// Analyzes dependencies in the most primitive way by checking simple read and
 /// write patterns.
 static LogicalResult
 verifyDependencies(ParallelOp firstPloop, ParallelOp secondPloop,
-                   const BlockAndValueMapping &firstToSecondPloopIndices) {
+                   const IRMapping &firstToSecondPloopIndices) {
   if (!haveNoReadsAfterWriteExceptSameIndex(firstPloop, secondPloop,
                                             firstToSecondPloopIndices))
     return failure();
 
-  BlockAndValueMapping secondToFirstPloopIndices;
+  IRMapping secondToFirstPloopIndices;
   secondToFirstPloopIndices.map(secondPloop.getBody()->getArguments(),
                                 firstPloop.getBody()->getArguments());
   return success(haveNoReadsAfterWriteExceptSameIndex(
       secondPloop, firstPloop, secondToFirstPloopIndices));
 }
 
-static bool
-isFusionLegal(ParallelOp firstPloop, ParallelOp secondPloop,
-              const BlockAndValueMapping &firstToSecondPloopIndices) {
+static bool isFusionLegal(ParallelOp firstPloop, ParallelOp secondPloop,
+                          const IRMapping &firstToSecondPloopIndices) {
   return !hasNestedParallelOp(firstPloop) &&
          !hasNestedParallelOp(secondPloop) &&
          equalIterationSpaces(firstPloop, secondPloop) &&
          succeeded(verifyDependencies(firstPloop, secondPloop,
                                       firstToSecondPloopIndices));
 }
 
 /// Prepends operations of firstPloop's body into secondPloop's body.
 static void fuseIfLegal(ParallelOp firstPloop, ParallelOp secondPloop,
                         OpBuilder b) {
-  BlockAndValueMapping firstToSecondPloopIndices;
+  IRMapping firstToSecondPloopIndices;
   firstToSecondPloopIndices.map(firstPloop.getBody()->getArguments(),
                                 secondPloop.getBody()->getArguments());
 
   if (!isFusionLegal(firstPloop, secondPloop, firstToSecondPloopIndices))
     return;
 
   b.setInsertionPointToStart(secondPloop.getBody());
   for (auto &op : firstPloop.getBody()->without_terminator())

mlir/lib/Dialect/SCF/Utils/Utils.cpp

 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/SCF/Utils/Utils.h"
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Interfaces/SideEffectInterfaces.h"
 #include "mlir/Support/MathExtras.h"
 #include "mlir/Transforms/RegionUtils.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/ADT/SmallVector.h"
Block *newBlock = rewriter.createBlock(
     llvm::append_range(callValues, outlinedValues);
     auto call = rewriter.create<func::CallOp>(loc, outlinedFunc, callValues);
     if (callOp)
       *callOp = call;
 
     // `originalTerminator` was moved to `outlinedFuncBody` and is still valid.
     // Clone `originalTerminator` to take the callOp results then erase it from
     // `outlinedFuncBody`.
-    BlockAndValueMapping bvm;
+    IRMapping bvm;
     bvm.map(originalTerminator->getOperands(), call->getResults());
     rewriter.clone(*originalTerminator, bvm);
     rewriter.eraseOp(originalTerminator);
   }
 
   // Lastly, explicit RAUW outlinedValues, only for uses within `outlinedFunc`.
   // Clone the `arith::ConstantIndexOp` at the start of `outlinedFuncBody`.
   for (auto it : llvm::zip(outlinedValues, outlinedFuncBlockArgs.take_back(
                                                outlinedValues.size()))) {
     Value orig = std::get<0>(it);
     Value repl = std::get<1>(it);
     {
       OpBuilder::InsertionGuard g(rewriter);
       rewriter.setInsertionPointToStart(outlinedFuncBody);
       if (Operation *cst = orig.getDefiningOp<arith::ConstantIndexOp>()) {
-        BlockAndValueMapping bvm;
+        IRMapping bvm;
         repl = rewriter.clone(*cst, bvm)->getResult(0);
       }
     }
     orig.replaceUsesWithIf(repl, [&](OpOperand &opOperand) {
       return outlinedFunc->isProperAncestor(opOperand.getOwner());
     });
   }
 
static void generateUnrolledLoop(
   // Keep a pointer to the last non-terminator operation in the original block
   // so that we know what to clone (since we are doing this in-place).
   Block::iterator srcBlockEnd = std::prev(loopBodyBlock->end(), 2);
 
   // Unroll the contents of 'forOp' (append unrollFactor - 1 additional copies).
   SmallVector<Value, 4> lastYielded(yieldedValues);
 
   for (unsigned i = 1; i < unrollFactor; i++) {
-    BlockAndValueMapping operandMap;
+    IRMapping operandMap;
 
     // Prepare operand map.
     operandMap.map(iterArgs, lastYielded);
 
     // If the induction variable is used, create a remapping to the value for
     // this unrolled instance.
     if (!forOpIV.use_empty()) {
       Value ivUnroll = ivRemapFn(i, forOpIV, builder);

mlir/lib/Dialect/SPIRV/IR/SPIRVDialect.cpp

struct SPIRVInlinerInterface : public DialectInlinerInterface {
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     return true;
   }
 
   /// Returns true if the given region 'src' can be inlined into the region
   /// 'dest' that is attached to an operation registered to the current dialect.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     // Return true here when inlining into spirv.func, spirv.mlir.selection, and
     // spirv.mlir.loop operations.
     auto *op = dest->getParentOp();
     return isa<spirv::FuncOp, spirv::SelectionOp, spirv::LoopOp>(op);
   }
 
   /// Returns true if the given operation 'op', that is registered to this
   /// dialect, can be inlined into the region 'dest' that is attached to an
   /// operation registered to the current dialect.
   bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     // TODO: Enable inlining structured control flows with return.
     if ((isa<spirv::SelectionOp, spirv::LoopOp>(op)) &&
         containsReturn(op->getRegion(0)))
       return false;
     // TODO: we need to filter OpKill here to avoid inlining it to
     // a loop continue construct:
     // https://github.com/KhronosGroup/SPIRV-Headers/issues/86
     // However OpKill is fragment shader specific and we don't support it yet.

mlir/lib/Dialect/Shape/IR/Shape.cpp

 namespace {
 /// This class defines the interface for inlining shape dialect ops.
 struct ShapeInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   // Returns true if the given region 'src' can be inlined into the region
   // 'dest' that is attached to an operation registered to the current dialect.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     return true;
   }
 
   // Returns true if the given operation 'op', that is registered to this
   // dialect, can be inlined into the region 'dest' that is attached to an
   // operation registered to the current dialect.
   bool isLegalToInline(Operation *op, Region *dest, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 void ShapeDialect::initialize() {
   addOperations<
 #define GET_OP_LIST

mlir/lib/Dialect/Shape/Transforms/OutlineShapeComputation.cpp

 //====----- OutlineShapeComputation.cpp -----------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/Shape/Analysis/ShapeMappingAnalysis.h"
 #include "mlir/Dialect/Shape/IR/Shape.h"
 #include "mlir/Dialect/Shape/Transforms/Passes.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/Support/Debug.h"
 #include <queue>
 #include <unordered_set>
createFuncFromCluster(OpBuilder &b, const SmallVector<Operation *, 8> &cluster,
   SmallVector<Value, 4> inputs = getInputsOfCluster(cluster);
   auto fnType =
       cluster.empty()
           ? b.getFunctionType(shape.getType(), shape.getType())
           : b.getFunctionType(ValueRange(inputs).getTypes(), shape.getType());
   shape::FuncOp fnOp = b.create<shape::FuncOp>(loc, fnName, fnType);
   Block *block = fnOp.addEntryBlock();
   b.setInsertionPoint(block, block->end());
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   if (cluster.empty()) {
     bvm.map(shape, fnOp.getArgument(0));
   } else {
     for (auto inputAndArg : llvm::zip(inputs, fnOp.getArguments()))
       bvm.map(std::get<0>(inputAndArg), std::get<1>(inputAndArg));
   }
 
   for (Operation *op : cluster)

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorRewriting.cpp

LogicalResult matchAndRewrite(GenericOp op,
     fusedIndexMaps.push_back(fusedIndexMaps.back()); // mimic other
     // Fuse producer and consumer into a new generic op.
     auto fusedOp = rewriter.create<GenericOp>(
         loc, op.getResult(0).getType(), inputOps, outputOps,
         rewriter.getAffineMapArrayAttr(fusedIndexMaps), prod.getIteratorTypes(),
         /*doc=*/nullptr, /*library_call=*/nullptr);
     Block &prodBlock = prod.getRegion().front();
     Block &consBlock = op.getRegion().front();
-    BlockAndValueMapping mapper;
+    IRMapping mapper;
     Block *fusedBlock = new Block();
     fusedOp.getRegion().push_back(fusedBlock);
     unsigned num = prodBlock.getNumArguments();
     for (unsigned i = 0; i < num - 1; i++)
       addArg(mapper, fusedBlock, prodBlock.getArgument(i));
     addArg(mapper, fusedBlock, consBlock.getArgument(1 - other));
     addArg(mapper, fusedBlock, prodBlock.getArgument(num - 1));
     // Clone bodies of the producer and consumer in new evaluation order.
LogicalResult matchAndRewrite(GenericOp op,
     // Replace consumer with fused operation. Old producer
     // and consumer ops will be removed by DCE.
     rewriter.replaceOp(op, fusedOp->getResults());
     return success();
   }
 
 private:
   // Helper to add argument and record the mapping.
-  static void addArg(BlockAndValueMapping &mapper, Block *b, BlockArgument a) {
+  static void addArg(IRMapping &mapper, Block *b, BlockArgument a) {
     mapper.map(a, b->addArgument(a.getType(), a.getLoc()));
   }
 };
 
 /// Sparse rewriting rule for sparse-to-sparse reshape operator.
 template <typename ReshapeOp>
 struct Sparse2SparseReshapeRewriter : public OpRewritePattern<ReshapeOp> {
 public:

mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp

Optional<unsigned> Merger::buildTensorExp(linalg::GenericOp op, Value v) {
   // Cannot build.
   return std::nullopt;
 }
 
 static Value insertYieldOp(RewriterBase &rewriter, Location loc, Region &region,
                            ValueRange vals) {
   // Make a clone of overlap region.
   Region tmpRegion;
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   region.cloneInto(&tmpRegion, tmpRegion.begin(), mapper);
   Block &clonedBlock = tmpRegion.front();
   YieldOp clonedYield = cast<YieldOp>(clonedBlock.getTerminator());
   // Merge cloned block and return yield value.
   Operation *placeholder = rewriter.create<arith::ConstantIndexOp>(loc, 0);
   rewriter.mergeBlockBefore(&tmpRegion.front(), placeholder, vals);
   Value val = clonedYield.getResult();
   rewriter.eraseOp(clonedYield);

mlir/lib/Dialect/Tensor/IR/TensorDialect.cpp

 //===----------------------------------------------------------------------===//
 // TensorDialect Dialect Interfaces
 //===----------------------------------------------------------------------===//
 
 namespace {
 struct TensorInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &valueMapping) const final {
+                       IRMapping &valueMapping) const final {
     return true;
   }
   bool isLegalToInline(Operation *, Region *, bool wouldBeCloned,
-                       BlockAndValueMapping &) const final {
+                       IRMapping &) const final {
     return true;
   }
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // TensorDialect Methods
 //===----------------------------------------------------------------------===//
 
 void TensorDialect::initialize() {
   addOperations<
 #define GET_OP_LIST
 #include "mlir/Dialect/Tensor/IR/TensorOps.cpp.inc"
       >();
   addInterfaces<TensorInlinerInterface>();
 }

mlir/lib/Dialect/Tensor/IR/TensorOps.cpp

 //===----------------------------------------------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Affine/IR/AffineOps.h"
 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Arith/Utils/Utils.h"
 #include "mlir/Dialect/Complex/IR/Complex.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Utils/IndexingUtils.h"
 #include "mlir/Dialect/Utils/ReshapeOpsUtils.h"
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributeInterfaces.h"
 #include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/OpDefinition.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Interfaces/DestinationStyleOpInterface.h"
 #include "mlir/Support/MathExtras.h"
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallBitVector.h"
struct ExtractFromTensorGenerate : public OpRewritePattern<tensor::ExtractOp> {
   using OpRewritePattern<tensor::ExtractOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(tensor::ExtractOp extract,
                                 PatternRewriter &rewriter) const final {
     auto tensorFromElements = extract.getTensor().getDefiningOp<GenerateOp>();
     if (!tensorFromElements || !wouldOpBeTriviallyDead(tensorFromElements))
       return failure();
 
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     Block *body = &tensorFromElements.getBody().front();
     mapping.map(body->getArguments(), extract.getIndices());
     for (auto &op : body->without_terminator())
       rewriter.clone(op, mapping);
 
     auto yield = cast<YieldOp>(body->getTerminator());
 
     rewriter.replaceOp(extract, mapping.lookupOrDefault(yield.getValue()));
if (newResultType == padTensorOp.getResultType()) {
         padTensorOp.getSourceMutable().assign(castOp.getSource());
       });
     } else {
       auto newOp = rewriter.create<PadOp>(
           padTensorOp->getLoc(), newResultType, padTensorOp.getSource(),
           padTensorOp.getLow(), padTensorOp.getHigh(),
           padTensorOp.getStaticLow(), padTensorOp.getStaticHigh(),
           padTensorOp.getNofold());
-      BlockAndValueMapping mapper;
+      IRMapping mapper;
       padTensorOp.getRegion().cloneInto(&newOp.getRegion(), mapper);
 
       rewriter.replaceOpWithNewOp<tensor::CastOp>(
           padTensorOp, padTensorOp.getResultType(), newOp);
     }
     return success();
   }
 };
if (!llvm::all_of(
             if (ShapedType::isDynamic(shape)) {
               // For the shape being dynamic when tile size is
               // specified, return true. In canonical form a constant
               // tile size should lead to constant shape of the tiled
               // dimension, but not needed for verification.
               return true;
             }
             return shape == constTileSize.value();
-           
           })) {
     return op->emitError("mismatch in inner tile sizes specified and shaped of "
                          "tiled dimension in the packed type");
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/Tensor/IR/TensorTilingInterfaceImpl.cpp

Operation *tensor::bubbleUpPadSlice(OpBuilder &b, tensor::PadOp padOp,
   auto createPadOfExtractSlice = [&]() {
     // Create pad(extract_slice(x)).
     auto newSliceOp = b.create<tensor::ExtractSliceOp>(
         loc, padOp.getSource(), newOffsets, newLengths, newStrides);
     auto newPadOp = b.create<PadOp>(loc, newSliceOp, staticNewLows,
                                     staticNewHighs, newLows, newHighs);
 
     // Copy region to new PadOp.
-    BlockAndValueMapping bvm;
+    IRMapping bvm;
     padOp.getRegion().cloneInto(&newPadOp.getRegion(), bvm);
 
     // Cast result and return.
     return castResult(newPadOp);
   };
 
   // Rewrite extract_slice(pad(x)) into a GenerateOp it is statically known that
   // the original data source x is not used.

mlir/lib/Dialect/Tosa/IR/TosaOps.cpp

struct TosaInlinerInterface : public DialectInlinerInterface {
   using DialectInlinerInterface::DialectInlinerInterface;
 
   //===--------------------------------------------------------------------===//
   // Analysis Hooks.
   //===--------------------------------------------------------------------===//
 
   /// All operations can be inlined by default.
   bool isLegalToInline(Operation *op, Region *region, bool wouldBeCloned,
-                       BlockAndValueMapping &map) const final {
+                       IRMapping &map) const final {
     return true;
   }
 
   /// All regions with If and While parent operators can be inlined.
   bool isLegalToInline(Region *dest, Region *src, bool wouldBeCloned,
-                       BlockAndValueMapping &map) const final {
+                       IRMapping &map) const final {
     return (isa<tosa::IfOp>(dest->getParentOp()) ||
             isa<tosa::WhileOp>(dest->getParentOp()));
   }
 };
 } // namespace
 
 //===----------------------------------------------------------------------===//
 // TOSA control flow support.

mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp

 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Tosa/Transforms/Passes.h"
 
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Tosa/IR/TosaOps.h"
 #include "mlir/Dialect/Tosa/Utils/ShapeUtils.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/Support/FormatVariadic.h"
 
 namespace mlir {
 namespace tosa {

mlir/lib/Dialect/Utils/StructuredOpsUtils.cpp

 //===- StructuredOpsUtils.cpp - Utilities used by structured ops ----------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Dialect/Utils/StructuredOpsUtils.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributes.h"
+#include "mlir/IR/IRMapping.h"
 
 #include "mlir/Dialect/Utils/DialectUtilsEnums.cpp.inc"
 
 using namespace mlir;
 
 bool mlir::isRowMajorMatmul(ArrayAttr indexingMaps) {
   if (indexingMaps.size() != 3)
     return false;
bool mlir::isRowMajorBatchMatmul(ArrayAttr indexingMaps) {
   auto mapB = AffineMapAttr::get(AffineMap::get(4, 0, {b, k, n}, context));
   auto mapC = AffineMapAttr::get(AffineMap::get(4, 0, {b, m, n}, context));
   auto maps = ArrayAttr::get(context, {mapA, mapB, mapC});
   return indexingMaps == maps;
 }
 
 Operation *mlir::clone(OpBuilder &b, Operation *op, TypeRange newResultTypes,
                        ValueRange newOperands) {
-  BlockAndValueMapping bvm;
+  IRMapping bvm;
   OperationState state(op->getLoc(), op->getName(), newOperands, newResultTypes,
                        op->getAttrs());
   for (Region &r : op->getRegions())
     r.cloneInto(state.addRegion(), bvm);
   return b.create(state);
 }
 
 Operation *mlir::cloneWithoutRegions(OpBuilder &b, Operation *op,
                                      TypeRange newResultTypes,
                                      ValueRange newOperands) {
   OperationState state(op->getLoc(), op->getName(), newOperands, newResultTypes,
                        op->getAttrs());
   for (size_t cnt = 0, e = op->getNumRegions(); cnt < e; ++cnt)
     state.addRegion();
   return b.create(state);
 }

mlir/lib/Dialect/Vector/IR/VectorOps.cpp

 #include "mlir/Dialect/Arith/IR/Arith.h"
 #include "mlir/Dialect/Arith/Utils/Utils.h"
 #include "mlir/Dialect/MemRef/IR/MemRef.h"
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Utils/IndexingUtils.h"
 #include "mlir/Dialect/Utils/StructuredOpsUtils.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinAttributes.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/DialectImplementation.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Support/LLVM.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringSet.h"
auto canonicalize = [&](Value maybeContraction,
           dyn_cast_or_null<vector::ContractionOp>(
               maybeContraction.getDefiningOp());
       if (!contractionOp)
         return vector::ContractionOp();
       if (auto maybeZero = dyn_cast_or_null<arith::ConstantOp>(
               contractionOp.getAcc().getDefiningOp())) {
         if (maybeZero.getValue() ==
             rewriter.getZeroAttr(contractionOp.getAcc().getType())) {
-          BlockAndValueMapping bvm;
+          IRMapping bvm;
           bvm.map(contractionOp.getAcc(), otherOperand);
           auto newContraction =
               cast<vector::ContractionOp>(rewriter.clone(*contractionOp, bvm));
           rewriter.replaceOp(addOp, newContraction.getResult());
           return newContraction;
         }
       }
       return vector::ContractionOp();

mlir/lib/Dialect/Vector/Transforms/VectorTransferSplitRewritePatterns.cpp

 static void createFullPartialVectorTransferWrite(RewriterBase &b,
                                                  vector::TransferWriteOp xferOp,
                                                  Value inBoundsCond,
                                                  Value alloc) {
   Location loc = xferOp.getLoc();
   auto notInBounds = b.create<arith::XOrIOp>(
       loc, inBoundsCond, b.create<arith::ConstantIntOp>(loc, true, 1));
   b.create<scf::IfOp>(loc, notInBounds, [&](OpBuilder &b, Location loc) {
-    BlockAndValueMapping mapping;
+    IRMapping mapping;
     Value load = b.create<memref::LoadOp>(
         loc,
         b.create<vector::TypeCastOp>(
             loc, MemRefType::get({}, xferOp.getVector().getType()), alloc));
     mapping.map(xferOp.getVector(), load);
     b.clone(*xferOp.getOperation(), mapping);
     b.create<scf::YieldOp>(loc, ValueRange{});
   });
LogicalResult mlir::vector::splitFullAndPartialTransfer(
 
   // Decide which location to write the entire vector to.
   auto memrefAndIndices = getLocationToWriteFullVec(
       b, xferWriteOp, returnTypes, inBoundsCond, compatibleMemRefType, alloc);
 
   // Do an in bounds write to either the output or the extra allocated buffer.
   // The operation is cloned to prevent deleting information needed for the
   // later IR creation.
-  BlockAndValueMapping mapping;
+  IRMapping mapping;
   mapping.map(xferWriteOp.getSource(), memrefAndIndices.front());
   mapping.map(xferWriteOp.getIndices(), memrefAndIndices.drop_front());
   auto *clone = b.clone(*xferWriteOp, mapping);
   clone->setAttr(xferWriteOp.getInBoundsAttrName(), inBoundsAttr);
 
   // Create a potential copy from the allocated buffer to the final output in
   // the slow path case.
   if (options.vectorTransferSplit == VectorTransferSplit::VectorTransfer)

mlir/lib/IR/Builders.cpp

 //===- Builders.cpp - Helpers for constructing MLIR Classes ---------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/AffineMap.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Dialect.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/IntegerSet.h"
 #include "mlir/IR/Matchers.h"
 #include "mlir/IR/SymbolTable.h"
 #include "llvm/Support/raw_ostream.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
LogicalResult OpBuilder::tryFold(Operation *op,
 
   // If we were successful, insert any generated constants.
   for (Operation *cst : generatedConstants)
     insert(cst);
 
   return success();
 }
 
-Operation *OpBuilder::clone(Operation &op, BlockAndValueMapping &mapper) {
+Operation *OpBuilder::clone(Operation &op, IRMapping &mapper) {
   Operation *newOp = op.clone(mapper);
   // The `insert` call below handles the notification for inserting `newOp`
   // itself. But if `newOp` has any regions, we need to notify the listener
   // about any ops that got inserted inside those regions as part of cloning.
   if (listener) {
     auto walkFn = [&](Operation *walkedOp) {
       listener->notifyOperationInserted(walkedOp);
     };
     for (Region &region : newOp->getRegions())
       region.walk(walkFn);
   }
   return insert(newOp);
 }
 
 Operation *OpBuilder::clone(Operation &op) {
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   return clone(op, mapper);
 }

mlir/lib/IR/BuiltinDialect.cpp

 //===- BuiltinDialect.cpp - MLIR Builtin Dialect --------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the Builtin dialect that contains all of the attributes,
 // operations, and types that are necessary for the validity of the IR.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/BuiltinDialect.h"
 #include "BuiltinDialectBytecode.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/DialectResourceBlobManager.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeRange.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
 // TableGen'erated dialect

mlir/lib/IR/Operation.cpp

 //===- Operation.cpp - Operation support code -----------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/Operation.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Dialect.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Interfaces/FoldInterfaces.h"
 #include "llvm/ADT/StringExtras.h"
 #include <numeric>
 
 using namespace mlir;
Operation::CloneOptions &Operation::CloneOptions::cloneOperands(bool enable) {
   cloneOperandsFlag = enable;
   return *this;
 }
 
 /// Create a deep copy of this operation but keep the operation regions empty.
 /// Operands are remapped using `mapper` (if present), and `mapper` is updated
 /// to contain the results. The `mapResults` flag specifies whether the results
 /// of the cloned operation should be added to the map.
-Operation *Operation::cloneWithoutRegions(BlockAndValueMapping &mapper) {
+Operation *Operation::cloneWithoutRegions(IRMapping &mapper) {
   return clone(mapper, CloneOptions::all().cloneRegions(false));
 }
 
 Operation *Operation::cloneWithoutRegions() {
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   return cloneWithoutRegions(mapper);
 }
 
 /// Create a deep copy of this operation, remapping any operands that use
 /// values outside of the operation using the map that is provided (leaving
 /// them alone if no entry is present).  Replaces references to cloned
 /// sub-operations to the corresponding operation that is copied, and adds
 /// those mappings to the map.
-Operation *Operation::clone(BlockAndValueMapping &mapper,
-                            CloneOptions options) {
+Operation *Operation::clone(IRMapping &mapper, CloneOptions options) {
   SmallVector<Value, 8> operands;
   SmallVector<Block *, 2> successors;
 
   // Remap the operands.
   if (options.shouldCloneOperands()) {
     operands.reserve(getNumOperands());
     for (auto opValue : getOperands())
       operands.push_back(mapper.lookupOrDefault(opValue));
   }
 
   // Remap the successors.
   successors.reserve(getNumSuccessors());
   for (Block *successor : getSuccessors())
     successors.push_back(mapper.lookupOrDefault(successor));
 
   // Create the new operation.
   auto *newOp = create(getLoc(), getName(), getResultTypes(), operands, attrs,
                        successors, getNumRegions());
+  mapper.map(this, newOp);
 
   // Clone the regions.
   if (options.shouldCloneRegions()) {
     for (unsigned i = 0; i != numRegions; ++i)
       getRegion(i).cloneInto(&newOp->getRegion(i), mapper);
   }
 
   // Remember the mapping of any results.
   for (unsigned i = 0, e = getNumResults(); i != e; ++i)
     mapper.map(getResult(i), newOp->getResult(i));
 
   return newOp;
 }
 
 Operation *Operation::clone(CloneOptions options) {
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   return clone(mapper, options);
 }
 
 //===----------------------------------------------------------------------===//
 // OpState trait class.
 //===----------------------------------------------------------------------===//
 
 // The fallback for the parser is to try for a dialect operation parser.

mlir/lib/IR/PatternMatch.cpp

 //===- PatternMatch.cpp - Base classes for pattern match ------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/PatternMatch.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
 // PatternBenefit
 //===----------------------------------------------------------------------===//
 
 PatternBenefit::PatternBenefit(unsigned benefit) : representation(benefit) {

 }
 
 /// Clone the blocks that belong to "region" before the given position in
 /// another region "parent". The two regions must be different. The caller is
 /// responsible for creating or updating the operation transferring flow of
 /// control to the region and passing it the correct block arguments.
 void RewriterBase::cloneRegionBefore(Region &region, Region &parent,
                                      Region::iterator before,
-                                     BlockAndValueMapping &mapping) {
+                                     IRMapping &mapping) {
   region.cloneInto(&parent, before, mapping);
 }
 void RewriterBase::cloneRegionBefore(Region &region, Region &parent,
                                      Region::iterator before) {
-  BlockAndValueMapping mapping;
+  IRMapping mapping;
   cloneRegionBefore(region, parent, before, mapping);
 }
 void RewriterBase::cloneRegionBefore(Region &region, Block *before) {
   cloneRegionBefore(region, *before->getParent(), before->getIterator());
 }

mlir/lib/IR/Region.cpp

 //===- Region.cpp - MLIR Region Class -------------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/Region.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 using namespace mlir;
 
 Region::Region(Operation *container) : container(container) {}
 
 Region::~Region() {
   // Operations may have cyclic references, which need to be dropped before we
   // can start deleting them.

 unsigned Region::getRegionNumber() {
   // Regions are always stored consecutively, so use pointer subtraction to
   // figure out what number this is.
   return this - &getParentOp()->getRegions()[0];
 }
 
 /// Clone the internal blocks from this region into `dest`. Any
 /// cloned blocks are appended to the back of dest.
-void Region::cloneInto(Region *dest, BlockAndValueMapping &mapper) {
+void Region::cloneInto(Region *dest, IRMapping &mapper) {
   assert(dest && "expected valid region to clone into");
   cloneInto(dest, dest->end(), mapper);
 }
 
 /// Clone this region into 'dest' before the given position in 'dest'.
 void Region::cloneInto(Region *dest, Region::iterator destPos,
-                       BlockAndValueMapping &mapper) {
+                       IRMapping &mapper) {
   assert(dest && "expected valid region to clone into");
   assert(this != dest && "cannot clone region into itself");
 
   // If the list is empty there is nothing to clone.
   if (empty())
     return;
 
   // The below clone implementation takes special care to be read only for the

mlir/lib/Reducer/ReductionNode.cpp

 // metadata for a specific generated variant within a reduction pass and are the
 // building blocks of the reduction tree structure. A reduction tree is used to
 // keep track of the different generated variants throughout a reduction pass in
 // the MLIR Reduce tool.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Reducer/ReductionNode.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "llvm/ADT/STLExtras.h"
 
 #include <algorithm>
 #include <limits>
 
 using namespace mlir;
 
 ReductionNode::ReductionNode(
     ReductionNode *parentNode, const std::vector<Range> &ranges,
     llvm::SpecificBumpPtrAllocator<ReductionNode> &allocator)
     /// Root node will have the parent pointer point to themselves.
     : parent(parentNode == nullptr ? this : parentNode),
       size(std::numeric_limits<size_t>::max()), ranges(ranges),
       startRanges(ranges), allocator(allocator) {
   if (parent != this)
     if (failed(initialize(parent->getModule(), parent->getRegion())))
       llvm_unreachable("unexpected initialization failure");
 }
 
 LogicalResult ReductionNode::initialize(ModuleOp parentModule,
                                         Region &targetRegion) {
   // Use the mapper help us find the corresponding region after module clone.
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   module = cast<ModuleOp>(parentModule->clone(mapper));
   // Use the first block of targetRegion to locate the cloned region.
   Block *block = mapper.lookup(&*targetRegion.begin());
   region = block->getParent();
   return success();
 }
 
 /// If we haven't explored any variants from this node, we will create N

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

 //===- OpenMPToLLVMIRTranslation.cpp - Translate OpenMP dialect to LLVM IR-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a translation between the MLIR OpenMP dialect and LLVM
 // IR.
 //
 //===----------------------------------------------------------------------===//
 #include "mlir/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.h"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/Support/LLVM.h"
 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
 
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/TypeSwitch.h"
 #include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
 #include "llvm/IR/DebugInfoMetadata.h"

mlir/lib/Target/SPIRV/Deserialization/Deserializer.cpp

 //===----------------------------------------------------------------------===//
 
 #include "Deserializer.h"
 
 #include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
 #include "mlir/Dialect/SPIRV/IR/SPIRVTypes.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Location.h"
 #include "mlir/Support/LogicalResult.h"
 #include "mlir/Target/SPIRV/SPIRVBinaryUtils.h"
 #include "llvm/ADT/STLExtras.h"
 #include "llvm/ADT/Sequence.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringExtras.h"
 #include "llvm/ADT/bit.h"
LogicalResult ControlFlowStructurizer::structurize() {
   } else {
     if (auto selectionOp = createSelectionOp(control))
       op = selectionOp.getOperation();
   }
   if (!op)
     return failure();
   Region &body = op->getRegion(0);
 
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   // All references to the old merge block should be directed to the
   // selection/loop merge block in the SelectionOp/LoopOp's region.
   mapper.map(mergeBlock, &body.back());
 
   collectBlocksInConstruct();
 
   // We've identified all blocks belonging to the selection/loop's region. Now
   // need to "move" them into the selection/loop. Instead of really moving the

mlir/lib/Transforms/Utils/DialectConversion.cpp

 //===- DialectConversion.cpp - MLIR dialect conversion generic pass -------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Transforms/DialectConversion.h"
 #include "mlir/IR/Block.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/BuiltinOps.h"
 #include "mlir/IR/FunctionInterfaces.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Rewrite/PatternApplicator.h"
 #include "llvm/ADT/ScopeExit.h"
 #include "llvm/ADT/SetVector.h"
 #include "llvm/ADT/SmallPtrSet.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/FormatVariadic.h"
 #include "llvm/Support/SaveAndRestore.h"
 #include "llvm/Support/ScopedPrinter.h"
static void logFailure(llvm::ScopedPrinter &os, StringRef fmt, Args &&...args) {
   });
 }
 
 //===----------------------------------------------------------------------===//
 // ConversionValueMapping
 //===----------------------------------------------------------------------===//
 
 namespace {
-/// This class wraps a BlockAndValueMapping to provide recursive lookup
+/// This class wraps a IRMapping to provide recursive lookup
 /// functionality, i.e. we will traverse if the mapped value also has a mapping.
 struct ConversionValueMapping {
   /// Lookup a mapped value within the map. If a mapping for the provided value
   /// does not exist then return the provided value. If `desiredType` is
   /// non-null, returns the most recently mapped value with that type. If an
   /// operand of that type does not exist, defaults to normal behavior.
   Value lookupOrDefault(Value from, Type desiredType = nullptr) const;
 
DenseMap<Value, SmallVector<Value>> getInverse() const {
     DenseMap<Value, SmallVector<Value>> inverse;
     for (auto &it : mapping.getValueMap())
       inverse[it.second].push_back(it.first);
     return inverse;
   }
 
 private:
   /// Current value mappings.
-  BlockAndValueMapping mapping;
+  IRMapping mapping;
 };
 } // namespace
 
 Value ConversionValueMapping::lookupOrDefault(Value from,
                                               Type desiredType) const {
   // If there was no desired type, simply find the leaf value.
   if (!desiredType) {
     // If this value had a valid mapping, unmap that value as well in the case

 
 void ConversionPatternRewriter::inlineRegionBefore(Region &region,
                                                    Region &parent,
                                                    Region::iterator before) {
   impl->notifyRegionIsBeingInlinedBefore(region, parent, before);
   PatternRewriter::inlineRegionBefore(region, parent, before);
 }
 
-void ConversionPatternRewriter::cloneRegionBefore(
-    Region &region, Region &parent, Region::iterator before,
-    BlockAndValueMapping &mapping) {
+void ConversionPatternRewriter::cloneRegionBefore(Region &region,
+                                                  Region &parent,
+                                                  Region::iterator before,
+                                                  IRMapping &mapping) {
   if (region.empty())
     return;
   PatternRewriter::cloneRegionBefore(region, parent, before, mapping);
 
   // Collect the range of the cloned blocks.
   auto clonedBeginIt = mapping.lookup(&region.front())->getIterator();
   auto clonedBlocks = llvm::make_range(clonedBeginIt, before);
   impl->notifyRegionWasClonedBefore(clonedBlocks, region.getLoc());

mlir/lib/Transforms/Utils/InliningUtils.cpp

 //===- InliningUtils.cpp ---- Misc utilities for inlining -----------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements miscellaneous inlining utilities.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Transforms/InliningUtils.h"
 
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Operation.h"
 #include "mlir/Interfaces/CallInterfaces.h"
 #include "llvm/ADT/MapVector.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
 
 #define DEBUG_TYPE "inlining"
 
auto remapOpLoc = [&](Operation *op) {
     }
     op->setLoc(it->second);
   };
   for (auto &block : inlinedBlocks)
     block.walk(remapOpLoc);
 }
 
 static void remapInlinedOperands(iterator_range<Region::iterator> inlinedBlocks,
-                                 BlockAndValueMapping &mapper) {
+                                 IRMapping &mapper) {
   auto remapOperands = [&](Operation *op) {
     for (auto &operand : op->getOpOperands())
       if (auto mappedOp = mapper.lookupOrNull(operand.get()))
         operand.set(mappedOp);
   };
   for (auto &block : inlinedBlocks)
     block.walk(remapOperands);
 }
 
 //===----------------------------------------------------------------------===//
 // InlinerInterface
 //===----------------------------------------------------------------------===//
 
 bool InlinerInterface::isLegalToInline(Operation *call, Operation *callable,
                                        bool wouldBeCloned) const {
   if (auto *handler = getInterfaceFor(call))
     return handler->isLegalToInline(call, callable, wouldBeCloned);
   return false;
 }
 
-bool InlinerInterface::isLegalToInline(
-    Region *dest, Region *src, bool wouldBeCloned,
-    BlockAndValueMapping &valueMapping) const {
+bool InlinerInterface::isLegalToInline(Region *dest, Region *src,
+                                       bool wouldBeCloned,
+                                       IRMapping &valueMapping) const {
   if (auto *handler = getInterfaceFor(dest->getParentOp()))
     return handler->isLegalToInline(dest, src, wouldBeCloned, valueMapping);
   return false;
 }
 
-bool InlinerInterface::isLegalToInline(
-    Operation *op, Region *dest, bool wouldBeCloned,
-    BlockAndValueMapping &valueMapping) const {
+bool InlinerInterface::isLegalToInline(Operation *op, Region *dest,
+                                       bool wouldBeCloned,
+                                       IRMapping &valueMapping) const {
   if (auto *handler = getInterfaceFor(op))
     return handler->isLegalToInline(op, dest, wouldBeCloned, valueMapping);
   return false;
 }
 
 bool InlinerInterface::shouldAnalyzeRecursively(Operation *op) const {
   auto *handler = getInterfaceFor(op);
   return handler ? handler->shouldAnalyzeRecursively(op) : true;
void InlinerInterface::processInlinedCallBlocks(
   auto *handler = getInterfaceFor(call);
   assert(handler && "expected valid dialect handler");
   handler->processInlinedCallBlocks(call, inlinedBlocks);
 }
 
 /// Utility to check that all of the operations within 'src' can be inlined.
 static bool isLegalToInline(InlinerInterface &interface, Region *src,
                             Region *insertRegion, bool shouldCloneInlinedRegion,
-                            BlockAndValueMapping &valueMapping) {
+                            IRMapping &valueMapping) {
   for (auto &block : *src) {
     for (auto &op : block) {
       // Check this operation.
       if (!interface.isLegalToInline(&op, insertRegion,
                                      shouldCloneInlinedRegion, valueMapping)) {
         LLVM_DEBUG({
           llvm::dbgs() << "* Illegal to inline because of op: ";
           op.dump();

 }
 
 //===----------------------------------------------------------------------===//
 // Inline Methods
 //===----------------------------------------------------------------------===//
 
 static LogicalResult
 inlineRegionImpl(InlinerInterface &interface, Region *src, Block *inlineBlock,
-                 Block::iterator inlinePoint, BlockAndValueMapping &mapper,
+                 Block::iterator inlinePoint, IRMapping &mapper,
                  ValueRange resultsToReplace, TypeRange regionResultTypes,
                  Optional<Location> inlineLoc, bool shouldCloneInlinedRegion,
                  Operation *call = nullptr) {
   assert(resultsToReplace.size() == regionResultTypes.size());
   // We expect the region to have at least one block.
   if (src->empty())
     return failure();
 
inlineRegionImpl(InlinerInterface &interface, Region *src, Block *inlineBlock,
   if (src->empty())
     return failure();
 
   auto *entryBlock = &src->front();
   if (inlinedOperands.size() != entryBlock->getNumArguments())
     return failure();
 
   // Map the provided call operands to the arguments of the region.
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   for (unsigned i = 0, e = inlinedOperands.size(); i != e; ++i) {
     // Verify that the types of the provided values match the function argument
     // types.
     BlockArgument regionArg = entryBlock->getArgument(i);
     if (inlinedOperands[i].getType() != regionArg.getType())
       return failure();
     mapper.map(regionArg, inlinedOperands[i]);
   }
 
   // Call into the main region inliner function.
   return inlineRegionImpl(interface, src, inlineBlock, inlinePoint, mapper,
                           resultsToReplace, resultsToReplace.getTypes(),
                           inlineLoc, shouldCloneInlinedRegion, call);
 }
 
 LogicalResult mlir::inlineRegion(InlinerInterface &interface, Region *src,
-                                 Operation *inlinePoint,
-                                 BlockAndValueMapping &mapper,
+                                 Operation *inlinePoint, IRMapping &mapper,
                                  ValueRange resultsToReplace,
                                  TypeRange regionResultTypes,
                                  Optional<Location> inlineLoc,
                                  bool shouldCloneInlinedRegion) {
   return inlineRegion(interface, src, inlinePoint->getBlock(),
                       ++inlinePoint->getIterator(), mapper, resultsToReplace,
                       regionResultTypes, inlineLoc, shouldCloneInlinedRegion);
 }
-LogicalResult
-mlir::inlineRegion(InlinerInterface &interface, Region *src, Block *inlineBlock,
-                   Block::iterator inlinePoint, BlockAndValueMapping &mapper,
-                   ValueRange resultsToReplace, TypeRange regionResultTypes,
+LogicalResult mlir::inlineRegion(InlinerInterface &interface, Region *src,
+                                 Block *inlineBlock,
+                                 Block::iterator inlinePoint, IRMapping &mapper,
+                                 ValueRange resultsToReplace,
+                                 TypeRange regionResultTypes,
                                  Optional<Location> inlineLoc,
                                  bool shouldCloneInlinedRegion) {
   return inlineRegionImpl(interface, src, inlineBlock, inlinePoint, mapper,
                           resultsToReplace, regionResultTypes, inlineLoc,
                           shouldCloneInlinedRegion);
 }
 
 LogicalResult mlir::inlineRegion(InlinerInterface &interface, Region *src,
                                  Operation *inlinePoint,
                                  ValueRange inlinedOperands,
LogicalResult mlir::inlineCall(InlinerInterface &interface,
   };
 
   // Builder used for any conversion operations that need to be materialized.
   OpBuilder castBuilder(call);
   Location castLoc = call.getLoc();
   const auto *callInterface = interface.getInterfaceFor(call->getDialect());
 
   // Map the provided call operands to the arguments of the region.
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   for (unsigned i = 0, e = callOperands.size(); i != e; ++i) {
     BlockArgument regionArg = entryBlock->getArgument(i);
     Value operand = callOperands[i];
 
     // If the call operand doesn't match the expected region argument, try to
     // generate a cast.
     Type regionArgType = regionArg.getType();
     if (operand.getType() != regionArgType) {

mlir/test/lib/Dialect/Test/TestDialect.cpp

struct TestInlinerInterface : public DialectInlinerInterface {
   // Analysis Hooks
   //===--------------------------------------------------------------------===//
 
   bool isLegalToInline(Operation *call, Operation *callable,
                        bool wouldBeCloned) const final {
     // Don't allow inlining calls that are marked `noinline`.
     return !call->hasAttr("noinline");
   }
-  bool isLegalToInline(Region *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Region *, Region *, bool, IRMapping &) const final {
     // Inlining into test dialect regions is legal.
     return true;
   }
-  bool isLegalToInline(Operation *, Region *, bool,
-                       BlockAndValueMapping &) const final {
+  bool isLegalToInline(Operation *, Region *, bool, IRMapping &) const final {
     return true;
   }
 
   bool shouldAnalyzeRecursively(Operation *op) const final {
     // Analyze recursively if this is not a functional region operation, it
     // froms a separate functional scope.
     return !isa<FunctionalRegionOp>(op);
   }

mlir/test/lib/IR/TestClone.cpp

void runOnOperation() override {
 
     Block &regionEntry = region.front();
     Operation *terminator = regionEntry.getTerminator();
 
     // Only handle functions whose returns match the inputs.
     if (terminator->getNumOperands() != regionEntry.getNumArguments())
       return;
 
-    BlockAndValueMapping map;
+    IRMapping map;
     for (auto tup :
          llvm::zip(terminator->getOperands(), regionEntry.getArguments())) {
       if (std::get<0>(tup).getType() != std::get<1>(tup).getType())
         return;
       map.map(std::get<1>(tup), std::get<0>(tup));
     }
 
     OpBuilder builder(op->getContext());

mlir/test/lib/IR/TestSlicing.cpp

 //===- TestSlicing.cpp - Testing slice functionality ----------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements a simple testing pass for slicing.
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/Analysis/SliceAnalysis.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
 #include "mlir/Dialect/Linalg/IR/Linalg.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Support/LLVM.h"
 
 using namespace mlir;
 
 /// Create a function with the same signature as the parent function of `op`
 /// with name being the function name and a `suffix`.
 static LogicalResult createBackwardSliceFunction(Operation *op,
                                                  StringRef suffix) {
   func::FuncOp parentFuncOp = op->getParentOfType<func::FuncOp>();
   OpBuilder builder(parentFuncOp);
   Location loc = op->getLoc();
   std::string clonedFuncOpName = parentFuncOp.getName().str() + suffix.str();
   func::FuncOp clonedFuncOp = builder.create<func::FuncOp>(
       loc, clonedFuncOpName, parentFuncOp.getFunctionType());
-  BlockAndValueMapping mapper;
+  IRMapping mapper;
   builder.setInsertionPointToEnd(clonedFuncOp.addEntryBlock());
   for (const auto &arg : enumerate(parentFuncOp.getArguments()))
     mapper.map(arg.value(), clonedFuncOp.getArgument(arg.index()));
   SetVector<Operation *> slice;
   getBackwardSlice(op, &slice);
   for (Operation *slicedOp : slice)
     builder.clone(*slicedOp, mapper);
   builder.create<func::ReturnOp>(loc);

mlir/test/lib/Transforms/TestInlining.cpp

 //===- TestInlining.cpp - Pass to inline calls in the test dialect --------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // TODO: This pass is only necessary because the main inlining pass
 // has not abstracted away the call+callee relationship. When the inlining
 // interface has this support, this pass should be removed.
 //
 //===----------------------------------------------------------------------===//
 
 #include "TestDialect.h"
 #include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/InliningUtils.h"
 #include "llvm/ADT/StringSet.h"
 
 using namespace mlir;
 using namespace test;
 
 namespace {

mlir/unittests/IR/CMakeLists.txt

 add_mlir_unittest(MLIRIRTests
   AttributeTest.cpp
-  BlockAndValueMapping.cpp
   DialectTest.cpp
   InterfaceTest.cpp
+  IRMapping.cpp
   InterfaceAttachmentTest.cpp
   OperationSupportTest.cpp
   PatternMatchTest.cpp
   ShapedTypeTest.cpp
   SubElementInterfaceTest.cpp
   TypeTest.cpp
 
   DEPENDS
   MLIRTestInterfaceIncGen
 )
 target_include_directories(MLIRIRTests PRIVATE "${MLIR_BINARY_DIR}/test/lib/Dialect/Test")
 target_link_libraries(MLIRIRTests
   PRIVATE
   MLIRIR
   MLIRTestDialect)

mlir/unittests/IR/IRMapping.cpp

-//===- BlockAndValueMapping.h -----------------------------------*- C++ -*-===//
+//===- IRMapping.cpp --------------------------------------------*- 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
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/IR/BlockAndValueMapping.h"
+#include "mlir/IR/IRMapping.h"
 #include "mlir/IR/Builders.h"
 #include "gtest/gtest.h"
 
 #include "../../test/lib/Dialect/Test/TestDialect.h"
 
 using namespace mlir;
 
-TEST(BlockAndValueMapping, TypedValue) {
+TEST(IRMapping, TypedValue) {
   MLIRContext context;
 
   context.loadDialect<test::TestDialect>();
 
   OpBuilder builder(&context);
   Location loc = builder.getUnknownLoc();
 
   Block block;
   builder.setInsertionPointToEnd(&block);
 
   Value i64Val = builder.create<test::TestOpConstant>(
       loc, builder.getI64Type(), builder.getI64IntegerAttr(0));
   Value f64Val = builder.create<test::TestOpConstant>(
       loc, builder.getF64Type(), builder.getF64FloatAttr(0.0));
 
-  BlockAndValueMapping mapping;
+  IRMapping mapping;
   mapping.map(i64Val, f64Val);
   TypedValue<IntegerType> typedI64Val = i64Val;
-  mapping.lookup(typedI64Val);
+  EXPECT_EQ(mapping.lookup(typedI64Val), f64Val);
+}
+
+TEST(IRMapping, OperationClone) {
+  MLIRContext ctx;
+  ctx.allowUnregisteredDialects();
+
+  OperationState state(UnknownLoc::get(&ctx), "no_results");
+  Operation *noResultsOp = Operation::create(state);
+
+  OperationState owner(UnknownLoc::get(&ctx), "owner");
+  owner.addRegion()->emplaceBlock().push_back(noResultsOp);
+  OwningOpRef<Operation *> ownerOp = Operation::create(owner);
+
+  IRMapping irMap;
+  OwningOpRef<Operation *> clonedOwnerOp = (*ownerOp)->clone(irMap);
+
+  EXPECT_EQ(irMap.lookupOrNull(*ownerOp), *clonedOwnerOp);
+  EXPECT_EQ(irMap.lookupOrNull(noResultsOp),
+            &(*clonedOwnerOp)->getRegion(0).front().front());
 }