Tracking Two TMA Loads with an mbarrier

T0 (producer)
TMA Engine
mbarrier
T1 (consumer)
Timeline
Working / active
Signaling operation
Dispatch to TMA
Blocked / waiting
Barrier state
Idle
Updates mbarrier state
Sync threads after init

Producer (T0)

The producer issues one TMA load for a 2,048-byte A tile and another for a 2,048-byte B tile. It then calls arrive.expect_tx(4096) to contribute one arrival and make the barrier wait for all 4,096 bytes.

TMA Engine

When each copy finishes, hardware reports a 2,048-byte complete-tx update to the barrier. After the second copy, tx-count reaches zero.

Consumer (T1)

The consumer calls TIRx's blocking try_wait(phase=0). It may read the A and B tiles only after the barrier completes that phase.