[Mlir-commits] [mlir] [MLIR][NVVM] Add support for dp4a instructions (PR #139043)

[Guray Ozen]

================
@@ -3444,6 +3444,54 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st"> {
   let hasVerifier = 1;
 }
 
+//===----------------------------------------------------------------------===//
+// NVVM dp4a Op
+//===----------------------------------------------------------------------===//
+
+def NVVM_Dp4aOp : NVVM_Op<"dp4a"> {
+  let summary = "Four-way byte dot product-accumulate instruction.";
+  let description = [{
+    Performs a four-way byte dot-product which is accumulated in a 32-bit
+    result.
+    Operand `a` and `b` can be passed either as packed 32-bit inputs holding
+    4 byte-inputs for the dot product, or as vectors of 4 i8 elements.
----------------
grypp wrote:

Passing packed 32bit isn't really type safe in MLIR side. Do you think do we really need to support that case? 

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