[Mlir-commits] [mlir] [MLIR][NVVM] Add SM version requirements to mbarrier Ops (PR #224105)
Dhairyashil R G
llvmlistbot at llvm.org
Wed Sep 16 11:02:10 PDT 2026
https://github.com/dhairyashilRG created https://github.com/llvm/llvm-project/pull/224105
The mbarrier Ops had no `NVVMRequiresSM` trait, so a `gpu.module` targeting an older chip (e.g. `sm_75`) accepted them, and the error only showed up later. This adds the trait to all 13 of them, same as #192257 and `mbarrier.check_layout` in #217252.
| floor | Ops |
|---|---|
| `sm_80` | `init`, `inval`, `arrive`, `arrive_drop`, `arrive.nocomplete`, `arrive_drop.nocomplete`, `test.wait` |
| `sm_90` | `expect_tx`, `complete_tx`, `arrive.expect_tx`, `arrive_drop.expect_tx`, `try_wait`, `try_wait.parity` |
The floors come from the Target ISA Notes in the [PTX ISA 8.5](https://docs.nvidia.com/cuda/pdf/ptx_isa_8.5.pdf) (section 9.7.12.15), and they match the backend: every Op that lowers to an intrinsic fails to compile one arch below its floor and compiles at it. `arrive_drop` at sm_80 relies on #222916.
Since the trait is per Op, some settings that need sm_90 still get through on sm_80: an explicit `count`, `scope = cluster`, a `shared::cluster` pointer, `relaxed = true`, or `layout` on `init`. Catching those would mean the Ops implementing `getRequiredMinSMVersion` themselves, which no Op does yet. Happy to add that if you'd like. PTX version isn't covered either (#69448).
One behaviour change: `gpu-lower-to-nvvm-pipeline` and `nvvm-attach-target` default to `sm_75`, so using mbarrier Ops with that default now fails earlier, with a clear message. `verifyTarget = false` turns the check off.
Tests are in `nvvm_check_target_sm.mlir`: one failing case per Op just below its floor, plus passing cases at the floor, above it, and on `sm_90a`.
Assisted-by: Claude
>From b1f7da1450dd6f4660bf9cb58aa01a548ee941bc Mon Sep 17 00:00:00 2001
From: Dhairyashil R G <dhairyashil25 at gmail.com>
Date: Wed, 16 Sep 2026 22:41:41 +0530
Subject: [PATCH] [MLIR][NVVM] Add SM version requirements to mbarrier Ops
Declare NVVMRequiresSM on the thirteen nvvm.mbarrier.* Ops that had no
SM requirement, so that a gpu.module with an #nvvm.target rejects them
on architectures that do not have the instruction, instead of failing
later during code generation or PTX assembly.
sm_80: mbarrier.init, mbarrier.inval, mbarrier.arrive,
mbarrier.arrive_drop, mbarrier.arrive.nocomplete,
mbarrier.arrive_drop.nocomplete, mbarrier.test.wait
sm_90: mbarrier.arrive.expect_tx, mbarrier.arrive_drop.expect_tx,
mbarrier.expect_tx, mbarrier.complete_tx, mbarrier.try_wait,
mbarrier.try_wait.parity
The floors follow the Target ISA Notes of each instruction in the PTX
ISA (Parallel Synchronization and Communication Instructions: mbarrier)
and match the Predicates already on the corresponding intrinsics in
NVPTXIntrinsics.td.
The trait is static per Op, so it declares the minimum over an Op's
configurations. Some configurations need sm_90 even when the Op's floor
is sm_80, and are not rejected by this change: an explicit count on
arrive/arrive_drop, scope = cluster, a shared::cluster pointer,
relaxed = true, and an explicit layout on init. The trait also
expresses only the SM half of a [PTX, SM] requirement.
Assisted-by: Claude
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 35 +--
.../Dialect/LLVMIR/nvvm_check_target_sm.mlir | 206 ++++++++++++++++++
2 files changed, 228 insertions(+), 13 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 88a6ee167bc74..38be955d1da34 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -645,7 +645,8 @@ def NVVM_PMEventOp : NVVM_VoidIntrinsicOp<"pmevent">,
//===----------------------------------------------------------------------===//
/// mbarrier.init instruction with generic pointer type
-def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
+def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init",
+ [NVVMRequiresSM<80>]>,
Arguments<(ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
I32:$count,
DefaultValuedAttr<ConfinedAttr<I32Attr,
@@ -697,7 +698,8 @@ def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
}];
}
-def NVVM_MBarrierInvalOp : NVVM_VoidIntrinsicOp<"mbarrier.inval">,
+def NVVM_MBarrierInvalOp : NVVM_VoidIntrinsicOp<"mbarrier.inval",
+ [NVVMRequiresSM<80>]>,
Arguments<(ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr)> {
let summary = "MBarrier Invalidation Operation";
let description = [{
@@ -753,7 +755,8 @@ def NVVM_MBarrierCheckLayoutOp :
"$addr (`layout` `=` $layout^)? attr-dict `:` type($addr) `->` type($res)";
}
-def NVVM_MBarrierExpectTxOp : NVVM_VoidIntrinsicOp<"mbarrier.expect_tx"> {
+def NVVM_MBarrierExpectTxOp : NVVM_VoidIntrinsicOp<"mbarrier.expect_tx",
+ [NVVMRequiresSM<90>]> {
let summary = "MBarrier expect-tx Operation";
let description = [{
The `nvvm.mbarrier.expect_tx` operation increases the transaction count
@@ -777,7 +780,8 @@ def NVVM_MBarrierExpectTxOp : NVVM_VoidIntrinsicOp<"mbarrier.expect_tx"> {
let hasVerifier = 1;
}
-def NVVM_MBarrierCompleteTxOp : NVVM_VoidIntrinsicOp<"mbarrier.complete_tx"> {
+def NVVM_MBarrierCompleteTxOp : NVVM_VoidIntrinsicOp<"mbarrier.complete_tx",
+ [NVVMRequiresSM<90>]> {
let summary = "MBarrier complete-tx Operation";
let description = [{
The `nvvm.mbarrier.complete_tx` operation decrements the transaction
@@ -803,7 +807,7 @@ def NVVM_MBarrierCompleteTxOp : NVVM_VoidIntrinsicOp<"mbarrier.complete_tx"> {
}
def NVVM_MBarrierArriveOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive",
- [InferTypeOpAdaptorWithIsCompatible]> {
+ [InferTypeOpAdaptorWithIsCompatible, NVVMRequiresSM<80>]> {
let summary = "MBarrier Arrive Operation";
let description = [{
The `nvvm.mbarrier.arrive` operation performs an arrive-on operation on the
@@ -856,7 +860,7 @@ def NVVM_MBarrierArriveOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive",
}
def NVVM_MBarrierArriveDropOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_drop",
- [InferTypeOpAdaptorWithIsCompatible]> {
+ [InferTypeOpAdaptorWithIsCompatible, NVVMRequiresSM<80>]> {
let summary = "MBarrier Arrive-Drop Operation";
let description = [{
The `nvvm.mbarrier.arrive_drop` operation decrements the expected arrival
@@ -885,7 +889,8 @@ def NVVM_MBarrierArriveDropOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_dr
let hasVerifier = 1;
}
-def NVVM_MBarrierArriveNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive.nocomplete">,
+def NVVM_MBarrierArriveNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive.nocomplete",
+ [NVVMRequiresSM<80>]>,
Results<(outs I64:$res)>,
Arguments<(ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
I32:$count)> {
@@ -919,7 +924,8 @@ def NVVM_MBarrierArriveNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arr
let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands) `->` type($res)";
}
-def NVVM_MBarrierArriveDropNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_drop.nocomplete">,
+def NVVM_MBarrierArriveDropNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_drop.nocomplete",
+ [NVVMRequiresSM<80>]>,
Results<(outs I64:$res)>,
Arguments<(ins AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
I32:$count)> {
@@ -937,7 +943,7 @@ def NVVM_MBarrierArriveDropNocompleteOp : NVVM_SingleResultIntrinsicOp<"mbarrier
}
def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_tx",
- [InferTypeOpAdaptorWithIsCompatible]> {
+ [InferTypeOpAdaptorWithIsCompatible, NVVMRequiresSM<90>]> {
let summary = "MBarrier Arrive with Expected Transaction Count";
let description = [{
The `nvvm.mbarrier.arrive.expect_tx` operation performs an expect-tx operation
@@ -1008,7 +1014,7 @@ def NVVM_MBarrierArriveExpectTxOp : NVVM_PTXBuilder_Op<"mbarrier.arrive.expect_t
}
def NVVM_MBarrierArriveDropExpectTxOp : NVVM_SingleResultIntrinsicOp<"mbarrier.arrive_drop.expect_tx",
- [InferTypeOpAdaptorWithIsCompatible]> {
+ [InferTypeOpAdaptorWithIsCompatible, NVVMRequiresSM<90>]> {
let summary = "MBarrier arrive_drop with expected transaction count";
let description = [{
The `nvvm.mbarrier.arrive_drop.expect_tx` operation is similar to the
@@ -1033,7 +1039,8 @@ def NVVM_MBarrierArriveDropExpectTxOp : NVVM_SingleResultIntrinsicOp<"mbarrier.a
let hasVerifier = 1;
}
-def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity">,
+def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity",
+ [NVVMRequiresSM<90>]>,
Arguments<(ins
AnyTypeOf<[LLVM_PointerGeneric, LLVM_PointerShared]>:$addr,
I32:$phase, I32:$ticks)> {
@@ -1091,7 +1098,8 @@ def NVVM_MBarrierTryWaitParityOp : NVVM_PTXBuilder_Op<"mbarrier.try_wait.parity"
let assemblyFormat = "$addr `,` $phase `,` $ticks attr-dict `:` type(operands)";
}
-def NVVM_MBarrierTestWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.test.wait"> {
+def NVVM_MBarrierTestWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.test.wait",
+ [NVVMRequiresSM<80>]> {
let summary = "MBarrier Non-Blocking Test Wait Operation";
let description = [{
The `nvvm.mbarrier.test.wait` operation performs a non-blocking test for the
@@ -1161,7 +1169,8 @@ def NVVM_MBarrierTestWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.test.wait">
let hasVerifier = 1;
}
-def NVVM_MBarrierTryWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.try_wait"> {
+def NVVM_MBarrierTryWaitOp : NVVM_SingleResultIntrinsicOp<"mbarrier.try_wait",
+ [NVVMRequiresSM<90>]> {
let summary = "MBarrier try wait on state or phase with an optional timelimit";
let description = [{
The `nvvm.mbarrier.try_wait` operation checks whether the specified
diff --git a/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir b/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
index 074af98ac4b74..a9b1e6d0fd3b5 100644
--- a/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir
@@ -227,3 +227,209 @@ gpu.module @tensormap_replace_sm90 [#nvvm.target<chip = "sm_90">] {
return
}
}
+
+// -----
+
+//===----------------------------------------------------------------------===//
+// mbarrier
+//===----------------------------------------------------------------------===//
+
+// Just check these don't emit errors: every op at exactly its floor.
+gpu.module @mbarrier_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_sm80(%addr: !llvm.ptr, %count: i32, %state: i64) {
+ nvvm.mbarrier.init %addr, %count : !llvm.ptr, i32
+ %0 = nvvm.mbarrier.arrive %addr : !llvm.ptr -> i64
+ %1 = nvvm.mbarrier.arrive_drop %addr : !llvm.ptr -> i64
+ %2 = nvvm.mbarrier.arrive.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ %3 = nvvm.mbarrier.arrive_drop.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ %4 = nvvm.mbarrier.test.wait %addr, %state : !llvm.ptr, i64 -> i1
+ nvvm.mbarrier.inval %addr : !llvm.ptr
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_sm90 [#nvvm.target<chip = "sm_90">] {
+ func.func @mbarrier_sm90(%addr: !llvm.ptr, %shared: !llvm.ptr<3>, %count: i32,
+ %state: i64, %phase: i32, %ticks: i32) {
+ %0 = nvvm.mbarrier.arrive.expect_tx %addr, %count : !llvm.ptr, i32 -> i64
+ %1 = nvvm.mbarrier.arrive_drop.expect_tx %addr, %count : !llvm.ptr, i32 -> i64
+ nvvm.mbarrier.expect_tx %shared, %count : !llvm.ptr<3>, i32
+ nvvm.mbarrier.complete_tx %shared, %count : !llvm.ptr<3>, i32
+ %2 = nvvm.mbarrier.try_wait %addr, %state : !llvm.ptr, i64 -> i1
+ nvvm.mbarrier.try_wait.parity %addr, %phase, %ticks : !llvm.ptr, i32, i32
+ return
+ }
+}
+
+// -----
+
+// Arch-accelerated targets satisfy a plain NVVMRequiresSM floor: sm_90a is
+// >= sm_80 and >= sm_90. This is the chip string real Hopper users write.
+gpu.module @mbarrier_sm90a [#nvvm.target<chip = "sm_90a">] {
+ func.func @mbarrier_sm90a(%addr: !llvm.ptr, %shared: !llvm.ptr<3>, %count: i32) {
+ nvvm.mbarrier.init %addr, %count : !llvm.ptr, i32
+ nvvm.mbarrier.expect_tx %shared, %count : !llvm.ptr<3>, i32
+ return
+ }
+}
+
+// -----
+
+// NVVMRequiresSM is a floor, not an equality: every sm_80 op stays valid on
+// sm_90. These positive modules pass before and after the traits were added;
+// they exist to catch a floor set too high, which the negatives cannot.
+gpu.module @mbarrier_sm80_ops_on_sm90 [#nvvm.target<chip = "sm_90">] {
+ func.func @mbarrier_sm80_ops_on_sm90(%addr: !llvm.ptr, %count: i32, %state: i64) {
+ nvvm.mbarrier.init %addr, %count : !llvm.ptr, i32
+ %0 = nvvm.mbarrier.arrive %addr : !llvm.ptr -> i64
+ %1 = nvvm.mbarrier.arrive_drop %addr : !llvm.ptr -> i64
+ %2 = nvvm.mbarrier.arrive.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ %3 = nvvm.mbarrier.arrive_drop.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ %4 = nvvm.mbarrier.test.wait %addr, %state : !llvm.ptr, i64 -> i1
+ nvvm.mbarrier.inval %addr : !llvm.ptr
+ return
+ }
+}
+
+// -----
+
+// The check stays opt-out, as for every other op in this file.
+gpu.module @mbarrier_verify_target_disabled
+ [#nvvm.target<chip = "sm_70", verifyTarget = false>] {
+ func.func @mbarrier_verify_target_disabled(%addr: !llvm.ptr, %count: i32) {
+ nvvm.mbarrier.init %addr, %count : !llvm.ptr, i32
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_init_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_init_sm70(%addr: !llvm.ptr, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.init' op is not supported on sm_70}}
+ nvvm.mbarrier.init %addr, %count : !llvm.ptr, i32
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_inval_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_inval_sm70(%addr: !llvm.ptr) {
+ // expected-error @below {{'nvvm.mbarrier.inval' op is not supported on sm_70}}
+ nvvm.mbarrier.inval %addr : !llvm.ptr
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_arrive_sm70(%addr: !llvm.ptr) {
+ // expected-error @below {{'nvvm.mbarrier.arrive' op is not supported on sm_70}}
+ %0 = nvvm.mbarrier.arrive %addr : !llvm.ptr -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_drop_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_arrive_drop_sm70(%addr: !llvm.ptr) {
+ // expected-error @below {{'nvvm.mbarrier.arrive_drop' op is not supported on sm_70}}
+ %0 = nvvm.mbarrier.arrive_drop %addr : !llvm.ptr -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_nocomplete_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_arrive_nocomplete_sm70(%addr: !llvm.ptr, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.arrive.nocomplete' op is not supported on sm_70}}
+ %0 = nvvm.mbarrier.arrive.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_drop_nocomplete_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_arrive_drop_nocomplete_sm70(%addr: !llvm.ptr, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.arrive_drop.nocomplete' op is not supported on sm_70}}
+ %0 = nvvm.mbarrier.arrive_drop.nocomplete %addr, %count : !llvm.ptr, i32 -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_test_wait_sm70 [#nvvm.target<chip = "sm_70">] {
+ func.func @mbarrier_test_wait_sm70(%addr: !llvm.ptr, %state: i64) {
+ // expected-error @below {{'nvvm.mbarrier.test.wait' op is not supported on sm_70}}
+ %0 = nvvm.mbarrier.test.wait %addr, %state : !llvm.ptr, i64 -> i1
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_expect_tx_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_arrive_expect_tx_sm80(%addr: !llvm.ptr, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.arrive.expect_tx' op is not supported on sm_80}}
+ %0 = nvvm.mbarrier.arrive.expect_tx %addr, %count : !llvm.ptr, i32 -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_arrive_drop_expect_tx_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_arrive_drop_expect_tx_sm80(%addr: !llvm.ptr, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.arrive_drop.expect_tx' op is not supported on sm_80}}
+ %0 = nvvm.mbarrier.arrive_drop.expect_tx %addr, %count : !llvm.ptr, i32 -> i64
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_expect_tx_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_expect_tx_sm80(%addr: !llvm.ptr<3>, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.expect_tx' op is not supported on sm_80}}
+ nvvm.mbarrier.expect_tx %addr, %count : !llvm.ptr<3>, i32
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_complete_tx_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_complete_tx_sm80(%addr: !llvm.ptr<3>, %count: i32) {
+ // expected-error @below {{'nvvm.mbarrier.complete_tx' op is not supported on sm_80}}
+ nvvm.mbarrier.complete_tx %addr, %count : !llvm.ptr<3>, i32
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_try_wait_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_try_wait_sm80(%addr: !llvm.ptr, %state: i64) {
+ // expected-error @below {{'nvvm.mbarrier.try_wait' op is not supported on sm_80}}
+ %0 = nvvm.mbarrier.try_wait %addr, %state : !llvm.ptr, i64 -> i1
+ return
+ }
+}
+
+// -----
+
+gpu.module @mbarrier_try_wait_parity_sm80 [#nvvm.target<chip = "sm_80">] {
+ func.func @mbarrier_try_wait_parity_sm80(%addr: !llvm.ptr, %phase: i32, %ticks: i32) {
+ // expected-error @below {{'nvvm.mbarrier.try_wait.parity' op is not supported on sm_80}}
+ nvvm.mbarrier.try_wait.parity %addr, %phase, %ticks : !llvm.ptr, i32, i32
+ return
+ }
+}
More information about the Mlir-commits
mailing list