[Mlir-commits] [mlir] [NVVM][MLIR] Fixed valgrind leak in MMAOp (PR #207221)

[Stefan Mada]

https://github.com/smada3 created https://github.com/llvm/llvm-project/pull/207221

MmaOp::getIntrinsicID can trigger an uninitialized read with valgrind enabled: intOverflowBehavior is optional, so when it's omitted it can be read without being initialized. It is set to a default value, which follows the PTX spec / existing lowering. Documentation updated to not be stale.

>From 603c5478a8b83d5e040cde10839be5663203d460 Mon Sep 17 00:00:00 2001
From: Stefan Mada <smada at nvidia.com>
Date: Tue, 30 Jun 2026 23:29:21 +0000
Subject: [PATCH] Fixed valgrind leak in MMAOp

---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td |  8 ++++----
 mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp  | 12 +++---------
 2 files changed, 7 insertions(+), 13 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 40f7f15b694cb..553219e5ed382 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3384,9 +3384,9 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
     `intOverflowBehavior` is only relevant when the `multiplicandType` attribute
     is one of `u8, s8, u4, s4`, this attribute describes how overflow is handled
     in the accumulator. When the attribute is `satfinite`, the accumulator values
-    are clamped in the int32 range on overflow. This is the default behavior.
-    Alternatively, accumulator behavior `wrapped` can also be specified, in
-    which case overflow wraps from one end of the range to the other.
+    are clamped in the int32 range on overflow. Alternatively, accumulator
+    behavior `wrapped` can be specified (this is the default), in which case
+    overflow wraps from one end of the range to the other.
 
     `layoutA` and `layoutB` are required and should generally be set to
     `#nvvm.mma_layout<row>` and `#nvvm.mma_layout<col>` respectively, but other
@@ -3431,7 +3431,7 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
   let results = (outs LLVM_AnyStruct:$res);
   let arguments = (ins NVVM_MMAShapeAttr:$shape,
              OptionalAttr<MMAB1OpAttr>:$b1Op,
-             OptionalAttr<MMAIntOverflowAttr>:$intOverflowBehavior,
+             DefaultValuedAttr<MMAIntOverflowAttr, "MMAIntOverflow::wrapped">:$intOverflowBehavior,
              MMALayoutAttr:$layoutA,
              MMALayoutAttr:$layoutB,
              OptionalAttr<MMATypesAttr>:$multiplicandAPtxType,
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 14406ab55f138..d9ef70dc1c12d 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -1129,15 +1129,9 @@ LogicalResult MmaOp::verify() {
                        " attribute");
   }
 
-  // Ensure int4/int8 MMA variants specify the accum overflow behavior
-  // attribute.
-  if (isInt4PtxType(*getMultiplicandAPtxType()) ||
-      isInt8PtxType(*getMultiplicandAPtxType())) {
-    if (!getIntOverflowBehavior())
-      return emitOpError("op requires " +
-                         getIntOverflowBehaviorAttrName().strref() +
-                         " attribute");
-  }
+  // `intOverflowBehavior` is a default-valued attribute (`wrapped`), so it
+  // always carries a value; no presence check is required for int4/int8
+  // variants.
 
   // Validate layout combinations. According to the operation description, most
   // MMA operations require layoutA=row and layoutB=col. Only m8n8k4 with f16