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

[llvmlistbot at llvm.org]

Author: Stefan Mada
Date: 2026-07-09T08:10:56-07:00
New Revision: 2e81babdf17d29d2b7e833137f92a598f76a8cbb

URL: https://github.com/llvm/llvm-project/commit/2e81babdf17d29d2b7e833137f92a598f76a8cbb
DIFF: https://github.com/llvm/llvm-project/commit/2e81babdf17d29d2b7e833137f92a598f76a8cbb.diff

LOG: [NVVM][MLIR] Fixed valgrind leak in MMAOp (#208063)

MmaOp::getIntrinsicID triggers an uninitialized-read warning under
Valgrind. `intOverflowBehavior` is optional; in the generated selector
its `has_value()` guarded access can be speculatively loaded before the
guard, reading the empty std::optional's uninitialized payload.

The read is benign: the payload is correctly stack-allocated but
uninitialized while empty, and the loaded value is discarded by the
guard, so this is not incorrect / undefined behavior.

This MR keeps the attribute as optional, but the lowering now resolves
it to a concrete value (wrapped when omitted) before calling
getIntrinsicID, so the selector always reads an initialized value.
`wrapped` matches the PTX spec and the existing lowering, where an
omitted attribute already selected the non-satfinite intrinsic;
intrinsic selection and printed IR are unchanged. Documentation updated
to correct the stale default.

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 89950d95e647ff..6dd32615451a7d 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,16 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 
   string llvmBuilder = [{
     auto operands = moduleTranslation.lookupValues(opInst.getOperands());
+    // `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();
     auto intId = mlir::NVVM::MmaOp::getIntrinsicID(
         $shape.getM(), $shape.getN(), $shape.getK(),
-        $b1Op, $intOverflowBehavior,
+        $b1Op, intOverflow,
         $layoutA, $layoutB,
         *$multiplicandAPtxType,
         *$multiplicandBPtxType,