[Mlir-commits] [mlir] [MLIR][NVVM] Add Ops for tcgen05 cp and shift (PR #127798)

[Guray Ozen]

================
@@ -1107,6 +1111,38 @@ LogicalResult NVVM::BarrierOp::verify() {
   return success();
 }
 
+LogicalResult NVVM::Tcgen05CpOp::verify() {
+  auto mc = getMulticast();
+
+  using SH = Tcgen05CpShape;
+  using MC = Tcgen05CpMulticast;
+  switch (getShape()) {
+  case SH::SHAPE_128x256b:
+  case SH::SHAPE_128x128b:
+  case SH::SHAPE_4x256b:
+    if (mc != MC::NONE)
+      return emitError("Invalid multicast type for tcgen05.cp Op");
+    break;
+  case SH::SHAPE_64x128b:
+    if (mc != MC::WARPX2_01_23 && mc != MC::WARPX2_02_13)
+      return emitError("Shape 64x128b requires multicast warpx2_01_23 or "
+                       "warpx2_02_13 for tcgen05.cp Op");
+    break;
+  case SH::SHAPE_32x128b:
+    if (mc != MC::WARPX4)
+      return emitError(
+          "Shape 32x128b requires multicast warpx4 for tcgen05.cp Op");
+    break;
+  default:
+    return emitError("Invalid shape for tcgen05.cp Op");
+  }
+  return success();
+}
+
+//===----------------------------------------------------------------------===//
+// NVVMDialect: getIntrinsicID/getIntrinsicIDAndArgs methods
----------------
grypp wrote:

ditto: remove NVVMDialect:

https://github.com/llvm/llvm-project/pull/127798