[Mlir-commits] [mlir] [MLIR] [NVVM] Update summary and description fields for Barrier Ops in NVVMOps.td (PR #156726)
Guray Ozen
llvmlistbot at llvm.org
Wed Sep 3 23:19:30 PDT 2025
================
@@ -494,6 +494,26 @@ def NVVM_PMEventOp : NVVM_PTXBuilder_Op<"pmevent">,
/// mbarrier.init instruction with generic pointer type
def NVVM_MBarrierInitOp : NVVM_PTXBuilder_Op<"mbarrier.init">,
Arguments<(ins LLVM_AnyPointer:$addr, I32:$count, PtxPredicate:$predicate)> {
+ let summary = "MBarrier Initialization Op";
+ let description = [{
+ The `nvvm.mbarrier.init` operation initializes an mbarrier object at the specified
+ memory location.
+
+ This operation initializes the mbarrier object with the following state:
+ - Current phase: 0
+ - Expected arrival count: `count`
+ - Pending arrival count: `count`
+ - Transaction count (tx-count): 0
+
+ The operation takes the following operands:
+ - `addr`: A pointer to the memory location of the mbarrier object. Uses generic
+ addressing, but the address must still be in the shared memory space.
+ - `count`: Integer specifying the number of threads that will participate in barrier
+ synchronization. Must be in the range [1, 2²⁰ - 1].
----------------
grypp wrote:
`2²⁰` might not render. can we double check
https://github.com/llvm/llvm-project/pull/156726
More information about the Mlir-commits
mailing list