Repository navigation
[Bug]: LDS grouped expert kernels skip groups beyond the launch grid #1908
Description
What happened
native_expert_grouped can route more expert groups than its launch has block rows. In STRATA_EXP_MODE=8, the fused gate/up kernel only visits blockIdx.y; the LDS down kernel used by modes 7/8 does the same. Groups at or beyond gridDim.y are silently skipped, leaving scratch/output stale.
Reproducer: a synthetic 512 x 256 expert layout with 11 groups, 32 entry capacity, varied group sizes and scattered output rows. Compare a full-capacity launch with cap_groups=1,3,8. Test mode 7's down pass separately and mode 8's complete pass. The same launch should produce identical output regardless of block-row count.
The regression fixture on W1nge:codex/grouped-stride-fix produces 120 mismatches per GPU against main's kernels; adding the group stride eliminates them. No model files or graph optimization are needed.
Strata version, GPU, OS
Main 61b3fb5d (0.1.42 development), Windows, CUDA 12.4 / MSVC; reproduced on RTX 2080 Ti (sm75) and P100 (sm60).
Activity
Niko1221 commentedon Oct 19, 2026
Thanks for taking the time to report this.
We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to read properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.
To reopen it takes about two minutes:
/reopen. A check runs and either reopens the issue or lists what is still missing.The rules are in CONTRIBUTING.md.