Skip to content

Fix a latent SMEM handoff race in the SM100 DSA indexer backward kernel #426

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
zkyue wants to merge 2 commits into
base: develop
Choose a base branch
from
+25 −20
Open

Fix a latent SMEM handoff race in the SM100 DSA indexer backward kernel #426

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
98a5a83
Fix indexer_backward_sm100 W_LOADED handoff: whole-warp mbarrier arrive
zkyue Jul 23, 2026
df68845
Elect a single lane for mbarrier init in indexer_backward_sm100
zkyue Jul 24, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
45 changes: 25 additions & 20 deletions python/cudnn/deepseek_sparse_attention/indexer_backward/indexer_backward_sm100.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -525,24 +525,25 @@ class SharedStorage:
cute.group_modes(gdQ, 0, 2),
)

# Init all barriers (warp 0)
# Init all barriers (single lane of warp 0)
if warp_idx == 0:
cute.arch.mbarrier_init(mbar + MBAR_S_FULL_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_S_FULL_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_DS_READY_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DS_READY_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DK_FULL_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_DK_FULL_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_DK_EMPTY_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DK_EMPTY_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_2, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_2, 1)
cute.arch.mbarrier_init(mbar + MBAR_W_LOADED, 1)
cute.arch.mbarrier_init(mbar + MBAR_DQ_DONE, 1)
with cute.arch.elect_one():
cute.arch.mbarrier_init(mbar + MBAR_S_FULL_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_S_FULL_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_DS_READY_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DS_READY_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DK_FULL_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_DK_FULL_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_DK_EMPTY_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DK_EMPTY_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_0, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_1, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_LOADED_2, self.WARPGROUP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_0, 1)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_1, 1)
cute.arch.mbarrier_init(mbar + MBAR_K_CONSUMED_2, 1)
cute.arch.mbarrier_init(mbar + MBAR_W_LOADED, self.WARP_SIZE)
cute.arch.mbarrier_init(mbar + MBAR_DQ_DONE, 1)
cute.arch.sync_threads()

# Pre-load topk indices into SMEM cooperatively (all 512 threads).
Expand Down Expand Up @@ -745,9 +746,13 @@ def _load_warp(
sW[idx] = mW[seq_idx, idx, batch_idx]

cute.arch.fence_view_async_shared()
# Signal W + grad_signal loaded for compute warpgroup
with cute.arch.elect_one():
cute.arch.mbarrier_arrive(mbar + MBAR_W_LOADED)
# Signal W + grad_signal loaded for compute warpgroup.
# All 32 lanes arrive (count = WARP_SIZE): mbarrier.arrive has release
# semantics for the *executing thread* only, so a single elected
# arrival would not order the other 31 lanes' sW/sGradSignal stores
# before the consumer's mbarrier_wait (racecheck flags exactly those
# 31 lanes). Whole-warp arrival closes the happens-before chain.
cute.arch.mbarrier_arrive(mbar + MBAR_W_LOADED)

# --- TMA Q load (dsa-next pattern: cute.copy with pre-partitioned tensors) ---
Q_producer.reset()
Expand Down