[all-commits] [llvm/llvm-project] 141ad3: [MLIR][NVVM] Add SM version requirements to mbarri...

[Dhairyashil R G via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 141ad3d5e3312eccc977b20af18c978e46fd9b40
      https://github.com/llvm/llvm-project/commit/141ad3d5e3312eccc977b20af18c978e46fd9b40
  Author: Dhairyashil R G <dhairyashil25 at gmail.com>
  Date:   2026-09-28 (Mon, 28 Sep 2026)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/test/Dialect/LLVMIR/nvvm_check_target_sm.mlir

  Log Message:
  -----------
  [MLIR][NVVM] Add SM version requirements to mbarrier Ops (#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



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications