[all-commits] [llvm/llvm-project] 9b734e: [mlir][NVVM] Add memory clobber support to inline_...

[Siyi Xu via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 9b734e25473bc9e299a4b58575b9b0e618925b78
      https://github.com/llvm/llvm-project/commit/9b734e25473bc9e299a4b58575b9b0e618925b78
  Author: Siyi Xu <siyix at nvidia.com>
  Date:   2026-07-31 (Fri, 31 Jul 2026)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/BasicPtxBuilderInterface.cpp
    M mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

  Log Message:
  -----------
  [mlir][NVVM] Add memory clobber support to inline_ptx and BasicPtxBui… (#212937)

PTX with memory side effects (stores, atomics, mbarrier operations with
acquire/release semantics) emitted through `nvvm.inline_ptx` lowers to
`llvm.inline_asm` with register constraints only. Without a `~{memory}`
clobber, LLVM may reorder memory accesses across the inline assembly.
There was no way to express the clobber through this op or the
`BasicPtxBuilderInterface` machinery.

This patch adds, as discussed with @grypp:
- A `hasMemoryClobber` interface method (default `false`) on
`BasicPtxBuilderOpInterface`; when it returns `true`, `PtxBuilder`
appends `~{memory}` to the constraints of the generated inline assembly,
after all register constraints and tied indices.
- An opt-in `memory_clobber` boolean attribute (default `false`) on
`nvvm.inline_ptx` exposing this. Existing behavior is unchanged unless
the attribute is set to `true`.

For example:
```mlir
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" ro(%p, %c : !llvm.ptr, i32) memory_clobber = true
```
now lowers to inline assembly with constraints `"l,r,~{memory}"`.

Tests: constraint emission with/without the attribute, combination with
predicate (`"l,r,b,~{memory}"`), read-write operands with tied indices
(`"=f,=f,r,r,0,1,~{memory}"`), and the no-operand case (`"~{memory}"`).
The op documentation's first example (mbarrier.init) is updated to show
correct usage.



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