mbarrier State and Operations

mbarrier object (64-bit, in shared memory)

phase parity

0 or 1

pending arrival
count

arrivals left

expected arrival
count

expected total

tx-count

bytes pending

phase parity The parity of the current phase. It alternates between 0 and 1 whenever the barrier advances.

pending arrival count The arrivals still required before the current phase can complete.

expected arrival count The number of arrivals required in each phase, set by init.

tx-count Bytes of asynchronous transfer still outstanding. arrive.expect_tx increases this count; hardware decreases it as transfers complete.

Phase Completion Condition

        pending_arrival_count == 0

        &&

        tx_count == 0

The current phase completes only after all expected arrivals have occurred and all registered asynchronous transfer bytes have completed. The barrier then advances atomically to the next phase and restores the pending arrival count from the expected arrival count.

Core APIs

Instruction	Role	Effect on mbarrier
mbarrier.init setup	Set the number of arrivals required in each phase	phase=0, expected_arrival_count=count, pending_arrival_count=count, tx_count=0
mbarrier.arrive producer	Thread reports an arrival to the barrier	pending_arrival_count -= 1
mbarrier.arrive.expect_tx participant	Arrive and declare the expected asynchronous transfer bytes	pending_arrival_count -= 1, tx_count += byte_count
tcgen05.commit producer	Track asynchronous tcgen05 operations previously issued by this thread	The system performs one arrive(count=1) after completion
T.ptx.mbarrier.try_wait consumer	Wait for the requested phase to complete	Does not modify barrier state; retries internally and returns after completion

Three update paths: (1) mbarrier.arrive directly reduces the pending arrival count. (2) arrive.expect_tx arrives and increases tx-count; when the associated TMA copy finishes, hardware reduces tx-count through complete_tx. (3) tcgen05.commit...arrive::one makes the system report one arrival after the preceding asynchronous tcgen05 operations complete.