[QST]Behavior of TMA Store and Wait Mechanism in CUTLASS

[ziyuhuang123]

In CUTLASS, there is a tma_store_wait function, which corresponds to cp.async.bulk.wait_group.read. Based on my observations while working with TMA, it seems that after completing a TMA-store operation, waiting is not necessary. It appears to behave like expect_tx, where the operation seems to complete automatically.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[hwu36]

@Junkai-Wu

[Junkai-Wu]

In epilogue, usually we have storing sequences like (r2s = store data from register to shared memory, s2g = tma store):

r2s
s2g
r2s
s2g
...

the second r2s must wait until the first s2g completes. If we don't wait until s2g completes, the second r2s will fill data in shared memory position that the first s2g is storing which will cause unexpected behaivour.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[github-actions]

This issue has been labeled inactive-90d due to no recent activity in the past 90 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.