Atlas-Inf/atlas · main @ 4df6e609 · plan approved 2026-09-03 · executed 2026-09-04
Seven PR-sized items, smallest-safe first, derived from the five defects hit on the GB10 bench and two independent code investigations (prod 27B and Flash-Next Q4) whose citations were re-verified against the tree. Each names the exact code, the change, the risk, the regression test and how it is verified on spark. All seven items are implemented on the fork branch assist/flashnext-longform-fixes (seven commits on 4df6e609); status per item is in the table. No upstream PR has been opened.
Two independent agentic investigations over the Atlas main tree with the same brief and tools (grep/read over the repo, the bench findings, and the PR #16 review): one driven by prod (abliterated Qwen3.8-27B, 112 steps, 120 tool calls, 85 min, 21 KB report) and one by Flash-Next Q4 on llama.cpp (config A, 105 steps, 126 tool calls, 92 min, 38 KB report). Where they disagreed, the tree decided; the second-opinion section at the end records what each got right. Every file:line below was then re-read by me against the tree; the two places the prod agent said it could not confirm were located by hand and are marked as such.
| # | Item | Fixes | Size | Risk | Needed for long-form? | Status |
|---|---|---|---|---|---|---|
| 1 | Derive the PLE token cap and n-gram row cache from the serve config | D-3a, D-3b | S | low | yes | landed e623e010 |
| 2 | Tie ATLAS_QSA_MAX_TOKENS to --max-seq-len; end with length, not 500 | D-5 | S | low | yes | landed 561be215 |
| 3 | Never launder an engine error into finish_reason=stop; log the cuBLASLt failure before the scalar GEMM fallback | D-2b, D-4 observability | S | low | yes (diagnosis) | landed 16af4d33 |
| 4 | Treat sticky CUDA errors as fatal: exit non-zero, /health → 503 | D-4 aftermath | S | low | yes (operability) | landed 35028544 · verified live |
| 5 | Watchdog rollback restores QSA/PLE aux state from the boundary snapshot, or declines | D-1 | M | medium | yes if the watchdog stays on | landed f4e81291 · cursor off by one at resume, open |
| 6 | Latch MTP off before a sequence can cross the QSA inert bound mid-verify | D-2a | S | low | no (speculative only) | landed 86bdf303 · verified live |
| 7 | Continuation-chunk prefill fault: sanitizer repro, typed row offsets, then fix the misaligned operand | D-4 | M–L | unknown until repro | yes, unless single-chunk prefill is accepted as the workaround | landed 97aa9fa4 · root cause found, verified |
PLE: 8193 tokens exceeds the 8192 this layer was sized for, then, with the cap raised, NgramRowCache: every one of 65536 slots is pinned by the batch in flight.crates/spark-server/src/main_modules/serve_phases/kv_cache.rs:114 sets max_batch_tokens = prefill_budget + max_batch_size, i.e. 8,193 by default; a prefill chunk is exactly that wide.
crates/spark-model/src/weight_loader/qwen4_exp.rs:197 derives the PLE floor from max_batch_tokens (correct), but an ATLAS_PLE_MAX_TOKENS override below it is honoured after a warning, and the refusal fires at crates/spark-model/src/layers/ple/layer.rs:267.
crates/spark-model/src/weight_loader/qwen4_exp/ple.rs:43-48: slots_from_env() defaults the row cache to 65,536 slots; a chunk pins tokens × ngram_heads rows (8,193 × 16 = 131,088), and crates/spark-storage/src/ngram_cache.rs:450 bails when nothing is evictable.
ATLAS_PLE_MAX_TOKENS below the derived floor as an error at load (or clamp up with a warning), not a warning that guarantees a later 500.max_batch_tokens × ngram_heads rounded up to a power of two (131,072 → 262,144 slots ≈ 42–84 MB pinned for the 160-byte FP8 rows), keeping ATLAS_PLE_CACHE_SLOTS as an override. Log the derived numbers next to the existing "PLE at MODEL LAYER 1 … served off NVMe with N cached slots" line.--max-prefill-tokens / --max-batch-size.ngram_cache test opening a cache sized from (8193, 16) and pinning 131,088 rows without a bail. Serve probe on spark: default flags, 9,000-token prompt, expect 200.ATLAS_QSA_MAX_TOKENS to --max-seq-len and end with finish_reason=lengthQSA: pos 32768 >= ATLAS_QSA_MAX_TOKENS despite --max-seq-len 65536.crates/spark-model/src/layers/qsa.rs:131-134 hard-codes the default 32,768 for the per-sequence raw-key ring and compressed-key cache; the guards are at qsa.rs:208-213 (prefill) and qsa.rs:297-300 (decode).max_seq_len into QsaIndexer::new and use it as the default (env stays as an override, validated ≥ max_seq_len at load). Cost at 65,536 tokens: about 2× today's 10.5 MB per sequence. Add an admission-time check (config A's addition): reject with HTTP 400 when prompt_tokens + max_tokens exceeds the indexer capacity, and when the client sent no max_tokens clamp the budget to capacity − prompt and log once. Keep the ensure!s in qsa.rs as last-resort bounds guards with a message naming max_seq_len. Per-sequence cost at 64K: about 240 MB across the 12 QSA layers; publish it in the startup log.max_seq_len = 40,000, assert decode_select at pos 39,999 succeeds and the cap equals the config. Serve probe: --max-seq-len 40000, 30K-token prompt, 5K completion, expect 200 with stop or length.stop after 1,679 tokens. (b) The continuation-chunk CUDA fault was preceded by a cuBLASLt failure that left no trace in the log.crates/spark-server/src/scheduler/verify_k2_step.rs:130-142: both verify arms do tracing::error!; a.finished = true; return; on Err, which the response layer then reports as a clean completion. (Located by hand; the prod agent could not open this file.)
crates/spark-model/src/layers/qwen3_ssm/trait_prefill_proj.rs:380-403: the GDN QKVZ projection tries ops::cublas_bf16_proj_dense(...) and on .is_err() silently launches the scalar ops::dense_gemm (gemm_dense.rs:90-103, grid ceil(N/16) × ceil(M/16), block 16×16). The faulting launch in our log, grid=[1024,70,1] block=[16,16,1], is exactly that fallback with N = 16,384 and M = 1,118 tokens.
finish_reason of error, or an HTTP 500 if nothing has been streamed) instead of a.finished = true. Keep the log line.warn before taking the scalar fallback, and add a debug assertion that normed and proj_dst are 16-byte aligned. Same for the other two cublas … .is_err() fallbacks in that file.stop will now see an error, which is the point.verify_k2_step: inject an Err from the verify and assert the finish reason is not stop. Serve probe: re-run the D-2 configuration and confirm a 500 or error instead of a truncated success./tokenize, returned 500s for everything else, and held ~100 GB until SIGKILL. Its own driver did not notice it had died.atlas::fault already recognises the sticky case ("a no-op synchronize issued afterwards also failed … the CUDA context is destroyed") and in one path logs "Exiting after a fatal GPU fault", but the fix-round path (phase_continue_prefills::run_standard) returned an error to the request and continued.GpuError::Fatal in the backend error mapping, propagate it as a shutdown signal from every scheduler phase, and have /health return 503 once a fatal error is recorded so supervisors and load balancers react.ATLAS_TEST_FAULT=1 that injects a sticky error on the next step; assert the process exits non-zero within one step and /health flips to 503 first.QSA: decode at pos 2864 but 2932 tokens ingested — the indexer cache lost sync.crates/spark-model/src/layers/qsa_aux.rs:92-110 rewind_verify(st, rejected) is a delta rewind (ingested -= rejected) whose comment at 77-91 states it assumes pairing with the verify that just ingested k rows.
That hook (crates/spark-model/src/layers/qwen3_attention/trait_impl.rs:122-147 → rewind_verify) is called from exactly one place, the speculative verify commit at crates/spark-model/src/model/trait_impl/async_chkpt.rs:266-277. The watchdog's rollback_to_boundary (crates/spark-server/src/scheduler/rollback.rs:203-290) is SSM-only: it picks a boundary that has a live SSM ring snapshot, restores it, lowers seq_len and rewinds the grammar FSM, and never touches aux state (no aux reference in the file). The aux snapshot/restore API that does exist (snapshot_aux/restore_aux, trait_impl.rs:149-188; collect_aux_states/apply_aux_states, model/trait_impl/mod.rs:897-935) is only used by prefill and checkpoint paths. So after a rollback of dropped tokens, seq_len is lower while ingested is unchanged, and the next decode_select fails pos == st.ingested at qsa.rs:291. Note table_len (qsa.rs:410-413) is a monotone high-water mark, so any arithmetic rewind must also reset it.
snapshot_aux in each SsmDecodeRing entry, restore them with restore_aux in rollback_to_boundary next to the SSM restore, and fold "entry carries aux" into the eligibility predicate with a new RollbackFallback::NoAuxSnapshot so a rollback that cannot be made exact declines into the existing hard stop instead of desyncing. A lighter arithmetic variant (ingested -= dropped; pooled = ingested/ratio; table_len = 0 plus PLE history truncation) is acceptable as a stopgap only behind a flag, since PLE conv carry cannot be rewound exactly that way.spark-server around rewind_buffers with a fake model asserting the aux restore is invoked once per rollback and that a ring entry without aux yields NoAuxSnapshot. Unit in qsa_tests.rs: ingest past the inert bound, roll back 3 tokens, decode_select must succeed and select the same blocks as a fresh sequence that never took them. Serve probe: stream:false (streaming declines rollback as unsafe), watchdog on, a prompt that loops, QSA active; expect 200 and zero "indexer cache lost sync" lines.--speculative --num-drafts 1, generation dies at position 2,051: QSA selection active for row 1 on the batched ms path; the dispatch gate should have routed this batch per-seq.The gate exists for plain decode: crates/spark-model/src/model/trait_impl/decode_a2.rs:127-136 computes qsa_active from index_topk + index_compress_ratio - 1 and forces the per-sequence loop. The batched multi-seq attention refuses loudly when selection turns out active: crates/spark-model/src/layers/qwen3_attention/trait_impl/multi_seq/mod.rs:247-288.
The MTP lane does have a gate, found by config A: crates/spark-server/src/scheduler/mod.rs:757-762 admits the batch only if every sequence has seq_len < verify_context_limit (2,051 here), with the in-tree comment "a verify error finishes the request — 10 of them died that way in a BFCL subset run before this gate". It is evaluated at dispatch time only: a sequence at 2,049 is admitted, the verify ingests k rows across the bound, and the next verify runs with selection active, which the batched attention refuses at multi_seq/mod.rs:284-288. That is why our run died at exactly 2,051 and why it looks intermittent.
seq_len + num_drafts + 1 ≥ limit, set the existing a.disable_mtp (already honoured by spec_dispatch_eligible at mod.rs:774) for the rest of the request, and log mtp declined: qsa_active once. Add a pre-mutation bail in the verify entry (VerifyUnsupportedWithActiveQsa) as a belt-and-braces so a future caller cannot mutate state before refusing. Making the verify QSA-aware (applying selection per row past the bound) is the larger follow-up and is not needed for correctness.seq_len ∈ {2048..2052} × num_drafts ∈ {1,2,3} asserting Serial whenever the latch condition holds. Serve probe: --speculative --num-drafts 1, prompt padded to bound − 8 tokens, 64-token completion, stream:false; expect finish_reason=length and token-for-token equality with the same request run with MTP off, which is the only check that also catches a silently dense verify.CUDA_ERROR_MISALIGNED_ADDRESS in chunk 2+ of a multi-chunk prefill, layer 0, in the scalar QKVZ GEMM fallback; chunk 0 never faults; not explained by chunk-width parity (1,117 and 1,118-wide last chunks both fault, 1,862 and 4,092 do not). After the fault the process keeps the dead context, holds ~100 GB and keeps answering /tokenize.Chunk widths after the first are max_prefill_tokens − 4 wide (4,096 → 4,092), so continuation chunks start at non-power-of-two token offsets (8,188, 12,280 …). The projection is launched over normed (the chunk's normalized hidden slice) into proj_dst with K = 10,240 and N = 16,384. cuBLASLt rejects the operands first (item 3 makes that visible), then the scalar kernel faults on them: the misalignment is in a pointer, not the kernel. Candidates are the arena sub-buffer offsets computed from the chunk's start token and the per-chunk staging of the GDN state, both in crates/spark-server/src/scheduler/phase_continue_prefills/run_standard.rs and the qwen3_ssm prefill entry.
compute-sanitizer --tool memcheck with item 3's logging in place: the 13,397-token prompt at --max-prefill-tokens 4095 faults in under a minute on spark. The sanitizer names the kernel and address; the log names the operand.RowOff element-count newtype for every DevicePtr::offset in the chunked paths, so token counts and byte offsets cannot be confused, plus debug_assert!(base % 16 == 0 && pitch % 16 == 0) on every pitched copy and GEMM operand.On the candidate operand the two investigations differ, and the log sides with the prod reading: the faulting launch shape matches the GDN QKVZ scalar GEMM (a layer-0 GDN layer has no QSA), so config A's suspects in the QSA prefill-ingest pitched copy (qsa.rs:231-239) and the block-table upload (qsa.rs:169) would show up on an attention layer, not layer 0. They stay on the list for the sanitizer run because chunk-2 asymmetry is documented there (qsa.rs:100-103).
--max-prefill-tokens 4095 and the default 8,192, all must prefill. Unit: a prefill-chunk planner test asserting every chunk's byte offsets are 256-byte aligned for the model's hidden size.ATLAS_QSA_MAX_TOKENS=<max-seq-len> ATLAS_PLE_CACHE_SLOTS=524288 ATLAS_PLE_MAX_TOKENS=<max-prefill-tokens + 20> \
spark serve … --max-prefill-tokens 20480 --content-loop-watchdog false --request-timeout 3600 # no --speculativeSingle-chunk prefill costs about 8 GB of arena at 20K tokens; with 48K context and 0.95 utilization the box ran with 2–5 GB to spare and once touched 0 GB of system memory. It is a bench workaround, not a deployment recommendation.
ATLAS_QSA_DISABLE) to dodge items 4–5: dense attention past the budget is "not the reference model", as the code comments say, and the benchmark would no longer measure the shipped architecture.gpu-memory-utilization in the recipe to make room: the registry maintainers report 0.90 freezing a GB10 hard; the memory should come from smaller PLE scratch and arena defaults instead.Both agents converged on defects 3 and 5 with identical citations. Where they differed:
rollback_aux_verify". Wrong: the tree shows that hook has exactly one caller (the verify commit) and the watchdog rollback is SSM-only. Config A got this right and its snapshot-based restore with a decline gate is what item 5 now proposes.scheduler/mod.rs:757-762 and diagnosed it as an instantaneous check that leaks at the boundary crossing; verified. Its one-field latch replaces the heavier re-routing prod proposed. Neither agent opened verify_k2_step.rs:130-142; that is where the error becomes a plain "finished" (located by hand).file:line either agent cited that this plan relies on was re-read against 4df6e609.Sources: investigation reports (reports/atlas-investigation-prod-2026-09-03.md, reports/atlas-investigation-flashnext-A-2026-09-03.md; spark ~/agentbench/out7/{P,A}/), bench findings (~/claude/atlas-flashnext-findings-2026-09-03.md), serve logs under spark:~/agentbench/atlas-td2-*.log, and the tree at spark:~/agentbench/atlas-work (main @ 4df6e609).
Items 1–7 are on the branch behind PR #23. Item 5's residual off-by-one is resolved: it was the token-buffer rewind, not the snapshot (the two watchdog sites sit on different sides of the step's push; the rewind now pops by fed tokens and the stale sampled token is dropped). Verified on the GB10 with five rollbacks and no desync.
Second batch, done: preset penalties opt-out for core requests (the presence 1.5 that killed long files), BF16 lm_head default for Flash-Next, --attn-proj-dtype bf16, the rollback rework. Measured neutral: spin-sync and PLE fault-thread depth. Deferred by request: the admission-time 400/clamp. Proposed next (code): device-side QSA top-k + lift the graph veto; PLE event instead of the per-layer barrier; single-launch hyper-connection sites; a code-aware content-loop watchdog (it false-positives on coordinate arrays).