[QST] Missing tcgen05.fence::before_thread_sync
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Based on the ptx docs, cross-thread ordering between async tcgen05 instructions is guaranteed by before_thread_sync and after_thread_sync. However, all the examples such as the one below only use
fence_view_async_tmem, which is tcgen05.wait.
https://github.com/NVIDIA/cutlass/blob/3f5bafb326e834b69b16fa6d10bc85b644dc1805/examples/python/CuTeDSL/blackwell/fmha_bwd.py#L2148
https://github.com/NVIDIA/cutlass/blob/3f5bafb326e834b69b16fa6d10bc85b644dc1805/examples/python/CuTeDSL/blackwell/fmha_bwd.py#L1684
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start with the PTX memory consistency documentation linked in the issue, then inspect the referenced examples/python/CuTeDSL/blackwell/fmha_bwd.py locations around lines 1684 and 2148. Compare their use of fence_view_async_tmem with the documented before_thread_sync and after_thread_sync ordering, and treat the issue as done when the missing support or usage is addressed and the affected examples follow the required synchronization pattern.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, python
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100