Change ASM Op printer to print the operation name in the framework instead of leaving it up to each individual operation

[mehdi_amini]

Summary

This aligns the printer with the parser contract: the operation isn't part of the user-controllable part of the syntax.

Depends On D108803

[rriddle]

mlir/lib/IR/Operation.cpp

650

Why printEscapedString? AFAIK we aren't going to be able to parse any escaped characters back in.

[mehdi_amini]

Fix for printOneResultOp

[lattner]

Thank you so much for doing this!

[mehdi_amini]

Fix parseOneResultSameOperandTypeOp to support parsing the generic form emitted by printOneResultOp

One consequence of the current patch is that using printGenericOp from a custom printer
requires supporting parsing back the generic form from the custom parser, since the op name
is already printed without quotes.

[mehdi_amini]

Rebase

[mehdi_amini]

Fix flang

[bondhugula]

mlir/lib/IR/AsmPrinter.cpp

385

Is this method named correctly? It's not printing the entire op.

2515

Nit: Print the rest of the op now.

2532

Nit: single quotes

[mehdi_amini]

mlir/lib/IR/AsmPrinter.cpp

385

With the default value printGenericOp(op) is still printing the entire op, I don't feel that having the ability to write printGenericOp(op, /*printOpName=*/false) really warrants a renaming: how would you rename it?

[mehdi_amini]

Address Uday's comments

mlir/examples/toy/Ch2/mlir/Dialect.cpp

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verifier for the constant operation. This corresponds to the `::verify(...)`
 /// in the op definition.
 static mlir::LogicalResult verify(ConstantOp op) {
   // If the return type of the constant is not an unranked tensor, the shape

mlir/examples/toy/Ch3/mlir/Dialect.cpp

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verifier for the constant operation. This corresponds to the `::verify(...)`
 /// in the op definition.
 static mlir::LogicalResult verify(ConstantOp op) {
   // If the return type of the constant is not an unranked tensor, the shape

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

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verifier for the constant operation. This corresponds to the `::verify(...)`
 /// in the op definition.
 static mlir::LogicalResult verify(ConstantOp op) {
   // If the return type of the constant is not an unranked tensor, the shape

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

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verifier for the constant operation. This corresponds to the `::verify(...)`
 /// in the op definition.
 static mlir::LogicalResult verify(ConstantOp op) {
   // If the return type of the constant is not an unranked tensor, the shape

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

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verifier for the constant operation. This corresponds to the `::verify(...)`
 /// in the op definition.
 static mlir::LogicalResult verify(ConstantOp op) {
   // If the return type of the constant is not an unranked tensor, the shape

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

if (parser.resolveOperands(operands, type, result.operands))
     return mlir::failure();
   result.addTypes(type);
   return mlir::success();
 }
 
 /// A generalized printer for binary operations. It prints in two different
 /// forms depending on if all of the types match.
 static void printBinaryOp(mlir::OpAsmPrinter &printer, mlir::Operation *op) {
-  printer << op->getName() << " " << op->getOperands();
+  printer << " " << op->getOperands();
   printer.printOptionalAttrDict(op->getAttrs());
   printer << " : ";
 
   // If all of the types are the same, print the type directly.
   Type resultType = *op->result_type_begin();
   if (llvm::all_of(op->getOperandTypes(),
                    [=](Type type) { return type == resultType; })) {
     printer << resultType;
static mlir::ParseResult parseConstantOp(mlir::OpAsmParser &parser,
 
   result.addTypes(value.getType());
   return success();
 }
 
 /// The 'OpAsmPrinter' class is a stream that allows for formatting
 /// strings, attributes, operands, types, etc.
 static void print(mlir::OpAsmPrinter &printer, ConstantOp op) {
-  printer << "toy.constant ";
+  printer << " ";
   printer.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
   printer << op.value();
 }
 
 /// Verify that the given attribute value is valid for the given type.
 static mlir::LogicalResult verifyConstantForType(mlir::Type type,
                                                  mlir::Attribute opaqueValue,
                                                  mlir::Operation *op) {

mlir/include/mlir/Dialect/GPU/GPUOps.td

def GPU_ReturnOp : GPU_Op<"return", [HasParent<"GPUFuncOp">, NoSideEffect,
   let description = [{
     A terminator operation for regions that appear in the body of  `gpu.func`
     functions. The operands to the `gpu.return` are the result values returned
     by an invocation of the `gpu.func`.
   }];
 
   let builders = [OpBuilder<(ins), [{ // empty}]>];
 
-  let parser = [{ return parseReturnOp(parser, result); }];
-  let printer = [{ p << getOperationName(); }];
+  let assemblyFormat = "attr-dict ($operands^ `:` type($operands))?";
   let verifier = [{ return ::verify(*this); }];
 }
 
 def GPU_TerminatorOp : GPU_Op<"terminator", [HasParent<"LaunchOp">,
                                              NoSideEffect, Terminator]>,
     Arguments<(ins)>, Results<(outs)> {
   let summary = "Terminator for GPU launch regions.";
   let description = [{
     A terminator operation for regions that appear in the body of `gpu.launch`
     operation.  These regions are not expected to return any value so the
     terminator takes no operands.
   }];
 
-  let parser = [{ return success(); }];
-  let printer = [{ p << getOperationName(); }];
+  let assemblyFormat = "attr-dict";
 }
 
 def GPU_YieldOp : GPU_Op<"yield", [NoSideEffect, Terminator]>,
     Arguments<(ins Variadic<AnyType>:$values)> {
   let summary = "GPU yield operation";
   let description = [{
     gpu.yield` is a special terminator operation for blocks inside regions
     in gpu ops. It returns values to the immediately enclosing gpu op.
let description = [{
     and all memory accesses made by these work items prior to the op are
     visible to all work items in the workgroup. Data hazards between work items
     accessing the same memory can be avoided by synchronizing work items
     in-between these accesses.
 
     Either none or all work items of a workgroup need to execute this op
     in convergence.
   }];
-  let parser = [{ return success(); }];
-  let printer = [{ p << getOperationName(); }];
+  let assemblyFormat = "attr-dict";
 }
 
 def GPU_GPUModuleOp : GPU_Op<"module", [
   DataLayoutOpInterface, HasDefaultDLTIDataLayout, IsolatedFromAbove,
   SymbolTable, Symbol,
   SingleBlockImplicitTerminator<"ModuleEndOp">
 ]> {
   let summary = "A top level compilation unit containing code to be run on a GPU.";

 def GPU_ModuleEndOp : GPU_Op<"module_end", [
   Terminator, HasParent<"GPUModuleOp">
 ]> {
   let summary = "A pseudo op that marks the end of a gpu.module.";
   let description = [{
     This op terminates the only block inside the only region of a `gpu.module`.
   }];
 
-  let parser = [{ return success(); }];
-  let printer = [{ p << getOperationName(); }];
+  let assemblyFormat = "attr-dict";
 }
 
 def GPU_HostRegisterOp : GPU_Op<"host_register">,
     Arguments<(ins AnyUnrankedMemRef:$value)> {
   let summary = "Registers a memref for access from device.";
   let description = [{
     This op maps the provided host buffer into the device address space.
 

mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td

let verifier = [{
     // No check for personality of function - landingpad op verifies it.
     return success();
   }];
 
   let assemblyFormat = "$value attr-dict `:` type($value)";
 }
 def LLVM_UnreachableOp : LLVM_TerminatorOp<"unreachable", []> {
   string llvmBuilder = [{ builder.CreateUnreachable(); }];
-  let parser = [{ return success(); }];
-  let printer = [{ p << getOperationName(); }];
+  let assemblyFormat = "attr-dict";
 }
 
 def LLVM_SwitchOp : LLVM_TerminatorOp<"switch",
     [AttrSizedOperandSegments, DeclareOpInterfaceMethods<BranchOpInterface>,
      NoSideEffect]> {
   let arguments = (ins
     I32:$value,
     Variadic<AnyType>:$defaultOperands,

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

def ROCDL_MubufLoadOp :
   string llvmBuilder = [{
       $res = createIntrinsicCall(builder,
           llvm::Intrinsic::amdgcn_buffer_load, {$rsrc, $vindex, $offset, $glc,
           $slc}, {$_resultType});
   }];
   let parser = [{ return parseROCDLMubufLoadOp(parser, result); }];
   let printer = [{
     Operation *op = this->getOperation();
-    p << op->getName() << " " << op->getOperands()
+    p << " " << op->getOperands()
       << " : " << op->getResultTypes();
   }];
 }
 
 def ROCDL_MubufStoreOp :
   ROCDL_Op<"buffer.store">,
   Arguments<(ins LLVM_Type:$vdata,
                  LLVM_Type:$rsrc,
                  LLVM_Type:$vindex,
                  LLVM_Type:$offset,
                  LLVM_Type:$glc,
                  LLVM_Type:$slc)>{
   string llvmBuilder = [{
     auto vdataType = moduleTranslation.convertType(op.vdata().getType());
     createIntrinsicCall(builder,
           llvm::Intrinsic::amdgcn_buffer_store, {$vdata, $rsrc, $vindex,
           $offset, $glc, $slc}, {vdataType});
   }];
   let parser = [{ return parseROCDLMubufStoreOp(parser, result); }];
   let printer = [{
     Operation *op = this->getOperation();
-    p << op->getName() << " " << op->getOperands()
+    p << " " << op->getOperands()
       << " : " << vdata().getType();
   }];
 }
 
 #endif // ROCDLIR_OPS

mlir/include/mlir/Dialect/Utils/ReshapeOpsUtils.h

 /// Parse a reshape-like op, i.e. linalg::(Tensor)ExpandShapeOp,
 /// linalg::(Tensor)CollapseShapeOp.
 ParseResult parseReshapeLikeOp(OpAsmParser &parser, OperationState &result);
 
 /// Print a reshape-like op, i.e. linalg::(Tensor)ExpandShapeOp,
 /// linalg::(Tensor)CollapseShapeOp.
 template <typename ReshapeLikeOp>
 void printReshapeOp(OpAsmPrinter &p, ReshapeLikeOp op) {
-  p << op.getOperationName() << ' ' << op.src() << " [";
+  p << ' ' << op.src() << " [";
 
   llvm::interleaveComma(op.reassociation(), p, [&](const Attribute &attr) {
     p << '[';
     auto arrayAttr = attr.template cast<ArrayAttr>();
     llvm::interleaveComma(arrayAttr, p, [&](const Attribute &attr) {
       p << attr.cast<IntegerAttr>().getInt();
     });
     p << ']';
LogicalResult matchAndRewrite(ReshapeOpTy reshapeOp,
                                 PatternRewriter &rewriter) const override {
     auto srcReshapeOp = reshapeOp.src().template getDefiningOp<ReshapeOpTy>();
     if (!srcReshapeOp)
       return failure();
 
     ShapedType resultType = reshapeOp.getResultType();
     Optional<SmallVector<ReassociationIndices>> reassociationIndices =
         composeReassociationIndices(srcReshapeOp.getReassociationIndices(),
-                                     reshapeOp.getReassociationIndices(),
-                                     rewriter.getContext());
     if (!reassociationIndices)
       return failure();
     rewriter.replaceOpWithNewOp<ReshapeOpTy>(
         reshapeOp, resultType, srcReshapeOp.src(), *reassociationIndices);
     return success();
   }
 };
 

mlir/include/mlir/IR/OpDefinition.h

protected:
 
   /// Unless overridden, the custom assembly form of an op is always rejected.
   /// Op implementations should implement this to return failure.
   /// On success, they should fill in result with the fields to use.
   static ParseResult parse(OpAsmParser &parser, OperationState &result);
 
   // The fallback for the printer is to print it the generic assembly form.
   static void print(Operation *op, OpAsmPrinter &p);
+  static void printOpName(Operation *op, OpAsmPrinter &p);
 
   /// Mutability management is handled by the OpWrapper/OpConstWrapper classes,
   /// so we can cast it away here.
   explicit OpState(Operation *state) : state(state) {}
 
 private:
   Operation *state;
 
static AbstractOperation::PrintAssemblyFn getPrintAssemblyFn() {
     return getPrintAssemblyFnImpl<ConcreteType>();
   }
   /// The internal implementation of `getPrintAssemblyFn` that is invoked when
   /// the concrete operation does not define a `print` method.
   template <typename ConcreteOpT>
   static std::enable_if_t<!detect_has_print<ConcreteOpT>::value,
                           AbstractOperation::PrintAssemblyFn>
   getPrintAssemblyFnImpl() {
-    return [](Operation *op, OpAsmPrinter &parser) {
-      return OpState::print(op, parser);
+    return [](Operation *op, OpAsmPrinter &printer) {
+      return OpState::print(op, printer);
     };
   }
   /// The internal implementation of `getPrintAssemblyFn` that is invoked when
   /// the concrete operation defines a `print` method.
   template <typename ConcreteOpT>
   static std::enable_if_t<detect_has_print<ConcreteOpT>::value,
                           AbstractOperation::PrintAssemblyFn>
   getPrintAssemblyFnImpl() {
     return &printAssembly;
   }
   static void printAssembly(Operation *op, OpAsmPrinter &p) {
+    OpState::printOpName(op, p);
     return cast<ConcreteType>(op).print(p);
   }
   /// Implementation of `VerifyInvariantsFn` AbstractOperation hook.
   static AbstractOperation::VerifyInvariantsFn getVerifyInvariantsFn() {
     return &verifyInvariants;
   }
 
   static constexpr bool hasNoDataMembers() {

mlir/include/mlir/IR/OpImplementation.h

public:
 
   /// If the specified operation has attributes, print out an attribute
   /// dictionary prefixed with 'attributes'.
   virtual void
   printOptionalAttrDictWithKeyword(ArrayRef<NamedAttribute> attrs,
                                    ArrayRef<StringRef> elidedAttrs = {}) = 0;
 
   /// Print the entire operation with the default generic assembly form.
-  virtual void printGenericOp(Operation *op) = 0;
+  /// If `printOpName` is true, then the operation name is printed (the default)
+  /// otherwise it is ommited and the print will start with the operand list.
+  virtual void printGenericOp(Operation *op, bool printOpName = true) = 0;
 
   /// Prints a region.
   /// If 'printEntryBlockArgs' is false, the arguments of the
   /// block are not printed. If 'printBlockTerminator' is false, the terminator
   /// operation of the block is not printed. If printEmptyBlock is true, then
   /// the block header is printed even if the block is empty.
   virtual void printRegion(Region &blocks, bool printEntryBlockArgs = true,
                            bool printBlockTerminators = true,

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

return parser.emitError(parser.getNameLoc(),
                             "dimension or symbol index mismatch");
   }
 
   result.types.append(map.getNumResults(), indexTy);
   return success();
 }
 
 static void print(OpAsmPrinter &p, AffineApplyOp op) {
-  p << AffineApplyOp::getOperationName() << " " << op.mapAttr();
+  p << " " << op.mapAttr();
   printDimAndSymbolList(op.operand_begin(), op.operand_end(),
                         op.getAffineMap().getNumDims(), p);
   p.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"map"});
 }
 
 static LogicalResult verify(AffineApplyOp op) {
   // Check input and output dimensions match.
   auto map = op.map();
void AffineDmaStartOp::build(OpBuilder &builder, OperationState &result,
   result.addOperands(tagIndices);
   result.addOperands(numElements);
   if (stride) {
     result.addOperands({stride, elementsPerStride});
   }
 }
 
 void AffineDmaStartOp::print(OpAsmPrinter &p) {
-  p << "affine.dma_start " << getSrcMemRef() << '[';
+  p << " " << getSrcMemRef() << '[';
   p.printAffineMapOfSSAIds(getSrcMapAttr(), getSrcIndices());
   p << "], " << getDstMemRef() << '[';
   p.printAffineMapOfSSAIds(getDstMapAttr(), getDstIndices());
   p << "], " << getTagMemRef() << '[';
   p.printAffineMapOfSSAIds(getTagMapAttr(), getTagIndices());
   p << "], " << getNumElements();
   if (isStrided()) {
     p << ", " << getStride();
void AffineDmaWaitOp::build(OpBuilder &builder, OperationState &result,
                             ValueRange tagIndices, Value numElements) {
   result.addOperands(tagMemRef);
   result.addAttribute(getTagMapAttrName(), AffineMapAttr::get(tagMap));
   result.addOperands(tagIndices);
   result.addOperands(numElements);
 }
 
 void AffineDmaWaitOp::print(OpAsmPrinter &p) {
-  p << "affine.dma_wait " << getTagMemRef() << '[';
+  p << " " << getTagMemRef() << '[';
   SmallVector<Value, 2> operands(getTagIndices());
   p.printAffineMapOfSSAIds(getTagMapAttr(), operands);
   p << "], ";
   p.printOperand(getNumElements());
   p << " : " << getTagMemRef().getType();
 }
 
 // Parse AffineDmaWaitOp.

 unsigned AffineForOp::getNumIterOperands() {
   AffineMap lbMap = getLowerBoundMapAttr().getValue();
   AffineMap ubMap = getUpperBoundMapAttr().getValue();
 
   return getNumOperands() - lbMap.getNumInputs() - ubMap.getNumInputs();
 }
 
 static void print(OpAsmPrinter &p, AffineForOp op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   p.printOperand(op.getBody()->getArgument(0));
   p << " = ";
   printBound(op.getLowerBoundMapAttr(), op.getLowerBoundOperands(), "max", p);
   p << " to ";
   printBound(op.getUpperBoundMapAttr(), op.getUpperBoundOperands(), "min", p);
 
   if (op.getStep() != 1)
     p << " step " << op.getStep();
if (parser.parseOptionalAttrDict(result.attributes))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, AffineIfOp op) {
   auto conditionAttr =
       op->getAttrOfType<IntegerSetAttr>(op.getConditionAttrName());
-  p << "affine.if " << conditionAttr;
+  p << " " << conditionAttr;
   printDimAndSymbolList(op.operand_begin(), op.operand_end(),
                         conditionAttr.getValue().getNumDims(), p);
   p.printOptionalArrowTypeList(op.getResultTypes());
   p.printRegion(op.thenRegion(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/op.getNumResults());
 
   // Print the 'else' regions if it has any blocks.
return failure(
       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));
 }
 
 static void print(OpAsmPrinter &p, AffineLoadOp op) {
-  p << "affine.load " << op.getMemRef() << '[';
+  p << " " << op.getMemRef() << '[';
   if (AffineMapAttr mapAttr =
           op->getAttrOfType<AffineMapAttr>(op.getMapAttrName()))
     p.printAffineMapOfSSAIds(mapAttr, op.getMapOperands());
   p << ']';
   p.printOptionalAttrDict(op->getAttrs(),
                           /*elidedAttrs=*/{op.getMapAttrName()});
   p << " : " << op.getMemRefType();
 }
return failure(parser.parseOperand(storeValueInfo) || parser.parseComma() ||
                  parser.parseColonType(type) ||
                  parser.resolveOperand(storeValueInfo, type.getElementType(),
                                        result.operands) ||
                  parser.resolveOperand(memrefInfo, type, result.operands) ||
                  parser.resolveOperands(mapOperands, indexTy, result.operands));
 }
 
 static void print(OpAsmPrinter &p, AffineStoreOp op) {
-  p << "affine.store " << op.getValueToStore();
+  p << " " << op.getValueToStore();
   p << ", " << op.getMemRef() << '[';
   if (AffineMapAttr mapAttr =
           op->getAttrOfType<AffineMapAttr>(op.getMapAttrName()))
     p.printAffineMapOfSSAIds(mapAttr, op.getMapOperands());
   p << ']';
   p.printOptionalAttrDict(op->getAttrs(),
                           /*elidedAttrs=*/{op.getMapAttrName()});
   p << " : " << op.getMemRefType();
static LogicalResult verifyAffineMinMaxOp(T op) {
   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.getOperationName() << ' ' << op->getAttr(T::getMapAttrName());
+  p << ' ' << op->getAttr(T::getMapAttrName());
   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()});
static ParseResult parseAffinePrefetchOp(OpAsmParser &parser,
   result.addAttribute(
       AffinePrefetchOp::getIsDataCacheAttrName(),
       parser.getBuilder().getBoolAttr(cacheType.equals("data")));
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, AffinePrefetchOp op) {
-  p << AffinePrefetchOp::getOperationName() << " " << op.memref() << '[';
+  p << " " << op.memref() << '[';
   AffineMapAttr mapAttr = op->getAttrOfType<AffineMapAttr>(op.getMapAttrName());
   if (mapAttr) {
     SmallVector<Value, 2> operands(op.getMapOperands());
     p.printAffineMapOfSSAIds(mapAttr, operands);
   }
   p << ']' << ", " << (op.isWrite() ? "write" : "read") << ", "
     << "locality<" << op.localityHint() << ">, "
     << (op.isDataCache() ? "data" : "instr");
if (size == 1) {
       p.printAffineMapOfSSAIds(AffineMapAttr::get(submap), operands);
       p << ')';
       start += size;
     }
   }
 }
 
 static void print(OpAsmPrinter &p, AffineParallelOp op) {
-  p << op.getOperationName() << " (" << op.getBody()->getArguments() << ") = (";
+  p << " (" << op.getBody()->getArguments() << ") = (";
   printMinMaxBound(p, op.lowerBoundsMapAttr(), op.lowerBoundsGroupsAttr(),
                    op.getLowerBoundsOperands(), "max");
   p << ") to (";
   printMinMaxBound(p, op.upperBoundsMapAttr(), op.upperBoundsGroupsAttr(),
                    op.getUpperBoundsOperands(), "min");
   p << ')';
   SmallVector<int64_t, 8> steps = op.getSteps();
   bool elideSteps = llvm::all_of(steps, [](int64_t step) { return step == 1; });
return failure(
       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));
 }
 
 static void print(OpAsmPrinter &p, AffineVectorLoadOp op) {
-  p << "affine.vector_load " << op.getMemRef() << '[';
+  p << " " << op.getMemRef() << '[';
   if (AffineMapAttr mapAttr =
           op->getAttrOfType<AffineMapAttr>(op.getMapAttrName()))
     p.printAffineMapOfSSAIds(mapAttr, op.getMapOperands());
   p << ']';
   p.printOptionalAttrDict(op->getAttrs(),
                           /*elidedAttrs=*/{op.getMapAttrName()});
   p << " : " << op.getMemRefType() << ", " << op.getType();
 }
return failure(
       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));
 }
 
 static void print(OpAsmPrinter &p, AffineVectorStoreOp op) {
-  p << "affine.vector_store " << op.getValueToStore();
+  p << " " << op.getValueToStore();
   p << ", " << op.getMemRef() << '[';
   if (AffineMapAttr mapAttr =
           op->getAttrOfType<AffineMapAttr>(op.getMapAttrName()))
     p.printAffineMapOfSSAIds(mapAttr, op.getMapOperands());
   p << ']';
   p.printOptionalAttrDict(op->getAttrs(),
                           /*elidedAttrs=*/{op.getMapAttrName()});
   p << " : " << op.getMemRefType() << ", " << op.getValueToStore().getType();

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

 
 //===----------------------------------------------------------------------===//
 // IncludeOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, IncludeOp &op) {
   bool standardInclude = op.is_standard_include();
 
-  p << IncludeOp::getOperationName() << " ";
+  p << " ";
   if (standardInclude)
     p << "<";
   p << "\"" << op.include() << "\"";
   if (standardInclude)
     p << ">";
 }
 
 static ParseResult parseIncludeOp(OpAsmParser &parser, OperationState &result) {

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

static LogicalResult verifyShuffleOp(gpu::ShuffleOp shuffleOp) {
   if (!type.isSignlessIntOrFloat() || type.getIntOrFloatBitWidth() != 32) {
     return shuffleOp.emitOpError()
            << "requires value operand type to be f32 or i32";
   }
   return success();
 }
 
 static void printShuffleOp(OpAsmPrinter &p, ShuffleOp op) {
-  p << ShuffleOp::getOperationName() << ' ' << op.getOperands() << ' '
-    << op.mode() << " : " << op.value().getType();
+  p << ' ' << op.getOperands() << ' ' << op.mode() << " : "
+    << op.value().getType();
 }
 
 static ParseResult parseShuffleOp(OpAsmParser &parser, OperationState &state) {
   SmallVector<OpAsmParser::OperandType, 3> operandInfo;
   if (parser.parseOperandList(operandInfo, 3))
     return failure();
 
   StringRef mode;
static void printSizeAssignment(OpAsmPrinter &p, KernelDim3 size,
   p << '(' << ids.x << ", " << ids.y << ", " << ids.z << ") in (";
   p << size.x << " = " << operands.x << ", ";
   p << size.y << " = " << operands.y << ", ";
   p << size.z << " = " << operands.z << ')';
 }
 
 static void printLaunchOp(OpAsmPrinter &p, LaunchOp op) {
   // Print the launch configuration.
-  p << LaunchOp::getOperationName() << ' ' << op.getBlocksKeyword();
+  p << ' ' << op.getBlocksKeyword();
   printSizeAssignment(p, op.getGridSize(), op.getGridSizeOperandValues(),
                       op.getBlockIds());
   p << ' ' << op.getThreadsKeyword();
   printSizeAssignment(p, op.getBlockSize(), op.getBlockSizeOperandValues(),
                       op.getThreadIds());
 
   p.printRegion(op.body(), /*printEntryBlockArgs=*/false);
   p.printOptionalAttrDict(op->getAttrs());
static void printAttributions(OpAsmPrinter &p, StringRef keyword,
   p << ' ' << keyword << '(';
   llvm::interleaveComma(
       values, p, [&p](BlockArgument v) { p << v << " : " << v.getType(); });
   p << ')';
 }
 
 /// Prints a GPU Func op.
 static void printGPUFuncOp(OpAsmPrinter &p, GPUFuncOp op) {
-  p << GPUFuncOp::getOperationName() << ' ';
+  p << ' ';
   p.printSymbolName(op.getName());
 
   FunctionType type = op.getType();
   function_like_impl::printFunctionSignature(
       p, op.getOperation(), type.getInputs(),
       /*isVariadic=*/false, type.getResults());
 
   printAttributions(p, op.getWorkgroupKeyword(), op.getWorkgroupAttributions());
LogicalResult GPUFuncOp::verifyBody() {
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ReturnOp
 //===----------------------------------------------------------------------===//
 
-static ParseResult parseReturnOp(OpAsmParser &parser, OperationState &result) {
-  llvm::SmallVector<OpAsmParser::OperandType, 4> operands;
-  llvm::SmallVector<Type, 4> types;
-  if (parser.parseOperandList(operands) ||
-      parser.parseOptionalColonTypeList(types) ||
-      parser.resolveOperands(operands, types, parser.getCurrentLocation(),
-                             result.operands))
-    return failure();
-
-  return success();
-}
-
 static LogicalResult verify(gpu::ReturnOp returnOp) {
   GPUFuncOp function = returnOp->getParentOfType<GPUFuncOp>();
 
   FunctionType funType = function.getType();
 
   if (funType.getNumResults() != returnOp.operands().size())
     return returnOp.emitOpError()
         .append("expected ", funType.getNumResults(), " result operands")
if (parser.parseRegion(*body, None, None))
     return failure();
 
   // Ensure that this module has a valid terminator.
   GPUModuleOp::ensureTerminator(*body, parser.getBuilder(), result.location);
   return success();
 }
 
 static void print(OpAsmPrinter &p, GPUModuleOp op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   p.printSymbolName(op.getName());
   p.printOptionalAttrDictWithKeyword(op->getAttrs(),
                                      {SymbolTable::getSymbolAttrName()});
   p.printRegion(op->getRegion(0), /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/false);
 }
 
 //===----------------------------------------------------------------------===//

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

static void printLLVMOpAttrs(OpAsmPrinter &printer, Operation *op,
                              DictionaryAttr attrs) {
   printer.printOptionalAttrDict(processFMFAttr(attrs.getValue()));
 }
 
 //===----------------------------------------------------------------------===//
 // Printing/parsing for LLVM::CmpOp.
 //===----------------------------------------------------------------------===//
 static void printICmpOp(OpAsmPrinter &p, ICmpOp &op) {
-  p << op.getOperationName() << " \"" << stringifyICmpPredicate(op.predicate())
-    << "\" " << op.getOperand(0) << ", " << op.getOperand(1);
+  p << " \"" << stringifyICmpPredicate(op.predicate()) << "\" "
+    << op.getOperand(0) << ", " << op.getOperand(1);
   p.printOptionalAttrDict(op->getAttrs(), {"predicate"});
   p << " : " << op.lhs().getType();
 }
 
 static void printFCmpOp(OpAsmPrinter &p, FCmpOp &op) {
-  p << op.getOperationName() << " \"" << stringifyFCmpPredicate(op.predicate())
-    << "\" " << op.getOperand(0) << ", " << op.getOperand(1);
+  p << " \"" << stringifyFCmpPredicate(op.predicate()) << "\" "
+    << op.getOperand(0) << ", " << op.getOperand(1);
   p.printOptionalAttrDict(processFMFAttr(op->getAttrs()), {"predicate"});
   p << " : " << op.lhs().getType();
 }
 
 // <operation> ::= `llvm.icmp` string-literal ssa-use `,` ssa-use
 //                 attribute-dict? `:` type
 // <operation> ::= `llvm.fcmp` string-literal ssa-use `,` ssa-use
 //                 attribute-dict? `:` type

 //===----------------------------------------------------------------------===//
 
 static void printAllocaOp(OpAsmPrinter &p, AllocaOp &op) {
   auto elemTy = op.getType().cast<LLVM::LLVMPointerType>().getElementType();
 
   auto funcTy = FunctionType::get(op.getContext(), {op.arraySize().getType()},
                                   {op.getType()});
 
-  p << op.getOperationName() << ' ' << op.arraySize() << " x " << elemTy;
+  p << ' ' << op.arraySize() << " x " << elemTy;
   if (op.alignment().hasValue() && *op.alignment() != 0)
     p.printOptionalAttrDict(op->getAttrs());
   else
     p.printOptionalAttrDict(op->getAttrs(), {"alignment"});
   p << " : " << funcTy;
 }
 
 // <operation> ::= `llvm.alloca` ssa-use `x` type attribute-dict?
if (isVolatile)
     result.addAttribute(kVolatileAttrName, builder.getUnitAttr());
   if (isNonTemporal)
     result.addAttribute(kNonTemporalAttrName, builder.getUnitAttr());
   if (alignment != 0)
     result.addAttribute("alignment", builder.getI64IntegerAttr(alignment));
 }
 
 static void printLoadOp(OpAsmPrinter &p, LoadOp &op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   if (op.volatile_())
     p << "volatile ";
   p << op.addr();
   p.printOptionalAttrDict(op->getAttrs(), {kVolatileAttrName});
   p << " : " << op.addr().getType();
 }
 
 // Extract the pointee type from the LLVM pointer type wrapped in MLIR.  Return
if (isVolatile)
     result.addAttribute(kVolatileAttrName, builder.getUnitAttr());
   if (isNonTemporal)
     result.addAttribute(kNonTemporalAttrName, builder.getUnitAttr());
   if (alignment != 0)
     result.addAttribute("alignment", builder.getI64IntegerAttr(alignment));
 }
 
 static void printStoreOp(OpAsmPrinter &p, StoreOp &op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   if (op.volatile_())
     p << "volatile ";
   p << op.value() << ", " << op.addr();
   p.printOptionalAttrDict(op->getAttrs(), {kVolatileAttrName});
   p << " : " << op.addr().getType();
 }
 
 // <operation> ::= `llvm.store` `volatile` ssa-use `,` ssa-use
static LogicalResult verify(InvokeOp op) {
 
   return success();
 }
 
 static void printInvokeOp(OpAsmPrinter &p, InvokeOp op) {
   auto callee = op.callee();
   bool isDirect = callee.hasValue();
 
-  p << op.getOperationName() << ' ';
+  p << ' ';
 
   // Either function name or pointer
   if (isDirect)
     p.printSymbolName(callee.getValue());
   else
     p << op.getOperand(0);
 
   p << '(' << op.getOperands().drop_front(isDirect ? 0 : 1) << ')';
if (isFilter) {
       return op.emitError("clause #")
              << idx << " is not a known constant - null, addressof, bitcast";
     }
   }
   return success();
 }
 
 static void printLandingpadOp(OpAsmPrinter &p, LandingpadOp &op) {
-  p << op.getOperationName() << (op.cleanup() ? " cleanup " : " ");
+  p << (op.cleanup() ? " cleanup " : " ");
 
   // Clauses
   for (auto value : op.getOperands()) {
     // Similar to llvm - if clause is an array type then it is filter
     // clause else catch clause
     bool isArrayTy = value.getType().isa<LLVMArrayType>();
     p << '(' << (isArrayTy ? "filter " : "catch ") << value << " : "
       << value.getType() << ") ";

 }
 
 static void printCallOp(OpAsmPrinter &p, CallOp &op) {
   auto callee = op.callee();
   bool isDirect = callee.hasValue();
 
   // Print the direct callee if present as a function attribute, or an indirect
   // callee (first operand) otherwise.
-  p << op.getOperationName() << ' ';
+  p << ' ';
   if (isDirect)
     p.printSymbolName(callee.getValue());
   else
     p << op.getOperand(0);
 
   auto args = op.getOperands().drop_front(isDirect ? 0 : 1);
   p << '(' << args << ')';
   p.printOptionalAttrDict(processFMFAttr(op->getAttrs()), {"callee"});
void LLVM::ExtractElementOp::build(OpBuilder &b, OperationState &result,
                                    ArrayRef<NamedAttribute> attrs) {
   auto vectorType = vector.getType();
   auto llvmType = LLVM::getVectorElementType(vectorType);
   build(b, result, llvmType, vector, position);
   result.addAttributes(attrs);
 }
 
 static void printExtractElementOp(OpAsmPrinter &p, ExtractElementOp &op) {
-  p << op.getOperationName() << ' ' << op.vector() << "[" << op.position()
-    << " : " << op.position().getType() << "]";
+  p << ' ' << op.vector() << "[" << op.position() << " : "
+    << op.position().getType() << "]";
   p.printOptionalAttrDict(op->getAttrs());
   p << " : " << op.vector().getType();
 }
 
 // <operation> ::= `llvm.extractelement` ssa-use `, ` ssa-use
 //                 attribute-dict? `:` type
 static ParseResult parseExtractElementOp(OpAsmParser &parser,
                                          OperationState &result) {
static LogicalResult verify(ExtractElementOp op) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Printing/parsing for LLVM::ExtractValueOp.
 //===----------------------------------------------------------------------===//
 
 static void printExtractValueOp(OpAsmPrinter &p, ExtractValueOp &op) {
-  p << op.getOperationName() << ' ' << op.container() << op.position();
+  p << ' ' << op.container() << op.position();
   p.printOptionalAttrDict(op->getAttrs(), {"position"});
   p << " : " << op.container().getType();
 }
 
 // Extract the type at `position` in the wrapped LLVM IR aggregate type
 // `containerType`.  Position is an integer array attribute where each value
 // is a zero-based position of the element in the aggregate type.  Return the
 // resulting type wrapped in MLIR, or nullptr on error.
static LogicalResult verify(ExtractValueOp op) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Printing/parsing for LLVM::InsertElementOp.
 //===----------------------------------------------------------------------===//
 
 static void printInsertElementOp(OpAsmPrinter &p, InsertElementOp &op) {
-  p << op.getOperationName() << ' ' << op.value() << ", " << op.vector() << "["
-    << op.position() << " : " << op.position().getType() << "]";
+  p << ' ' << op.value() << ", " << op.vector() << "[" << op.position() << " : "
+    << op.position().getType() << "]";
   p.printOptionalAttrDict(op->getAttrs());
   p << " : " << op.vector().getType();
 }
 
 // <operation> ::= `llvm.insertelement` ssa-use `,` ssa-use `,` ssa-use
 //                 attribute-dict? `:` type
 static ParseResult parseInsertElementOp(OpAsmParser &parser,
                                         OperationState &result) {
return op.emitOpError()
            << " into " << op.vector().getType();
   return success();
 }
 //===----------------------------------------------------------------------===//
 // Printing/parsing for LLVM::InsertValueOp.
 //===----------------------------------------------------------------------===//
 
 static void printInsertValueOp(OpAsmPrinter &p, InsertValueOp &op) {
-  p << op.getOperationName() << ' ' << op.value() << ", " << op.container()
-    << op.position();
+  p << ' ' << op.value() << ", " << op.container() << op.position();
   p.printOptionalAttrDict(op->getAttrs(), {"position"});
   p << " : " << op.container().getType();
 }
 
 // <operation> ::= `llvm.insertvaluevalue` ssa-use `,` ssa-use
 //                 `[` integer-literal (`,` integer-literal)* `]`
 //                 attribute-dict? `:` type
 static ParseResult parseInsertValueOp(OpAsmParser &parser,
static LogicalResult verify(InsertValueOp op) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Printing, parsing and verification for LLVM::ReturnOp.
 //===----------------------------------------------------------------------===//
 
 static void printReturnOp(OpAsmPrinter &p, ReturnOp op) {
-  p << op.getOperationName();
   p.printOptionalAttrDict(op->getAttrs());
   assert(op.getNumOperands() <= 1);
 
   if (op.getNumOperands() == 0)
     return;
 
   p << ' ' << op.getOperand(0) << " : " << op.getOperand(0).getType();
 }
result.addAttribute(getLinkageAttrName(),
                       builder.getI64IntegerAttr(static_cast<int64_t>(linkage)));
   if (addrSpace != 0)
     result.addAttribute("addr_space", builder.getI32IntegerAttr(addrSpace));
   result.attributes.append(attrs.begin(), attrs.end());
   result.addRegion();
 }
 
 static void printGlobalOp(OpAsmPrinter &p, GlobalOp op) {
-  p << op.getOperationName() << ' ' << stringifyLinkage(op.linkage()) << ' ';
+  p << ' ' << stringifyLinkage(op.linkage()) << ' ';
   if (op.unnamed_addr())
     p << stringifyUnnamedAddr(*op.unnamed_addr()) << ' ';
   if (op.constant())
     p << "constant ";
   p.printSymbolName(op.sym_name());
   p << '(';
   if (auto value = op.getValueOrNull())
     p.printAttribute(value);
void LLVM::ShuffleVectorOp::build(OpBuilder &b, OperationState &result,
   auto containerType = v1.getType();
   auto vType = LLVM::getFixedVectorType(
       LLVM::getVectorElementType(containerType), mask.size());
   build(b, result, vType, v1, v2, mask);
   result.addAttributes(attrs);
 }
 
 static void printShuffleVectorOp(OpAsmPrinter &p, ShuffleVectorOp &op) {
-  p << op.getOperationName() << ' ' << op.v1() << ", " << op.v2() << " "
-    << op.mask();
+  p << ' ' << op.v1() << ", " << op.v2() << " " << op.mask();
   p.printOptionalAttrDict(op->getAttrs(), {"mask"});
   p << " : " << op.v1().getType() << ", " << op.v2().getType();
 }
 
 // <operation> ::= `llvm.shufflevector` ssa-use `, ` ssa-use
 //                 `[` integer-literal (`,` integer-literal)* `]`
 //                 attribute-dict? `:` type
 static ParseResult parseShuffleVectorOp(OpAsmParser &parser,
OptionalParseResult parseResult = parser.parseOptionalRegion(
       *body, entryArgs, entryArgs.empty() ? ArrayRef<Type>() : argTypes);
   return failure(parseResult.hasValue() && failed(*parseResult));
 }
 
 // Print the LLVMFuncOp. Collects argument and result types and passes them to
 // helper functions. Drops "void" result since it cannot be parsed back. Skips
 // the external linkage since it is the default value.
 static void printLLVMFuncOp(OpAsmPrinter &p, LLVMFuncOp op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   if (op.linkage() != LLVM::Linkage::External)
     p << stringifyLinkage(op.linkage()) << ' ';
   p.printSymbolName(op.getName());
 
   LLVMFunctionType fnType = op.getType();
   SmallVector<Type, 8> argTypes;
   SmallVector<Type, 1> resTypes;
   argTypes.reserve(fnType.getNumParams());
static ParseResult parseAtomicOrdering(OpAsmParser &parser,
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Printer, parser and verifier for LLVM::AtomicRMWOp.
 //===----------------------------------------------------------------------===//
 
 static void printAtomicRMWOp(OpAsmPrinter &p, AtomicRMWOp &op) {
-  p << op.getOperationName() << ' ' << stringifyAtomicBinOp(op.bin_op()) << ' '
-    << op.ptr() << ", " << op.val() << ' '
-    << stringifyAtomicOrdering(op.ordering()) << ' ';
+  p << ' ' << stringifyAtomicBinOp(op.bin_op()) << ' ' << op.ptr() << ", "
+    << op.val() << ' ' << stringifyAtomicOrdering(op.ordering()) << ' ';
   p.printOptionalAttrDict(op->getAttrs(), {"bin_op", "ordering"});
   p << " : " << op.res().getType();
 }
 
 // <operation> ::= `llvm.atomicrmw` keyword ssa-use `,` ssa-use keyword
 //                 attribute-dict? `:` type
 static ParseResult parseAtomicRMWOp(OpAsmParser &parser,
                                     OperationState &result) {
static LogicalResult verify(AtomicRMWOp op) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Printer, parser and verifier for LLVM::AtomicCmpXchgOp.
 //===----------------------------------------------------------------------===//
 
 static void printAtomicCmpXchgOp(OpAsmPrinter &p, AtomicCmpXchgOp &op) {
-  p << op.getOperationName() << ' ' << op.ptr() << ", " << op.cmp() << ", "
-    << op.val() << ' ' << stringifyAtomicOrdering(op.success_ordering()) << ' '
+  p << ' ' << op.ptr() << ", " << op.cmp() << ", " << op.val() << ' '
+    << stringifyAtomicOrdering(op.success_ordering()) << ' '
     << stringifyAtomicOrdering(op.failure_ordering());
   p.printOptionalAttrDict(op->getAttrs(),
                           {"success_ordering", "failure_ordering"});
   p << " : " << op.val().getType();
 }
 
 // <operation> ::= `llvm.cmpxchg` ssa-use `,` ssa-use `,` ssa-use
 //                 keyword keyword attribute-dict? `:` type
static ParseResult parseFenceOp(OpAsmParser &parser, OperationState &result) {
   if (parseAtomicOrdering(parser, result, "ordering") ||
       parser.parseOptionalAttrDict(result.attributes))
     return failure();
   return success();
 }
 
 static void printFenceOp(OpAsmPrinter &p, FenceOp &op) {
   StringRef syncscopeKeyword = "syncscope";
-  p << op.getOperationName() << ' ';
+  p << ' ';
   if (!op->getAttr(syncscopeKeyword).cast<StringAttr>().getValue().empty())
     p << "syncscope(" << op->getAttr(syncscopeKeyword) << ") ";
   p << stringifyAtomicOrdering(op.ordering());
 }
 
 static LogicalResult verify(FenceOp &op) {
   if (op.ordering() == AtomicOrdering::not_atomic ||
       op.ordering() == AtomicOrdering::unordered ||

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

 
 #include "mlir/Dialect/LLVMIR/NVVMOpsDialect.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // Printing/parsing for NVVM ops
 //===----------------------------------------------------------------------===//
 
 static void printNVVMIntrinsicOp(OpAsmPrinter &p, Operation *op) {
-  p << op->getName() << " " << op->getOperands();
+  p << " " << op->getOperands();
   if (op->getNumResults() > 0)
     p << " : " << op->getResultTypes();
 }
 
 // <operation> ::=
 //     `llvm.nvvm.shfl.sync.bfly %dst, %val, %offset, %clamp_and_mask`
 //      ({return_value_and_is_valid})? : result_type
 static ParseResult parseNVVMShflSyncBflyOp(OpAsmParser &parser,
static LogicalResult parseWMMAMmaF16F16M16N16K16Op(OpAsmParser &parser,
   if (parser.resolveOperands(operands, operandTypes, operandsLoc,
                              result.operands))
     return failure();
   return success();
 }
 
 static void printWMMAMmaF16F16M16N16K16Op(OpAsmPrinter &p,
                                           WMMAMmaF16F16M16N16K16Op &op) {
-  p << op.getOperationName();
   p << ' ';
   p << op.args();
   p.printOptionalAttrDict(op->getAttrs(), {});
   p << " : ";
   p << op->getOperand(0).getType();
   p << ' ' << "->";
   p << ' ';
   p << ::llvm::ArrayRef<::mlir::Type>(op.res().getType());

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

void GenericOp::build(
     ArrayRef<NamedAttribute> attributes) {
   build(builder, result, resultTensorTypes, inputs, outputs, indexingMaps,
         iteratorTypes,
         /*doc=*/"",
         /*libraryCall=*/"", bodyBuild, attributes);
 }
 
 static void print(OpAsmPrinter &p, GenericOp op) {
-  p << op.getOperationName() << " ";
+  p << " ";
 
   // Print extra attributes.
   auto genericAttrNames = op.linalgTraitAttrNames();
 
   llvm::StringSet<> genericAttrNamesSet;
   genericAttrNamesSet.insert(genericAttrNames.begin(), genericAttrNames.end());
   SmallVector<NamedAttribute, 8> genericAttrs;
   for (auto attr : op->getAttrs())
LogicalResult TensorCollapseShapeOp::reifyResultShapes(
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // YieldOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, linalg::YieldOp op) {
-  p << op.getOperationName();
   if (op.getNumOperands() > 0)
     p << ' ' << op.getOperands();
   p.printOptionalAttrDict(op->getAttrs());
   if (op.getNumOperands() > 0)
     p << " : " << op.getOperandTypes();
 }
 
 static ParseResult parseYieldOp(OpAsmParser &parser, OperationState &result) {
bodyBuilderFn(builder, result.location,
                   bodyBlock->getArguments().take_front(numIVs),
                   bodyBlock->getArguments().slice(numIVs, inputs.size()),
                   bodyBlock->getArguments().take_back(outputs.size()));
     TiledLoopOp::ensureTerminator(*bodyRegion, builder, result.location);
   }
 }
 
 static void print(OpAsmPrinter &p, TiledLoopOp op) {
-  p << op.getOperationName() << " (" << op.getInductionVars() << ") = ("
-    << op.lowerBound() << ") to (" << op.upperBound() << ") step (" << op.step()
-    << ")";
+  p << " (" << op.getInductionVars() << ") = (" << op.lowerBound() << ") to ("
+    << op.upperBound() << ") step (" << op.step() << ")";
 
   if (!op.inputs().empty()) {
     p << " ins (";
     llvm::interleaveComma(llvm::zip(op.getRegionInputArgs(), op.inputs()), p,
                           [&](auto it) {
                             p << std::get<0>(it) << " = " << std::get<1>(it)
                               << ": " << std::get<1>(it).getType();
                           });

 struct DimOfTiledLoopInsOutsFolder : public OpRewritePattern<OpTy> {
   using OpRewritePattern<OpTy>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(OpTy dimOp,
                                 PatternRewriter &rewriter) const final {
     auto src = dimOp.source().template dyn_cast<BlockArgument>();
     if (!src)
       return failure();
-    auto loopOp = dyn_cast<TiledLoopOp>(
-        src.getOwner()->getParent()->getParentOp());
+    auto loopOp =
+        dyn_cast<TiledLoopOp>(src.getOwner()->getParent()->getParentOp());
     if (!loopOp)
       return failure();
 
     auto inputArgs = loopOp.getRegionInputArgs();
     auto it1 = llvm::find(inputArgs, src);
     if (it1 != inputArgs.end()) {
       rewriter.updateRootInPlace(dimOp, [&] {
         dimOp.sourceMutable().assign(loopOp.inputs()[it1 - inputArgs.begin()]);
static void printNamedStructuredOpResults(OpAsmPrinter &p,
                                           TypeRange resultTypes) {
   if (resultTypes.empty())
     return;
   p.printOptionalArrowTypeList(resultTypes);
 }
 
 template <typename NamedStructuredOpType>
 static void printNamedStructuredOp(OpAsmPrinter &p, NamedStructuredOpType op) {
-  p << op.getOperationName();
   p.printOptionalAttrDict(
       op->getAttrs(),
       /*elidedAttrs=*/{"operand_segment_sizes",
                        // See generated code in mlir-linalg-yaml-gen.cpp
                        "linalg.memoized_indexing_maps"});
 
   // Printing is shared with generic ops, except for the region and
   // attributes.

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

 
 //===----------------------------------------------------------------------===//
 // AllocaScopeOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, AllocaScopeOp &op) {
   bool printBlockTerminators = false;
 
-  p << AllocaScopeOp::getOperationName() << " ";
+  p << " ";
   if (!op.results().empty()) {
     p << " -> (" << op.getResultTypes() << ")";
     printBlockTerminators = true;
   }
   p.printRegion(op.bodyRegion(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/printBlockTerminators);
   p.printOptionalAttrDict(op->getAttrs());
void DmaStartOp::build(OpBuilder &builder, OperationState &result,
   result.addOperands(destIndices);
   result.addOperands({numElements, tagMemRef});
   result.addOperands(tagIndices);
   if (stride)
     result.addOperands({stride, elementsPerStride});
 }
 
 void DmaStartOp::print(OpAsmPrinter &p) {
-  p << getOperationName() << " " << getSrcMemRef() << '[' << getSrcIndices()
-    << "], " << getDstMemRef() << '[' << getDstIndices() << "], "
-    << getNumElements() << ", " << getTagMemRef() << '[' << getTagIndices()
-    << ']';
+  p << " " << getSrcMemRef() << '[' << getSrcIndices() << "], "
+    << getDstMemRef() << '[' << getDstIndices() << "], " << getNumElements()
+    << ", " << getTagMemRef() << '[' << getTagIndices() << ']';
   if (isStrided())
     p << ", " << getStride() << ", " << getNumElementsPerStride();
 
   p.printOptionalAttrDict((*this)->getAttrs());
   p << " : " << getSrcMemRef().getType() << ", " << getDstMemRef().getType()
     << ", " << getTagMemRef().getType();
 }
 
void DmaWaitOp::build(OpBuilder &builder, OperationState &result,
                       Value tagMemRef, ValueRange tagIndices,
                       Value numElements) {
   result.addOperands(tagMemRef);
   result.addOperands(tagIndices);
   result.addOperands(numElements);
 }
 
 void DmaWaitOp::print(OpAsmPrinter &p) {
-  p << getOperationName() << " " << getTagMemRef() << '[' << getTagIndices()
-    << "], " << getNumElements();
+  p << " " << getTagMemRef() << '[' << getTagIndices() << "], "
+    << getNumElements();
   p.printOptionalAttrDict((*this)->getAttrs());
   p << " : " << getTagMemRef().getType();
 }
 
 // Parse DmaWaitOp.
 // Eg:
 //   dma_wait %tag[%index], %num_elements : memref<1 x i32, (d0) -> (d0), 4>
 //
void LoadOp::getCanonicalizationPatterns(RewritePatternSet &results,
   results.add<LoadOfBufferCast>(context);
 }
 
 //===----------------------------------------------------------------------===//
 // PrefetchOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, PrefetchOp op) {
-  p << PrefetchOp::getOperationName() << " " << op.memref() << '[';
+  p << " " << op.memref() << '[';
   p.printOperands(op.indices());
   p << ']' << ", " << (op.isWrite() ? "write" : "read");
   p << ", locality<" << op.localityHint();
   p << ">, " << (op.isDataCache() ? "data" : "instr");
   p.printOptionalAttrDict(
       op->getAttrs(),
       /*elidedAttrs=*/{"localityHint", "isWrite", "isDataCache"});
   p << " : " << op.getMemRefType();
void TransposeOp::build(OpBuilder &b, OperationState &result, Value in,
   MemRefType resultType = inferTransposeResultType(memRefType, permutationMap);
 
   build(b, result, resultType, in, attrs);
   result.addAttribute(TransposeOp::getPermutationAttrName(), permutation);
 }
 
 // transpose $in $permutation attr-dict : type($in) `to` type(results)
 static void print(OpAsmPrinter &p, TransposeOp op) {
-  p << "memref.transpose " << op.in() << " " << op.permutation();
+  p << " " << op.in() << " " << op.permutation();
   p.printOptionalAttrDict(op->getAttrs(),
                           {TransposeOp::getPermutationAttrName()});
   p << " : " << op.in().getType() << " to " << op.getType();
 }
 
 static ParseResult parseTransposeOp(OpAsmParser &parser,
                                     OperationState &result) {
   OpAsmParser::OperandType in;
return failure(
       parser.resolveOperand(srcInfo, srcType, result.operands) ||
       parser.resolveOperands(offsetInfo, indexType, result.operands) ||
       parser.resolveOperands(sizesInfo, indexType, result.operands) ||
       parser.parseKeywordType("to", dstType) ||
       parser.addTypeToList(dstType, result.types));
 }
 
 static void print(OpAsmPrinter &p, ViewOp op) {
-  p << op.getOperationName() << ' ' << op.getOperand(0) << '[';
+  p << ' ' << op.getOperand(0) << '[';
   p.printOperand(op.byte_shift());
   p << "][" << op.sizes() << ']';
   p.printOptionalAttrDict(op->getAttrs());
   p << " : " << op.getOperand(0).getType() << " to " << op.getType();
 }
 
 static LogicalResult verify(ViewOp op) {
   auto baseType = op.getOperand(0).getType().cast<MemRefType>();

mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp

static ParseResult parseParallelOp(OpAsmParser &parser,
   // Additional attributes
   if (failed(parser.parseOptionalAttrDictWithKeyword(result.attributes)))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &printer, ParallelOp &op) {
-  printer << ParallelOp::getOperationName();
-
   // async()?
   if (Value async = op.async())
     printer << " " << ParallelOp::getAsyncKeyword() << "(" << async << ": "
             << async.getType() << ")";
 
   // wait()?
   printOperandList(op.waitOperands(), ParallelOp::getWaitKeyword(), printer);
 
static ParseResult parseLoopOp(OpAsmParser &parser, OperationState &result) {
 
   if (failed(parser.parseOptionalAttrDictWithKeyword(result.attributes)))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &printer, LoopOp &op) {
-  printer << LoopOp::getOperationName();
-
   unsigned execMapping = op.exec_mapping();
   if (execMapping & OpenACCExecMapping::GANG) {
     printer << " " << LoopOp::getGangKeyword();
     Value gangNum = op.gangNum();
     Value gangStatic = op.gangStatic();
 
     // Print optional gang operands
     if (gangNum || gangStatic) {

mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp

 static LogicalResult verifyParallelOp(ParallelOp op) {
   if (op.allocate_vars().size() != op.allocators_vars().size())
     return op.emitError(
         "expected equal sizes for allocate and allocator variables");
   return success();
 }
 
 static void printParallelOp(OpAsmPrinter &p, ParallelOp op) {
-  p << "omp.parallel";
-
   if (auto ifCond = op.if_expr_var())
     p << " if(" << ifCond << " : " << ifCond.getType() << ")";
 
   if (auto threads = op.num_threads_var())
     p << " num_threads(" << threads << " : " << threads.getType() << ")";
 
   // Print private, firstprivate, shared and copyin parameters
   auto printDataVars = [&p](StringRef name, OperandRange vars) {
static ParseResult parseWsLoopOp(OpAsmParser &parser, OperationState &result) {
   SmallVector<OpAsmParser::OperandType> blockArgs(ivs);
   if (parser.parseRegion(*body, blockArgs, ivTypes))
     return failure();
   return success();
 }
 
 static void printWsLoopOp(OpAsmPrinter &p, WsLoopOp op) {
   auto args = op.getRegion().front().getArguments();
-  p << op.getOperationName() << " (" << args << ") : " << args[0].getType()
-    << " = (" << op.lowerBound() << ") to (" << op.upperBound() << ") step ("
-    << op.step() << ")";
+  p << " (" << args << ") : " << args[0].getType() << " = (" << op.lowerBound()
+    << ") to (" << op.upperBound() << ") step (" << op.step() << ")";
 
   // Print private, firstprivate, shared and copyin parameters
   auto printDataVars = [&p](StringRef name, OperandRange vars) {
     if (vars.empty())
       return;
 
     p << " " << name << "(";
     llvm::interleaveComma(

mlir/lib/Dialect/SCF/SCF.cpp

static ParseResult parseExecuteRegionOp(OpAsmParser &parser,
   if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{}) ||
       parser.parseOptionalAttrDict(result.attributes))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, ExecuteRegionOp op) {
-  p << ExecuteRegionOp::getOperationName();
   p.printOptionalArrowTypeList(op.getResultTypes());
 
   p.printRegion(op.region(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/true);
 
   p.printOptionalAttrDict(op->getAttrs());
 }
static void printInitializationList(OpAsmPrinter &p,
   p << prefix << '(';
   llvm::interleaveComma(llvm::zip(blocksArgs, initializers), p, [&](auto it) {
     p << std::get<0>(it) << " = " << std::get<1>(it);
   });
   p << ")";
 }
 
 static void print(OpAsmPrinter &p, ForOp op) {
-  p << op.getOperationName() << " " << op.getInductionVar() << " = "
-    << op.lowerBound() << " to " << op.upperBound() << " step " << op.step();
+  p << " " << op.getInductionVar() << " = " << op.lowerBound() << " to "
+    << op.upperBound() << " step " << op.step();
 
   printInitializationList(p, op.getRegionIterArgs(), op.getIterOperands(),
                           " iter_args");
   if (!op.getIterOperands().empty())
     p << " -> (" << op.getIterOperands().getTypes() << ')';
   p.printRegion(op.region(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/op.hasIterOperands());
static ParseResult parseIfOp(OpAsmParser &parser, OperationState &result) {
   if (parser.parseOptionalAttrDict(result.attributes))
     return failure();
   return success();
 }
 
 static void print(OpAsmPrinter &p, IfOp op) {
   bool printBlockTerminators = false;
 
-  p << IfOp::getOperationName() << " " << op.condition();
+  p << " " << op.condition();
   if (!op.results().empty()) {
     p << " -> (" << op.getResultTypes() << ")";
     // Print yield explicitly if the op defines values.
     printBlockTerminators = true;
   }
   p.printRegion(op.thenRegion(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/printBlockTerminators);
parser.resolveOperands(initVals, result.types, parser.getNameLoc(),
                            result.operands);
   // Add a terminator if none was parsed.
   ForOp::ensureTerminator(*body, builder, result.location);
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, ParallelOp op) {
-  p << op.getOperationName() << " (" << op.getBody()->getArguments() << ") = ("
-    << op.lowerBound() << ") to (" << op.upperBound() << ") step (" << op.step()
-    << ")";
+  p << " (" << op.getBody()->getArguments() << ") = (" << op.lowerBound()
+    << ") to (" << op.upperBound() << ") step (" << op.step() << ")";
   if (!op.initVals().empty())
     p << " init (" << op.initVals() << ")";
   p.printOptionalArrowTypeList(op.getResultTypes());
   p.printRegion(op.region(), /*printEntryBlockArgs=*/false);
   p.printOptionalAttrDict(
       op->getAttrs(), /*elidedAttrs=*/ParallelOp::getOperandSegmentSizeAttr());
 }
 
static ParseResult parseReduceOp(OpAsmParser &parser, OperationState &result) {
   Region *body = result.addRegion();
   if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{}))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, ReduceOp op) {
-  p << op.getOperationName() << "(" << op.operand() << ") ";
+  p << "(" << op.operand() << ") ";
   p << " : " << op.operand().getType();
   p.printRegion(op.reductionOperator());
 }
 
 //===----------------------------------------------------------------------===//
 // ReduceReturnOp
 //===----------------------------------------------------------------------===//
 
static ParseResult parseWhileOp(OpAsmParser &parser, OperationState &result) {
   return failure(
       parser.parseRegion(*before, regionArgs, functionType.getInputs()) ||
       parser.parseKeyword("do") || parser.parseRegion(*after) ||
       parser.parseOptionalAttrDictWithKeyword(result.attributes));
 }
 
 /// Prints a `while` op.
 static void print(OpAsmPrinter &p, scf::WhileOp op) {
-  p << op.getOperationName();
   printInitializationList(p, op.before().front().getArguments(), op.inits(),
                           " ");
   p << " : ";
   p.printFunctionalType(op.inits().getTypes(), op.results().getTypes());
   p.printRegion(op.before(), /*printEntryBlockArgs=*/false);
   p << " do";
   p.printRegion(op.after());
   p.printOptionalAttrDictWithKeyword(op->getAttrs());

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

static ParseResult parseAtomicUpdateOp(OpAsmParser &parser,
   if (parser.resolveOperands(operandInfo, operandTypes, parser.getNameLoc(),
                              state.operands))
     return failure();
   return parser.addTypeToList(ptrType.getPointeeType(), state.types);
 }
 
 // Prints an atomic update op.
 static void printAtomicUpdateOp(Operation *op, OpAsmPrinter &printer) {
-  printer << op->getName() << " \"";
+  printer << " \"";
   auto scopeAttr = op->getAttrOfType<IntegerAttr>(kMemoryScopeAttrName);
   printer << spirv::stringifyScope(
                  static_cast<spirv::Scope>(scopeAttr.getInt()))
           << "\" \"";
   auto memorySemanticsAttr = op->getAttrOfType<IntegerAttr>(kSemanticsAttrName);
   printer << spirv::stringifyMemorySemantics(
                  static_cast<spirv::MemorySemantics>(
                      memorySemanticsAttr.getInt()))
if (parser.resolveOperand(*clusterSizeInfo, i32Type, state.operands))
       return failure();
   }
 
   return parser.addTypeToList(resultType, state.types);
 }
 
 static void printGroupNonUniformArithmeticOp(Operation *groupOp,
                                              OpAsmPrinter &printer) {
-  printer << groupOp->getName() << " \""
+  printer << " \""
           << stringifyScope(static_cast<spirv::Scope>(
                  groupOp->getAttrOfType<IntegerAttr>(kExecutionScopeAttrName)
                      .getInt()))
           << "\" \""
           << stringifyGroupOperation(static_cast<spirv::GroupOperation>(
                  groupOp->getAttrOfType<IntegerAttr>(kGroupOperationAttrName)
                      .getInt()))
           << "\" " << groupOp->getOperand(0);
if (parser.parseOperand(operandInfo) || parser.parseColonType(type) ||
       parser.resolveOperands(operandInfo, type, state.operands)) {
     return failure();
   }
   state.addTypes(type);
   return success();
 }
 
 static void printUnaryOp(Operation *unaryOp, OpAsmPrinter &printer) {
-  printer << unaryOp->getName() << ' ' << unaryOp->getOperand(0) << " : "
+  printer << ' ' << unaryOp->getOperand(0) << " : "
           << unaryOp->getOperand(0).getType();
 }
 
 /// Result of a logical op must be a scalar or vector of boolean type.
 static Type getUnaryOpResultType(Builder &builder, Type operandType) {
   Type resultType = builder.getIntegerType(1);
   if (auto vecType = operandType.dyn_cast<VectorType>()) {
     return VectorType::get(vecType.getNumElements(), resultType);
if (parser.parseOperandList(ops, 2) || parser.parseColonType(type) ||
       parser.resolveOperands(ops, type, result.operands)) {
     return failure();
   }
   result.addTypes(getUnaryOpResultType(parser.getBuilder(), type));
   return success();
 }
 
 static void printLogicalOp(Operation *logicalOp, OpAsmPrinter &printer) {
-  printer << logicalOp->getName() << ' ' << logicalOp->getOperands() << " : "
+  printer << ' ' << logicalOp->getOperands() << " : "
           << logicalOp->getOperand(0).getType();
 }
 
 static ParseResult parseShiftOp(OpAsmParser &parser, OperationState &state) {
   SmallVector<OpAsmParser::OperandType, 2> operandInfo;
   Type baseType;
   Type shiftType;
   auto loc = parser.getCurrentLocation();
 
   if (parser.parseOperandList(operandInfo, 2) || parser.parseColon() ||
       parser.parseType(baseType) || parser.parseComma() ||
       parser.parseType(shiftType) ||
       parser.resolveOperands(operandInfo, {baseType, shiftType}, loc,
                              state.operands)) {
     return failure();
   }
   state.addTypes(baseType);
   return success();
 }
 
 static void printShiftOp(Operation *op, OpAsmPrinter &printer) {
   Value base = op->getOperand(0);
   Value shift = op->getOperand(1);
-  printer << op->getName() << ' ' << base << ", " << shift << " : "
-          << base.getType() << ", " << shift.getType();
+  printer << ' ' << base << ", " << shift << " : " << base.getType() << ", "
+          << shift.getType();
 }
 
 static LogicalResult verifyShiftOp(Operation *op) {
   if (op->getOperand(0).getType() != op->getResult(0).getType()) {
     return op->emitError("expected the same type for the first operand and "
                          "result, but provided ")
            << op->getOperand(0).getType() << " and "
            << op->getResult(0).getType();
static ParseResult parseAccessChainOp(OpAsmParser &parser,
   }
 
   state.addTypes(resultType);
   return success();
 }
 
 template <typename Op>
 static void printAccessChain(Op op, ValueRange indices, OpAsmPrinter &printer) {
-  printer << Op::getOperationName() << ' ' << op.base_ptr() << '[' << indices
+  printer << ' ' << op.base_ptr() << '[' << indices
           << "] : " << op.base_ptr().getType() << ", " << indices.getTypes();
 }
 
 static void print(spirv::AccessChainOp op, OpAsmPrinter &printer) {
   printAccessChain(op, op.indices(), printer);
 }
 
 template <typename Op>
if (parser.resolveOperands(
           parser.getNameLoc(), state.operands))
     return failure();
 
   return parser.addTypeToList(ptrType.getPointeeType(), state.types);
 }
 
 static void print(spirv::AtomicCompareExchangeWeakOp atomOp,
                   OpAsmPrinter &printer) {
-  printer << spirv::AtomicCompareExchangeWeakOp::getOperationName() << " \""
-          << stringifyScope(atomOp.memory_scope()) << "\" \""
+  printer << " \"" << stringifyScope(atomOp.memory_scope()) << "\" \""
           << stringifyMemorySemantics(atomOp.equal_semantics()) << "\" \""
           << stringifyMemorySemantics(atomOp.unequal_semantics()) << "\" "
           << atomOp.getOperands() << " : " << atomOp.pointer().getType();
 }
 
 static LogicalResult verify(spirv::AtomicCompareExchangeWeakOp atomOp) {
   // According to the spec:
   // "The type of Value must be the same as Result Type. The type of the value
state.addAttribute(
       spirv::BranchConditionalOp::getOperandSegmentSizeAttr(),
       builder.getI32VectorAttr({1, static_cast<int32_t>(trueOperands.size()),
                                 static_cast<int32_t>(falseOperands.size())}));
 
   return success();
 }
 
 static void print(spirv::BranchConditionalOp branchOp, OpAsmPrinter &printer) {
-  printer << spirv::BranchConditionalOp::getOperationName() << ' '
-          << branchOp.condition();
+  printer << ' ' << branchOp.condition();
 
   if (auto weights = branchOp.branch_weights()) {
     printer << " [";
     llvm::interleaveComma(weights->getValue(), printer, [&](Attribute a) {
       printer << a.cast<IntegerAttr>().getInt();
     });
     printer << "]";
   }
for (auto index : llvm::seq<uint32_t>(0, operands.size())) {
     elementTypes.push_back(cType.getElementType(index));
   }
   state.addTypes(type);
   return parser.resolveOperands(operands, elementTypes, loc, state.operands);
 }
 
 static void print(spirv::CompositeConstructOp compositeConstructOp,
                   OpAsmPrinter &printer) {
-  printer << spirv::CompositeConstructOp::getOperationName() << " "
-          << compositeConstructOp.constituents() << " : "
+  printer << " " << compositeConstructOp.constituents() << " : "
           << compositeConstructOp.getResult().getType();
 }
 
 static LogicalResult verify(spirv::CompositeConstructOp compositeConstructOp) {
   auto cType = compositeConstructOp.getType().cast<spirv::CompositeType>();
   SmallVector<Value, 4> constituents(compositeConstructOp.constituents());
 
   if (cType.isa<spirv::CooperativeMatrixNVType>()) {
if (!resultType) {
     return failure();
   }
   state.addTypes(resultType);
   return success();
 }
 
 static void print(spirv::CompositeExtractOp compositeExtractOp,
                   OpAsmPrinter &printer) {
-  printer << spirv::CompositeExtractOp::getOperationName() << ' '
-          << compositeExtractOp.composite() << compositeExtractOp.indices()
-          << " : " << compositeExtractOp.composite().getType();
+  printer << ' ' << compositeExtractOp.composite()
+          << compositeExtractOp.indices() << " : "
+          << compositeExtractOp.composite().getType();
 }
 
 static LogicalResult verify(spirv::CompositeExtractOp compExOp) {
   auto indicesArrayAttr = compExOp.indices().dyn_cast<ArrayAttr>();
   auto resultType = getElementType(compExOp.composite().getType(),
                                    indicesArrayAttr, compExOp.getLoc());
   if (!resultType)
     return failure();
return compositeInsertOp.emitOpError("result type should be the same as "
            << compositeInsertOp.getType();
   }
 
   return success();
 }
 
 static void print(spirv::CompositeInsertOp compositeInsertOp,
                   OpAsmPrinter &printer) {
-  printer << spirv::CompositeInsertOp::getOperationName() << " "
-          << compositeInsertOp.object() << ", " << compositeInsertOp.composite()
-          << compositeInsertOp.indices() << " : "
-          << compositeInsertOp.object().getType() << " into "
+  printer << " " << compositeInsertOp.object() << ", "
+          << compositeInsertOp.composite() << compositeInsertOp.indices()
+          << " : " << compositeInsertOp.object().getType() << " into "
           << compositeInsertOp.composite().getType();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.Constant
 //===----------------------------------------------------------------------===//
 
 static ParseResult parseConstantOp(OpAsmParser &parser, OperationState &state) {
   Attribute value;
   if (parser.parseAttribute(value, kValueAttrName, state.attributes))
     return failure();
 
   Type type = value.getType();
   if (type.isa<NoneType, TensorType>()) {
     if (parser.parseColonType(type))
       return failure();
   }
 
   return parser.addTypeToList(type, state.types);
 }
 
 static void print(spirv::ConstantOp constOp, OpAsmPrinter &printer) {
-  printer << spirv::ConstantOp::getOperationName() << ' ' << constOp.value();
+  printer << ' ' << constOp.value();
   if (constOp.getType().isa<spirv::ArrayType>())
     printer << " : " << constOp.getType();
 }
 
 static LogicalResult verify(spirv::ConstantOp constOp) {
   auto opType = constOp.getType();
   auto value = constOp.value();
   auto valueType = value.getType();
if (!parser.parseOptionalComma()) {
     } while (!parser.parseOptionalComma());
   }
   state.addAttribute(kInterfaceAttrName,
                      parser.getBuilder().getArrayAttr(interfaceVars));
   return success();
 }
 
 static void print(spirv::EntryPointOp entryPointOp, OpAsmPrinter &printer) {
-  printer << spirv::EntryPointOp::getOperationName() << " \""
-          << stringifyExecutionModel(entryPointOp.execution_model()) << "\" ";
+  printer << " \"" << stringifyExecutionModel(entryPointOp.execution_model())
+          << "\" ";
   printer.printSymbolName(entryPointOp.fn());
   auto interfaceVars = entryPointOp.interface().getValue();
   if (!interfaceVars.empty()) {
     printer << ", ";
     llvm::interleaveComma(interfaceVars, printer);
   }
 }
 
while (!parser.parseOptionalComma()) {
     values.push_back(value.cast<IntegerAttr>().getInt());
   }
   state.addAttribute(kValuesAttrName,
                      parser.getBuilder().getI32ArrayAttr(values));
   return success();
 }
 
 static void print(spirv::ExecutionModeOp execModeOp, OpAsmPrinter &printer) {
-  printer << spirv::ExecutionModeOp::getOperationName() << " ";
+  printer << " ";
   printer.printSymbolName(execModeOp.fn());
   printer << " \"" << stringifyExecutionMode(execModeOp.execution_mode())
           << "\"";
   auto values = execModeOp.values();
   if (!values.size())
     return;
   printer << ", ";
   llvm::interleaveComma(values, printer, [&](Attribute a) {
static ParseResult parseFuncOp(OpAsmParser &parser, OperationState &state) {
   auto *body = state.addRegion();
   OptionalParseResult result = parser.parseOptionalRegion(
       *body, entryArgs, entryArgs.empty() ? ArrayRef<Type>() : argTypes);
   return failure(result.hasValue() && failed(*result));
 }
 
 static void print(spirv::FuncOp fnOp, OpAsmPrinter &printer) {
   // Print function name, signature, and control.
-  printer << spirv::FuncOp::getOperationName() << " ";
+  printer << " ";
   printer.printSymbolName(fnOp.sym_name());
   auto fnType = fnOp.getType();
   function_like_impl::printFunctionSignature(printer, fnOp, fnType.getInputs(),
                                              /*isVariadic=*/false,
                                              fnType.getResults());
   printer << " \"" << spirv::stringifyFunctionControl(fnOp.function_control())
           << "\"";
   function_like_impl::printFunctionAttributes(
static ParseResult parseGlobalVariableOp(OpAsmParser &parser,
 
   return success();
 }
 
 static void print(spirv::GlobalVariableOp varOp, OpAsmPrinter &printer) {
   auto *op = varOp.getOperation();
   SmallVector<StringRef, 4> elidedAttrs{
       spirv::attributeName<spirv::StorageClass>()};
-  printer << spirv::GlobalVariableOp::getOperationName();
 
   // Print variable name.
   printer << ' ';
   printer.printSymbolName(varOp.sym_name());
   elidedAttrs.push_back(SymbolTable::getSymbolAttrName());
 
   // Print optional initializer
   if (auto initializer = varOp.initializer()) {
static ParseResult parseSubgroupBlockReadINTELOp(OpAsmParser &parser,
 
   state.addTypes(elementType);
   return success();
 }
 
 static void print(spirv::SubgroupBlockReadINTELOp blockReadOp,
                   OpAsmPrinter &printer) {
   SmallVector<StringRef, 4> elidedAttrs;
-  printer << spirv::SubgroupBlockReadINTELOp::getOperationName() << " "
-          << blockReadOp.ptr();
+  printer << " " << blockReadOp.ptr();
   printer << " : " << blockReadOp.getType();
 }
 
 static LogicalResult verify(spirv::SubgroupBlockReadINTELOp blockReadOp) {
   if (failed(verifyBlockReadWritePtrAndValTypes(blockReadOp, blockReadOp.ptr(),
                                                 blockReadOp.value())))
     return failure();
 
if (parser.resolveOperands(operandInfo, {ptrType, elementType}, loc,
     return failure();
   }
   return success();
 }
 
 static void print(spirv::SubgroupBlockWriteINTELOp blockWriteOp,
                   OpAsmPrinter &printer) {
   SmallVector<StringRef, 4> elidedAttrs;
-  printer << spirv::SubgroupBlockWriteINTELOp::getOperationName() << " "
-          << blockWriteOp.ptr() << ", " << blockWriteOp.value();
+  printer << " " << blockWriteOp.ptr() << ", " << blockWriteOp.value();
   printer << " : " << blockWriteOp.value().getType();
 }
 
 static LogicalResult verify(spirv::SubgroupBlockWriteINTELOp blockWriteOp) {
   if (failed(verifyBlockReadWritePtrAndValTypes(
           blockWriteOp, blockWriteOp.ptr(), blockWriteOp.value())))
     return failure();
 
static ParseResult parseLoadOp(OpAsmParser &parser, OperationState &state) {
   return success();
 }
 
 static void print(spirv::LoadOp loadOp, OpAsmPrinter &printer) {
   auto *op = loadOp.getOperation();
   SmallVector<StringRef, 4> elidedAttrs;
   StringRef sc = stringifyStorageClass(
       loadOp.ptr().getType().cast<spirv::PointerType>().getStorageClass());
-  printer << spirv::LoadOp::getOperationName() << " \"" << sc << "\" "
-          << loadOp.ptr();
+  printer << " \"" << sc << "\" " << loadOp.ptr();
 
   printMemoryAccessAttribute(loadOp, printer, elidedAttrs);
 
   printer.printOptionalAttrDict(op->getAttrs(), elidedAttrs);
   printer << " : " << loadOp.getType();
 }
 
 static LogicalResult verify(spirv::LoadOp loadOp) {
if (parseControlAttribute<spirv::LoopControl>(parser, state))
     return failure();
   return parser.parseRegion(*state.addRegion(), /*arguments=*/{},
                             /*argTypes=*/{});
 }
 
 static void print(spirv::LoopOp loopOp, OpAsmPrinter &printer) {
   auto *op = loopOp.getOperation();
 
-  printer << spirv::LoopOp::getOperationName();
   auto control = loopOp.loop_control();
   if (control != spirv::LoopControl::None)
     printer << " control(" << spirv::stringifyLoopControl(control) << ")";
   printer.printRegion(op->getRegion(0), /*printEntryBlockArgs=*/false,
                       /*printBlockTerminators=*/true);
 }
 
 /// Returns true if the given `srcBlock` contains only one `spv.Branch` to the
static ParseResult parseModuleOp(OpAsmParser &parser, OperationState &state) {
   // Make sure we have at least one block.
   if (body->empty())
     body->push_back(new Block());
 
   return success();
 }
 
 static void print(spirv::ModuleOp moduleOp, OpAsmPrinter &printer) {
-  printer << spirv::ModuleOp::getOperationName();
-
   if (Optional<StringRef> name = moduleOp.getName()) {
     printer << ' ';
     printer.printSymbolName(*name);
   }
 
   SmallVector<StringRef, 2> elidedAttrs;
 
   printer << " " << spirv::stringifyAddressingModel(moduleOp.addressing_model())
static ParseResult parseSelectionOp(OpAsmParser &parser,
   if (parseControlAttribute<spirv::SelectionControl>(parser, state))
     return failure();
   return parser.parseRegion(*state.addRegion(), /*arguments=*/{},
                             /*argTypes=*/{});
 }
 
 static void print(spirv::SelectionOp selectionOp, OpAsmPrinter &printer) {
   auto *op = selectionOp.getOperation();
-
-  printer << spirv::SelectionOp::getOperationName();
   auto control = selectionOp.selection_control();
   if (control != spirv::SelectionControl::None)
     printer << " control(" << spirv::stringifySelectionControl(control) << ")";
   printer.printRegion(op->getRegion(0), /*printEntryBlockArgs=*/false,
                       /*printBlockTerminators=*/true);
 }
 
 static LogicalResult verify(spirv::SelectionOp selectionOp) {
static ParseResult parseSpecConstantOp(OpAsmParser &parser,
   if (parser.parseEqual() ||
       parser.parseAttribute(valueAttr, kDefaultValueAttrName, state.attributes))
     return failure();
 
   return success();
 }
 
 static void print(spirv::SpecConstantOp constOp, OpAsmPrinter &printer) {
-  printer << spirv::SpecConstantOp::getOperationName() << ' ';
+  printer << ' ';
   printer.printSymbolName(constOp.sym_name());
   if (auto specID = constOp->getAttrOfType<IntegerAttr>(kSpecIdAttrName))
     printer << ' ' << kSpecIdAttrName << '(' << specID.getInt() << ')';
   printer << " = " << constOp.default_value();
 }
 
 static LogicalResult verify(spirv::SpecConstantOp constOp) {
   if (auto specID = constOp->getAttrOfType<IntegerAttr>(kSpecIdAttrName))
static ParseResult parseStoreOp(OpAsmParser &parser, OperationState &state) {
   return success();
 }
 
 static void print(spirv::StoreOp storeOp, OpAsmPrinter &printer) {
   auto *op = storeOp.getOperation();
   SmallVector<StringRef, 4> elidedAttrs;
   StringRef sc = stringifyStorageClass(
       storeOp.ptr().getType().cast<spirv::PointerType>().getStorageClass());
-  printer << spirv::StoreOp::getOperationName() << " \"" << sc << "\" "
-          << storeOp.ptr() << ", " << storeOp.value();
+  printer << " \"" << sc << "\" " << storeOp.ptr() << ", " << storeOp.value();
 
   printMemoryAccessAttribute(storeOp, printer, elidedAttrs);
 
   printer << " : " << storeOp.value().getType();
   printer.printOptionalAttrDict(op->getAttrs(), elidedAttrs);
 }
 
 static LogicalResult verify(spirv::StoreOp storeOp) {
static ParseResult parseVariableOp(OpAsmParser &parser, OperationState &state) {
   state.addAttribute(spirv::attributeName<spirv::StorageClass>(), attr);
 
   return success();
 }
 
 static void print(spirv::VariableOp varOp, OpAsmPrinter &printer) {
   SmallVector<StringRef, 4> elidedAttrs{
       spirv::attributeName<spirv::StorageClass>()};
-  printer << spirv::VariableOp::getOperationName();
-
   // Print optional initializer
   if (varOp.getNumOperands() != 0)
     printer << " init(" << varOp.initializer() << ")";
 
   printVariableDecorations(varOp, printer, elidedAttrs);
   printer << " : " << varOp.getType();
 }
 
if (parser.resolveOperands(operandInfo,
     return failure();
   }
 
   state.addTypes(elementType);
   return success();
 }
 
 static void print(spirv::CooperativeMatrixLoadNVOp M, OpAsmPrinter &printer) {
-  printer << spirv::CooperativeMatrixLoadNVOp::getOperationName() << " "
-          << M.pointer() << ", " << M.stride() << ", " << M.columnmajor();
+  printer << " " << M.pointer() << ", " << M.stride() << ", "
+          << M.columnmajor();
   // Print optional memory access attribute.
   if (auto memAccess = M.memory_access())
     printer << " [\"" << stringifyMemoryAccess(*memAccess) << "\"]";
   printer << " : " << M.pointer().getType() << " as " << M.getType();
 }
 
 static LogicalResult verifyPointerAndCoopMatrixType(Operation *op, Type pointer,
                                                     Type coopMatrix) {
if (parser.resolveOperands(
     return failure();
   }
 
   return success();
 }
 
 static void print(spirv::CooperativeMatrixStoreNVOp coopMatrix,
                   OpAsmPrinter &printer) {
-  printer << spirv::CooperativeMatrixStoreNVOp::getOperationName() << " "
-          << coopMatrix.pointer() << ", " << coopMatrix.object() << ", "
+  printer << " " << coopMatrix.pointer() << ", " << coopMatrix.object() << ", "
           << coopMatrix.stride() << ", " << coopMatrix.columnmajor();
   // Print optional memory access attribute.
   if (auto memAccess = coopMatrix.memory_access())
     printer << " [\"" << stringifyMemoryAccess(*memAccess) << "\"]";
   printer << " : " << coopMatrix.pointer().getType() << ", "
           << coopMatrix.getOperand(1).getType();
 }
 

 }
 
 //===----------------------------------------------------------------------===//
 // spv.CopyMemory
 //===----------------------------------------------------------------------===//
 
 static void print(spirv::CopyMemoryOp copyMemory, OpAsmPrinter &printer) {
   auto *op = copyMemory.getOperation();
-  printer << spirv::CopyMemoryOp::getOperationName() << ' ';
+  printer << ' ';
 
   StringRef targetStorageClass =
       stringifyStorageClass(copyMemory.target()
                                 .getType()
                                 .cast<spirv::PointerType>()
                                 .getStorageClass());
   printer << " \"" << targetStorageClass << "\" " << copyMemory.target()
           << ", ";
if (parser.parseColonType(type))
     return failure();
 
   state.addAttribute(kTypeAttrName, TypeAttr::get(type));
 
   return success();
 }
 
 static void print(spirv::SpecConstantCompositeOp op, OpAsmPrinter &printer) {
-  printer << spirv::SpecConstantCompositeOp::getOperationName() << " ";
+  printer << " ";
   printer.printSymbolName(op.sym_name());
   printer << " (";
   auto constituents = op.constituents().getValue();
 
   if (!constituents.empty())
     llvm::interleaveComma(constituents, printer);
 
   printer << ") : " << op.type();
static ParseResult parseSpecConstantOperationOp(OpAsmParser &parser,
 
   if (parser.parseOptionalAttrDict(state.attributes))
     return failure();
 
   return success();
 }
 
 static void print(spirv::SpecConstantOperationOp op, OpAsmPrinter &printer) {
-  printer << op.getOperationName() << " wraps ";
+  printer << " wraps ";
   printer.printGenericOp(&op.body().front().front());
 }
 
 static LogicalResult verify(spirv::SpecConstantOperationOp constOp) {
   Block &block = constOp.getRegion().getBlocks().front();
 
   if (block.getOperations().size() != 2)
     return constOp.emitOpError("expected exactly 2 nested ops");

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

static ParseResult parseAssumingOp(OpAsmParser &parser,
   if (parser.parseOptionalAttrDict(result.attributes))
     return failure();
   return success();
 }
 
 static void print(OpAsmPrinter &p, AssumingOp op) {
   bool yieldsResults = !op.results().empty();
 
-  p << AssumingOp::getOperationName() << " " << op.witness();
+  p << " " << op.witness();
   if (yieldsResults) {
     p << " -> (" << op.getResultTypes() << ")";
   }
   p.printRegion(op.doRegion(),
                 /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/yieldsResults);
   p.printOptionalAttrDict(op->getAttrs());
 }
OpFoldResult ConcatOp::fold(ArrayRef<Attribute> operands) {
   return builder.getIndexTensorAttr(resultShape);
 }
 
 //===----------------------------------------------------------------------===//
 // ConstShapeOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, ConstShapeOp &op) {
-  p << "shape.const_shape ";
+  p << " ";
   p.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"shape"});
   p << "[";
   interleaveComma(op.shape().getValues<int64_t>(), p,
                   [&](int64_t i) { p << i; });
   p << "] : ";
   p.printType(op.getType());
 }
 
ParseResult parseFunctionLibraryOp(OpAsmParser &parser,
   if (parser.parseAttribute(mappingAttr,
                             parser.getBuilder().getType<NoneType>(), "mapping",
                             result.attributes))
     return failure();
   return success();
 }
 
 void print(OpAsmPrinter &p, FunctionLibraryOp op) {
-  p << op.getOperationName() << ' ';
+  p << ' ';
   p.printSymbolName(op.getName());
   p.printOptionalAttrDictWithKeyword(
       op->getAttrs(), {SymbolTable::getSymbolAttrName(), "mapping"});
   p.printRegion(op.getOperation()->getRegion(0), /*printEntryBlockArgs=*/false,
                 /*printBlockTerminators=*/false);
   p << " mapping ";
   p.printAttributeWithoutType(op.mappingAttr());
 }
static ParseResult parseReduceOp(OpAsmParser &parser, OperationState &result) {
   // Parse attributes.
   if (parser.parseOptionalAttrDict(result.attributes))
     return failure();
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, ReduceOp op) {
-  p << op.getOperationName() << '(' << op.shape() << ", " << op.initVals()
+  p << '(' << op.shape() << ", " << op.initVals()
     << ") : " << op.shape().getType();
   p.printOptionalArrowTypeList(op.getResultTypes());
   p.printRegion(op.region());
   p.printOptionalAttrDict(op->getAttrs());
 }
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/Shape/IR/ShapeOps.cpp.inc"

mlir/lib/Dialect/StandardOps/IR/Ops.cpp

 //===----------------------------------------------------------------------===//
 
 /// A custom unary operation printer that omits the "std." prefix from the
 /// operation names.
 static void printStandardUnaryOp(Operation *op, OpAsmPrinter &p) {
   assert(op->getNumOperands() == 1 && "unary op should have one operand");
   assert(op->getNumResults() == 1 && "unary op should have one result");
 
-  int stdDotLen = StandardOpsDialect::getDialectNamespace().size() + 1;
-  p << op->getName().getStringRef().drop_front(stdDotLen) << ' '
-    << op->getOperand(0);
+  p << ' ' << op->getOperand(0);
   p.printOptionalAttrDict(op->getAttrs());
   p << " : " << op->getOperand(0).getType();
 }
 
 /// A custom binary operation printer that omits the "std." prefix from the
 /// operation names.
 static void printStandardBinaryOp(Operation *op, OpAsmPrinter &p) {
   assert(op->getNumOperands() == 2 && "binary op should have two operands");
   assert(op->getNumResults() == 1 && "binary op should have one result");
 
   // If not all the operand and result types are the same, just use the
   // generic assembly form to avoid omitting information in printing.
   auto resultType = op->getResult(0).getType();
   if (op->getOperand(0).getType() != resultType ||
       op->getOperand(1).getType() != resultType) {
     p.printGenericOp(op);
     return;
   }
 
-  int stdDotLen = StandardOpsDialect::getDialectNamespace().size() + 1;
-  p << op->getName().getStringRef().drop_front(stdDotLen) << ' '
-    << op->getOperand(0) << ", " << op->getOperand(1);
+  p << ' ' << op->getOperand(0) << ", " << op->getOperand(1);
   p.printOptionalAttrDict(op->getAttrs());
 
   // Now we can output only one type for all operands and the result.
   p << " : " << op->getResult(0).getType();
 }
 
 /// A custom ternary operation printer that omits the "std." prefix from the
 /// operation names.
 static void printStandardTernaryOp(Operation *op, OpAsmPrinter &p) {
   assert(op->getNumOperands() == 3 && "ternary op should have three operands");
   assert(op->getNumResults() == 1 && "ternary op should have one result");
 
   // If not all the operand and result types are the same, just use the
   // generic assembly form to avoid omitting information in printing.
   auto resultType = op->getResult(0).getType();
   if (op->getOperand(0).getType() != resultType ||
       op->getOperand(1).getType() != resultType ||
       op->getOperand(2).getType() != resultType) {
     p.printGenericOp(op);
     return;
   }
 
-  int stdDotLen = StandardOpsDialect::getDialectNamespace().size() + 1;
-  p << op->getName().getStringRef().drop_front(stdDotLen) << ' '
-    << op->getOperand(0) << ", " << op->getOperand(1) << ", "
+  p << ' ' << op->getOperand(0) << ", " << op->getOperand(1) << ", "
     << op->getOperand(2);
   p.printOptionalAttrDict(op->getAttrs());
 
   // Now we can output only one type for all operands and the result.
   p << " : " << op->getResult(0).getType();
 }
 
 /// A custom cast operation printer that omits the "std." prefix from the
 /// operation names.
 static void printStandardCastOp(Operation *op, OpAsmPrinter &p) {
-  int stdDotLen = StandardOpsDialect::getDialectNamespace().size() + 1;
-  p << op->getName().getStringRef().drop_front(stdDotLen) << ' '
-    << op->getOperand(0) << " : " << op->getOperand(0).getType() << " to "
-    << op->getResult(0).getType();
+  p << ' ' << op->getOperand(0) << " : " << op->getOperand(0).getType()
+    << " to " << op->getResult(0).getType();
 }
 
 void StandardOpsDialect::initialize() {
   addOperations<
 #define GET_OP_LIST
 #include "mlir/Dialect/StandardOps/IR/Ops.cpp.inc"
       >();
   addInterfaces<StdInlinerInterface>();
static ParseResult parseGenericAtomicRMWOp(OpAsmParser &parser,
   if (parser.parseRegion(*body, llvm::None, llvm::None) ||
       parser.parseOptionalAttrDict(result.attributes))
     return failure();
   result.types.push_back(memrefType.cast<MemRefType>().getElementType());
   return success();
 }
 
 static void print(OpAsmPrinter &p, GenericAtomicRMWOp op) {
-  p << op.getOperationName() << ' ' << op.memref() << "[" << op.indices()
+  p << ' ' << op.memref() << "[" << op.indices()
     << "] : " << op.memref().getType();
   p.printRegion(op.body());
   p.printOptionalAttrDict(op->getAttrs());
 }
 
 //===----------------------------------------------------------------------===//
 // AtomicYieldOp
 //===----------------------------------------------------------------------===//
Block *CondBranchOp::getSuccessorForOperands(ArrayRef<Attribute> operands) {
   return nullptr;
 }
 
 //===----------------------------------------------------------------------===//
 // Constant*Op
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, ConstantOp &op) {
-  p << "constant ";
+  p << " ";
   p.printOptionalAttrDict(op->getAttrs(), /*elidedAttrs=*/{"value"});
 
   if (op->getAttrs().size() > 1)
     p << ' ';
   p << op.getValue();
 
   // If the value is a symbol reference or Array, print a trailing type.
   if (op.getValue().isa<SymbolRefAttr, ArrayAttr>())
if (pred == mlir::CmpIPredicate::eq || pred == mlir::CmpIPredicate::ne) {
           (cmpRhs == trueVal && cmpLhs == falseVal))
         return pred == mlir::CmpIPredicate::ne ? trueVal : falseVal;
     }
   }
   return nullptr;
 }
 
 static void print(OpAsmPrinter &p, SelectOp op) {
-  p << "select " << op.getOperands();
+  p << " " << op.getOperands();
   p.printOptionalAttrDict(op->getAttrs());
   p << " : ";
   if (ShapedType condType = op.getCondition().getType().dyn_cast<ShapedType>())
     p << condType << ", ";
   p << op.getType();
 }
 
 static ParseResult parseSelectOp(OpAsmParser &parser, OperationState &result) {

mlir/lib/Dialect/Vector/VectorOps.cpp

if (parser.parseAttribute(attr, "kind", result.attributes) ||
     return failure();
   if (operandsInfo.size() < 1 || operandsInfo.size() > 2)
     return parser.emitError(parser.getNameLoc(),
                             "unsupported number of operands");
   return success();
 }
 
 static void print(OpAsmPrinter &p, ReductionOp op) {
-  p << op.getOperationName() << " \"" << op.kind() << "\", " << op.vector();
+  p << " \"" << op.kind() << "\", " << op.vector();
   if (!op.acc().empty())
     p << ", " << op.acc();
   p << " : " << op.vector().getType() << " into " << op.dest().getType();
 }
 
 Value mlir::vector::getVectorReductionOp(AtomicRMWKind op, OpBuilder &builder,
                                          Location loc, Value vector) {
   Type scalarType = vector.getType().cast<ShapedType>().getElementType();
static void print(OpAsmPrinter &p, ContractionOp op) {
   llvm::StringSet<> traitAttrsSet;
   traitAttrsSet.insert(attrNames.begin(), attrNames.end());
   SmallVector<NamedAttribute, 8> attrs;
   for (auto attr : op->getAttrs())
     if (traitAttrsSet.count(attr.first.strref()) > 0)
       attrs.push_back(attr);
 
   auto dictAttr = DictionaryAttr::get(op.getContext(), attrs);
-  p << op.getOperationName() << " " << dictAttr << " " << op.lhs() << ", ";
+  p << " " << dictAttr << " " << op.lhs() << ", ";
   p << op.rhs() << ", " << op.acc();
   if (op.masks().size() == 2)
     p << ", " << op.masks();
 
   p.printOptionalAttrDict(op->getAttrs(), attrNames);
   p << " : " << op.lhs().getType() << ", " << op.rhs().getType() << " into "
     << op.getResultType();
 }
void vector::ExtractOp::build(OpBuilder &builder, OperationState &result,
   SmallVector<int64_t, 4> positionConstants =
       llvm::to_vector<4>(llvm::map_range(position, [](Value pos) {
         return pos.getDefiningOp<ConstantIndexOp>().getValue();
       }));
   build(builder, result, source, positionConstants);
 }
 
 static void print(OpAsmPrinter &p, vector::ExtractOp op) {
-  p << op.getOperationName() << " " << op.vector() << op.position();
+  p << " " << op.vector() << op.position();
   p.printOptionalAttrDict(op->getAttrs(), {"position"});
   p << " : " << op.vector().getType();
 }
 
 static ParseResult parseExtractOp(OpAsmParser &parser, OperationState &result) {
   llvm::SMLoc attributeLoc, typeLoc;
   NamedAttrList attrs;
   OpAsmParser::OperandType vector;
void ShuffleOp::build(OpBuilder &builder, OperationState &result, Value v1,
                       Value v2, ArrayRef<int64_t> mask) {
   result.addOperands({v1, v2});
   auto maskAttr = getVectorSubscriptAttr(builder, mask);
   result.addTypes(v1.getType());
   result.addAttribute(getMaskAttrName(), maskAttr);
 }
 
 static void print(OpAsmPrinter &p, ShuffleOp op) {
-  p << op.getOperationName() << " " << op.v1() << ", " << op.v2() << " "
-    << op.mask();
+  p << " " << op.v1() << ", " << op.v2() << " " << op.mask();
   p.printOptionalAttrDict(op->getAttrs(), {ShuffleOp::getMaskAttrName()});
   p << " : " << op.v1().getType() << ", " << op.v2().getType();
 }
 
 static LogicalResult verify(ShuffleOp op) {
   VectorType resultType = op.getVectorType();
   VectorType v1Type = op.getV1VectorType();
   VectorType v2Type = op.getV2VectorType();

 /// Build an op without mask, use the type of `acc` as the return type.
 void OuterProductOp::build(OpBuilder &builder, OperationState &result,
                            Value lhs, Value rhs, Value acc) {
   result.addOperands({lhs, rhs, acc});
   result.addTypes(acc.getType());
 }
 
 static void print(OpAsmPrinter &p, OuterProductOp op) {
-  p << op.getOperationName() << " " << op.lhs() << ", " << op.rhs();
+  p << " " << op.lhs() << ", " << op.rhs();
   if (!op.acc().empty()) {
     p << ", " << op.acc();
     p.printOptionalAttrDict(op->getAttrs());
   }
   p << " : " << op.lhs().getType() << ", " << op.rhs().getType();
 }
 
 static ParseResult parseOuterProductOp(OpAsmParser &parser,
return op->emitOpError("requires a permutation_map with input dims of the "
                            "same rank as the source type");
 
   if (inBounds) {
     if (permutationMap.getNumResults() != static_cast<int64_t>(inBounds.size()))
       return op->emitOpError("expects the optional in_bounds attr of same rank "
                              "as permutation_map results: ")
              << AffineMapAttr::get(permutationMap);
     for (unsigned int i = 0; i < permutationMap.getNumResults(); ++i)
-      if (permutationMap.getResult(i).isa<AffineConstantExpr>()
-          && !inBounds.getValue()[i].cast<BoolAttr>().getValue())
+      if (permutationMap.getResult(i).isa<AffineConstantExpr>() &&
+          !inBounds.getValue()[i].cast<BoolAttr>().getValue())
         return op->emitOpError("requires broadcast dimensions to be in-bounds");
   }
 
   return success();
 }
 
 /// Builder that sets padding to zero.
 void TransferReadOp::build(OpBuilder &builder, OperationState &result,
if (auto inBounds = op.in_bounds()) {
     }
   }
   if (elideInBounds)
     elidedAttrs.push_back(op.getInBoundsAttrName());
   p.printOptionalAttrDict(op->getAttrs(), elidedAttrs);
 }
 
 static void print(OpAsmPrinter &p, TransferReadOp op) {
-  p << op.getOperationName() << " " << op.source() << "[" << op.indices()
-    << "], " << op.padding();
+  p << " " << op.source() << "[" << op.indices() << "], " << op.padding();
   if (op.mask())
     p << ", " << op.mask();
   printTransferAttrs(p, cast<VectorTransferOpInterface>(op.getOperation()));
   p << " : " << op.getShapedType() << ", " << op.getVectorType();
 }
 
 static ParseResult parseTransferReadOp(OpAsmParser &parser,
                                        OperationState &result) {
for (unsigned i = 0; i < op.getTransferRank(); ++i) {
     if (op.isDimInBounds(i)) {
       newInBounds.push_back(true);
       continue;
     }
     // Currently out-of-bounds, check whether we can statically determine it is
     // inBounds.
     auto dimExpr = permutationMap.getResult(i).dyn_cast<AffineDimExpr>();
     assert(dimExpr && "Broadcast dims must be in-bounds");
-    auto inBounds = isInBounds(
-        op, /*resultIdx=*/i, /*indicesIdx=*/dimExpr.getPosition());
+    auto inBounds =
+        isInBounds(op, /*resultIdx=*/i, /*indicesIdx=*/dimExpr.getPosition());
     newInBounds.push_back(inBounds);
     // We commit the pattern if it is "more inbounds".
     changed |= inBounds;
   }
   if (!changed)
     return failure();
   // OpBuilder is only used as a helper to build an I64ArrayAttr.
   OpBuilder b(op.getContext());
result.addAttribute(
       TransferWriteOp::getOperandSegmentSizeAttr(),
       builder.getI32VectorAttr({1, 1, static_cast<int32_t>(indexInfo.size()),
                                 static_cast<int32_t>(hasMask.succeeded())}));
   return failure(shapedType.isa<RankedTensorType>() &&
                  parser.addTypeToList(shapedType, result.types));
 }
 
 static void print(OpAsmPrinter &p, TransferWriteOp op) {
-  p << op.getOperationName() << " " << op.vector() << ", " << op.source() << "["
-    << op.indices() << "]";
+  p << " " << op.vector() << ", " << op.source() << "[" << op.indices() << "]";
   if (op.mask())
     p << ", " << op.mask();
   printTransferAttrs(p, cast<VectorTransferOpInterface>(op.getOperation()));
   p << " : " << op.getVectorType() << ", " << op.getShapedType();
 }
 
 static LogicalResult verify(TransferWriteOp op) {
   // Consistency of elemental types in shape and vector.

mlir/lib/IR/AsmPrinter.cpp

void print(Operation *op) {
     }
 
     // Otherwise print with the generic assembly form.
     printGenericOp(op);
   }
 
 private:
   /// Print the given operation in the generic form.
-  void printGenericOp(Operation *op) override {
+  void printGenericOp(Operation *op, bool printOpName = true) override {
     // Consider nested operations for aliases.
     if (op->getNumRegions() != 0) {
       for (Region &region : op->getRegions())
         printRegion(region, /*printEntryBlockArgs=*/true,
                     /*printBlockTerminators=*/true);
     }
 
     // Visit all the types used in the operation.
public:
   /// Print the given top-level operation.
   void printTopLevelOperation(Operation *op);
 
   /// Print the given operation with its indent and location.
   void print(Operation *op);
   /// Print the bare location, not including indentation/location/etc.
   void printOperation(Operation *op);
   /// Print the given operation in the generic form.
-  void printGenericOp(Operation *op) override;
+  void printGenericOp(Operation *op, bool printOpName) override;
 
   /// Print the name of the given block.
   void printBlockName(Block *block);
 
   /// Print the given block. If 'printBlockArgs' is false, the arguments of the
   /// block are not printed. If 'printBlockTerminator' is false, the terminator
   /// operation of the block is not printed.
   void print(Block *block, bool printBlockArgs = true,
if (!printerFlags.shouldPrintGenericOpForm()) {
     // custom printer hook.
     if (auto *opInfo = op->getAbstractOperation()) {
       opInfo->printAssembly(op, *this);
       return;
     }
     // Otherwise try to dispatch to the dialect, if available.
     if (Dialect *dialect = op->getDialect()) {
       if (auto opPrinter = dialect->getOperationPrinter(op)) {
+        // Print the op name first.
+        StringRef name = op->getName().getStringRef();
+        printEscapedString(name, os);
+        // Print the op now.
         opPrinter(op, *this);
         return;
       }
     }
   }
 
   // Otherwise print with the generic assembly form.
-  printGenericOp(op);
+  printGenericOp(op, /*printOpName=*/true);
 }
 
-void OperationPrinter::printGenericOp(Operation *op) {
+void OperationPrinter::printGenericOp(Operation *op, bool printOpName) {
+  if (printOpName) {
     os << '"';
     printEscapedString(op->getName().getStringRef(), os);
-  os << "\"(";
+    os << '"';
+  }
+  os << "(";
   interleaveComma(op->getOperands(), [&](Value value) { printValueID(value); });
   os << ')';
 
   // For terminators, print the list of successors and their operands.
   if (op->getNumSuccessors() != 0) {
     os << '[';
     interleaveComma(op->getSuccessors(),
                     [&](Block *successor) { printBlockName(successor); });

mlir/lib/IR/FunctionImplementation.cpp

void mlir::function_like_impl::printFunctionLikeOp(OpAsmPrinter &p,
                                                    Operation *op,
                                                    ArrayRef<Type> argTypes,
                                                    bool isVariadic,
                                                    ArrayRef<Type> resultTypes) {
   // Print the operation and the function name.
   auto funcName =
       op->getAttrOfType<StringAttr>(SymbolTable::getSymbolAttrName())
           .getValue();
-  p << op->getName() << ' ';
+  p << ' ';
 
   StringRef visibilityAttrName = SymbolTable::getVisibilityAttrName();
   if (auto visibility = op->getAttrOfType<StringAttr>(visibilityAttrName))
     p << visibility.getValue() << ' ';
   p.printSymbolName(funcName);
 
   printFunctionSignature(p, op, argTypes, isVariadic, resultTypes);
   printFunctionAttributes(p, op, argTypes.size(), resultTypes.size(),
                           {visibilityAttrName});
   // Print the body if this is not an external function.
   Region &body = op->getRegion(0);
   if (!body.empty())
     p.printRegion(body, /*printEntryBlockArgs=*/false,
                   /*printBlockTerminators=*/true);
 }

mlir/lib/IR/Operation.cpp

 //===- Operation.cpp - Operation support code -----------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "mlir/IR/Operation.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpImplementation.h"
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/IR/TypeUtilities.h"
 #include "mlir/Interfaces/FoldInterfaces.h"
+#include "llvm/ADT/StringExtras.h"
 #include <numeric>
 
 using namespace mlir;
 
 OpAsmParser::~OpAsmParser() {}
 
 //===----------------------------------------------------------------------===//
 // OperationName

 
 // The fallback for the parser is to reject the custom assembly form.
 ParseResult OpState::parse(OpAsmParser &parser, OperationState &result) {
   return parser.emitError(parser.getNameLoc(), "has no custom assembly form");
 }
 
 // The fallback for the printer is to print in the generic assembly form.
 void OpState::print(Operation *op, OpAsmPrinter &p) { p.printGenericOp(op); }
+// The fallback for the printer is to print in the generic assembly form.
+void OpState::printOpName(Operation *op, OpAsmPrinter &p) {
+  StringRef name = op->getName().getStringRef();
+  if (name.startswith("std."))
+    name = name.drop_front(4);
+  llvm::printEscapedString(name, p.getStream());
+}
 
 /// Emit an error about fatal conditions with this operation, reporting up to
 /// any diagnostic handlers that may be listening.
 InFlightDiagnostic OpState::emitError(const Twine &message) {
   return getOperation()->emitError(message);
 }
 
 /// Emit an error with the op name prefixed, like "'dim' op " which is

 void impl::printOneResultOp(Operation *op, OpAsmPrinter &p) {
   assert(op->getNumResults() == 1 && "op should have one result");
 
   // If not all the operand and result types are the same, just use the
   // generic assembly form to avoid omitting information in printing.
   auto resultType = op->getResult(0).getType();
   if (llvm::any_of(op->getOperandTypes(),
                    [&](Type type) { return type != resultType; })) {
-    p.printGenericOp(op);
+    p.printGenericOp(op, /*printOpName=*/false);
     return;
   }
 
-  p << op->getName() << ' ';
+  p << ' ';
   p.printOperands(op->getOperands());
   p.printOptionalAttrDict(op->getAttrs());
   // Now we can output only one type for all operands and the result.
   p << " : " << resultType;
 }
 
 //===----------------------------------------------------------------------===//
 // CastOp implementation
return failure(parser.parseOperand(srcInfo) ||
                  parser.parseOptionalAttrDict(result.attributes) ||
                  parser.parseColonType(srcType) ||
                  parser.resolveOperand(srcInfo, srcType, result.operands) ||
                  parser.parseKeywordType("to", dstType) ||
                  parser.addTypeToList(dstType, result.types));
 }
 
 void impl::printCastOp(Operation *op, OpAsmPrinter &p) {
-  p << op->getName() << ' ' << op->getOperand(0);
+  p << ' ' << op->getOperand(0);
   p.printOptionalAttrDict(op->getAttrs());
   p << " : " << op->getOperand(0).getType() << " to "
     << op->getResult(0).getType();
 }
 
 Value impl::foldCastOp(Operation *op) {
   // Identity cast
   if (op->getOperand(0).getType() == op->getResult(0).getType())

mlir/test/Dialect/Standard/expand-ops.mlir

 // RUN: mlir-opt -std-expand %s -split-input-file | FileCheck %s
 
 // CHECK-LABEL: func @atomic_rmw_to_generic
 // CHECK-SAME: ([[F:%.*]]: memref<10xf32>, [[f:%.*]]: f32, [[i:%.*]]: index)
 func @atomic_rmw_to_generic(%F: memref<10xf32>, %f: f32, %i: index) -> f32 {
   %x = atomic_rmw maxf %f, %F[%i] : (f32, memref<10xf32>) -> f32
   return %x : f32
 }
-// CHECK: %0 = std.generic_atomic_rmw %arg0[%arg2] : memref<10xf32> {
+// CHECK: %0 = generic_atomic_rmw %arg0[%arg2] : memref<10xf32> {
 // CHECK: ^bb0([[CUR_VAL:%.*]]: f32):
 // CHECK:   [[CMP:%.*]] = cmpf ogt, [[CUR_VAL]], [[f]] : f32
 // CHECK:   [[SELECT:%.*]] = select [[CMP]], [[CUR_VAL]], [[f]] : f32
 // CHECK:   atomic_yield [[SELECT]] : f32
 // CHECK: }
 // CHECK: return %0 : f32
 
 // -----

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

TestDialect::getParseOperationHook(StringRef opName) const {
   return None;
 }
 
 function_ref<void(Operation *, OpAsmPrinter &)>
 TestDialect::getOperationPrinter(Operation *op) const {
   StringRef opName = op->getName().getStringRef();
   if (opName == "test.dialect_custom_printer") {
     return [](Operation *op, OpAsmPrinter &printer) {
-      printer.getStream() << op->getName().getStringRef() << " custom_format";
+      printer.getStream() << " custom_format";
     };
   }
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // TestBranchOp
 //===----------------------------------------------------------------------===//
static ParseResult parseParseIntegerLiteralOp(OpAsmParser &parser,
 
   IndexType type = parser.getBuilder().getIndexType();
   for (unsigned i = 0; i < numResults; ++i)
     result.addTypes(type);
   return success();
 }
 
 static void print(OpAsmPrinter &p, ParseIntegerLiteralOp op) {
-  p << ParseIntegerLiteralOp::getOperationName();
   if (unsigned numResults = op->getNumResults())
     p << " : " << numResults;
 }
 
 static ParseResult parseParseWrappedKeywordOp(OpAsmParser &parser,
                                               OperationState &result) {
   StringRef keyword;
   if (parser.parseKeyword(&keyword))
     return failure();
   result.addAttribute("keyword", parser.getBuilder().getStringAttr(keyword));
   return success();
 }
 
 static void print(OpAsmPrinter &p, ParseWrappedKeywordOp op) {
-  p << ParseWrappedKeywordOp::getOperationName() << " " << op.keyword();
+  p << " " << op.keyword();
 }
 
 //===----------------------------------------------------------------------===//
 // Test WrapRegionOp - wrapping op exercising `parseGenericOperation()`.
 
 static ParseResult parseWrappingRegionOp(OpAsmParser &parser,
                                          OperationState &result) {
   if (parser.parseKeyword("wraps"))
static ParseResult parseWrappingRegionOp(OpAsmParser &parser,
 
   // Use the location of the wrapped op for the "test.wrapping_region" op.
   result.location = wrapped_op->getLoc();
 
   return success();
 }
 
 static void print(OpAsmPrinter &p, WrappingRegionOp op) {
-  p << op.getOperationName() << " wraps ";
+  p << " wraps ";
   p.printGenericOp(&op.region().front().front());
 }
 
 //===----------------------------------------------------------------------===//
 // Test PolyForOp - parse list of region arguments.
 //===----------------------------------------------------------------------===//
 
 static ParseResult parsePolyForOp(OpAsmParser &parser, OperationState &result) {
static ParseResult parseStringAttrPrettyNameOp(OpAsmParser &parser,
   }
 
   auto namesAttr = parser.getBuilder().getStrArrayAttr(names);
   result.attributes.push_back({Identifier::get("names", context), namesAttr});
   return success();
 }
 
 static void print(OpAsmPrinter &p, StringAttrPrettyNameOp op) {
-  p << "test.string_attr_pretty_name";
-
   // Note that we only need to print the "name" attribute if the asmprinter
   // result name disagrees with it.  This can happen in strange cases, e.g.
   // when there are conflicts.
   bool namesDisagree = op.names().size() != op.getNumResults();
 
   SmallString<32> resultNameStr;
   for (size_t i = 0, e = op.getNumResults(); i != e && !namesDisagree; ++i) {
     resultNameStr.clear();
if (auto str = value[i].dyn_cast<StringAttr>())
         setNameFn(getResult(i), str.getValue());
 }
 
 //===----------------------------------------------------------------------===//
 // RegionIfOp
 //===----------------------------------------------------------------------===//
 
 static void print(OpAsmPrinter &p, RegionIfOp op) {
-  p << RegionIfOp::getOperationName() << " ";
+  p << " ";
   p.printOperands(op.getOperands());
   p << ": " << op.getOperandTypes();
   p.printArrowTypeList(op.getResultTypes());
   p << " then";
   p.printRegion(op.thenRegion(),
                 /*printEntryBlockArgs=*/true,
                 /*printBlockTerminators=*/true);
   p << " else";
static ParseResult parseSingleNoTerminatorCustomAsmOp(OpAsmParser &parser,
                                                       OperationState &state) {
   Region *body = state.addRegion();
   if (parser.parseRegion(*body, /*arguments=*/{}, /*argTypes=*/{}))
     return failure();
   return success();
 }
 
 static void print(SingleNoTerminatorCustomAsmOp op, OpAsmPrinter &printer) {
-  printer << op.getOperationName();
   printer.printRegion(
       op.getRegion(), /*printEntryBlockArgs=*/false,
       // This op has a single block without terminators. But explicitly mark
       // as not printing block terminators for testing.
       /*printBlockTerminators=*/false);
 }
 
 #include "TestOpEnums.cpp.inc"
 #include "TestOpInterfaces.cpp.inc"
 #include "TestOpStructs.cpp.inc"
 #include "TestTypeInterfaces.cpp.inc"
 
 #define GET_OP_CLASSES
 #include "TestOps.cpp.inc"

mlir/tools/mlir-tblgen/OpFormatGen.cpp

void OperationFormat::genElementPrinter(Element *element, OpMethodBody &body,
   }
 }
 
 void OperationFormat::genPrinter(Operator &op, OpClass &opClass) {
   auto *method =
       opClass.addMethodAndPrune("void", "print", "::mlir::OpAsmPrinter &p");
   auto &body = method->body();
 
-  // Emit the operation name, trimming the prefix if this is the standard
-  // dialect.
-  body << "  p << \"";
-  std::string opName = op.getOperationName();
-  if (op.getDialectName() == "std")
-    body << StringRef(opName).drop_front(4);
-  else
-    body << opName;
-  body << "\";\n";
-
   // Flags for if we should emit a space, and if the last element was
   // punctuation.
   bool shouldEmitSpace = true, lastWasPunctuation = false;
   for (auto &element : elements)
     genElementPrinter(element.get(), body, op, shouldEmitSpace,
                       lastWasPunctuation);
 }