[QST]Synchronizing Threads Between Loading Q/K and V in WASP

[ziyuhuang123]

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?

[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.