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

[Guray Ozen]

================
@@ -3444,6 +3444,71 @@ def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st"> {
   let hasVerifier = 1;
 }
 
+//===----------------------------------------------------------------------===//
+// NVVM dp4a Op
+//===----------------------------------------------------------------------===//
+
+def DP4aS8 : I32EnumAttrCase<"S8", 1, "s8">;
+def DP4aU8 : I32EnumAttrCase<"U8", 0, "u8">;
+
+def DP4aType : I32EnumAttr<"DP4aType", "NVVM DP4aType",
+  [DP4aS8, DP4aU8]> {
+  let cppNamespace = "::mlir::NVVM";
+  let genSpecializedAttr = 0;
+}
+
+def DP4aTypeAttr : EnumAttr<NVVM_Dialect, DP4aType, "dp4a_type"> {
+  let assemblyFormat = "`<` $value `>`";
+}
+
+
+def NVVM_Dp4aOp : NVVM_Op<"dot.accumulate.4way"> {
+  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` are vectors of 4 bytes between which the dot product is 
+    computed.
+    The `a_type` and `b_type` attributes specify the type of the elements in `a`
+    and `b` respectively.
+    If `a_type` or `b_type` is `s8`, then the elements in the corresponding 
+    vector are sign-extended to 32-bit before the dot product is computed.
+    If `a_type` or `b_type` is `u8`, then the elements in the corresponding 
+    vector are zero-extended to 32-bit instead.
+    Operand `c` is a 32-bit integer to which the result is accumulated. It is
+    treated as holding a signed integer if any of `a_type` or `b_type` is `s8`.
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#integer-arithmetic-instructions-dp4a)
+  }];
+  
+  let arguments = (ins
+    VectorOfLengthAndType<[4], [I8]>:$a,
+    DP4aTypeAttr:$a_type,
+    VectorOfLengthAndType<[4], [I8]>:$b,
+    DP4aTypeAttr:$b_type,
+    I32:$c,
+    DefaultValuedAttr<UnitAttr, "false">:$a_siext,
+    DefaultValuedAttr<UnitAttr, "false">:$b_siext
----------------
grypp wrote:

I believe we don't need them anymore

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