[mlir] Update various operations to declaratively specify their assembly format.

[rriddle]

Summary

This revision switches over many operations to use the declarative methods for defining the assembly specification. This updates operations in the NVVM, ROCDL, Standard, and VectorOps dialects.

Depends On D73406

[merge_guards_bot]

Unit tests: pass. 62197 tests passed, 0 failed and 815 were skipped.

clang-tidy: pass.

clang-format: fail. Please format your changes with clang-format by running git-clang-format HEAD^ or applying this patch.

Build artifacts: diff.json, clang-tidy.txt, clang-format.patch, CMakeCache.txt, console-log.txt, test-results.xml

Pre-merge checks is in beta. Report issue. Please join beta or enable it for your project.

[rriddle]

Rebase

[merge_guards_bot]

Unit tests: pass. 62252 tests passed, 0 failed and 816 were skipped.

clang-tidy: pass.

clang-format: fail. Please format your changes with clang-format by running git-clang-format HEAD^ or applying this patch.

Build artifacts: diff.json, clang-tidy.txt, clang-format.patch, CMakeCache.txt, console-log.txt, test-results.xml

Pre-merge checks is in beta. Report issue. Please join beta or enable it for your project.

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

 //===----------------------------------------------------------------------===//
 
 class NVVM_SpecialRegisterOp<string mnemonic,
     list<OpTrait> traits = []> :
   NVVM_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
   Results<(outs LLVM_Type:$res)>, Arguments<(ins)> {
   string llvmBuilder = "$res = createIntrinsicCall(builder,"
     # "llvm::Intrinsic::nvvm_" # !subst(".","_", mnemonic) # ");";
-  let parser = [{ return parseNVVMSpecialRegisterOp(parser, result); }];
-  let printer = [{ printNVVMIntrinsicOp(p, this->getOperation()); }];
+  let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 //===----------------------------------------------------------------------===//
 // Lane index and range
 def NVVM_LaneIdOp : NVVM_SpecialRegisterOp<"read.ptx.sreg.laneid">;
 def NVVM_WarpSizeOp : NVVM_SpecialRegisterOp<"read.ptx.sreg.warpsize">;
 
 //===----------------------------------------------------------------------===//

 //===----------------------------------------------------------------------===//
 // NVVM synchronization op definitions
 //===----------------------------------------------------------------------===//
 
 def NVVM_Barrier0Op : NVVM_Op<"barrier0"> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0);
   }];
-  let parser = [{ return success(); }];
-  let printer = [{ printNVVMIntrinsicOp(p, this->getOperation()); }];
+  let assemblyFormat = "attr-dict";
 }
 
 def NVVM_ShflBflyOp :
   NVVM_Op<"shfl.sync.bfly">,
   Results<(outs LLVM_Type:$res)>,
   Arguments<(ins LLVM_Type:$dst,
                  LLVM_Type:$val,
                  LLVM_Type:$offset,

 def NVVM_MmaOp :
   NVVM_Op<"mma.sync">,
   Results<(outs LLVM_Type:$res)>,
   Arguments<(ins Variadic<LLVM_Type>:$args)> {
   string llvmBuilder = [{
     $res = createIntrinsicCall(
         builder, llvm::Intrinsic::nvvm_mma_m8n8k4_row_row_f32_f32, $args);
   }];
-  let parser = [{ return parseNVVMMmaOp(parser, result); }];
-  let printer = [{ printNVVMMmaOp(p, *this); }];
+  let assemblyFormat = "$args attr-dict `:` functional-type($args, $res)";
   let verifier = [{ return ::verify(*this); }];
 }
 
 #endif // NVVMIR_OPS

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

 //===----------------------------------------------------------------------===//
 
 class ROCDL_SpecialRegisterOp<string mnemonic,
     list<OpTrait> traits = []> :
   ROCDL_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
   Results<(outs LLVM_Type:$res)>, Arguments<(ins)> {
   string llvmBuilder = "$res = createIntrinsicCall(builder,"
     # "llvm::Intrinsic::amdgcn_" # !subst(".","_", mnemonic) # ");";
-  let parser = [{ return parseROCDLOp(parser, result); }];
-  let printer = [{ printROCDLOp(p, this->getOperation()); }];
+  let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 class ROCDL_DeviceFunctionOp<string mnemonic, string device_function,
                              int parameter, list<OpTrait> traits = []> :
   ROCDL_Op<mnemonic, !listconcat(traits, [NoSideEffect])>,
   Results<(outs LLVM_Type:$res)>, Arguments<(ins)> {
   string llvmBuilder = "$res = createDeviceFunctionCall(builder, \""
   # device_function # "\", " # parameter # ");";
-  let parser = [{ return parseROCDLOp(parser, result); }];
-  let printer = [{ printROCDLOp(p, this->getOperation()); }];
+  let assemblyFormat = "attr-dict `:` type($res)";
 }
 
 //===----------------------------------------------------------------------===//
 // Thread index and Block index
 
 def ROCDL_ThreadIdXOp : ROCDL_SpecialRegisterOp<"workitem.id.x">;
 def ROCDL_ThreadIdYOp : ROCDL_SpecialRegisterOp<"workitem.id.y">;
 def ROCDL_ThreadIdZOp : ROCDL_SpecialRegisterOp<"workitem.id.z">;

mlir/include/mlir/Dialect/StandardOps/Ops.td

let extraClassDeclaration = [{
     operand_iterator arg_operand_begin() { return operand_begin(); }
     operand_iterator arg_operand_end() { return operand_end(); }
 
     /// Return the callee of this operation.
     CallInterfaceCallable getCallableForCallee() {
       return getAttrOfType<SymbolRefAttr>("callee");
     }
   }];
+
+  let assemblyFormat = [{
+    $callee `(` $operands `)` attr-dict `:` functional-type($operands, results)
+  }];
 }
 
 def CallIndirectOp : Std_Op<"call_indirect", [CallOpInterface]> {
   let summary = "indirect call operation";
   let description = [{
     The "call_indirect" operation represents an indirect call to a value of
     function type.  Functions are first class types in MLIR, and may be passed
     as arguments and merged together with block arguments.  The operands
operations).
       %0 = alloc() : memref<8x64xf32, (d0, d1) -> (d0, d1), 1>
       dealloc %0 : memref<8x64xf32, (d0, d1) -> (d0, d1), 1>
   }];
 
   let arguments = (ins AnyMemRef:$memref);
 
   let hasCanonicalizer = 1;
   let hasFolder = 1;
+  let assemblyFormat = "$memref attr-dict `:` type($memref)";
 }
 
 def DimOp : Std_Op<"dim", [NoSideEffect]> {
   let summary = "dimension index operation";
   let description = [{
     The "dim" operation takes a memref or tensor operand and returns an "index".
     It requires a single integer attribute named "index". It returns the size
     of the specified dimension. For example:
def RankOp : Std_Op<"rank", [NoSideEffect]> {
 
   let builders = [OpBuilder<
     "Builder *builder, OperationState &result, Value tensor", [{
       auto indexType = builder->getIndexType();
       build(builder, result, indexType, tensor);
     }]>];
 
   let hasFolder = 1;
+  let assemblyFormat = "operands attr-dict `:` type(operands)";
 }
 
 def RemFOp : FloatArithmeticOp<"remf"> {
   let summary = "floating point division remainder operation";
 }
 
 def SignedRemIOp : IntArithmeticOp<"remi_signed"> {
   let summary = "signed integer division remainder operation";

mlir/include/mlir/Dialect/VectorOps/VectorOps.td

let description = [{
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
   }];
+  let assemblyFormat = "$source attr-dict `:` type($source) `to` type($vector)";
 }
 
 def Vector_ShuffleOp :
   Vector_Op<"shuffle", [NoSideEffect,
      PredOpTrait<"first operand v1 and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>,
      PredOpTrait<"second operand v2 and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 1>>]>,
let extraClassDeclaration = [{
     TupleType getResultTupleType() {
       return getResult().getType().cast<TupleType>();
     }
     void getSizes(SmallVectorImpl<int64_t> &results);
     void getStrides(SmallVectorImpl<int64_t> &results);
     static StringRef getSizesAttrName() { return "sizes"; }
     static StringRef getStridesAttrName() { return "strides"; }
   }];
+  let assemblyFormat = [{
+    $vector `,` $sizes `,` $strides attr-dict `:` type($vector) `into`
+    type(results)
+  }];
 }
 
 def Vector_InsertElementOp :
   Vector_Op<"insertelement", [NoSideEffect,
      PredOpTrait<"source operand and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>,
      PredOpTrait<"dest operand and result have same type",
                  TCresIsSameAsOpBase<0, 1>>]>,
let extraClassDeclaration = [{
     VectorType getResultVectorType() {
       return getResult().getType().cast<VectorType>();
     }
     void getSizes(SmallVectorImpl<int64_t> &results);
     void getStrides(SmallVectorImpl<int64_t> &results);
     static StringRef getSizesAttrName() { return "sizes"; }
     static StringRef getStridesAttrName() { return "strides"; }
   }];
+  let assemblyFormat = [{
+    $vectors `,` $sizes `,` $strides attr-dict `:` type($vectors) `into`
+    type(results)
+  }];
 }
 
 def Vector_InsertStridedSliceOp :
   Vector_Op<"insert_strided_slice", [NoSideEffect,
     PredOpTrait<"operand #0 and result have same element type",
                  TCresVTEtIsSameAsOpBase<0, 0>>,
     PredOpTrait<"dest operand and result have same type",
                  TCresIsSameAsOpBase<0, 1>>]>,
def Vector_StridedSliceOp :
   let extraClassDeclaration = [{
     static StringRef getOffsetsAttrName() { return "offsets"; }
     static StringRef getSizesAttrName() { return "sizes"; }
     static StringRef getStridesAttrName() { return "strides"; }
     VectorType getVectorType(){ return vector().getType().cast<VectorType>(); }
     void getOffsets(SmallVectorImpl<int64_t> &results);
   }];
   let hasCanonicalizer = 1;
+  let assemblyFormat = "$vector attr-dict `:` type($vector) `to` type(results)";
 }
 
 def Vector_TransferReadOp :
   Vector_Op<"transfer_read">,
     Arguments<(ins AnyMemRef:$memref, Variadic<Index>:$indices,
                AffineMapAttr:$permutation_map, AnyType:$padding)>,
     Results<(outs AnyVector:$vector)> {
 
def Vector_TransferWriteOp :
   let extraClassDeclaration = [{
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
     }
     MemRefType getMemRefType() {
       return memref().getType().cast<MemRefType>();
     }
   }];
+  let assemblyFormat = [{
+    $vector `,` $memref `[` $indices `]` attr-dict `:` type($vector) `,`
+    type($memref)
+  }];
 }
 
 def Vector_TypeCastOp :
   Vector_Op<"type_cast", [NoSideEffect]>,
     Arguments<(ins StaticShapeMemRefOf<[AnyType]>:$memref)>,
     Results<(outs AnyMemRef)> {
   let summary = "type_cast op converts a scalar memref to a vector memref";
   let description = [{
Example: create a constant vector mask of size 4x3xi1 with elements in range
         rows  1 | 1    1    0
               2 | 1    1    0
               3 | 0    0    0
   }];
 
   let extraClassDeclaration = [{
     static StringRef getMaskDimSizesAttrName() { return "mask_dim_sizes"; }
   }];
+  let assemblyFormat = "$mask_dim_sizes attr-dict `:` type(results)";
 }
 
 def Vector_CreateMaskOp :
   Vector_Op<"create_mask", [NoSideEffect]>,
     Arguments<(ins Variadic<Index>:$operands)>, Results<(outs VectorOf<[I1]>)> {
   let summary = "creates a vector mask";
   let description = [{
     Creates and returns a vector mask where elements of the result vector
Example: create a vector mask of size 4x3xi1 where elements in range
                 |------------
               0 | 1    1    0
         rows  1 | 1    1    0
               2 | 1    1    0
               3 | 0    0    0
   }];
 
   let hasCanonicalizer = 1;
+  let assemblyFormat = "$operands attr-dict `:` type(results)";
 }
 
 def Vector_TupleOp :
   Vector_Op<"tuple", [NoSideEffect]>,
     Arguments<(ins Variadic<AnyVector>:$vectors)>,
     Results<(outs TupleOf<[AnyVector]>)> {
   let summary = "make tuple of vectors operation";
   let description = [{
let description = [{
     ```
   }];
   let verifier = ?;
   let extraClassDeclaration = [{
     Type getPrintType() {
       return source().getType();
     }
   }];
+  let assemblyFormat = "$source attr-dict `:` type($source)";
 }
 
 #endif // VECTOR_OPS

mlir/include/mlir/IR/OpBase.td

!if(!eq(description, ""),
         StrJoin<!foreach(t, allowedTypes, t.description), " or ">.result,
         description)>;
 
 // Integer types.
 // Any integer type irrespective of its width.
 def AnyInteger : Type<CPred<"$_self.isa<IntegerType>()">, "integer">;
 
 // Index type.
-def Index : Type<CPred<"$_self.isa<IndexType>()">, "index">;
+def Index : Type<CPred<"$_self.isa<IndexType>()">, "index">,
+            BuildableType<"getIndexType()">;
 
 // Integer type of a specific width.
 class I<int width>
     : Type<CPred<"$_self.isInteger(" # width # ")">,
                   width # "-bit integer">,
       BuildableType<"getIntegerType(" # width # ")"> {
   int bitwidth = width;
 }

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

 //===----------------------------------------------------------------------===//
 
 static void printNVVMIntrinsicOp(OpAsmPrinter &p, Operation *op) {
   p << op->getName() << " " << op->getOperands();
   if (op->getNumResults() > 0)
     p << " : " << op->getResultTypes();
 }
 
-// <operation> ::= `llvm.nvvm.XYZ` : type
-static ParseResult parseNVVMSpecialRegisterOp(OpAsmParser &parser,
-                                              OperationState &result) {
-  Type type;
-  if (parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonType(type))
-    return failure();
-
-  result.addTypes(type);
-  return success();
-}
-
 static LLVM::LLVMDialect *getLlvmDialect(OpAsmParser &parser) {
   return parser.getBuilder()
       .getContext()
       ->getRegisteredDialect<LLVM::LLVMDialect>();
 }
 
 // <operation> ::=
 //     `llvm.nvvm.shfl.sync.bfly %dst, %val, %offset, %clamp_and_mask`
static ParseResult parseNVVMVoteBallotOp(OpAsmParser &parser,
   return failure(parser.parseOperandList(ops) ||
                  parser.parseOptionalAttrDict(result.attributes) ||
                  parser.parseColonType(type) ||
                  parser.addTypeToList(type, result.types) ||
                  parser.resolveOperands(ops, {int32Ty, int1Ty},
                                         parser.getNameLoc(), result.operands));
 }
 
-// <operation> ::= `llvm.nvvm.mma.sync %lhs... %rhs... %acc...`
-//                 : signature_type
-static ParseResult parseNVVMMmaOp(OpAsmParser &parser, OperationState &result) {
-  SmallVector<OpAsmParser::OperandType, 12> ops;
-  Type type;
-  llvm::SMLoc typeLoc;
-  if (parser.parseOperandList(ops) ||
-      parser.parseOptionalAttrDict(result.attributes) || parser.parseColon() ||
-      parser.getCurrentLocation(&typeLoc) || parser.parseType(type)) {
-    return failure();
-  }
-
-  auto signature = type.dyn_cast<FunctionType>();
-  if (!signature) {
-    return parser.emitError(
-        typeLoc, "expected the type to be the full list of input and output");
-  }
-
-  if (signature.getNumResults() != 1) {
-    return parser.emitError(typeLoc, "expected single result");
-  }
-
-  return failure(parser.addTypeToList(signature.getResult(0), result.types) ||
-                 parser.resolveOperands(ops, signature.getInputs(),
-                                        parser.getNameLoc(), result.operands));
-}
-
-static void printNVVMMmaOp(OpAsmPrinter &p, MmaOp &op) {
-  p << op.getOperationName() << " " << op.getOperands();
-  p.printOptionalAttrDict(op.getAttrs());
-  p << " : "
-    << FunctionType::get(llvm::to_vector<12>(op.getOperandTypes()),
-                         op.getType(), op.getContext());
-}
-
 static LogicalResult verify(MmaOp op) {
   auto dialect = op.getContext()->getRegisteredDialect<LLVM::LLVMDialect>();
   auto f16Ty = LLVM::LLVMType::getHalfTy(dialect);
   auto f16x2Ty = LLVM::LLVMType::getVectorTy(f16Ty, 2);
   auto f32Ty = LLVM::LLVMType::getFloatTy(dialect);
   auto f16x2x4StructTy = LLVM::LLVMType::getStructTy(
       dialect, {f16x2Ty, f16x2Ty, f16x2Ty, f16x2Ty});
   auto f32x8StructTy = LLVM::LLVMType::getStructTy(

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

 #include "llvm/IR/Function.h"
 #include "llvm/IR/Type.h"
 #include "llvm/Support/SourceMgr.h"
 
 using namespace mlir;
 using namespace ROCDL;
 
 //===----------------------------------------------------------------------===//
-// Printing/parsing for ROCDL ops
-//===----------------------------------------------------------------------===//
-
-static void printROCDLOp(OpAsmPrinter &p, Operation *op) {
-  p << op->getName() << " " << op->getOperands();
-  if (op->getNumResults() > 0)
-    p << " : " << op->getResultTypes();
-}
-
-// <operation> ::= `rocdl.XYZ` : type
-static ParseResult parseROCDLOp(OpAsmParser &parser, OperationState &result) {
-  Type type;
-  return failure(parser.parseOptionalAttrDict(result.attributes) ||
-                 parser.parseColonType(type) ||
-                 parser.addTypeToList(type, result.types));
-}
-
-//===----------------------------------------------------------------------===//
 // ROCDLDialect initialization, type parsing, and registration.
 //===----------------------------------------------------------------------===//
 
 // TODO(herhut): This should be the llvm.rocdl dialect once this is supported.
 ROCDLDialect::ROCDLDialect(MLIRContext *context) : Dialect("rocdl", context) {
   addOperations<
 #define GET_OP_LIST
 #include "mlir/Dialect/LLVMIR/ROCDLOps.cpp.inc"

mlir/lib/Dialect/StandardOps/Ops.cpp

void BranchOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                            MLIRContext *context) {
   results.insert<SimplifyBrToBlockWithSinglePred>(context);
 }
 
 //===----------------------------------------------------------------------===//
 // CallOp
 //===----------------------------------------------------------------------===//
 
-static ParseResult parseCallOp(OpAsmParser &parser, OperationState &result) {
-  FlatSymbolRefAttr calleeAttr;
-  FunctionType calleeType;
-  SmallVector<OpAsmParser::OperandType, 4> operands;
-  auto calleeLoc = parser.getNameLoc();
-  if (parser.parseAttribute(calleeAttr, "callee", result.attributes) ||
-      parser.parseOperandList(operands, OpAsmParser::Delimiter::Paren) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonType(calleeType) ||
-      parser.addTypesToList(calleeType.getResults(), result.types) ||
-      parser.resolveOperands(operands, calleeType.getInputs(), calleeLoc,
-                             result.operands))
-    return failure();
-
-  return success();
-}
-
-static void print(OpAsmPrinter &p, CallOp op) {
-  p << "call " << op.getAttr("callee") << '(' << op.getOperands() << ')';
-  p.printOptionalAttrDict(op.getAttrs(), /*elidedAttrs=*/{"callee"});
-  p << " : " << op.getCalleeType();
-}
-
 static LogicalResult verify(CallOp op) {
   // Check that the callee attribute was specified.
   auto fnAttr = op.getAttrOfType<FlatSymbolRefAttr>("callee");
   if (!fnAttr)
     return op.emitOpError("requires a 'callee' symbol reference attribute");
   auto fn =
       op.getParentOfType<ModuleOp>().lookupSymbol<FuncOp>(fnAttr.getValue());
   if (!fn)
PatternMatchResult matchAndRewrite(DeallocOp dealloc,
 
     // Erase the dealloc operation.
     rewriter.eraseOp(dealloc);
     return matchSuccess();
   }
 };
 } // end anonymous namespace.
 
-static void print(OpAsmPrinter &p, DeallocOp op) {
-  p << "dealloc " << op.memref() << " : " << op.memref().getType();
-}
-
-static ParseResult parseDeallocOp(OpAsmParser &parser, OperationState &result) {
-  OpAsmParser::OperandType memrefInfo;
-  MemRefType type;
-
-  return failure(parser.parseOperand(memrefInfo) ||
-                 parser.parseColonType(type) ||
-                 parser.resolveOperand(memrefInfo, type, result.operands));
-}
-
 static LogicalResult verify(DeallocOp op) {
   if (!op.memref().getType().isa<MemRefType>())
     return op.emitOpError("operand must be a memref");
   return success();
 }
 
 void DeallocOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
                                             MLIRContext *context) {
LogicalResult PrefetchOp::fold(ArrayRef<Attribute> cstOperands,
   // prefetch(memrefcast) -> prefetch
   return foldMemRefCast(*this);
 }
 
 //===----------------------------------------------------------------------===//
 // RankOp
 //===----------------------------------------------------------------------===//
 
-static void print(OpAsmPrinter &p, RankOp op) {
-  p << "rank " << op.getOperand() << " : " << op.getOperand().getType();
-}
-
-static ParseResult parseRankOp(OpAsmParser &parser, OperationState &result) {
-  OpAsmParser::OperandType operandInfo;
-  Type type;
-  Type indexType = parser.getBuilder().getIndexType();
-  return failure(parser.parseOperand(operandInfo) ||
-                 parser.parseColonType(type) ||
-                 parser.resolveOperand(operandInfo, type, result.operands) ||
-                 parser.addTypeToList(indexType, result.types));
-}
-
 OpFoldResult RankOp::fold(ArrayRef<Attribute> operands) {
   // Constant fold rank when the rank of the tensor is known.
   auto type = getOperand().getType();
   if (auto tensorType = type.dyn_cast<RankedTensorType>())
     return IntegerAttr::get(IndexType::get(getContext()), tensorType.getRank());
   return IntegerAttr();
 }
 

mlir/lib/Dialect/VectorOps/VectorOps.cpp

void ExtractSlicesOp::build(Builder *builder, OperationState &result,
   result.addOperands(vector);
   auto sizesAttr = getVectorSubscriptAttr(*builder, sizes);
   auto stridesAttr = getVectorSubscriptAttr(*builder, strides);
   result.addTypes(tupleType);
   result.addAttribute(getSizesAttrName(), sizesAttr);
   result.addAttribute(getStridesAttrName(), stridesAttr);
 }
 
-static ParseResult parseExtractSlicesOp(OpAsmParser &parser,
-                                        OperationState &result) {
-  OpAsmParser::OperandType operandInfo;
-  ArrayAttr sizesAttr;
-  StringRef sizesAttrName = ExtractSlicesOp::getSizesAttrName();
-  ArrayAttr stridesAttr;
-  StringRef stridesAttrName = ExtractSlicesOp::getStridesAttrName();
-  VectorType vectorType;
-  TupleType resultTupleType;
-  return failure(
-      parser.parseOperand(operandInfo) || parser.parseComma() ||
-      parser.parseAttribute(sizesAttr, sizesAttrName, result.attributes) ||
-      parser.parseComma() ||
-      parser.parseAttribute(stridesAttr, stridesAttrName, result.attributes) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonType(vectorType) ||
-      parser.parseKeywordType("into", resultTupleType) ||
-      parser.resolveOperand(operandInfo, vectorType, result.operands) ||
-      parser.addTypeToList(resultTupleType, result.types));
-}
-
-static void print(OpAsmPrinter &p, ExtractSlicesOp op) {
-  p << op.getOperationName() << ' ' << op.vector() << ", ";
-  p << op.sizes() << ", " << op.strides();
-  p.printOptionalAttrDict(
-      op.getAttrs(),
-      /*elidedAttrs=*/{ExtractSlicesOp::getSizesAttrName(),
-                       ExtractSlicesOp::getStridesAttrName()});
-  p << " : " << op.vector().getType();
-  p << " into " << op.getResultTupleType();
-}
-
 static LogicalResult
 isValidExtractOrInsertSlicesType(Operation *op, VectorType vectorType,
                                  TupleType tupleType, ArrayRef<int64_t> sizes,
                                  ArrayRef<int64_t> strides) {
   // Check for non-unit strides.
   // TODO(b/144845578) Support non-1 strides.
   if (llvm::any_of(strides, [](int64_t s) { return s != 1; }))
     return op->emitError("requires unit strides");

 void ExtractSlicesOp::getStrides(SmallVectorImpl<int64_t> &results) {
   populateFromInt64AttrArray(strides(), results);
 }
 
 //===----------------------------------------------------------------------===//
 // BroadcastOp
 //===----------------------------------------------------------------------===//
 
-static void print(OpAsmPrinter &p, BroadcastOp op) {
-  p << op.getOperationName() << " " << op.source() << " : "
-    << op.getSourceType() << " to " << op.getVectorType();
-}
-
 static LogicalResult verify(BroadcastOp op) {
   VectorType srcVectorType = op.getSourceType().dyn_cast<VectorType>();
   VectorType dstVectorType = op.getVectorType();
   // Scalar to vector broadcast is always valid. A vector
   // to vector broadcast needs some additional checking.
   if (srcVectorType) {
     int64_t srcRank = srcVectorType.getRank();
     int64_t dstRank = dstVectorType.getRank();
     if (srcRank > dstRank)
       return op.emitOpError("source rank higher than destination rank");
     // Source has an exact match or singleton value for all trailing dimensions
     // (all leading dimensions are simply duplicated).
     int64_t lead = dstRank - srcRank;
     for (int64_t r = 0; r < srcRank; ++r) {
       int64_t srcDim = srcVectorType.getDimSize(r);
       int64_t dstDim = dstVectorType.getDimSize(lead + r);
       if (srcDim != 1 && srcDim != dstDim)
         return op.emitOpError("dimension mismatch (")
                << srcDim << " vs. " << dstDim << ")";
     }
   }
   return success();
 }
 
-static ParseResult parseBroadcastOp(OpAsmParser &parser,
-                                    OperationState &result) {
-  OpAsmParser::OperandType source;
-  Type sourceType;
-  VectorType vectorType;
-  return failure(parser.parseOperand(source) ||
-                 parser.parseColonType(sourceType) ||
-                 parser.parseKeywordType("to", vectorType) ||
-                 parser.resolveOperand(source, sourceType, result.operands) ||
-                 parser.addTypeToList(vectorType, result.types));
-}
-
 //===----------------------------------------------------------------------===//
 // ShuffleOp
 //===----------------------------------------------------------------------===//
 
 void ShuffleOp::build(Builder *builder, OperationState &result, Value v1,
                       Value v2, ArrayRef<int64_t> mask) {
   result.addOperands({v1, v2});
   auto maskAttr = getVectorSubscriptAttr(*builder, mask);
static LogicalResult verify(InsertOp op) {
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // InsertSlicesOp
 //===----------------------------------------------------------------------===//
 
-static ParseResult parseInsertSlicesOp(OpAsmParser &parser,
-                                       OperationState &result) {
-  OpAsmParser::OperandType operandInfo;
-  ArrayAttr sizesAttr;
-  StringRef sizesAttrName = InsertSlicesOp::getSizesAttrName();
-  ArrayAttr stridesAttr;
-  StringRef stridesAttrName = InsertSlicesOp::getStridesAttrName();
-  TupleType tupleType;
-  VectorType resultVectorType;
-  return failure(
-      parser.parseOperand(operandInfo) || parser.parseComma() ||
-      parser.parseAttribute(sizesAttr, sizesAttrName, result.attributes) ||
-      parser.parseComma() ||
-      parser.parseAttribute(stridesAttr, stridesAttrName, result.attributes) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonType(tupleType) ||
-      parser.parseKeywordType("into", resultVectorType) ||
-      parser.resolveOperand(operandInfo, tupleType, result.operands) ||
-      parser.addTypeToList(resultVectorType, result.types));
-}
-
-static void print(OpAsmPrinter &p, InsertSlicesOp op) {
-  p << op.getOperationName() << ' ' << op.vectors() << ", ";
-  p << op.sizes() << ", " << op.strides();
-  p.printOptionalAttrDict(
-      op.getAttrs(),
-      /*elidedAttrs=*/{InsertSlicesOp::getSizesAttrName(),
-                       InsertSlicesOp::getStridesAttrName()});
-  p << " : " << op.vectors().getType();
-  p << " into " << op.getResultVectorType();
-}
-
 static LogicalResult verify(InsertSlicesOp op) {
   SmallVector<int64_t, 4> sizes;
   op.getSizes(sizes);
   SmallVector<int64_t, 4> strides;
   op.getStrides(strides);
   return isValidExtractOrInsertSlicesType(
       op.getOperation(), op.getResultVectorType(), op.getSourceTupleType(),
       sizes, strides);
void StridedSliceOp::build(Builder *builder, OperationState &result,
   result.addTypes(
       inferStridedSliceOpResultType(source.getType().cast<VectorType>(),
                                     offsetsAttr, sizesAttr, stridesAttr));
   result.addAttribute(getOffsetsAttrName(), offsetsAttr);
   result.addAttribute(getSizesAttrName(), sizesAttr);
   result.addAttribute(getStridesAttrName(), stridesAttr);
 }
 
-static void print(OpAsmPrinter &p, StridedSliceOp op) {
-  p << op.getOperationName() << " " << op.vector();
-  p.printOptionalAttrDict(op.getAttrs());
-  p << " : " << op.vector().getType() << " to " << op.getResult().getType();
-}
-
-static ParseResult parseStridedSliceOp(OpAsmParser &parser,
-                                       OperationState &result) {
-  llvm::SMLoc attributeLoc, typeLoc;
-  OpAsmParser::OperandType vector;
-  VectorType vectorType, resultVectorType;
-  return failure(parser.parseOperand(vector) ||
-                 parser.getCurrentLocation(&attributeLoc) ||
-                 parser.parseOptionalAttrDict(result.attributes) ||
-                 parser.getCurrentLocation(&typeLoc) ||
-                 parser.parseColonType(vectorType) ||
-                 parser.parseKeywordType("to", resultVectorType) ||
-                 parser.resolveOperand(vector, vectorType, result.operands) ||
-                 parser.addTypeToList(resultVectorType, result.types));
-}
-
 static LogicalResult verify(StridedSliceOp op) {
   auto type = op.getVectorType();
   auto offsets = op.offsets();
   auto sizes = op.sizes();
   auto strides = op.strides();
   if (offsets.size() != sizes.size() || offsets.size() != strides.size()) {
     op.emitOpError(
         "expected offsets, sizes and strides attributes of same size");
static LogicalResult verify(TransferReadOp op) {
 
   return verifyPermutationMap(permutationMap,
                               [&op](Twine t) { return op.emitOpError(t); });
 }
 
 //===----------------------------------------------------------------------===//
 // TransferWriteOp
 //===----------------------------------------------------------------------===//
-static void print(OpAsmPrinter &p, TransferWriteOp op) {
-  p << op.getOperationName() << " " << op.vector() << ", " << op.memref() << "["
-    << op.indices() << "]";
-  p.printOptionalAttrDict(op.getAttrs());
-  p << " : " << op.getVectorType() << ", " << op.getMemRefType();
-}
-
-static ParseResult parseTransferWriteOp(OpAsmParser &parser,
-                                        OperationState &result) {
-  llvm::SMLoc typesLoc;
-  OpAsmParser::OperandType storeValueInfo;
-  OpAsmParser::OperandType memRefInfo;
-  SmallVector<OpAsmParser::OperandType, 4> indexInfo;
-  SmallVector<Type, 2> types;
-  if (parser.parseOperand(storeValueInfo) || parser.parseComma() ||
-      parser.parseOperand(memRefInfo) ||
-      parser.parseOperandList(indexInfo, OpAsmParser::Delimiter::Square) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.getCurrentLocation(&typesLoc) || parser.parseColonTypeList(types))
-    return failure();
-  if (types.size() != 2)
-    return parser.emitError(typesLoc, "two types required");
-  auto indexType = parser.getBuilder().getIndexType();
-  Type vectorType = types[0], memRefType = types[1];
-  return failure(
-      parser.resolveOperand(storeValueInfo, vectorType, result.operands) ||
-      parser.resolveOperand(memRefInfo, memRefType, result.operands) ||
-      parser.resolveOperands(indexInfo, indexType, result.operands));
-}
 
 static LogicalResult verify(TransferWriteOp op) {
   // Consistency of elemental types in memref and vector.
   MemRefType memrefType = op.getMemRefType();
   VectorType vectorType = op.getVectorType();
   auto permutationMap = op.permutation_map();
 
   if (llvm::size(op.indices()) != memrefType.getRank())
if (auto tupleOp = dyn_cast_or_null<TupleOp>(getOperand().getDefiningOp()))
     return tupleOp.getOperand(getIndex());
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // ConstantMaskOp
 //===----------------------------------------------------------------------===//
 
-static ParseResult parseConstantMaskOp(OpAsmParser &parser,
-                                       OperationState &result) {
-  Type resultType;
-  ArrayAttr maskDimSizesAttr;
-  StringRef attrName = ConstantMaskOp::getMaskDimSizesAttrName();
-  return failure(
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseAttribute(maskDimSizesAttr, attrName, result.attributes) ||
-      parser.parseColonType(resultType) ||
-      parser.addTypeToList(resultType, result.types));
-}
-
-static void print(OpAsmPrinter &p, ConstantMaskOp op) {
-  p << op.getOperationName() << ' ' << op.mask_dim_sizes() << " : "
-    << op.getResult().getType();
-}
-
 static LogicalResult verify(ConstantMaskOp &op) {
   // Verify that array attr size matches the rank of the vector result.
   auto resultType = op.getResult().getType().cast<VectorType>();
   if (static_cast<int64_t>(op.mask_dim_sizes().size()) != resultType.getRank())
     return op.emitOpError(
         "must specify array attr of size equal vector result rank");
   // Verify that each array attr element is in bounds of corresponding vector
   // result dimension size.
return op.emitOpError("expected all mask dim sizes to be zeros, "
                           "as a result of conjunction with zero mask dim");
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // CreateMaskOp
 //===----------------------------------------------------------------------===//
 
-static ParseResult parseCreateMaskOp(OpAsmParser &parser,
-                                     OperationState &result) {
-  auto indexType = parser.getBuilder().getIndexType();
-  Type resultType;
-  SmallVector<OpAsmParser::OperandType, 4> operandInfo;
-  return failure(
-      parser.parseOperandList(operandInfo) ||
-      parser.parseOptionalAttrDict(result.attributes) ||
-      parser.parseColonType(resultType) ||
-      parser.resolveOperands(operandInfo, indexType, result.operands) ||
-      parser.addTypeToList(resultType, result.types));
-}
-
-static void print(OpAsmPrinter &p, CreateMaskOp op) {
-  p << op.getOperationName() << ' ' << op.operands() << " : " << op.getType();
-}
-
 static LogicalResult verify(CreateMaskOp op) {
   // Verify that an operand was specified for each result vector each dimension.
   if (op.getNumOperands() !=
       op.getResult().getType().cast<VectorType>().getRank())
     return op.emitOpError(
         "must specify an operand for each result vector dimension");
   return success();
 }
 
-//===----------------------------------------------------------------------===//
-// PrintOp
-//===----------------------------------------------------------------------===//
-
-static ParseResult parsePrintOp(OpAsmParser &parser, OperationState &result) {
-  OpAsmParser::OperandType source;
-  Type sourceType;
-  return failure(parser.parseOperand(source) ||
-                 parser.parseColonType(sourceType) ||
-                 parser.resolveOperand(source, sourceType, result.operands));
-}
-
-static void print(OpAsmPrinter &p, PrintOp op) {
-  p << op.getOperationName() << ' ' << op.source() << " : "
-    << op.getPrintType();
-}
-
 namespace {
 
 // Pattern to rewrite a CreateMaskOp with a ConstantMaskOp.
 class CreateMaskFolder final : public OpRewritePattern<CreateMaskOp> {
 public:
   using OpRewritePattern<CreateMaskOp>::OpRewritePattern;
 
   PatternMatchResult matchAndRewrite(CreateMaskOp createMaskOp,

mlir/test/Dialect/LLVMIR/invalid.mlir

 
 // -----
 
 // CHECK-LABEL: @nvvm_invalid_mma_6
 func @nvvm_invalid_mma_6(%a0 : !llvm<"<2 x half>">, %a1 : !llvm<"<2 x half>">,
                          %b0 : !llvm<"<2 x half>">, %b1 : !llvm<"<2 x half>">,
                          %c0 : !llvm.float, %c1 : !llvm.float, %c2 : !llvm.float, %c3 : !llvm.float,
                          %c4 : !llvm.float, %c5 : !llvm.float, %c6 : !llvm.float, %c7 : !llvm.float) {
-  // expected-error@+1 {{expected the type to be the full list of input and output}}
+  // expected-error@+1 {{invalid kind of type specified}}
   %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : !llvm<"{ float, float, float, float, float, float, float, float }">
   llvm.return %0 : !llvm<"{ float, float, float, float, float, float, float, float }">
 }
 
 // -----
 
 // CHECK-LABEL: @nvvm_invalid_mma_7
 func @nvvm_invalid_mma_7(%a0 : !llvm<"<2 x half>">, %a1 : !llvm<"<2 x half>">,
                          %b0 : !llvm<"<2 x half>">, %b1 : !llvm<"<2 x half>">,
                          %c0 : !llvm.float, %c1 : !llvm.float, %c2 : !llvm.float, %c3 : !llvm.float,
                          %c4 : !llvm.float, %c5 : !llvm.float, %c6 : !llvm.float, %c7 : !llvm.float) {
-  // expected-error@+1 {{expected single result}}
-  %0 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : (!llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float) -> (!llvm<"{ float, float, float, float, float, float, float, float }">, !llvm.i32)
-  llvm.return %0 : (!llvm<"{ float, float, float, float, float, float, float, float }">, !llvm.i32)
+  // expected-error@+1 {{op requires one result}}
+  %0:2 = nvvm.mma.sync %a0, %a1, %b0, %b1, %c0, %c1, %c2, %c3, %c4, %c5, %c6, %c7 {alayout="col", blayout="row"} : (!llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm<"<2 x half>">, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float, !llvm.float) -> (!llvm<"{ float, float, float, float, float, float, float, float }">, !llvm.i32)
+  llvm.return %0#0 : !llvm<"{ float, float, float, float, float, float, float, float }">
 }
 
 // -----
 
 // FIXME: the LLVM-IR dialect should parse mutually recursive types
 // CHECK-LABEL: @recursive_type
 // expected-error@+1 {{expected end of string}}
 llvm.func @recursive_type(%a : !llvm<"%a = type { %a* }">) ->