[Mlir-commits] [mlir] [MLIR][NVVM] Propagate verification failure for unsupported SM targets (PR #170001)

[llvmlistbot at llvm.org]

https://github.com/Men-cotton created https://github.com/llvm/llvm-project/pull/170001

Fixes: https://github.com/llvm/llvm-project/issues/169113

Correctly propagate verification failure when `NVVM::RequiresSMInterface` check fails during `gpu.module` verification.
Previously, the walk was interrupted but the function returned `success()`, causing a mismatch between the emitted diagnostic and the return status. This led to assertion failures in Python bindings which expect `failure()` when diagnostics are emitted.

CC: @grypp

>From 92a0353ca20c264c84280735e4d5cebbca4ab93b Mon Sep 17 00:00:00 2001
From: mencotton <mencotton0410 at gmail.com>
Date: Sun, 30 Nov 2025 01:15:51 +0900
Subject: [PATCH] [MLIR][NVVM] Propagate verification failure for unsupported
 SM targets

---
 mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp    | 24 +++++++++++--------
 .../Dialect/LLVMIR/nvvm-target-invalid.mlir   | 11 +++++++++
 2 files changed, 25 insertions(+), 10 deletions(-)
 create mode 100644 mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir

diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index d3c305555fde8..b98f15cfe6d75 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -4707,16 +4707,20 @@ LogicalResult NVVMTargetAttr::verifyTarget(Operation *gpuModule) {
                      "Minimum NVVM target SM version is sm_20");
   }
 
-  gpuModuleOp->walk([&](Operation *op) {
-    if (auto reqOp = llvm::dyn_cast<NVVM::RequiresSMInterface>(op)) {
-      const NVVMCheckSMVersion requirement = reqOp.getRequiredMinSMVersion();
-      if (!requirement.isCompatibleWith(targetSMVersion)) {
-        op->emitOpError() << "is not supported on " << getChip();
-        return WalkResult::interrupt();
-      }
-    }
-    return WalkResult::advance();
-  });
+  if (gpuModuleOp
+          ->walk([&](Operation *op) {
+            if (auto reqOp = llvm::dyn_cast<NVVM::RequiresSMInterface>(op)) {
+              const NVVMCheckSMVersion requirement =
+                  reqOp.getRequiredMinSMVersion();
+              if (!requirement.isCompatibleWith(targetSMVersion)) {
+                op->emitOpError() << "is not supported on " << getChip();
+                return WalkResult::interrupt();
+              }
+            }
+            return WalkResult::advance();
+          })
+          .wasInterrupted())
+    return failure();
 
   return success();
 }
diff --git a/mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir b/mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir
new file mode 100644
index 0000000000000..c2cfa7689978b
--- /dev/null
+++ b/mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir
@@ -0,0 +1,11 @@
+// RUN: not mlir-opt %s 2>&1 | FileCheck %s
+// CHECK: 'nvvm.tcgen05.alloc' op is not supported on sm_90
+
+module {
+    gpu.module @mod [#nvvm.target<chip = "sm_90">] {
+        func.func @tcgen05_alloc(%arg0: !llvm.ptr<7>, %arg1: i32) {
+             nvvm.tcgen05.alloc %arg0, %arg1 : !llvm.ptr<7>, i32
+             return
+        }
+    }
+}