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

[Guray Ozen]

================
@@ -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;
----------------
grypp wrote:

This attribute is integer-only. Making it `DefaultValuedAttr` would leak an integer-overflow flag onto float ops during round-trip, which is odd.

Current solution looks good ot me

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