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

[Durgadoss R]

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

ok, I see what you mentioned about "leaking this attr on float Ops" too. So, `DefaultValuedAttr` is not the right choice here. I think `DefaultValuedOptionalAttr` will help us prevent the round-trip leak but I see that this needs a different set of Verifier checks now (w/o buying us much).

So, this change LGTM as is. Thanks, Guray for clarifying!

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