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

[Srinivasa Ravi]

================
@@ -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.
+    The `a_signed` and `b_signed` unit attributes specify whether the
+    individual byte inputs in operands `a` and `b` are signed or unsigned
+    respectively.
+    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` or `b` are signed.
+    
+    [For more information, see PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#integer-arithmetic-instructions-dp4a)
+  }];
+  
+  let arguments = (ins
+    AnyTypeOf<[I32, VectorOfLengthAndType<[4], [I8]>]>:$a,
+    AnyTypeOf<[I32, VectorOfLengthAndType<[4], [I8]>]>:$b,
+    I32:$c,
+    DefaultValuedAttr<UnitAttr, "false">:$a_signed,
+    DefaultValuedAttr<UnitAttr, "false">:$b_signed
+  );
+
+  let results = (outs I32:$res);
+
+  let assemblyFormat = "$a `,` $b `,` $c attr-dict `:` type($a) `,` type($b)";
+  
+  let extraClassDeclaration = [{
+    static llvm::Intrinsic::ID getIntrinsicID(bool a_signed, bool b_signed);
+  }];
+
+  string llvmBuilder = [{
+    auto id = NVVM::Dp4aOp::getIntrinsicID($a_signed, $b_signed);
+    llvm::Value* argA = $a;
+    llvm::Value* argB = $b;
+    if (!op.getA().getType().isInteger(32))
+      argA = builder.CreateBitCast(argA, llvm::Type::getInt32Ty(builder.getContext()));
+    if (!op.getB().getType().isInteger(32))
+      argB = builder.CreateBitCast(argB, llvm::Type::getInt32Ty(builder.getContext()));
+    $res = createIntrinsicCall(builder, id, {argA, argB, $c});
----------------
Wolfram70 wrote:

Moved it to `NVVMDialect.cpp` in the latest revision, thanks!

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