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

[Stefan Mada]

https://github.com/smada3 updated https://github.com/llvm/llvm-project/pull/208063

>From 5458e4a6a9d12bab08793042addb0b87cf91117e Mon Sep 17 00:00:00 2001
From: Stefan Mada <smada at nvidia.com>
Date: Tue, 7 Jul 2026 18:20:31 +0000
Subject: [PATCH 1/2] [NVVM][MLIR] Fixed valgrind leak in MMAOp

---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 22 +++++++++++++++++----
 1 file changed, 18 insertions(+), 4 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 40f7f15b694cb..1f5738589a6a0 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
@@ -3477,9 +3477,23 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 
   string llvmBuilder = [{
     auto operands = moduleTranslation.lookupValues(opInst.getOperands());
+    // `intOverflowBehavior` is optional; an omitted attribute selects the
+    // default `wrapped` behavior. Resolve it to a fully-initialized value here
+    // so the intrinsic selector always inspects an engaged optional and never
+    // reads an empty optional's uninitialized payload (a benign read the
+    // optimizer can hoist out of the guarding `has_value()` check under -O2).
+    // An empty std::optional still reserves storage for its payload but leaves
+    // it uninitialized, and at -O2 the optimizer may if-convert the selector's
+    // `has_value() ? ... *sat ... : true` guard and speculatively load `*sat`
+    // before the guard discards it. This is not UB -- the read touches validly
+    // allocated storage and its result is never used -- but Valgrind still
+    // reports the load of uninitialized bytes.
+    std::optional<mlir::NVVM::MMAIntOverflow> intOverflow = mlir::NVVM::MMAIntOverflow::wrapped;
+    if (mlir::NVVM::MMAIntOverflowAttr satAttr = op.getIntOverflowBehaviorAttr())
+      intOverflow = satAttr.getValue();
     auto intId = mlir::NVVM::MmaOp::getIntrinsicID(
         $shape.getM(), $shape.getN(), $shape.getK(),
-        $b1Op, $intOverflowBehavior,
+        $b1Op, intOverflow,
         $layoutA, $layoutB,
         *$multiplicandAPtxType,
         *$multiplicandBPtxType,

>From 964544ea817f23d0a055d14fcfc72c5493755475 Mon Sep 17 00:00:00 2001
From: Stefan Mada <smada at nvidia.com>
Date: Wed, 8 Jul 2026 15:31:25 +0000
Subject: [PATCH 2/2] Shortened comment on mmaop valgrind fix

---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 15 ++++-----------
 1 file changed, 4 insertions(+), 11 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 1f5738589a6a0..627d9a861d780 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -3477,17 +3477,10 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 
   string llvmBuilder = [{
     auto operands = moduleTranslation.lookupValues(opInst.getOperands());
-    // `intOverflowBehavior` is optional; an omitted attribute selects the
-    // default `wrapped` behavior. Resolve it to a fully-initialized value here
-    // so the intrinsic selector always inspects an engaged optional and never
-    // reads an empty optional's uninitialized payload (a benign read the
-    // optimizer can hoist out of the guarding `has_value()` check under -O2).
-    // An empty std::optional still reserves storage for its payload but leaves
-    // it uninitialized, and at -O2 the optimizer may if-convert the selector's
-    // `has_value() ? ... *sat ... : true` guard and speculatively load `*sat`
-    // before the guard discards it. This is not UB -- the read touches validly
-    // allocated storage and its result is never used -- but Valgrind still
-    // reports the load of uninitialized bytes.
+    // `intOverflowBehavior` is optional and defaults to `wrapped` when omitted.
+    // We initialize it to a concrete value so that the intrinsic selector never
+    // speculatively reads an empty optional's uninitialized payload (benign,
+    // but flagged by Valgrind).
     std::optional<mlir::NVVM::MMAIntOverflow> intOverflow = mlir::NVVM::MMAIntOverflow::wrapped;
     if (mlir::NVVM::MMAIntOverflowAttr satAttr = op.getIntOverflowBehaviorAttr())
       intOverflow = satAttr.getValue();