用 mbarrier 追踪两次 TMA load

T0 (producer)
TMA Engine
mbarrier
T1 (consumer)
时间线
正在执行
发出信号的操作
提交给 TMA
阻塞 / 等待
barrier 状态
空闲
更新 mbarrier 状态
初始化后进行 CTA 同步

producer(T0)

producer 分别用一次 TMA load 搬运 2048 bytes 的 A tile 和 2048 bytes 的 B tile。随后调用 arrive.expect_tx(4096),报告一次 arrival,并让 barrier 等待两次 load 共 4096 bytes 的传输。

TMA engine

每次 copy 完成后,硬件都会向 barrier 报告 2048 bytes 的 complete-tx,并相应扣减 tx-count。第二次完成后,tx-count 回到 0。

consumer(T1)

调用 TIRx 的阻塞式 try_wait(phase=0);barrier 完成这一 phase 后,才能读取 A、B tiles。