[all-commits] [llvm/llvm-project] 46ef57: [MLIR][NVVM] Add mbarrier.try_wait Op (#170285)

[Durgadoss R via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 46ef57a78bec4c8f8c6837e665becf785f0c0290
      https://github.com/llvm/llvm-project/commit/46ef57a78bec4c8f8c6837e665becf785f0c0290
  Author: Durgadoss R <durgadossr at nvidia.com>
  Date:   2025-12-04 (Thu, 04 Dec 2025)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
    M mlir/test/Target/LLVMIR/nvvm/mbar_invalid.mlir
    A mlir/test/Target/LLVMIR/nvvm/mbar_try_wait.mlir

  Log Message:
  -----------
  [MLIR][NVVM] Add mbarrier.try_wait Op (#170285)

This patch adds an Op for mbarrier.try_wait operation which lowers
to the corresponding intrinsics. This Op has support for an optional
time-limit, state-or-phase as well as relaxed memory semantics,
completing the features on this Op up to Blackwell.

Unlike the existing `nvvm.mbarrier.try_wait.parity` Op, this Op
does not provide a _blocking_ implementation. We intend to
add looping around this at NVGPU in a subsequent PR
(and deprecate the inline-asm based Op here).

lit tests are added to verify the lowering to the intrinsics.

Signed-off-by: Durgadoss R <durgadossr at nvidia.com>



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