[MLIR][NVVM] Propagate verification failure for unsupported SM targets

[Men-cotton]

Fixes: #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

[llvmbot]

@llvm/pr-subscribers-mlir-llvm

@llvm/pr-subscribers-mlir

Author: Men-cotton (Men-cotton)

Changes

Fixes: #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

---

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

2 Files Affected:

• (modified) mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp (+14-10)
• (added) mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir (+11)

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
+        }
+    }
+}

[Wolfram70]

LGTM, thanks for the fix!

[Men-cotton]

@Wolfram70
I don't have write access. Could you please merge this?