[QST]Synchronizing Threads Between Loading Q/K and V in WASP
Open
Nobody has claimed this yet.
? - Needs Triage
inactive-30d
inactive-90d
question
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
In WASP, we can't use __syncthreads(), but my requirement is: first, load Q and K, then I want to sync, and finally load V. How should I write this sync?
Is this correct?
cutlass::arch::NamedBarrier::sync(NumMmaThreads + cutlass::NumThreadsPerWarp, static_cast<int>(FwdNamedBarriers::QueryEmpty) /*id*/);
I don't understand the concept of NamedBarrier. Could anyone explain if this is the right approach?
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 WASP synchronization requirement and the referenced cutlass::arch::NamedBarrier::sync call, including FwdNamedBarriers::QueryEmpty. Determine whether this barrier sequence correctly separates loading Q/K from loading V, and document the approved approach; the issue provides no file or test to run.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100