mbarrier：状态与操作

mbarrier 概念状态（64-bit，内部编码不透明）

phase parity
0 或 1
pending arrival
count
剩余 arrival
expected arrival
count
预期总数
tx-count
待完成字节数
phase parity 当前 phase 的奇偶值。phase 前进一轮后，parity 会在 0 和 1 之间翻转。
pending arrival count 当前 phase 完成前还需要多少次 arrival。
expected arrival count 每个 phase 需要等待的 arrival 数量，由 init 设置。
tx-count 尚未完成的异步传输字节数。arrive.expect_tx 会增加它；传输完成后硬件会减少它。

phase 完成条件

pending_arrival_count == 0
&&
tx_count == 0
当 预期的 arrival 都已发生，并且 所有登记的异步传输字节都已完成 时，当前 phase 才算完成。 随后 barrier 会原子地进入下一 phase，并将 pending arrival count 恢复为 expected arrival count。

核心 API

操作作用对 mbarrier 的影响
mbarrier.init setup 设置每个 phase 需要等待的 arrival 数量 phase=0, expected_arrival_count=count, pending_arrival_count=count, tx_count=0
mbarrier.arrive participant thread 向 barrier 发出 arrival 信号 pending_arrival_count -= 1
mbarrier.arrive.expect_tx participant arrive，同时声明预期的异步传输字节数 pending_arrival_count -= 1, tx_count += byte_count
tcgen05.commit...mbarrier::arrive::one issuer 让 barrier 追踪当前 thread 此前发出的异步 tcgen05 操作 完成后由 system 执行一次 arrive(count=1)
T.ptx.mbarrier.try_wait consumer 等待指定 phase 完成 不修改 barrier 状态；未完成时在内部重试，完成后返回
三种更新路径： (1) mbarrier.arrive 直接减少 pending arrival count。 (2) arrive.expect_tx 同时 arrive 并增加 tx-count；关联的 TMA copy 完成时，硬件通过 complete_tx 扣减 tx-count。 (3) tcgen05.commit...arrive::one 在此前的异步 tcgen05 操作完成后，由 system 报告一次 arrival。