[mlir] Start migrating more dialects to prefixed form

[jpienaar]

Summary

Marked all dialects that could be (reasonably) easily flipped to _Both
prefix. Updating the accessors to prefixed form will happen in follow
up, this was to flush out conflicts and to mark all dialects explicitly
as I plan to flip OpBase default to _Prefixed to avoid needing to
migrate new dialects.

Except for Standalone example which got flipped to _Prefixed.

[Mogball]

It begins

mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td

924

These should really just be deleted.

[jsetoain]

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

197–198

Should this change be part of this patch? Looks unrelated to all the other changes.

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

689–690

This doesn't look like it belongs in this patch either.

mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp

476–477

Same here. If the purpose is to add type prefixes to accessor methods, this doesn't look like part of that.

[Mogball]

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

197–198

Some of the member functions now do different things because they were hidden before. E.g. getBody now regions the region body and hides the auto-generated function getBody for ops with single block regions

[jpienaar]

mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td

924

Agreed, I kept it mechanical so that folks can just do regex as part of this. I think we should improve this generically rather than these, but I'll leave them as is for now and make it easy to flip. But then yes we should do something better here, these are inconsistently added and feels out of place (we have the variant already that requires a context, perhaps that's an alternative to conciliate around that).

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

197–198

Yes, so Jeff explained the getBody instances. And here the type of steps change to match getSteps (as these are now the same function and getSteps is actually the convenience accessor with the type most used). If you look inside getSteps it's doing the casting internally that was done here. With Mehdi's upcoming change we could simplify this more.

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

689–690

It does as Jeff mentioned above.

mlir/examples/standalone/include/Standalone/StandaloneDialect.td

def Standalone_Dialect : Dialect {
     let name = "standalone";
     let summary = "A standalone out-of-tree MLIR dialect.";
     let description = [{
         This dialect is an example of an out-of-tree MLIR dialect designed to
         illustrate the basic setup required to develop MLIR-based tools without
         working inside of the LLVM source tree.
     }];
     let cppNamespace = "::mlir::standalone";
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed;
 }
 
 //===----------------------------------------------------------------------===//
 // Base standalone operation definition.
 //===----------------------------------------------------------------------===//
 
 class Standalone_Op<string mnemonic, list<Trait> traits = []> :
         Op<Standalone_Dialect, mnemonic, traits>;
 
 #endif // STANDALONE_DIALECT

mlir/include/mlir/Dialect/AMDGPU/AMDGPU.td

def AMDGPU_Dialect : Dialect {
   let name = "amdgpu";
   let cppNamespace = "::mlir::amdgpu";
   let description = [{
     The `AMDGPU` dialect provides wrappers around AMD-specific functionality
     and LLVM intrinsics. These wrappers should be used in conjunction with
     more generic dialects, such as `gpu` and `vector`, when generating LLVM IR
     that will eventually be executed on AMD hardware.
   }];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // AMDGPU Op definitions
 //===----------------------------------------------------------------------===//
 
 class AMDGPU_Op<string mnemonic, list<Trait> traits = []> :
   Op<AMDGPU_Dialect, mnemonic, traits> {}

mlir/include/mlir/Dialect/AMX/AMX.td

let description = [{
 
     Note that since configuration changes (implicit at dialect level) are
     costly, it is highly recommended to use the AMX dialect on same-shaped
     vectors, at least within a single method.
 
     For details, see the Intel documentation:
     https://software.intel.com/content/www/us/en/develop/articles/intel-sdm.html
   }];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // AMX Op and IntrOp definitions.
 //===----------------------------------------------------------------------===//
 
 class AMX_Op<string mnemonic, list<Trait> traits = []> :
   Op<AMX_Dialect, mnemonic, traits> {}

mlir/include/mlir/Dialect/Affine/IR/AffineMemoryOpInterfaces.td

InterfaceMethod<
       /*retTy=*/"NamedAttribute",
       /*methodName=*/"getAffineMapAttrForMemRef",
       /*args=*/(ins "Value":$memref),
       /*methodBody=*/[{}],
       /*defaultImplementation=*/[{
         ConcreteOp op = cast<ConcreteOp>(this->getOperation());
         assert(memref == op.getMemRef() &&
                "Expected memref argument to match memref operand");
-        return {StringAttr::get(op.getContext(), op.getMapAttrName()),
+        return {StringAttr::get(op.getContext(), op.getMapAttrStrName()),
                 op.getAffineMapAttr()};
       }]
     >,
   ];
 }
 
 #endif // AFFINEMEMORYOPINTERFACES

mlir/include/mlir/Dialect/Affine/IR/AffineOps.h

public:
   }
 
   /// Returns the rank (number of indices) of the source MemRefType.
   unsigned getSrcMemRefRank() { return getSrcMemRefType().getRank(); }
 
   /// Returns the affine map used to access the source memref.
   AffineMap getSrcMap() { return getSrcMapAttr().getValue(); }
   AffineMapAttr getSrcMapAttr() {
-    return (*this)->getAttr(getSrcMapAttrName()).cast<AffineMapAttr>();
+    return (*this)->getAttr(getSrcMapAttrStrName()).cast<AffineMapAttr>();
   }
 
   /// Returns the source memref affine map indices for this DMA operation.
   operand_range getSrcIndices() {
     return {operand_begin() + getSrcMemRefOperandIndex() + 1,
             operand_begin() + getSrcMemRefOperandIndex() + 1 +
                 getSrcMap().getNumInputs()};
   }
public:
   /// Returns the memory space of the source memref.
   unsigned getDstMemorySpace() {
     return getDstMemRef().getType().cast<MemRefType>().getMemorySpaceAsInt();
   }
 
   /// Returns the affine map used to access the destination memref.
   AffineMap getDstMap() { return getDstMapAttr().getValue(); }
   AffineMapAttr getDstMapAttr() {
-    return (*this)->getAttr(getDstMapAttrName()).cast<AffineMapAttr>();
+    return (*this)->getAttr(getDstMapAttrStrName()).cast<AffineMapAttr>();
   }
 
   /// Returns the destination memref indices for this DMA operation.
   operand_range getDstIndices() {
     return {operand_begin() + getDstMemRefOperandIndex() + 1,
             operand_begin() + getDstMemRefOperandIndex() + 1 +
                 getDstMap().getNumInputs()};
   }
public:
   /// Returns the rank (number of indices) of the tag MemRefType.
   unsigned getTagMemRefRank() {
     return getTagMemRef().getType().cast<MemRefType>().getRank();
   }
 
   /// Returns the affine map used to access the tag memref.
   AffineMap getTagMap() { return getTagMapAttr().getValue(); }
   AffineMapAttr getTagMapAttr() {
-    return (*this)->getAttr(getTagMapAttrName()).cast<AffineMapAttr>();
+    return (*this)->getAttr(getTagMapAttrStrName()).cast<AffineMapAttr>();
   }
 
   /// Returns the tag memref indices for this DMA operation.
   operand_range getTagIndices() {
     return {operand_begin() + getTagMemRefOperandIndex() + 1,
             operand_begin() + getTagMemRefOperandIndex() + 1 +
                 getTagMap().getNumInputs()};
   }
 
   /// Returns the number of elements being transferred by this DMA operation.
   Value getNumElements() {
     return getOperand(getTagMemRefOperandIndex() + 1 +
                       getTagMap().getNumInputs());
   }
 
   /// Impelements the AffineMapAccessInterface.
   /// Returns the AffineMapAttr associated with 'memref'.
   NamedAttribute getAffineMapAttrForMemRef(Value memref) {
     if (memref == getSrcMemRef())
-      return {StringAttr::get(getContext(), getSrcMapAttrName()),
+      return {StringAttr::get(getContext(), getSrcMapAttrStrName()),
               getSrcMapAttr()};
     if (memref == getDstMemRef())
-      return {StringAttr::get(getContext(), getDstMapAttrName()),
+      return {StringAttr::get(getContext(), getDstMapAttrStrName()),
               getDstMapAttr()};
     assert(memref == getTagMemRef() &&
            "DmaStartOp expected source, destination or tag memref");
-    return {StringAttr::get(getContext(), getTagMapAttrName()),
+    return {StringAttr::get(getContext(), getTagMapAttrStrName()),
             getTagMapAttr()};
   }
 
   /// Returns true if this is a DMA from a faster memory space to a slower one.
   bool isDestMemorySpaceFaster() {
     return (getSrcMemorySpace() < getDstMemorySpace());
   }
 
   /// Returns true if this is a DMA from a slower memory space to a faster one.
   bool isSrcMemorySpaceFaster() {
     // Assumes that a lower number is for a slower memory space.
     return (getDstMemorySpace() < getSrcMemorySpace());
   }
 
   /// Given a DMA start operation, returns the operand position of either the
   /// source or destination memref depending on the one that is at the higher
   /// level of the memory hierarchy. Asserts failure if neither is true.
   unsigned getFasterMemPos() {
     assert(isSrcMemorySpaceFaster() || isDestMemorySpaceFaster());
     return isSrcMemorySpaceFaster() ? 0 : getDstMemRefOperandIndex();
   }
 
-  static StringRef getSrcMapAttrName() { return "src_map"; }
-  static StringRef getDstMapAttrName() { return "dst_map"; }
-  static StringRef getTagMapAttrName() { return "tag_map"; }
+  static StringRef getSrcMapAttrStrName() { return "src_map"; }
+  static StringRef getDstMapAttrStrName() { return "dst_map"; }
+  static StringRef getTagMapAttrStrName() { return "tag_map"; }
 
   static StringRef getOperationName() { return "affine.dma_start"; }
   static ParseResult parse(OpAsmParser &parser, OperationState &result);
   void print(OpAsmPrinter &p);
   LogicalResult verifyInvariantsImpl();
   LogicalResult verifyInvariants() { return verifyInvariantsImpl(); }
   LogicalResult fold(ArrayRef<Attribute> cstOperands,
                      SmallVectorImpl<OpFoldResult> &results);
public:
   Value getTagMemRef() { return getOperand(0); }
   MemRefType getTagMemRefType() {
     return getTagMemRef().getType().cast<MemRefType>();
   }
 
   /// Returns the affine map used to access the tag memref.
   AffineMap getTagMap() { return getTagMapAttr().getValue(); }
   AffineMapAttr getTagMapAttr() {
-    return (*this)->getAttr(getTagMapAttrName()).cast<AffineMapAttr>();
+    return (*this)->getAttr(getTagMapAttrStrName()).cast<AffineMapAttr>();
   }
 
   /// Returns the tag memref index for this DMA operation.
   operand_range getTagIndices() {
     return {operand_begin() + 1,
             operand_begin() + 1 + getTagMap().getNumInputs()};
   }
 
   /// Returns the rank (number of indices) of the tag memref.
   unsigned getTagMemRefRank() {
     return getTagMemRef().getType().cast<MemRefType>().getRank();
   }
 
   /// Impelements the AffineMapAccessInterface. Returns the AffineMapAttr
   /// associated with 'memref'.
   NamedAttribute getAffineMapAttrForMemRef(Value memref) {
     assert(memref == getTagMemRef());
-    return {StringAttr::get(getContext(), getTagMapAttrName()),
+    return {StringAttr::get(getContext(), getTagMapAttrStrName()),
             getTagMapAttr()};
   }
 
   /// Returns the number of elements transferred by the associated DMA op.
   Value getNumElements() { return getOperand(1 + getTagMap().getNumInputs()); }
 
-  static StringRef getTagMapAttrName() { return "tag_map"; }
+  static StringRef getTagMapAttrStrName() { return "tag_map"; }
   static ParseResult parse(OpAsmParser &parser, OperationState &result);
   void print(OpAsmPrinter &p);
   LogicalResult verifyInvariantsImpl();
   LogicalResult verifyInvariants() { return verifyInvariantsImpl(); }
   LogicalResult fold(ArrayRef<Attribute> cstOperands,
                      SmallVectorImpl<OpFoldResult> &results);
 };
 

mlir/include/mlir/Dialect/Affine/IR/AffineOps.td

 include "mlir/Interfaces/LoopLikeInterface.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 def Affine_Dialect : Dialect {
   let name = "affine";
   let cppNamespace = "mlir";
   let hasConstantMaterializer = 1;
   let dependentDialects = ["arith::ArithmeticDialect"];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 // Base class for Affine dialect ops.
 class Affine_Op<string mnemonic, list<Trait> traits = []> :
     Op<Affine_Dialect, mnemonic, traits>;
 
 // Require regions to have affine.yield.
 def ImplicitAffineTerminator
let extraClassDeclaration = [{
 
     /// Returns true if the result of this operation is a symbol in the region
     /// of the closest surrounding op that has the trait AffineScope.
     bool isValidSymbol();
 
     /// Returns true if the result of this operation is a symbol for all its
     /// uses in `region`.
     bool isValidSymbol(Region *region);
-
-    operand_range getMapOperands() { return getOperands(); }
   }];
 
   let hasCanonicalizer = 1;
   let hasCustomAssemblyFormat = 1;
   let hasFolder = 1;
   let hasVerifier = 1;
 }
 
OpBuilder<(ins "ValueRange":$lbOperands, "AffineMap":$lbMap,
            "nullptr">:$bodyBuilder)>
   ];
 
   let extraClassDeclaration = [{
     /// Defining the function type we use for building the body of affine.for.
     using BodyBuilderFn =
         function_ref<void(OpBuilder &, Location, Value, ValueRange)>;
 
-    static StringRef getStepAttrName() { return "step"; }
-    static StringRef getLowerBoundAttrName() { return "lower_bound"; }
-    static StringRef getUpperBoundAttrName() { return "upper_bound"; }
+    static StringRef getStepAttrStrName() { return "step"; }
+    static StringRef getLowerBoundAttrStrName() { return "lower_bound"; }
+    static StringRef getUpperBoundAttrStrName() { return "upper_bound"; }
 
     BlockArgument getInductionVar() { return getBody()->getArgument(0); }
     Block::BlockArgListType getRegionIterArgs() {
       return getBody()->getArguments().drop_front();
     }
     Operation::operand_range getIterOperands() {
       return getOperands().drop_front(getNumControlOperands());
     }
let extraClassDeclaration = [{
     /// Returns information about the lower bound as a single object.
     AffineBound getLowerBound();
 
     /// Returns information about the upper bound as a single object.
     AffineBound getUpperBound();
 
     /// Returns loop step.
     int64_t getStep() {
-      return (*this)->getAttr(getStepAttrName()).cast<IntegerAttr>().getInt();
+      return (*this)->getAttr(getStepAttrStrName()).cast<IntegerAttr>().getInt();
     }
 
     /// Returns affine map for the lower bound.
     AffineMap getLowerBoundMap() { return getLowerBoundMapAttr().getValue(); }
     AffineMapAttr getLowerBoundMapAttr() {
-      return (*this)->getAttr(getLowerBoundAttrName()).cast<AffineMapAttr>();
+      return (*this)->getAttr(getLowerBoundAttrStrName()).cast<AffineMapAttr>();
     }
     /// Returns affine map for the upper bound. The upper bound is exclusive.
     AffineMap getUpperBoundMap() { return getUpperBoundMapAttr().getValue(); }
     AffineMapAttr getUpperBoundMapAttr() {
-      return (*this)->getAttr(getUpperBoundAttrName()).cast<AffineMapAttr>();
+      return (*this)->getAttr(getUpperBoundAttrStrName()).cast<AffineMapAttr>();
     }
 
     /// Set lower bound. The new bound must have the same number of operands as
     /// the current bound map. Otherwise, 'replaceForLowerBound' should be used.
     void setLowerBound(ValueRange operands, AffineMap map);
     /// Set upper bound. The new bound must not have more operands than the
     /// current bound map. Otherwise, 'replaceForUpperBound' should be used.
     void setUpperBound(ValueRange operands, AffineMap map);
 
     /// Set the lower bound map without changing operands.
     void setLowerBoundMap(AffineMap map);
 
     /// Set the upper bound map without changing operands.
     void setUpperBoundMap(AffineMap map);
 
     /// Set loop step.
     void setStep(int64_t step) {
       assert(step > 0 && "step has to be a positive integer constant");
       auto *context = getLowerBoundMap().getContext();
-      (*this)->setAttr(StringAttr::get(context, getStepAttrName()),
+      (*this)->setAttr(StringAttr::get(context, getStepAttrStrName()),
                        IntegerAttr::get(IndexType::get(context), step));
     }
 
     /// Returns number of region arguments for loop-carried values.
     unsigned getNumRegionIterArgs() {
       return getBody()->getNumArguments() - 1;
     }
 
def AffineIfOp : Affine_Op<"if",
   let builders = [
     OpBuilder<(ins "IntegerSet":$set, "ValueRange":$args,
       "bool":$withElseRegion)>,
     OpBuilder<(ins "TypeRange":$resultTypes, "IntegerSet":$set,
       "ValueRange":$args, "bool":$withElseRegion)>,
   ];
 
   let extraClassDeclaration = [{
-    static StringRef getConditionAttrName() { return "condition"; }
+    static StringRef getConditionAttrStrName() { return "condition"; }
 
     IntegerSet getIntegerSet();
     void setIntegerSet(IntegerSet newSet);
 
     /// Sets the integer set with its operands.
     void setConditional(IntegerSet set, ValueRange operands);
 
     /// Returns true if an else block exists.
class AffineLoadOpBase<string mnemonic, list<Trait> traits = []> :
   code extraClassDeclarationBase = [{
     /// Returns the operand index of the memref.
     unsigned getMemRefOperandIndex() { return 0; }
 
     void setMemRef(Value value) { setOperand(getMemRefOperandIndex(), value); }
 
     /// Returns the affine map used to index the memref for this operation.
     AffineMapAttr getAffineMapAttr() {
-      return (*this)->getAttr(getMapAttrName()).cast<AffineMapAttr>();
+      return (*this)->getAttr(getMapAttrStrName()).cast<AffineMapAttr>();
     }
 
-    static StringRef getMapAttrName() { return "map"; }
-
-    // TODO: Remove once prefixing is flipped.
-    operand_range getIndices() { return indices(); }
+    static StringRef getMapAttrStrName() { return "map"; }
   }];
 }
 
 def AffineLoadOp : AffineLoadOpBase<"load"> {
   let summary = "affine load operation";
   let description = [{
     The "affine.load" op reads an element from a memref, where the index
     for each memref dimension is an affine expression of loop induction
class AffineMinMaxOpBase<string mnemonic, list<Trait> traits = []> :
   let builders = [
     OpBuilder<(ins "AffineMap":$affineMap, "ValueRange":$mapOperands),
     [{
       build($_builder, $_state, $_builder.getIndexType(), affineMap, mapOperands);
     }]>
   ];
 
   let extraClassDeclaration = [{
-    static StringRef getMapAttrName() { return "map"; }
+    static StringRef getMapAttrStrName() { return "map"; }
     AffineMap getAffineMap() { return map(); }
     ValueRange getMapOperands() { return operands(); }
     ValueRange getDimOperands() {
       return OperandRange{operands().begin(),
                         operands().begin() + map().getNumDims()};
     }
     ValueRange getSymbolOperands() {
       return OperandRange{operands().begin() + map().getNumDims(),
def AffineParallelOp : Affine_Op<"parallel",
   }];
 
   let arguments = (ins
      TypedArrayAttrBase<AtomicRMWKindAttr, "Reduction ops">:$reductions,
      AffineMapAttr:$lowerBoundsMap,
      I32ElementsAttr:$lowerBoundsGroups,
      AffineMapAttr:$upperBoundsMap,
      I32ElementsAttr:$upperBoundsGroups,
-     I64ArrayAttr:$steps,
+     I64SmallVectorArrayAttr:$steps,
      Variadic<Index>:$mapOperands);
   let results = (outs Variadic<AnyType>:$results);
   let regions = (region SizedRegion<1>:$region);
 
   let builders = [
     OpBuilder<(ins "TypeRange":$resultTypes,
       "ArrayRef<arith::AtomicRMWKind>":$reductions, "ArrayRef<int64_t>":$ranges)>,
     OpBuilder<(ins "TypeRange":$resultTypes,
let extraClassDeclaration = [{
     AffineMap getUpperBoundMap(unsigned pos);
     operand_range getUpperBoundsOperands();
     AffineValueMap getUpperBoundsValueMap();
 
     /// Sets elements fo the loop upper bound.
     void setUpperBounds(ValueRange operands, AffineMap map);
     void setUpperBoundsMap(AffineMap map);
 
-    SmallVector<int64_t, 8> getSteps();
     void setSteps(ArrayRef<int64_t> newSteps);
 
     /// Returns attribute names to use in op construction. Not expected to be
     /// used directly.
-    static StringRef getReductionsAttrName() { return "reductions"; }
-    static StringRef getLowerBoundsMapAttrName() { return "lowerBoundsMap"; }
-    static StringRef getLowerBoundsGroupsAttrName() {
+    static StringRef getReductionsAttrStrName() { return "reductions"; }
+    static StringRef getLowerBoundsMapAttrStrName() { return "lowerBoundsMap"; }
+    static StringRef getLowerBoundsGroupsAttrStrName() {
       return "lowerBoundsGroups";
     }
-    static StringRef getUpperBoundsMapAttrName() { return "upperBoundsMap"; }
-    static StringRef getUpperBoundsGroupsAttrName() {
+    static StringRef getUpperBoundsMapAttrStrName() { return "upperBoundsMap"; }
+    static StringRef getUpperBoundsGroupsAttrStrName() {
       return "upperBoundsGroups";
     }
-    static StringRef getStepsAttrName() { return "steps"; }
+    static StringRef getStepsAttrStrName() { return "steps"; }
 
     /// Returns `true` if the loop bounds have min/max expressions.
     bool hasMinMaxBounds() {
       return lowerBoundsMap().getNumResults() != getNumDims() ||
              upperBoundsMap().getNumResults() != getNumDims();
     }
   }];
 
OpBuilder<(ins "Value":$memref, "AffineMap":$map,
       "bool":$isDataCache),
     [{
       assert(map.getNumInputs() == mapOperands.size()
              && "inconsistent index info");
       auto localityHintAttr = $_builder.getI32IntegerAttr(localityHint);
       auto isWriteAttr = $_builder.getBoolAttr(isWrite);
       auto isDataCacheAttr = $_builder.getBoolAttr(isDataCache);
       $_state.addOperands(memref);
-      $_state.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+      $_state.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
       $_state.addOperands(mapOperands);
-      $_state.addAttribute(getLocalityHintAttrName(), localityHintAttr);
-      $_state.addAttribute(getIsWriteAttrName(), isWriteAttr);
-      $_state.addAttribute(getIsDataCacheAttrName(), isDataCacheAttr);
+      $_state.addAttribute(getLocalityHintAttrStrName(), localityHintAttr);
+      $_state.addAttribute(getIsWriteAttrStrName(), isWriteAttr);
+      $_state.addAttribute(getIsDataCacheAttrStrName(), isDataCacheAttr);
     }]>];
 
   let extraClassDeclaration = [{
     MemRefType getMemRefType() {
       return memref().getType().cast<MemRefType>();
     }
 
     /// Returns the affine map used to index the memref for this operation.
     AffineMap getAffineMap() { return getAffineMapAttr().getValue(); }
     AffineMapAttr getAffineMapAttr() {
-      return (*this)->getAttr(getMapAttrName()).cast<AffineMapAttr>();
+      return (*this)->getAttr(getMapAttrStrName()).cast<AffineMapAttr>();
     }
 
     /// Impelements the AffineMapAccessInterface.
     /// Returns the AffineMapAttr associated with 'memref'.
     NamedAttribute getAffineMapAttrForMemRef(Value mref) {
       assert(mref == memref() &&
              "Expected mref argument to match memref operand");
-      return {StringAttr::get(getContext(), getMapAttrName()),
+      return {StringAttr::get(getContext(), getMapAttrStrName()),
         getAffineMapAttr()};
     }
 
     /// Get affine map operands.
     operand_range getMapOperands() {
       return {operand_begin() + 1, operand_end()};
     }
 
-    static StringRef getMapAttrName() { return "map"; }
-    static StringRef getLocalityHintAttrName() { return "localityHint"; }
-    static StringRef getIsWriteAttrName() { return "isWrite"; }
-    static StringRef getIsDataCacheAttrName() { return "isDataCache"; }
+    static StringRef getMapAttrStrName() { return "map"; }
+    static StringRef getLocalityHintAttrStrName() { return "localityHint"; }
+    static StringRef getIsWriteAttrStrName() { return "isWrite"; }
+    static StringRef getIsDataCacheAttrStrName() { return "isDataCache"; }
   }];
 
   let hasCanonicalizer = 1;
   let hasCustomAssemblyFormat = 1;
   let hasFolder = 1;
   let hasVerifier = 1;
 }
 
 class AffineStoreOpBase<string mnemonic, list<Trait> traits = []> :
     Affine_Op<mnemonic, !listconcat(traits,
     [DeclareOpInterfaceMethods<AffineWriteOpInterface>,
     DeclareOpInterfaceMethods<AffineMapAccessInterface>,
     MemRefsNormalizable])> {
   code extraClassDeclarationBase = [{
     /// Returns the operand index of the value to be stored.
     unsigned getStoredValOperandIndex() { return 0; }
 
     /// Returns the operand index of the memref.
     unsigned getMemRefOperandIndex() { return 1; }
 
     void setMemRef(Value value) { setOperand(getMemRefOperandIndex(), value); }
 
     /// Returns the affine map used to index the memref for this operation.
     AffineMapAttr getAffineMapAttr() {
-      return (*this)->getAttr(getMapAttrName()).cast<AffineMapAttr>();
+      return (*this)->getAttr(getMapAttrStrName()).cast<AffineMapAttr>();
     }
 
-    static StringRef getMapAttrName() { return "map"; }
-
-    // TODO: Remove once prefixing is flipped.
-    operand_range getIndices() { return indices(); }
+    static StringRef getMapAttrStrName() { return "map"; }
   }];
 }
 
 def AffineStoreOp : AffineStoreOpBase<"store"> {
   let summary = "affine store operation";
   let description = [{
     The "affine.store" op writes an element to a memref, where the index
     for each memref dimension is an affine expression of loop induction

mlir/include/mlir/Dialect/ArmNeon/ArmNeon.td

 
 def ArmNeon_Dialect : Dialect {
   let name = "arm_neon";
   let cppNamespace = "::mlir::arm_neon";
 
   // Note: this does not need to depend on LLVMDialect as long as functions in
   // this dialect (such as canonicalization) do not produce entities belonging
   // to the LLVMDialect (ops or types).
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // ArmNeon op definitions
 //===----------------------------------------------------------------------===//
 
 // ArmNeon dialect op that corresponds (and is convertible to) an LLVM IR
 // intrinsic.

mlir/include/mlir/Dialect/ArmSVE/ArmSVE.td

 def ArmSVE_Dialect : Dialect {
   let name = "arm_sve";
   let cppNamespace = "::mlir::arm_sve";
   let summary = "Basic dialect to target Arm SVE architectures";
   let description = [{
     This dialect contains the definitions necessary to target specific Arm SVE
     scalable vector operations.
   }];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // ArmSVE op definitions
 //===----------------------------------------------------------------------===//
 
 class ArmSVE_Op<string mnemonic, list<Trait> traits = []> :
   Op<ArmSVE_Dialect, mnemonic, traits> {}

mlir/include/mlir/Dialect/Async/IR/AsyncDialect.td

let extraClassDeclaration = [{
     /// The name of a unit attribute on funcs that are allowed to have a
     /// blocking async.runtime.await ops. Only useful in combination with
     /// 'eliminate-blocking-await-ops' option, which in absence of this
     /// attribute might convert a func to a coroutine.
     static constexpr StringRef kAllowedToBlockAttrName =
         "async.allowed_to_block";
   }];
 
+  // TODO: Prefixed form overlaps with generated names, update before flipping.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 #endif // ASYNC_DIALECT_TD

mlir/include/mlir/Dialect/Bufferization/IR/BufferizationBase.td

constexpr const static ::llvm::StringLiteral
         kWritableAttrName = "bufferization.writable";
 
     /// Attribute name used to mark the bufferization layout for region
     /// arguments during One-Shot Module Bufferize.
     constexpr const static ::llvm::StringLiteral
         kBufferLayoutAttrName = "bufferization.buffer_layout";
   }];
   let hasOperationAttrVerify = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 #endif // BUFFERIZATION_BASE

mlir/include/mlir/Dialect/DLTI/DLTIBase.td

let extraClassDeclaration = [{
     constexpr const static ::llvm::StringLiteral
     kDataLayoutEndiannessBig = "big";
 
     constexpr const static ::llvm::StringLiteral
     kDataLayoutEndiannessLittle = "little";
   }];
 
   let useDefaultAttributePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 def DLTI_DataLayoutEntryAttr : DialectAttr<
     DLTI_Dialect,
     CPred<"$_self.isa<::mlir::DataLayoutEntryAttr>()">,
     "Target data layout entry"> {
   let storageType = "::mlir::DataLayoutEntryAttr";
   let returnType = "::mlir::DataLayoutEntryAttr";

mlir/include/mlir/Dialect/EmitC/IR/EmitCBase.td

def EmitC_Dialect : Dialect {
   let summary = "Dialect to generate C/C++ from MLIR.";
   let description = [{
     [include "Dialects/emitc.md"]
   }];
 
   let hasConstantMaterializer = 1;
   let useDefaultTypePrinterParser = 1;
   let useDefaultAttributePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 #endif // MLIR_DIALECT_EMITC_IR_EMITCBASE

mlir/include/mlir/Dialect/GPU/IR/GPUBase.td

let extraClassDeclaration = [{
     /// Returns the numeric value used to identify the private memory address
     /// space.
     static unsigned getPrivateAddressSpace() { return 5; }
   }];
 
   let dependentDialects = ["arith::ArithmeticDialect"];
   let useDefaultAttributePrinterParser = 1;
   let useDefaultTypePrinterParser = 1;
+
+  // TODO: This has overlapping accessors with generated when switched to
+  // prefixed. Fix and update to _Both & then _Prefixed.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 def GPU_AsyncToken : DialectType<
   GPU_Dialect, CPred<"$_self.isa<::mlir::gpu::AsyncTokenType>()">, "async token type">,
              BuildableType<"mlir::gpu::AsyncTokenType::get($_builder.getContext())">;
 
 // Predicat to check if type is gpu::MMAMatrixType.
 def IsMMAMatrixTypePred : CPred<"$_self.isa<::mlir::gpu::MMAMatrixType>()">;

mlir/include/mlir/Dialect/LLVMIR/LLVMIntrinsicOps.td

 def LLVM_VPFPToSIOp : LLVM_VPCastF<"fptosi">;
 
 def LLVM_VPUIToFPOp : LLVM_VPCastI<"uitofp">;
 def LLVM_VPSIToFPOp : LLVM_VPCastI<"sitofp">;
 
 def LLVM_VPPtrToIntOp : LLVM_VPCastPtr<"ptrtoint">;
 def LLVM_VPIntToPtrOp : LLVM_VPCastI<"inttoptr">;
 
-#endif // LLVM_INTRINSIC_OPS
+#endif // LLVM_INTRINSIC_OP
 

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

def NVVM_Dialect : Dialect {
 
   let extraClassDeclaration = [{
     /// Get the name of the attribute used to annotate external kernel
     /// functions.
     static StringRef getKernelFuncAttrName() { return "nvvm.kernel"; }
   }];
 
   let useDefaultAttributePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // NVVM op definitions
 //===----------------------------------------------------------------------===//
 
 class NVVM_Op<string mnemonic, list<Trait> traits = []> :
   LLVM_OpBase<NVVM_Dialect, mnemonic, traits> {

mlir/include/mlir/Dialect/LLVMIR/ROCDLOps.td

def ROCDL_Dialect : Dialect {
   let dependentDialects = ["LLVM::LLVMDialect"];
   let hasOperationAttrVerify = 1;
 
   let extraClassDeclaration = [{
     /// Get the name of the attribute used to annotate external kernel
     /// functions.
     static StringRef getKernelFuncAttrName() { return "rocdl.kernel"; }
   }];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // ROCDL op definitions
 //===----------------------------------------------------------------------===//
 
 class ROCDL_Op<string mnemonic, list<Trait> traits = []> :
   LLVM_OpBase<ROCDL_Dialect, mnemonic, traits> {

mlir/include/mlir/Dialect/Linalg/IR/LinalgBase.td

let extraClassDeclaration = [{
     using RegionBuilderFunType = llvm::function_ref<
       void(ImplicitLocOpBuilder &b, Block &, ArrayRef<NamedAttribute>)>;
     RegionBuilderFunType getRegionBuilder(StringRef name) {
       return namedStructuredOpRegionBuilders.lookup(name);
     }
     private:
       llvm::StringMap<RegionBuilderFunType> namedStructuredOpRegionBuilders;
   }];
+
+  // TODO: This has overlapping accessors with generated when switched to
+  // prefixed. Fix and update to _Both & then _Prefixed.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 // Define the function attribute enums matching the OpDSL functions.
 def UnaryFn : I32EnumAttr<"UnaryFn", "", [
   I32EnumAttrCase<"exp", 0>,
   I32EnumAttrCase<"log", 1>,
   I32EnumAttrCase<"abs", 2>,
   I32EnumAttrCase<"ceil", 3>,

mlir/include/mlir/Dialect/MemRef/IR/MemRefBase.td

def MemRef_Dialect : Dialect {
   let cppNamespace = "::mlir::memref";
   let description = [{
     The `memref` dialect is intended to hold core memref creation and
     manipulation ops, which are not strongly associated with any particular
     other dialect or domain abstraction.
   }];
   let dependentDialects = ["arith::ArithmeticDialect"];
   let hasConstantMaterializer = 1;
+
+  // TODO: This has overlapping accessors with generated when switched to
+  // prefixed. Fix and update to _Both & then _Prefixed.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 #endif // MEMREF_BASE

mlir/include/mlir/Dialect/NVGPU/NVGPU.td

def NVGPU_Dialect : Dialect {
   let description = [{
     This `NVGPU` dialect provides a bridge between the target agnostic GPU and
     Vector dialects and the lower level LLVM IR based NVVM dialect. This allow
     representing PTX specific operations while using MLIR high level concepts
     like memref and 2-D vector.
   }];
 
   let useDefaultTypePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 /// Device-side synchronization token.
 def NVGPU_DeviceAsyncToken : DialectType<
   NVGPU_Dialect, CPred<"$_self.isa<::mlir::nvgpu::DeviceAsyncTokenType>()">,
    "device async token type">,
    BuildableType<
      "mlir::nvgpu::DeviceAsyncTokenType::get($_builder.getContext())">;

mlir/include/mlir/Dialect/PDL/IR/PDLDialect.td

def PDL_Dialect : Dialect {
 
   let name = "pdl";
   let cppNamespace = "::mlir::pdl";
 
   let useDefaultTypePrinterParser = 1;
   let extraClassDeclaration = [{
     void registerTypes();
   }];
+
+  // FIXME: Prefixed accessors overlap with builtin Operation members. Flip
+  // once resolved.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 #endif // MLIR_DIALECT_PDL_IR_PDLDIALECT

mlir/include/mlir/Dialect/Quant/QuantOpsBase.td

 
 include "mlir/IR/OpBase.td"
 
 def Quantization_Dialect : Dialect {
   let name = "quant";
   let cppNamespace = "::mlir::quant";
 
   let useDefaultTypePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // Quantization type definitions
 //===----------------------------------------------------------------------===//
 
 class quant_TypedPrimitiveOrContainer<Type etype> :
     Type<Or<[etype.predicate,

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVBase.td

let extraClassDeclaration = [{
 
     /// Dialect attribute parsing hook.
     Attribute parseAttribute(
         DialectAsmParser &parser, Type type) const override;
     /// Dialect attribute printing hook.
     void printAttribute(
         Attribute attr, DialectAsmPrinter &printer) const override;
   }];
+
+  // TODO: This has overlapping accessors with generated when switched to
+  // prefixed. Fix and update to _Both & then _Prefixed.
+  let emitAccessorPrefix = kEmitAccessorPrefix_Raw;
 }
 
 //===----------------------------------------------------------------------===//
 // Utility definitions
 //===----------------------------------------------------------------------===//
 
 // A predicate that checks whether `$_self` is a known enum case for the
 // enum class with `name`.

mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorBase.td

let description = [{
     and Saman Amarasinghe. Tensor Algebra Compilation with Workspaces,
     Proceedings of the IEEE/ACM International Symposium on Code Generation
     and Optimization, 2019.
     * [Kjolstad20] Fredrik Berg Kjolstad. Sparse Tensor Algebra Compilation.
     PhD thesis, MIT, February, 2020.
   }];
 
   let useDefaultAttributePrinterParser = 1;
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 #endif // SPARSETENSOR_BASE

mlir/include/mlir/Dialect/Tensor/IR/TensorBase.td

def Tensor_Dialect : Dialect {
 
   }];
 
   let hasConstantMaterializer = 1;
   let dependentDialects = [
     "arith::ArithmeticDialect",
     "complex::ComplexDialect",
   ];
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 #endif // TENSOR_BASE

mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td

 //===----------------------------------------------------------------------===//
 
 class Tensor_ReassociativeReshapeOp<string mnemonic, list<Trait> traits = []> :
     Tensor_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
     Arguments<(ins AnyTensor:$src, IndexListArrayAttr:$reassociation)>,
     Results<(outs AnyTensor:$result)> {
 
   code commonExtraClassDeclaration = [{
-    static StringRef getReassociationAttrName() { return "reassociation"; }
+    static StringRef getReassociationAttrStrName() { return "reassociation"; }
     SmallVector<AffineMap, 4> getReassociationMaps();
     SmallVector<ReassociationExprs, 4> getReassociationExprs();
     SmallVector<ReassociationIndices, 4> getReassociationIndices() {
       SmallVector<ReassociationIndices, 4> reassociationIndices;
       for (auto attr : reassociation())
         reassociationIndices.push_back(llvm::to_vector<2>(
             llvm::map_range(attr.cast<ArrayAttr>(), [&](Attribute indexAttr) {
               return indexAttr.cast<IntegerAttr>().getInt();
let assemblyFormat = [{
     $source
     (`nofold` $nofold^)?
     `low` `` custom<OperandsOrIntegersSizesList>($low, $static_low)
     `high` `` custom<OperandsOrIntegersSizesList>($high, $static_high)
     $region attr-dict `:` type($source) `to` type($result)
   }];
 
   let extraClassDeclaration = [{
-    static StringRef getStaticLowAttrName() {
+    static StringRef getStaticLowAttrStrName() {
       return "static_low";
     }
 
-    static StringRef getStaticHighAttrName() {
+    static StringRef getStaticHighAttrStrName() {
       return "static_high";
     }
 
     RankedTensorType getSourceType() {
       return source().getType().cast<RankedTensorType>();
     }
     RankedTensorType getResultType() {
       return getResult().getType().cast<RankedTensorType>();

mlir/include/mlir/Dialect/X86Vector/X86Vector.td

 
 //===----------------------------------------------------------------------===//
 // X86Vector dialect definition
 //===----------------------------------------------------------------------===//
 
 def X86Vector_Dialect : Dialect {
   let name = "x86vector";
   let cppNamespace = "::mlir::x86vector";
+
+  let emitAccessorPrefix = kEmitAccessorPrefix_Both;
 }
 
 //===----------------------------------------------------------------------===//
 // AVX512 op definitions
 //===----------------------------------------------------------------------===//
 
 // Operation that is part of the input dialect.
 class AVX512_Op<string mnemonic, list<Trait> traits = []> :

mlir/include/mlir/IR/OpBase.td

 def I32ArrayAttr : TypedArrayAttrBase<I32Attr,
                                       "32-bit integer array attribute"> {
   let constBuilderCall = "$_builder.getI32ArrayAttr($0)";
 }
 def I64ArrayAttr : TypedArrayAttrBase<I64Attr,
                                       "64-bit integer array attribute"> {
   let constBuilderCall = "$_builder.getI64ArrayAttr($0)";
 }
+// Variant of I64ArrayAttr whose user accessor is SmallVector<in64_t>.
+def I64SmallVectorArrayAttr :
+    TypedArrayAttrBase<I64Attr, "64-bit integer array attribute"> {
+  let returnType = [{ ::llvm::SmallVector<int64_t, 8> }];
+  let convertFromStorage = [{
+    llvm::to_vector<4>(
+      llvm::map_range($_self.getAsRange<mlir::IntegerAttr>(),
+      [](IntegerAttr attr) { return attr.getInt(); }));
+  }];
+  let constBuilderCall = "$_builder.getI64ArrayAttr($0)";
+}
 def F32ArrayAttr : TypedArrayAttrBase<F32Attr, "32-bit float array attribute"> {
   let constBuilderCall = "$_builder.getF32ArrayAttr($0)";
 }
 def F64ArrayAttr : TypedArrayAttrBase<F64Attr, "64-bit float array attribute"> {
   let constBuilderCall = "$_builder.getF64ArrayAttr($0)";
 }
 def StrArrayAttr : TypedArrayAttrBase<StrAttr, "string array attribute"> {
   let constBuilderCall = "$_builder.getStrArrayAttr($0)";

mlir/lib/Conversion/AffineToStandard/AffineToStandard.cpp

for (unsigned i = 0, e = op.getNumDims(); i < e; ++i) {
 
       Value upper = lowerAffineMapMin(rewriter, loc, op.getUpperBoundMap(i),
                                       op.getUpperBoundsOperands());
       if (!upper)
         return rewriter.notifyMatchFailure(op, "couldn't convert upper bounds");
       upperBoundTuple.push_back(upper);
     }
     steps.reserve(op.steps().size());
-    for (Attribute step : op.steps())
-      steps.push_back(rewriter.create<arith::ConstantIndexOp>(
-          loc, step.cast<IntegerAttr>().getInt()));
+    for (int64_t step : op.steps())
+      steps.push_back(rewriter.create<arith::ConstantIndexOp>(loc, step));
 
     // Get the terminator op.
     Operation *affineParOpTerminator = op.getBody()->getTerminator();
     scf::ParallelOp parOp;
     if (op.results().empty()) {
       // Case with no reduction operations/return values.
       parOp = rewriter.create<scf::ParallelOp>(loc, lowerBoundTuple,
                                                upperBoundTuple, steps,

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

 void AffineDmaStartOp::build(OpBuilder &builder, OperationState &result,
                              Value srcMemRef, AffineMap srcMap,
                              ValueRange srcIndices, Value destMemRef,
                              AffineMap dstMap, ValueRange destIndices,
                              Value tagMemRef, AffineMap tagMap,
                              ValueRange tagIndices, Value numElements,
                              Value stride, Value elementsPerStride) {
   result.addOperands(srcMemRef);
-  result.addAttribute(getSrcMapAttrName(), AffineMapAttr::get(srcMap));
+  result.addAttribute(getSrcMapAttrStrName(), AffineMapAttr::get(srcMap));
   result.addOperands(srcIndices);
   result.addOperands(destMemRef);
-  result.addAttribute(getDstMapAttrName(), AffineMapAttr::get(dstMap));
+  result.addAttribute(getDstMapAttrStrName(), AffineMapAttr::get(dstMap));
   result.addOperands(destIndices);
   result.addOperands(tagMemRef);
-  result.addAttribute(getTagMapAttrName(), AffineMapAttr::get(tagMap));
+  result.addAttribute(getTagMapAttrStrName(), AffineMapAttr::get(tagMap));
   result.addOperands(tagIndices);
   result.addOperands(numElements);
   if (stride) {
     result.addOperands({stride, elementsPerStride});
   }
 }
 
 void AffineDmaStartOp::print(OpAsmPrinter &p) {
ParseResult AffineDmaStartOp::parse(OpAsmParser &parser,
 
   // Parse and resolve the following list of operands:
   // *) dst memref followed by its affine maps operands (in square brackets).
   // *) src memref followed by its affine map operands (in square brackets).
   // *) tag memref followed by its affine map operands (in square brackets).
   // *) number of elements transferred by DMA operation.
   if (parser.parseOperand(srcMemRefInfo) ||
       parser.parseAffineMapOfSSAIds(srcMapOperands, srcMapAttr,
-                                    getSrcMapAttrName(), result.attributes) ||
+                                    getSrcMapAttrStrName(),
+                                    result.attributes) ||
       parser.parseComma() || parser.parseOperand(dstMemRefInfo) ||
       parser.parseAffineMapOfSSAIds(dstMapOperands, dstMapAttr,
-                                    getDstMapAttrName(), result.attributes) ||
+                                    getDstMapAttrStrName(),
+                                    result.attributes) ||
       parser.parseComma() || parser.parseOperand(tagMemRefInfo) ||
       parser.parseAffineMapOfSSAIds(tagMapOperands, tagMapAttr,
-                                    getTagMapAttrName(), result.attributes) ||
+                                    getTagMapAttrStrName(),
+                                    result.attributes) ||
       parser.parseComma() || parser.parseOperand(numElementsInfo))
     return failure();
 
   // Parse optional stride and elements per stride.
   if (parser.parseTrailingOperandList(strideInfo))
     return failure();
 
   if (!strideInfo.empty() && strideInfo.size() != 2) {

 // AffineDmaWaitOp
 //===----------------------------------------------------------------------===//
 
 // TODO: Check that map operands are loop IVs or symbols.
 void AffineDmaWaitOp::build(OpBuilder &builder, OperationState &result,
                             Value tagMemRef, AffineMap tagMap,
                             ValueRange tagIndices, Value numElements) {
   result.addOperands(tagMemRef);
-  result.addAttribute(getTagMapAttrName(), AffineMapAttr::get(tagMap));
+  result.addAttribute(getTagMapAttrStrName(), AffineMapAttr::get(tagMap));
   result.addOperands(tagIndices);
   result.addOperands(numElements);
 }
 
 void AffineDmaWaitOp::print(OpAsmPrinter &p) {
   p << " " << getTagMemRef() << '[';
   SmallVector<Value, 2> operands(getTagIndices());
   p.printAffineMapOfSSAIds(getTagMapAttr(), operands);
ParseResult AffineDmaWaitOp::parse(OpAsmParser &parser,
   SmallVector<OpAsmParser::UnresolvedOperand, 2> tagMapOperands;
   Type type;
   auto indexType = parser.getBuilder().getIndexType();
   OpAsmParser::UnresolvedOperand numElementsInfo;
 
   // Parse tag memref, its map operands, and dma size.
   if (parser.parseOperand(tagMemRefInfo) ||
       parser.parseAffineMapOfSSAIds(tagMapOperands, tagMapAttr,
-                                    getTagMapAttrName(), result.attributes) ||
+                                    getTagMapAttrStrName(),
+                                    result.attributes) ||
       parser.parseComma() || parser.parseOperand(numElementsInfo) ||
       parser.parseColonType(type) ||
       parser.resolveOperand(tagMemRefInfo, type, result.operands) ||
       parser.resolveOperands(tagMapOperands, indexType, result.operands) ||
       parser.resolveOperand(numElementsInfo, indexType, result.operands))
     return failure();
 
   if (!type.isa<MemRefType>())
assert(((!ubMap && ubOperands.empty()) ||
           ubOperands.size() == ubMap.getNumInputs()) &&
          "upper bound operand count does not match the affine map");
   assert(step > 0 && "step has to be a positive integer constant");
 
   for (Value val : iterArgs)
     result.addTypes(val.getType());
 
   // Add an attribute for the step.
-  result.addAttribute(getStepAttrName(),
+  result.addAttribute(getStepAttrStrName(),
                       builder.getIntegerAttr(builder.getIndexType(), step));
 
   // Add the lower bound.
-  result.addAttribute(getLowerBoundAttrName(), AffineMapAttr::get(lbMap));
+  result.addAttribute(getLowerBoundAttrStrName(), AffineMapAttr::get(lbMap));
   result.addOperands(lbOperands);
 
   // Add the upper bound.
-  result.addAttribute(getUpperBoundAttrName(), AffineMapAttr::get(ubMap));
+  result.addAttribute(getUpperBoundAttrStrName(), AffineMapAttr::get(ubMap));
   result.addOperands(ubOperands);
 
   result.addOperands(iterArgs);
   // Create a region and a block for the body.  The argument of the region is
   // the loop induction variable.
   Region *bodyRegion = result.addRegion();
   bodyRegion->push_back(new Block);
   Block &bodyBlock = bodyRegion->front();

 static ParseResult parseBound(bool isLower, OperationState &result,
                               OpAsmParser &p) {
   // 'min' / 'max' prefixes are generally syntactic sugar, but are required if
   // the map has multiple results.
   bool failedToParsedMinMax =
       failed(p.parseOptionalKeyword(isLower ? "max" : "min"));
 
   auto &builder = p.getBuilder();
-  auto boundAttrName = isLower ? AffineForOp::getLowerBoundAttrName()
-                               : AffineForOp::getUpperBoundAttrName();
+  auto boundAttrStrName = isLower ? AffineForOp::getLowerBoundAttrStrName()
+                                  : AffineForOp::getUpperBoundAttrStrName();
 
   // Parse ssa-id as identity map.
   SmallVector<OpAsmParser::UnresolvedOperand, 1> boundOpInfos;
   if (p.parseOperandList(boundOpInfos))
     return failure();
 
   if (!boundOpInfos.empty()) {
     // Check that only one operand was parsed.
     if (boundOpInfos.size() > 1)
       return p.emitError(p.getNameLoc(),
                          "expected only one loop bound operand");
 
     // TODO: improve error message when SSA value is not of index type.
     // Currently it is 'use of value ... expects different type than prior uses'
     if (p.resolveOperand(boundOpInfos.front(), builder.getIndexType(),
                          result.operands))
       return failure();
 
     // Create an identity map using symbol id. This representation is optimized
     // for storage. Analysis passes may expand it into a multi-dimensional map
     // if desired.
     AffineMap map = builder.getSymbolIdentityMap();
-    result.addAttribute(boundAttrName, AffineMapAttr::get(map));
+    result.addAttribute(boundAttrStrName, AffineMapAttr::get(map));
     return success();
   }
 
   // Get the attribute location.
   SMLoc attrLoc = p.getCurrentLocation();
 
   Attribute boundAttr;
-  if (p.parseAttribute(boundAttr, builder.getIndexType(), boundAttrName,
+  if (p.parseAttribute(boundAttr, builder.getIndexType(), boundAttrStrName,
                        result.attributes))
     return failure();
 
   // Parse full form - affine map followed by dim and symbol list.
   if (auto affineMapAttr = boundAttr.dyn_cast<AffineMapAttr>()) {
     unsigned currentNumOperands = result.operands.size();
     unsigned numDims;
     if (parseDimAndSymbolList(p, result.operands, numDims))
if (auto affineMapAttr = boundAttr.dyn_cast<AffineMapAttr>()) {
     }
     return success();
   }
 
   // Parse custom assembly form.
   if (auto integerAttr = boundAttr.dyn_cast<IntegerAttr>()) {
     result.attributes.pop_back();
     result.addAttribute(
-        boundAttrName,
+        boundAttrStrName,
         AffineMapAttr::get(builder.getConstantAffineMap(integerAttr.getInt())));
     return success();
   }
 
   return p.emitError(
       p.getNameLoc(),
       "expected valid affine map representation for loop bounds");
 }
ParseResult AffineForOp::parse(OpAsmParser &parser, OperationState &result) {
   if (parseBound(/*isLower=*/true, result, parser) ||
       parser.parseKeyword("to", " between bounds") ||
       parseBound(/*isLower=*/false, result, parser))
     return failure();
 
   // Parse the optional loop step, we default to 1 if one is not present.
   if (parser.parseOptionalKeyword("step")) {
     result.addAttribute(
-        AffineForOp::getStepAttrName(),
+        AffineForOp::getStepAttrStrName(),
         builder.getIntegerAttr(builder.getIndexType(), /*value=*/1));
   } else {
     SMLoc stepLoc = parser.getCurrentLocation();
     IntegerAttr stepAttr;
     if (parser.parseAttribute(stepAttr, builder.getIndexType(),
-                              AffineForOp::getStepAttrName().data(),
+                              AffineForOp::getStepAttrStrName().data(),
                               result.attributes))
       return failure();
 
     if (stepAttr.getValue().getSExtValue() < 0)
       return parser.emitError(
           stepLoc,
           "expected step to be representable as a positive signed integer");
   }
if (getNumIterOperands() > 0) {
     });
     p << ") -> (" << getResultTypes() << ")";
     printBlockTerminators = true;
   }
 
   p << ' ';
   p.printRegion(region(), /*printEntryBlockArgs=*/false, printBlockTerminators);
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/{getLowerBoundAttrName(),
-                                           getUpperBoundAttrName(),
-                                           getStepAttrName()});
+                          /*elidedAttrs=*/{getLowerBoundAttrStrName(),
+                                           getUpperBoundAttrStrName(),
+                                           getStepAttrStrName()});
 }
 
 /// Fold the constant bounds of a loop.
 static LogicalResult foldLoopBounds(AffineForOp forOp) {
   auto foldLowerOrUpperBound = [&forOp](bool lower) {
     // Check to see if each of the operands is the result of a constant.  If
     // so, get the value.  If not, ignore it.
     SmallVector<Attribute, 8> operandConstants;
void AffineForOp::setLowerBound(ValueRange lbOperands, AffineMap map) {
   SmallVector<Value, 4> newOperands(lbOperands.begin(), lbOperands.end());
 
   auto ubOperands = getUpperBoundOperands();
   newOperands.append(ubOperands.begin(), ubOperands.end());
   auto iterOperands = getIterOperands();
   newOperands.append(iterOperands.begin(), iterOperands.end());
   (*this)->setOperands(newOperands);
 
-  (*this)->setAttr(getLowerBoundAttrName(), AffineMapAttr::get(map));
+  (*this)->setAttr(getLowerBoundAttrStrName(), AffineMapAttr::get(map));
 }
 
 void AffineForOp::setUpperBound(ValueRange ubOperands, AffineMap map) {
   assert(ubOperands.size() == map.getNumInputs());
   assert(map.getNumResults() >= 1 && "bound map has at least one result");
 
   SmallVector<Value, 4> newOperands(getLowerBoundOperands());
   newOperands.append(ubOperands.begin(), ubOperands.end());
   auto iterOperands = getIterOperands();
   newOperands.append(iterOperands.begin(), iterOperands.end());
   (*this)->setOperands(newOperands);
 
-  (*this)->setAttr(getUpperBoundAttrName(), AffineMapAttr::get(map));
+  (*this)->setAttr(getUpperBoundAttrStrName(), AffineMapAttr::get(map));
 }
 
 void AffineForOp::setLowerBoundMap(AffineMap map) {
   auto lbMap = getLowerBoundMap();
   assert(lbMap.getNumDims() == map.getNumDims() &&
          lbMap.getNumSymbols() == map.getNumSymbols());
   assert(map.getNumResults() >= 1 && "bound map has at least one result");
   (void)lbMap;
-  (*this)->setAttr(getLowerBoundAttrName(), AffineMapAttr::get(map));
+  (*this)->setAttr(getLowerBoundAttrStrName(), AffineMapAttr::get(map));
 }
 
 void AffineForOp::setUpperBoundMap(AffineMap map) {
   auto ubMap = getUpperBoundMap();
   assert(ubMap.getNumDims() == map.getNumDims() &&
          ubMap.getNumSymbols() == map.getNumSymbols());
   assert(map.getNumResults() >= 1 && "bound map has at least one result");
   (void)ubMap;
-  (*this)->setAttr(getUpperBoundAttrName(), AffineMapAttr::get(map));
+  (*this)->setAttr(getUpperBoundAttrStrName(), AffineMapAttr::get(map));
 }
 
 bool AffineForOp::hasConstantLowerBound() {
   return getLowerBoundMap().isSingleConstant();
 }
 
 bool AffineForOp::hasConstantUpperBound() {
   return getUpperBoundMap().isSingleConstant();
struct AlwaysTrueOrFalseIf : public OpRewritePattern<AffineIfOp> {
   }
 };
 } // namespace
 
 LogicalResult AffineIfOp::verify() {
   // Verify that we have a condition attribute.
   // FIXME: This should be specified in the arguments list in ODS.
   auto conditionAttr =
-      (*this)->getAttrOfType<IntegerSetAttr>(getConditionAttrName());
+      (*this)->getAttrOfType<IntegerSetAttr>(getConditionAttrStrName());
   if (!conditionAttr)
     return emitOpError("requires an integer set attribute named 'condition'");
 
   // Verify that there are enough operands for the condition.
   IntegerSet condition = conditionAttr.getValue();
   if (getNumOperands() != condition.getNumInputs())
     return emitOpError("operand count and condition integer set dimension and "
                        "symbol count must match");
 
   // Verify that the operands are valid dimension/symbols.
   if (failed(verifyDimAndSymbolIdentifiers(*this, getOperands(),
                                            condition.getNumDims())))
     return failure();
 
   return success();
 }
 
 ParseResult AffineIfOp::parse(OpAsmParser &parser, OperationState &result) {
   // Parse the condition attribute set.
   IntegerSetAttr conditionAttr;
   unsigned numDims;
-  if (parser.parseAttribute(conditionAttr, AffineIfOp::getConditionAttrName(),
+  if (parser.parseAttribute(conditionAttr,
+                            AffineIfOp::getConditionAttrStrName(),
                             result.attributes) ||
       parseDimAndSymbolList(parser, result.operands, numDims))
     return failure();
 
   // Verify the condition operands.
   auto set = conditionAttr.getValue();
   if (set.getNumDims() != numDims)
     return parser.emitError(
ParseResult AffineIfOp::parse(OpAsmParser &parser, OperationState &result) {
   if (parser.parseOptionalAttrDict(result.attributes))
     return failure();
 
   return success();
 }
 
 void AffineIfOp::print(OpAsmPrinter &p) {
   auto conditionAttr =
-      (*this)->getAttrOfType<IntegerSetAttr>(getConditionAttrName());
+      (*this)->getAttrOfType<IntegerSetAttr>(getConditionAttrStrName());
   p << " " << conditionAttr;
   printDimAndSymbolList(operand_begin(), operand_end(),
                         conditionAttr.getValue().getNumDims(), p);
   p.printOptionalArrowTypeList(getResultTypes());
   p << ' ';
   p.printRegion(thenRegion(), /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/getNumResults());
 
   // Print the 'else' regions if it has any blocks.
   auto &elseRegion = this->elseRegion();
   if (!elseRegion.empty()) {
     p << " else ";
     p.printRegion(elseRegion,
                   /*printEntryBlockArgs=*/false,
                   /*printBlockTerminators=*/getNumResults());
   }
 
   // Print the attribute list.
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/getConditionAttrName());
+                          /*elidedAttrs=*/getConditionAttrStrName());
 }
 
 IntegerSet AffineIfOp::getIntegerSet() {
   return (*this)
-      ->getAttrOfType<IntegerSetAttr>(getConditionAttrName())
+      ->getAttrOfType<IntegerSetAttr>(getConditionAttrStrName())
       .getValue();
 }
 
 void AffineIfOp::setIntegerSet(IntegerSet newSet) {
-  (*this)->setAttr(getConditionAttrName(), IntegerSetAttr::get(newSet));
+  (*this)->setAttr(getConditionAttrStrName(), IntegerSetAttr::get(newSet));
 }
 
 void AffineIfOp::setConditional(IntegerSet set, ValueRange operands) {
   setIntegerSet(set);
   (*this)->setOperands(operands);
 }
 
 void AffineIfOp::build(OpBuilder &builder, OperationState &result,
                        TypeRange resultTypes, IntegerSet set, ValueRange args,
                        bool withElseRegion) {
   assert(resultTypes.empty() || withElseRegion);
   result.addTypes(resultTypes);
   result.addOperands(args);
-  result.addAttribute(getConditionAttrName(), IntegerSetAttr::get(set));
+  result.addAttribute(getConditionAttrStrName(), IntegerSetAttr::get(set));
 
   Region *thenRegion = result.addRegion();
   thenRegion->push_back(new Block());
   if (resultTypes.empty())
     AffineIfOp::ensureTerminator(*thenRegion, builder, result.location);
 
   Region *elseRegion = result.addRegion();
   if (withElseRegion) {

 // AffineLoadOp
 //===----------------------------------------------------------------------===//
 
 void AffineLoadOp::build(OpBuilder &builder, OperationState &result,
                          AffineMap map, ValueRange operands) {
   assert(operands.size() == 1 + map.getNumInputs() && "inconsistent operands");
   result.addOperands(operands);
   if (map)
-    result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+    result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
   auto memrefType = operands[0].getType().cast<MemRefType>();
   result.types.push_back(memrefType.getElementType());
 }
 
 void AffineLoadOp::build(OpBuilder &builder, OperationState &result,
                          Value memref, AffineMap map, ValueRange mapOperands) {
   assert(map.getNumInputs() == mapOperands.size() && "inconsistent index info");
   result.addOperands(memref);
   result.addOperands(mapOperands);
   auto memrefType = memref.getType().cast<MemRefType>();
-  result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+  result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
   result.types.push_back(memrefType.getElementType());
 }
 
 void AffineLoadOp::build(OpBuilder &builder, OperationState &result,
                          Value memref, ValueRange indices) {
   auto memrefType = memref.getType().cast<MemRefType>();
   int64_t rank = memrefType.getRank();
   // Create identity map for memrefs with at least one dimension or () -> ()
ParseResult AffineLoadOp::parse(OpAsmParser &parser, OperationState &result) {
 
   MemRefType type;
   OpAsmParser::UnresolvedOperand memrefInfo;
   AffineMapAttr mapAttr;
   SmallVector<OpAsmParser::UnresolvedOperand, 1> mapOperands;
   return failure(
       parser.parseOperand(memrefInfo) ||
       parser.parseAffineMapOfSSAIds(mapOperands, mapAttr,
-                                    AffineLoadOp::getMapAttrName(),
+                                    AffineLoadOp::getMapAttrStrName(),
                                     result.attributes) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(type) ||
       parser.resolveOperand(memrefInfo, type, result.operands) ||
       parser.resolveOperands(mapOperands, indexTy, result.operands) ||
       parser.addTypeToList(type.getElementType(), result.types));
 }
 
 void AffineLoadOp::print(OpAsmPrinter &p) {
   p << " " << getMemRef() << '[';
   if (AffineMapAttr mapAttr =
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()))
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()))
     p.printAffineMapOfSSAIds(mapAttr, getMapOperands());
   p << ']';
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/{getMapAttrName()});
+                          /*elidedAttrs=*/{getMapAttrStrName()});
   p << " : " << getMemRefType();
 }
 
 /// Verify common indexing invariants of affine.load, affine.store,
 /// affine.vector_load and affine.vector_store.
 static LogicalResult
 verifyMemoryOpIndexing(Operation *op, AffineMapAttr mapAttr,
                        Operation::operand_range mapOperands,

 
 LogicalResult AffineLoadOp::verify() {
   auto memrefType = getMemRefType();
   if (getType() != memrefType.getElementType())
     return emitOpError("result type must match element type of memref");
 
   if (failed(verifyMemoryOpIndexing(
           getOperation(),
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()),
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()),
           getMapOperands(), memrefType,
           /*numIndexOperands=*/getNumOperands() - 1)))
     return failure();
 
   return success();
 }
 
 void AffineLoadOp::getCanonicalizationPatterns(RewritePatternSet &results,

 
 void AffineStoreOp::build(OpBuilder &builder, OperationState &result,
                           Value valueToStore, Value memref, AffineMap map,
                           ValueRange mapOperands) {
   assert(map.getNumInputs() == mapOperands.size() && "inconsistent index info");
   result.addOperands(valueToStore);
   result.addOperands(memref);
   result.addOperands(mapOperands);
-  result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+  result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
 }
 
 // Use identity map.
 void AffineStoreOp::build(OpBuilder &builder, OperationState &result,
                           Value valueToStore, Value memref,
                           ValueRange indices) {
   auto memrefType = memref.getType().cast<MemRefType>();
   int64_t rank = memrefType.getRank();
ParseResult AffineStoreOp::parse(OpAsmParser &parser, OperationState &result) {
 
   MemRefType type;
   OpAsmParser::UnresolvedOperand storeValueInfo;
   OpAsmParser::UnresolvedOperand memrefInfo;
   AffineMapAttr mapAttr;
   SmallVector<OpAsmParser::UnresolvedOperand, 1> mapOperands;
   return failure(parser.parseOperand(storeValueInfo) || parser.parseComma() ||
                  parser.parseOperand(memrefInfo) ||
-                 parser.parseAffineMapOfSSAIds(mapOperands, mapAttr,
-                                               AffineStoreOp::getMapAttrName(),
-                                               result.attributes) ||
+                 parser.parseAffineMapOfSSAIds(
+                     mapOperands, mapAttr, AffineStoreOp::getMapAttrStrName(),
                  parser.parseOptionalAttrDict(result.attributes) ||
                  parser.parseColonType(type) ||
                  parser.resolveOperand(storeValueInfo, type.getElementType(),
                                        result.operands) ||
                  parser.resolveOperand(memrefInfo, type, result.operands) ||
                  parser.resolveOperands(mapOperands, indexTy, result.operands));
 }
 
 void AffineStoreOp::print(OpAsmPrinter &p) {
   p << " " << getValueToStore();
   p << ", " << getMemRef() << '[';
   if (AffineMapAttr mapAttr =
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()))
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()))
     p.printAffineMapOfSSAIds(mapAttr, getMapOperands());
   p << ']';
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/{getMapAttrName()});
+                          /*elidedAttrs=*/{getMapAttrStrName()});
   p << " : " << getMemRefType();
 }
 
 LogicalResult AffineStoreOp::verify() {
   // The value to store must have the same type as memref element type.
   auto memrefType = getMemRefType();
   if (getValueToStore().getType() != memrefType.getElementType())
     return emitOpError(
         "value to store must have the same type as memref element type");
 
   if (failed(verifyMemoryOpIndexing(
           getOperation(),
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()),
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()),
           getMapOperands(), memrefType,
           /*numIndexOperands=*/getNumOperands() - 2)))
     return failure();
 
   return success();
 }
 
 void AffineStoreOp::getCanonicalizationPatterns(RewritePatternSet &results,
template <typename T> static LogicalResult verifyAffineMinMaxOp(T op) {
   // Verify that operand count matches affine map dimension and symbol count.
   if (op.getNumOperands() != op.map().getNumDims() + op.map().getNumSymbols())
     return op.emitOpError(
         "operand count and affine map dimension and symbol count must match");
   return success();
 }
 
 template <typename T> static void printAffineMinMaxOp(OpAsmPrinter &p, T op) {
-  p << ' ' << op->getAttr(T::getMapAttrName());
+  p << ' ' << op->getAttr(T::getMapAttrStrName());
   auto operands = op.getOperands();
   unsigned numDims = op.map().getNumDims();
   p << '(' << operands.take_front(numDims) << ')';
 
   if (operands.size() != numDims)
     p << '[' << operands.drop_front(numDims) << ']';
   p.printOptionalAttrDict(op->getAttrs(),
-                          /*elidedAttrs=*/{T::getMapAttrName()});
+                          /*elidedAttrs=*/{T::getMapAttrStrName()});
 }
 
 template <typename T>
 static ParseResult parseAffineMinMaxOp(OpAsmParser &parser,
                                        OperationState &result) {
   auto &builder = parser.getBuilder();
   auto indexType = builder.getIndexType();
   SmallVector<OpAsmParser::UnresolvedOperand, 8> dimInfos;
   SmallVector<OpAsmParser::UnresolvedOperand, 8> symInfos;
   AffineMapAttr mapAttr;
   return failure(
-      parser.parseAttribute(mapAttr, T::getMapAttrName(), result.attributes) ||
+      parser.parseAttribute(mapAttr, T::getMapAttrStrName(),
+                            result.attributes) ||
       parser.parseOperandList(dimInfos, OpAsmParser::Delimiter::Paren) ||
       parser.parseOperandList(symInfos,
                               OpAsmParser::Delimiter::OptionalSquare) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.resolveOperands(dimInfos, indexType, result.operands) ||
       parser.resolveOperands(symInfos, indexType, result.operands) ||
       parser.addTypeToList(indexType, result.types));
 }
ParseResult AffinePrefetchOp::parse(OpAsmParser &parser,
   IntegerAttr hintInfo;
   auto i32Type = parser.getBuilder().getIntegerType(32);
   StringRef readOrWrite, cacheType;
 
   AffineMapAttr mapAttr;
   SmallVector<OpAsmParser::UnresolvedOperand, 1> mapOperands;
   if (parser.parseOperand(memrefInfo) ||
       parser.parseAffineMapOfSSAIds(mapOperands, mapAttr,
-                                    AffinePrefetchOp::getMapAttrName(),
+                                    AffinePrefetchOp::getMapAttrStrName(),
                                     result.attributes) ||
       parser.parseComma() || parser.parseKeyword(&readOrWrite) ||
       parser.parseComma() || parser.parseKeyword("locality") ||
       parser.parseLess() ||
       parser.parseAttribute(hintInfo, i32Type,
-                            AffinePrefetchOp::getLocalityHintAttrName(),
+                            AffinePrefetchOp::getLocalityHintAttrStrName(),
                             result.attributes) ||
       parser.parseGreater() || parser.parseComma() ||
       parser.parseKeyword(&cacheType) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(type) ||
       parser.resolveOperand(memrefInfo, type, result.operands) ||
       parser.resolveOperands(mapOperands, indexTy, result.operands))
     return failure();
 
   if (!readOrWrite.equals("read") && !readOrWrite.equals("write"))
     return parser.emitError(parser.getNameLoc(),
                             "rw specifier has to be 'read' or 'write'");
   result.addAttribute(
-      AffinePrefetchOp::getIsWriteAttrName(),
+      AffinePrefetchOp::getIsWriteAttrStrName(),
       parser.getBuilder().getBoolAttr(readOrWrite.equals("write")));
 
   if (!cacheType.equals("data") && !cacheType.equals("instr"))
     return parser.emitError(parser.getNameLoc(),
                             "cache type has to be 'data' or 'instr'");
 
   result.addAttribute(
-      AffinePrefetchOp::getIsDataCacheAttrName(),
+      AffinePrefetchOp::getIsDataCacheAttrStrName(),
       parser.getBuilder().getBoolAttr(cacheType.equals("data")));
 
   return success();
 }
 
 void AffinePrefetchOp::print(OpAsmPrinter &p) {
   p << " " << memref() << '[';
   AffineMapAttr mapAttr =
-      (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName());
+      (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName());
   if (mapAttr)
     p.printAffineMapOfSSAIds(mapAttr, getMapOperands());
   p << ']' << ", " << (isWrite() ? "write" : "read") << ", "
     << "locality<" << localityHint() << ">, "
     << (isDataCache() ? "data" : "instr");
   p.printOptionalAttrDict(
       (*this)->getAttrs(),
-      /*elidedAttrs=*/{getMapAttrName(), getLocalityHintAttrName(),
-                       getIsDataCacheAttrName(), getIsWriteAttrName()});
+      /*elidedAttrs=*/{getMapAttrStrName(), getLocalityHintAttrStrName(),
+                       getIsDataCacheAttrStrName(), getIsWriteAttrStrName()});
   p << " : " << getMemRefType();
 }
 
 LogicalResult AffinePrefetchOp::verify() {
-  auto mapAttr = (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName());
+  auto mapAttr = (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName());
   if (mapAttr) {
     AffineMap map = mapAttr.getValue();
     if (map.getNumResults() != getMemRefType().getRank())
       return emitOpError("affine.prefetch affine map num results must equal"
                          " memref rank");
     if (map.getNumInputs() + 1 != getNumOperands())
       return emitOpError("too few operands");
   } else {
void AffineParallelOp::build(OpBuilder &builder, OperationState &result,
 
   result.addTypes(resultTypes);
 
   // Convert the reductions to integer attributes.
   SmallVector<Attribute, 4> reductionAttrs;
   for (arith::AtomicRMWKind reduction : reductions)
     reductionAttrs.push_back(
         builder.getI64IntegerAttr(static_cast<int64_t>(reduction)));
-  result.addAttribute(getReductionsAttrName(),
+  result.addAttribute(getReductionsAttrStrName(),
                       builder.getArrayAttr(reductionAttrs));
 
   // Concatenates maps defined in the same input space (same dimensions and
   // symbols), assumes there is at least one map.
   auto concatMapsSameInput = [&builder](ArrayRef<AffineMap> maps,
                                         SmallVectorImpl<int32_t> &groups) {
     if (maps.empty())
       return AffineMap::get(builder.getContext());
     SmallVector<AffineExpr> exprs;
     groups.reserve(groups.size() + maps.size());
     exprs.reserve(maps.size());
     for (AffineMap m : maps) {
       llvm::append_range(exprs, m.getResults());
       groups.push_back(m.getNumResults());
     }
     return AffineMap::get(maps[0].getNumDims(), maps[0].getNumSymbols(), exprs,
                           maps[0].getContext());
   };
 
   // Set up the bounds.
   SmallVector<int32_t> lbGroups, ubGroups;
   AffineMap lbMap = concatMapsSameInput(lbMaps, lbGroups);
   AffineMap ubMap = concatMapsSameInput(ubMaps, ubGroups);
-  result.addAttribute(getLowerBoundsMapAttrName(), AffineMapAttr::get(lbMap));
-  result.addAttribute(getLowerBoundsGroupsAttrName(),
+  result.addAttribute(getLowerBoundsMapAttrStrName(),
+                      AffineMapAttr::get(lbMap));
+  result.addAttribute(getLowerBoundsGroupsAttrStrName(),
                       builder.getI32TensorAttr(lbGroups));
-  result.addAttribute(getUpperBoundsMapAttrName(), AffineMapAttr::get(ubMap));
-  result.addAttribute(getUpperBoundsGroupsAttrName(),
+  result.addAttribute(getUpperBoundsMapAttrStrName(),
+                      AffineMapAttr::get(ubMap));
+  result.addAttribute(getUpperBoundsGroupsAttrStrName(),
                       builder.getI32TensorAttr(ubGroups));
-  result.addAttribute(getStepsAttrName(), builder.getI64ArrayAttr(steps));
+  result.addAttribute(getStepsAttrStrName(), builder.getI64ArrayAttr(steps));
   result.addOperands(lbArgs);
   result.addOperands(ubArgs);
 
   // Create a region and a block for the body.
   auto *bodyRegion = result.addRegion();
   auto *body = new Block();
   // Add all the block arguments.
   for (unsigned i = 0, e = steps.size(); i < e; ++i)

 void AffineParallelOp::setUpperBoundsMap(AffineMap map) {
   AffineMap ubMap = upperBoundsMap();
   assert(ubMap.getNumDims() == map.getNumDims() &&
          ubMap.getNumSymbols() == map.getNumSymbols());
   (void)ubMap;
   upperBoundsMapAttr(AffineMapAttr::get(map));
 }
 
-SmallVector<int64_t, 8> AffineParallelOp::getSteps() {
-  SmallVector<int64_t, 8> result;
-  for (Attribute attr : steps()) {
-    result.push_back(attr.cast<IntegerAttr>().getInt());
-  }
-  return result;
-}
-
 void AffineParallelOp::setSteps(ArrayRef<int64_t> newSteps) {
   stepsAttr(getBodyBuilder().getI64ArrayAttr(newSteps));
 }
 
 LogicalResult AffineParallelOp::verify() {
   auto numDims = getNumDims();
   if (lowerBoundsGroups().getNumElements() != numDims ||
       upperBoundsGroups().getNumElements() != numDims ||
if (getNumResults()) {
     p << ") -> (" << getResultTypes() << ")";
   }
 
   p << ' ';
   p.printRegion(region(), /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/getNumResults());
   p.printOptionalAttrDict(
       (*this)->getAttrs(),
-      /*elidedAttrs=*/{AffineParallelOp::getReductionsAttrName(),
-                       AffineParallelOp::getLowerBoundsMapAttrName(),
-                       AffineParallelOp::getLowerBoundsGroupsAttrName(),
-                       AffineParallelOp::getUpperBoundsMapAttrName(),
-                       AffineParallelOp::getUpperBoundsGroupsAttrName(),
-                       AffineParallelOp::getStepsAttrName()});
+      /*elidedAttrs=*/{AffineParallelOp::getReductionsAttrStrName(),
+                       AffineParallelOp::getLowerBoundsMapAttrStrName(),
+                       AffineParallelOp::getLowerBoundsGroupsAttrStrName(),
+                       AffineParallelOp::getUpperBoundsMapAttrStrName(),
+                       AffineParallelOp::getUpperBoundsGroupsAttrStrName(),
+                       AffineParallelOp::getStepsAttrStrName()});
 }
 
 /// Given a list of lists of parsed operands, populates `uniqueOperands` with
 /// unique operands. Also populates `replacements with affine expressions of
 /// `kind` that can be used to update affine maps previously accepting a
 /// `operands` to accept `uniqueOperands` instead.
 static ParseResult deduplicateAndResolveOperands(
     OpAsmParser &parser,

 /// expr-of-ssa-ids-list ::= expr-of-ssa-ids (`,` expr-of-ssa-id-list)?
 ///
 /// Examples:
 ///   (%0, min(%1 + %2, %3), %4, min(%5 floordiv 32, %6))
 ///   (%0, max(%1 - 2 * %2))
 static ParseResult parseAffineMapWithMinMax(OpAsmParser &parser,
                                             OperationState &result,
                                             MinMaxKind kind) {
-  constexpr llvm::StringLiteral tmpAttrName = "__pseudo_bound_map";
+  constexpr llvm::StringLiteral tmpAttrStrName = "__pseudo_bound_map";
 
   StringRef mapName = kind == MinMaxKind::Min
-                          ? AffineParallelOp::getUpperBoundsMapAttrName()
-                          : AffineParallelOp::getLowerBoundsMapAttrName();
-  StringRef groupsName = kind == MinMaxKind::Min
-                             ? AffineParallelOp::getUpperBoundsGroupsAttrName()
-                             : AffineParallelOp::getLowerBoundsGroupsAttrName();
+                          ? AffineParallelOp::getUpperBoundsMapAttrStrName()
+                          : AffineParallelOp::getLowerBoundsMapAttrStrName();
+  StringRef groupsName =
+      kind == MinMaxKind::Min
+          ? AffineParallelOp::getUpperBoundsGroupsAttrStrName()
+          : AffineParallelOp::getLowerBoundsGroupsAttrStrName();
 
   if (failed(parser.parseLParen()))
     return failure();
 
   if (succeeded(parser.parseOptionalRParen())) {
     result.addAttribute(
         mapName, AffineMapAttr::get(parser.getBuilder().getEmptyAffineMap()));
     result.addAttribute(groupsName, parser.getBuilder().getI32TensorAttr({}));
     return success();
   }
 
   SmallVector<AffineExpr> flatExprs;
   SmallVector<SmallVector<OpAsmParser::UnresolvedOperand>> flatDimOperands;
   SmallVector<SmallVector<OpAsmParser::UnresolvedOperand>> flatSymOperands;
   SmallVector<int32_t> numMapsPerGroup;
   SmallVector<OpAsmParser::UnresolvedOperand> mapOperands;
   auto parseOperands = [&]() {
     if (succeeded(parser.parseOptionalKeyword(
             kind == MinMaxKind::Min ? "min" : "max"))) {
       mapOperands.clear();
       AffineMapAttr map;
-      if (failed(parser.parseAffineMapOfSSAIds(mapOperands, map, tmpAttrName,
+      if (failed(parser.parseAffineMapOfSSAIds(mapOperands, map, tmpAttrStrName,
                                                result.attributes,
                                                OpAsmParser::Delimiter::Paren)))
         return failure();
-      result.attributes.erase(tmpAttrName);
+      result.attributes.erase(tmpAttrStrName);
       llvm::append_range(flatExprs, map.getValue().getResults());
       auto operandsRef = llvm::makeArrayRef(mapOperands);
       auto dimsRef = operandsRef.take_front(map.getValue().getNumDims());
       SmallVector<OpAsmParser::UnresolvedOperand> dims(dimsRef.begin(),
                                                        dimsRef.end());
       auto symsRef = operandsRef.drop_front(map.getValue().getNumDims());
       SmallVector<OpAsmParser::UnresolvedOperand> syms(symsRef.begin(),
                                                        symsRef.end());
if (parser.parseArgumentList(ivs, OpAsmParser::Delimiter::Paren) ||
       parseAffineMapWithMinMax(parser, result, MinMaxKind::Min))
     return failure();
 
   AffineMapAttr stepsMapAttr;
   NamedAttrList stepsAttrs;
   SmallVector<OpAsmParser::UnresolvedOperand, 4> stepsMapOperands;
   if (failed(parser.parseOptionalKeyword("step"))) {
     SmallVector<int64_t, 4> steps(ivs.size(), 1);
-    result.addAttribute(AffineParallelOp::getStepsAttrName(),
+    result.addAttribute(AffineParallelOp::getStepsAttrStrName(),
                         builder.getI64ArrayAttr(steps));
   } else {
     if (parser.parseAffineMapOfSSAIds(stepsMapOperands, stepsMapAttr,
-                                      AffineParallelOp::getStepsAttrName(),
+                                      AffineParallelOp::getStepsAttrStrName(),
                                       stepsAttrs,
                                       OpAsmParser::Delimiter::Paren))
       return failure();
 
     // Convert steps from an AffineMap into an I64ArrayAttr.
     SmallVector<int64_t, 4> steps;
     auto stepsMap = stepsMapAttr.getValue();
     for (const auto &result : stepsMap.getResults()) {
       auto constExpr = result.dyn_cast<AffineConstantExpr>();
       if (!constExpr)
         return parser.emitError(parser.getNameLoc(),
                                 "steps must be constant integers");
       steps.push_back(constExpr.getValue());
     }
-    result.addAttribute(AffineParallelOp::getStepsAttrName(),
+    result.addAttribute(AffineParallelOp::getStepsAttrStrName(),
                         builder.getI64ArrayAttr(steps));
   }
 
   // Parse optional clause of the form: `reduce ("addf", "maxf")`, where the
   // quoted strings are a member of the enum AtomicRMWKind.
   SmallVector<Attribute, 4> reductions;
   if (succeeded(parser.parseOptionalKeyword("reduce"))) {
     if (parser.parseLParen())
auto parseAttributes = [&]() -> ParseResult {
       reductions.push_back(builder.getI64IntegerAttr(
           static_cast<int64_t>(reduction.getValue())));
       // While we keep getting commas, keep parsing.
       return success();
     };
     if (parser.parseCommaSeparatedList(parseAttributes) || parser.parseRParen())
       return failure();
   }
-  result.addAttribute(AffineParallelOp::getReductionsAttrName(),
+  result.addAttribute(AffineParallelOp::getReductionsAttrStrName(),
                       builder.getArrayAttr(reductions));
 
   // Parse return types of reductions (if any)
   if (parser.parseOptionalArrowTypeList(result.types))
     return failure();
 
   // Now parse the body.
   Region *body = result.addRegion();

 //===----------------------------------------------------------------------===//
 
 void AffineVectorLoadOp::build(OpBuilder &builder, OperationState &result,
                                VectorType resultType, AffineMap map,
                                ValueRange operands) {
   assert(operands.size() == 1 + map.getNumInputs() && "inconsistent operands");
   result.addOperands(operands);
   if (map)
-    result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+    result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
   result.types.push_back(resultType);
 }
 
 void AffineVectorLoadOp::build(OpBuilder &builder, OperationState &result,
                                VectorType resultType, Value memref,
                                AffineMap map, ValueRange mapOperands) {
   assert(map.getNumInputs() == mapOperands.size() && "inconsistent index info");
   result.addOperands(memref);
   result.addOperands(mapOperands);
-  result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+  result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
   result.types.push_back(resultType);
 }
 
 void AffineVectorLoadOp::build(OpBuilder &builder, OperationState &result,
                                VectorType resultType, Value memref,
                                ValueRange indices) {
   auto memrefType = memref.getType().cast<MemRefType>();
   int64_t rank = memrefType.getRank();
ParseResult AffineVectorLoadOp::parse(OpAsmParser &parser,
   MemRefType memrefType;
   VectorType resultType;
   OpAsmParser::UnresolvedOperand memrefInfo;
   AffineMapAttr mapAttr;
   SmallVector<OpAsmParser::UnresolvedOperand, 1> mapOperands;
   return failure(
       parser.parseOperand(memrefInfo) ||
       parser.parseAffineMapOfSSAIds(mapOperands, mapAttr,
-                                    AffineVectorLoadOp::getMapAttrName(),
+                                    AffineVectorLoadOp::getMapAttrStrName(),
                                     result.attributes) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(memrefType) || parser.parseComma() ||
       parser.parseType(resultType) ||
       parser.resolveOperand(memrefInfo, memrefType, result.operands) ||
       parser.resolveOperands(mapOperands, indexTy, result.operands) ||
       parser.addTypeToList(resultType, result.types));
 }
 
 void AffineVectorLoadOp::print(OpAsmPrinter &p) {
   p << " " << getMemRef() << '[';
   if (AffineMapAttr mapAttr =
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()))
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()))
     p.printAffineMapOfSSAIds(mapAttr, getMapOperands());
   p << ']';
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/{getMapAttrName()});
+                          /*elidedAttrs=*/{getMapAttrStrName()});
   p << " : " << getMemRefType() << ", " << getType();
 }
 
 /// Verify common invariants of affine.vector_load and affine.vector_store.
 static LogicalResult verifyVectorMemoryOp(Operation *op, MemRefType memrefType,
                                           VectorType vectorType) {
   // Check that memref and vector element types match.
   if (memrefType.getElementType() != vectorType.getElementType())
     return op->emitOpError(
         "requires memref and vector types of the same elemental type");
   return success();
 }
 
 LogicalResult AffineVectorLoadOp::verify() {
   MemRefType memrefType = getMemRefType();
   if (failed(verifyMemoryOpIndexing(
           getOperation(),
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()),
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()),
           getMapOperands(), memrefType,
           /*numIndexOperands=*/getNumOperands() - 1)))
     return failure();
 
   if (failed(verifyVectorMemoryOp(getOperation(), memrefType, getVectorType())))
     return failure();
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // AffineVectorStoreOp
 //===----------------------------------------------------------------------===//
 
 void AffineVectorStoreOp::build(OpBuilder &builder, OperationState &result,
                                 Value valueToStore, Value memref, AffineMap map,
                                 ValueRange mapOperands) {
   assert(map.getNumInputs() == mapOperands.size() && "inconsistent index info");
   result.addOperands(valueToStore);
   result.addOperands(memref);
   result.addOperands(mapOperands);
-  result.addAttribute(getMapAttrName(), AffineMapAttr::get(map));
+  result.addAttribute(getMapAttrStrName(), AffineMapAttr::get(map));
 }
 
 // Use identity map.
 void AffineVectorStoreOp::build(OpBuilder &builder, OperationState &result,
                                 Value valueToStore, Value memref,
                                 ValueRange indices) {
   auto memrefType = memref.getType().cast<MemRefType>();
   int64_t rank = memrefType.getRank();
ParseResult AffineVectorStoreOp::parse(OpAsmParser &parser,
   OpAsmParser::UnresolvedOperand storeValueInfo;
   OpAsmParser::UnresolvedOperand memrefInfo;
   AffineMapAttr mapAttr;
   SmallVector<OpAsmParser::UnresolvedOperand, 1> mapOperands;
   return failure(
       parser.parseOperand(storeValueInfo) || parser.parseComma() ||
       parser.parseOperand(memrefInfo) ||
       parser.parseAffineMapOfSSAIds(mapOperands, mapAttr,
-                                    AffineVectorStoreOp::getMapAttrName(),
+                                    AffineVectorStoreOp::getMapAttrStrName(),
                                     result.attributes) ||
       parser.parseOptionalAttrDict(result.attributes) ||
       parser.parseColonType(memrefType) || parser.parseComma() ||
       parser.parseType(resultType) ||
       parser.resolveOperand(storeValueInfo, resultType, result.operands) ||
       parser.resolveOperand(memrefInfo, memrefType, result.operands) ||
       parser.resolveOperands(mapOperands, indexTy, result.operands));
 }
 
 void AffineVectorStoreOp::print(OpAsmPrinter &p) {
   p << " " << getValueToStore();
   p << ", " << getMemRef() << '[';
   if (AffineMapAttr mapAttr =
-          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()))
+          (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()))
     p.printAffineMapOfSSAIds(mapAttr, getMapOperands());
   p << ']';
   p.printOptionalAttrDict((*this)->getAttrs(),
-                          /*elidedAttrs=*/{getMapAttrName()});
+                          /*elidedAttrs=*/{getMapAttrStrName()});
   p << " : " << getMemRefType() << ", " << getValueToStore().getType();
 }
 
 LogicalResult AffineVectorStoreOp::verify() {
   MemRefType memrefType = getMemRefType();
   if (failed(verifyMemoryOpIndexing(
-          *this, (*this)->getAttrOfType<AffineMapAttr>(getMapAttrName()),
+          *this, (*this)->getAttrOfType<AffineMapAttr>(getMapAttrStrName()),
           getMapOperands(), memrefType,
           /*numIndexOperands=*/getNumOperands() - 2)))
     return failure();
 
   if (failed(verifyVectorMemoryOp(*this, memrefType, getVectorType())))
     return failure();
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // TableGen'd op method definitions
 //===----------------------------------------------------------------------===//
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Affine/IR/AffineOps.cpp.inc"

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

struct ExtractFromTensorGenerate : public OpRewritePattern<tensor::ExtractOp> {
   using OpRewritePattern<tensor::ExtractOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(tensor::ExtractOp extract,
                                 PatternRewriter &rewriter) const final {
     auto tensorFromElements = extract.tensor().getDefiningOp<GenerateOp>();
     if (!tensorFromElements || !wouldOpBeTriviallyDead(tensorFromElements))
       return failure();
 
     BlockAndValueMapping mapping;
-    Block *body = tensorFromElements.getBody();
+    Block *body = &tensorFromElements.getBody().front();
     mapping.map(body->getArguments(), extract.indices());
     for (auto &op : body->without_terminator())
       rewriter.clone(op, mapping);
 
     auto yield = cast<YieldOp>(body->getTerminator());
 
     rewriter.replaceOp(extract, mapping.lookupOrDefault(yield.value()));
     return success();

 void CollapseShapeOp::build(OpBuilder &b, OperationState &result, Value src,
                             ArrayRef<ReassociationIndices> reassociation,
                             ArrayRef<NamedAttribute> attrs) {
   auto resultType = computeTensorReshapeCollapsedType(
       src.getType().cast<RankedTensorType>(),
       getSymbolLessAffineMaps(
           convertReassociationIndicesToExprs(b.getContext(), reassociation)));
   build(b, result, resultType, src, attrs);
-  result.addAttribute(getReassociationAttrName(),
+  result.addAttribute(getReassociationAttrStrName(),
                       getReassociationIndicesAttribute(b, reassociation));
 }
 
 template <typename TensorReshapeOp, bool isExpansion = std::is_same<
                                         TensorReshapeOp, ExpandShapeOp>::value>
 static LogicalResult verifyTensorReshapeOp(TensorReshapeOp op,
                                            RankedTensorType expandedType,
                                            RankedTensorType collapsedType) {

mlir/lib/Dialect/Tensor/Transforms/BufferizableOpInterfaceImpl.cpp

for (int i = 0; i < rank; i++) {
       upperBounds.push_back(upperBound);
     }
 
     // Generate tensor elements with a parallel loop that stores into
     // each element of the resulting memref. We use mergeBlockBefore to "move"
     // this op's body into the scf.parallel's body.
     auto parallel =
         rewriter.create<scf::ParallelOp>(loc, lowerBounds, upperBounds, steps);
     Block *parallelBody = parallel.getBody();
-    rewriter.mergeBlockBefore(generateOp.getBody(),
+    rewriter.mergeBlockBefore(&generateOp.getBody().front(),
                               parallelBody->getTerminator(),
                               parallelBody->getArguments());
     // Replace the inlined yield op with a store op. The scf.parallel's builder
     // already populated an scf.yield at the end, so we don't need to worry
     // about creating that.
     Operation *elementYield = parallelBody->getTerminator()->getPrevNode();
     rewriter.setInsertionPointAfter(elementYield);
     rewriter.replaceOpWithNewOp<memref::StoreOp>(