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

[llvmlistbot at llvm.org]

Author: Men-cotton
Date: 2025-12-01T09:50:13+05:30
New Revision: 81c5d468cf00d6e41112fba6c89d6c40013bcbda

URL: https://github.com/llvm/llvm-project/commit/81c5d468cf00d6e41112fba6c89d6c40013bcbda
DIFF: https://github.com/llvm/llvm-project/commit/81c5d468cf00d6e41112fba6c89d6c40013bcbda.diff

LOG: [MLIR][NVVM] Propagate verification failure for unsupported SM targets (#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

Added: 
    mlir/test/Dialect/LLVMIR/nvvm-target-invalid.mlir

Modified: 
    mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

Removed: 
    


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