#1908 [Bug]: LDS grouped expert kernels skip groups beyond the launch grid

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