[NVPTX][MLIR] Updated MmaOp verification and negative MMA tests. PR156040.

Author: kvederni

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

@@ -621,6 +621,9 @@ LogicalResult MmaOp::verify() {
     expectedB.emplace_back(unitB, multiplicandFragType);
     allowedShapes.push_back({16, 8, kFactor});
     allowedShapes.push_back({16, 8, kFactor * 2});
+
+    if (resultPtxType() != accumPtxType())
+      return emitOpError("ctype does not match dtype");
   }
 
   // In the M=8 case, there is only 1 possible case per data type.

mlir/test/Dialect/LLVMIR/invalid.mlir

@@ -749,6 +749,7 @@ llvm.func @nvvm_mma_m16n8k16_f32_f16(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
                                      %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                                      %c0 : vector<2xf16>, %c1 : vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32)> {
   // C and D should have the same type according to PTX ISA
+  // expected-error@+1 {{'nvvm.mma.sync' op ctype does not match dtype}}
   %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1]
     {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,
      shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(f32, f32, f32, f32)>
@@ -763,6 +764,7 @@ llvm.func @nvvm_mma_m16n8k16_f16_f32(%a0 : vector<2xf16>, %a1 : vector<2xf16>,
                                      %b0 : vector<2xf16>, %b1 : vector<2xf16>,
                                      %c0 : f32, %c1 : f32, %c2 : f32, %c3 : f32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> {
   // C and D should have the same type according to PTX ISA
+  // expected-error@+1 {{'nvvm.mma.sync' op ctype does not match dtype}}
   %0 = nvvm.mma.sync A[%a0, %a1, %a2, %a3] B[%b0, %b1] C[%c0, %c1, %c2, %c3]
     {layoutA = #nvvm.mma_layout<row>, layoutB = #nvvm.mma_layout<col>,
      shape = #nvvm.shape<m = 16, n = 8, k = 16>} : (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>