[Mlir-commits] [mlir] [MLIR][NVVM] Add support for dp4a instructions (PR #139043)
Srinivasa Ravi
llvmlistbot at llvm.org
Thu May 8 03:25:37 PDT 2025
================
@@ -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.
----------------
Wolfram70 wrote:
That makes sense. I have removed the support for the packed `I32` and retained only the vector type for the inputs in the latest revision since it is just a bitcast away anyways, thanks!
https://github.com/llvm/llvm-project/pull/139043
More information about the Mlir-commits
mailing list