mega_moe: reserve 2 SMs of headroom (grid = num_sms - 2) — original yingz/mega-sm-headroom branch - #14
mega_moe: reserve 2 SMs of headroom (grid = num_sms - 2) — original yingz/mega-sm-headroom branch#14mayinghan wants to merge 3 commits into
Conversation
The SM100 FP8/FP4 mega MoE kernel launches a cooperative grid of NUM_SMS CTAs and uses grid_sync / nvlink_barrier that require ALL NUM_SMS CTAs resident. In disaggregated serving, a concurrent NCCL ncclDevKernel_SendRecv (KV transfer) can grab an SM the mega grid needs, leaving it 1-2 CTAs short so the grid/NVLink barrier never completes -> cross-host barrier timeout (observed as DeepGEMM "signal=4 target=8" crash on GB300 pf8 deployments). Launch with grid = get_num_sms() - 2 so a co-resident SendRecv has headroom and cannot starve the cooperative grid. Stays even for the 2-CTA cluster. Co-authored-by: Cursor <cursoragent@cursor.com>
Runtime cuda-gdb capture showed num_sms-2 was insufficient: 2 concurrent KV SendRecv kernels left the mega cooperative grid at 148/150 CTAs on a generator GPU, hanging the cross-host NVLink barrier (signal=4 target=8). Bump headroom to 16 to test whether ample room removes the deadlock (vs a cooperative-grid placement issue that headroom cannot fix). Co-authored-by: Cursor <cursoragent@cursor.com>
The -16 debug headroom masked a *stuck* KV-transfer SendRecv holding an SM forever; that root cause (cross-generator-group head-of-line blocking) is now fixed in fireworks nccl_kv_cache_transfer.py. KV SendRecv is pinned to 1 CTA per communicator (max_ctas=1) and ~1 co-resident for CONSECUTIVE_SUBRANGE, so 2 SMs headroom is ample. NIXL transfer (RDMA) needs none. Co-authored-by: Cursor <cursoragent@cursor.com>
INC-1291: the same defect, reproduced end-to-end in trainingPosting the evidence that brought me to this branch, in case it's useful for deciding how far to The production failureFigma RLOR job then NCCL watchdogs failing with It is not hardware. DCGM showed Xid 0 on all 8 GPUs right up to the crash and flipped to 43 only Local reproduction2× B300, the production image, a 2-rank EP group via torchrun, plus a synthetic SM-occupancy kernel
One occupied SM on one rank is enough, and it kills both ranks. That is the cross-rank cascade: The part that surprised me: full occupancy is benign
When the whole grid is queued, nobody is spinning, so no Why I think reservation is mitigation rather than fixThe kernel semantically needs gang scheduling — all Reserving 2 SMs narrows that window — a later-arriving comm kernel lands in the headroom instead And cross-rank, patience doesn't help at all: a rank whose grid is merely queued is locally benign, Structural options, if you want to go further
Repro scripts and the full RCA are in fw-ai/fireworks#47752 |
There was a problem hiding this comment.
Risk: medium. Left a non-blocking comment; no reviewers were assigned. Human review is needed: this changes cooperative mega-kernel launch occupancy (get_num_sms() - 2) on a production-critical path, which is above the low-risk auto-approve threshold.
Sent by Cursor Approval Agent: Pull Request Approver
The SM100 FP8/FP4 and BF16 mega MoE kernels (forward + backward) launch
cooperative grids whose grid_sync / nvlink_barrier (comm/barrier.cuh)
require ALL grid CTAs simultaneously resident, with a hard 60s timeout that
traps and takes down the process (SIGABRT) and the GPU (Xid 43 fallout).
An ordinary launch gives no gang-scheduling guarantee: CUDA places CTAs
greedily, so a concurrent kernel holding some (but not all) SMs leaves part
of the grid queued while resident CTAs spin against the deadline and never
yield their slots -- mutual deadlock.
Reproduced locally (train-firetitan:0.638.337, 2x B300 -- same SKU as the
production node): a single concurrent smem-heavy kernel on one rank's GPU
delays that rank's grid; its own CTAs time out at barrier.cuh:39 ('Grid sync
timeout') and peer ranks die at barrier.cuh:80 ('NVLink barrier timeout') --
the exact INC-1291 signature (figma RLOR xy98ftpyg7kovb4m, AP_MALAYSIA_2,
2026-09-08).
Reserve headroom so a co-resident comm kernel has somewhere to land other
than an SM the grid needs, via a shared helper used
by the fp8_fp4 forward, all four fp8_fp4 backward launches, and the bf16
default (the absolute DG_BF16_MEGA_MOE_NUM_SMS override is retained). The
amount is tunable with DG_MEGA_MOE_SM_HEADROOM and is rounded up to an even
value, since these are 2-CTA cluster launches and several sites assert
num_sms % 2 == 0.
The default of 2 is Ying Zhang's value from yingz/mega-sm-headroom (95046f0,
see #14), validated for disaggregated serving where the competitor is NCCL
SendRecv. It is NOT calibrated for training, where the competitors are FSDP2
all-gather / reduce-scatter on dedicated comm streams (an 8-GPU NVLink comm
pool is up to 32 channels == 32 CTAs, measured). Making it configurable so
the training default can be set from measurement rather than inherited from
a different deployment shape.
Known gap: reservation narrows the window but does not close it -- a
preemptive long-lived kernel occupying SMs before the mega launch still
starves the grid. A structural fix needs residency-agnostic synchronization
or driver-validated cooperative launch.
Co-Authored-By: Claude Code <noreply@anthropic.com>
|
Follow-up evidence (final): the production Chronosphere logs for INC-1291 were retrieved
Also reproduced with the production BF16 kernel (not just fp8_fp4) and real funcol One note for the training case: this branch's fixed |
…_MOE_SM_HEADROOM)
The SM100 FP8/FP4 and BF16 mega MoE kernels (forward + backward) launch
cooperative grids whose grid_sync / nvlink_barrier (comm/barrier.cuh)
require ALL grid CTAs simultaneously resident, with a hard 60s timeout that
traps and takes down the process (SIGABRT) and the GPU (Xid 43 fallout).
An ordinary launch gives no gang-scheduling guarantee: CUDA places CTAs
greedily, so a concurrent kernel holding some (but not all) SMs leaves part
of the grid queued while resident CTAs spin against the deadline and never
yield their slots -- mutual deadlock.
Reproduced locally (train-firetitan:0.638.337, 2x B300 -- same SKU as the
production node): a single concurrent smem-heavy kernel on one rank's GPU
delays that rank's grid; its own CTAs time out at barrier.cuh:39 ('Grid sync
timeout') and peer ranks die at barrier.cuh:80 ('NVLink barrier timeout') --
the exact INC-1291 signature (figma RLOR xy98ftpyg7kovb4m, AP_MALAYSIA_2,
2026-09-08, forward dispatch barrier tag=1 per production Chronosphere logs).
With the production BF16 kernel and real FSDP-shaped collectives, all three
collective types (reduce-scatter / all-gather / HSDP all-reduce) abort the
stock build; none abort with headroom + pinned NCCL channels.
Reserve headroom via a shared get_mega_moe_num_sms() helper used by the
fp8_fp4 forward, all four fp8_fp4 backward launches, and the bf16 default
(the absolute DG_BF16_MEGA_MOE_NUM_SMS override is retained). The amount is
controlled exclusively by DG_MEGA_MOE_SM_HEADROOM; it has NO built-in
default -- unset means 0, the historical full-device grid -- because the
right value is deployment-specific (serving SendRecv: 2; training FSDP2
collectives with 32 NCCL channels: 8 alongside NCCL_MAX_NCHANNELS=8). The
value is rounded up to even, since these are 2-CTA cluster launches and
several sites assert num_sms % 2 == 0.
Measured sizing for reference (8x B300, production BF16 kernel):
DG_MEGA_MOE_SM_HEADROOM=8 + NCCL_MAX_NCHANNELS=8 + NCCL_MIN_NCHANNELS=8
+ NCCL_NVLS_ENABLE=0: zero aborts across 90 mixed + 200 e2e iterations,
no measurable slowdown (collectives ~9% faster at 64 MiB payloads).
Forward-path credit: Ying Zhang's yingz/mega-sm-headroom (95046f0, opened as
#14) -- same mechanism, fixed value 2; this PR generalizes it to all
cooperative sites and makes it a deployment knob.
Known gap: reservation narrows the window but does not close it -- a
preemptive long-lived kernel occupying SMs before the mega launch still
starves the grid. A structural fix needs residency-agnostic synchronization
or driver-validated cooperative launch.
Co-Authored-By: Claude Code <noreply@anthropic.com>
…_MOE_SM_HEADROOM) (#13) The SM100 FP8/FP4 and BF16 mega MoE kernels (forward + backward) launch cooperative grids whose grid_sync / nvlink_barrier (comm/barrier.cuh) require ALL grid CTAs simultaneously resident, with a hard 60s timeout that traps and takes down the process (SIGABRT) and the GPU (Xid 43 fallout). An ordinary launch gives no gang-scheduling guarantee: CUDA places CTAs greedily, so a concurrent kernel holding some (but not all) SMs leaves part of the grid queued while resident CTAs spin against the deadline and never yield their slots -- mutual deadlock. Reproduced locally (train-firetitan:0.638.337, 2x B300 -- same SKU as the production node): a single concurrent smem-heavy kernel on one rank's GPU delays that rank's grid; its own CTAs time out at barrier.cuh:39 ('Grid sync timeout') and peer ranks die at barrier.cuh:80 ('NVLink barrier timeout') -- the exact INC-1291 signature (figma RLOR xy98ftpyg7kovb4m, AP_MALAYSIA_2, 2026-09-08, forward dispatch barrier tag=1 per production Chronosphere logs). With the production BF16 kernel and real FSDP-shaped collectives, all three collective types (reduce-scatter / all-gather / HSDP all-reduce) abort the stock build; none abort with headroom + pinned NCCL channels. Reserve headroom via a shared get_mega_moe_num_sms() helper used by the fp8_fp4 forward, all four fp8_fp4 backward launches, and the bf16 default (the absolute DG_BF16_MEGA_MOE_NUM_SMS override is retained). The amount is controlled exclusively by DG_MEGA_MOE_SM_HEADROOM; it has NO built-in default -- unset means 0, the historical full-device grid -- because the right value is deployment-specific (serving SendRecv: 2; training FSDP2 collectives with 32 NCCL channels: 8 alongside NCCL_MAX_NCHANNELS=8). The value is rounded up to even, since these are 2-CTA cluster launches and several sites assert num_sms % 2 == 0. Measured sizing for reference (8x B300, production BF16 kernel): DG_MEGA_MOE_SM_HEADROOM=8 + NCCL_MAX_NCHANNELS=8 + NCCL_MIN_NCHANNELS=8 + NCCL_NVLS_ENABLE=0: zero aborts across 90 mixed + 200 e2e iterations, no measurable slowdown (collectives ~9% faster at 64 MiB payloads). Forward-path credit: Ying Zhang's yingz/mega-sm-headroom (95046f0, opened as #14) -- same mechanism, fixed value 2; this PR generalizes it to all cooperative sites and makes it a deployment knob. Known gap: reservation narrows the window but does not close it -- a preemptive long-lived kernel occupying SMs before the mega launch still starves the grid. A structural fix needs residency-agnostic synchronization or driver-validated cooperative launch. Co-authored-by: Claude Code <noreply@anthropic.com>


What this is
This opens Ying Zhang's existing
yingz/mega-sm-headroombranch as a PR. The work and the codeare Ying's (commit
95046f0, 2026-06-27); I am only surfacing it for review because the samedefect just hit us again in training. @ipiszy — please take this over / correct anything below.
Net diff against
mainis one file, +11/-1, and merges cleanly (the intermediateDEBUG reserve 16 SMscommit is reverted by the last commit, so the branch nets out to the 2-SMreservation on the fp8_fp4 forward launch).
Why it matters now
The fix was written for disaggregated serving — a concurrent NCCL
ncclDevKernel_SendRecv(KVtransfer) grabbing an SM the cooperative mega grid needs, leaving it 1-2 CTAs short so the
grid/NVLink barrier never completes (the "signal=4 target=8" crash on GB300 pf8).
It was never opened as a PR and never merged, so it is absent from
affab4d— the DeepGEMMcommit shipped in the production trainer image. On 2026-09-08 the same defect class took down an
RLOR training job (INC-1291): DeepGEMM
barrier.cuh:39"Grid sync timeout" +barrier.cuh:80"NVLink barrier timeout" mid-
forward_backward, then SIGABRT and Xid 43 on all 8 GPUs.I reproduced it end-to-end on healthy hardware (including with the production BF16 kernel and
real FSDP-shaped collectives — all three collective types are fatal; see the comment below for
the full analysis and the production-log decode). Full RCA in fw-ai/fireworks#47752.
Relationship to #13
#13 re-applies the same one-line change on the current base, makes the amount configurable
(
DG_MEGA_MOE_SM_HEADROOM, default 2), and extends it to the launches this branch does notcover (the 4 fp8_fp4 backward launches + the bf16 default).
Suggested order: merge this PR first so the original fix lands with correct provenance, then
rebase #13 down to only the extension. Happy to do that rebase, or to close #13 entirely if you'd
rather fold the extension in here.