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

[llvmlistbot at llvm.org]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Stefan Mada (smada3)

<details>
<summary>Changes</summary>

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.

---
Full diff: https://github.com/llvm/llvm-project/pull/207221.diff


2 Files Affected:

- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+4-4) 
- (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+3-9) 


``````````diff
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

``````````

</details>


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