After a 36,689-token request at 155k, every later request answers one repeated token until the engine reloads

Summary

One request generated 36,689 tokens at a 155,763-token prompt. After it, the server stayed up, /health stayed loaded: true, decode stayed at 157–183 tok/s — and every subsequent request answered with a single character repeated, including a 20-token question. Sampling parameters make no difference at all: six requests with temperature 0 / absent / 1.0, repetition_penalty 1.05 / 1.3 and penalty_last_n 64 / 4096 produced byte-identical output (!!!!…, 1,500 of them). Only stopping the service and letting the engine reload its experts fixed it.

This looks like in-process state that survives the end of a request (here: a request the client aborted), keeps feeding a constant token to later requests, and is invisible to every health check we have.

Environment

engine 0.1.35 (prebuilt Linux / CUDA 13), serve/server.py from the same tag
RTX 5090 D 32 GB, driver 595.91.07; Ryzen 9 9950X3D; 64 GB kit (OS sees 59.2 GiB)
resident low-RAM mode, --resident-experts: 39.82 GiB of experts page-locked, 0 blob reads from the file
model: Qwen3.8-Flash-Next IQ3_S (original, GSQ-RCO), --max-context 224000 --kv int8 --prefill auto:32768 --spec 3 --spec-min-p 0.5 --mtp …/rt --expert-cache auto --vision --vram-reserve-mib 700 --api-key …
run config sampling block: temperature 0.6, top_p 0.9, top_k 30, min_p 0.05, repetition_penalty 1.05
one slot in play, --conversation-cache-mib 0 (off)

How it unfolded (serve log, the same process throughout)

prompt 153,171 … 876  generated in 6,095 ms (143.7 tok/s), drafts accepted 596 of 688
prompt 154,112 … 958  generated in 8,309 ms (115.3),         drafts accepted 445 of 643   hit rate 92.9%
prompt 155,121 … 594  generated in 3,961 ms (150.0),         drafts accepted 429 of 472   hit rate 88.6%
prompt 155,763 = 155,715 reused + 48 read, 36,689 generated in 232,175 ms (158.0),
                       drafts accepted 18,441 of 18,551                                   hit rate 99.5%
        suffix drafts: 18,069 windows, 18,296 of 18,316 drafts accepted
prompt 156,132 … 780  generated in 4,957 ms (157.3),         drafts accepted 388 of 388   hit rate 99.6%
prompt 156,019 … 348  generated in 2,210 ms (157.4),         drafts accepted 172 of 172   hit rate 100.0%
        resident RAM: 39.82 GiB in RAM, exchanged with the VRAM tier: +3,039 → +85 → +8 per request

The client aborted that 232-second request (finish recorded as a disconnect). From then on everything collapsed, at any prompt size:

prompt	request	output
23 tokens	plain question, sampling default (0.6)	1,500 × `!`
23 tokens	`temperature: 0`	1,500 × `!` — identical text
23 tokens	`temperature: 1.0`	1,500 × `!` — identical text
23 tokens	`repetition_penalty: 1.3, penalty_last_n: 4096`	1,500 × `!` — identical text
8,192 / 65,536 tokens	same task	1,500 × `!`

After systemctl --user restart (engine reloads 40 GiB of experts): the same 23-token, 8k and 64k prompts all answer normally, drafts accepted back to 57–84%, longest same-character run 1–2.

The tell I would like your read on

Going from temperature: 0 to temperature: 1.0 cannot produce byte-identical text unless the sampler is no longer what picks the token. So whatever is stuck sits upstream of sampling — yet it is not per-request, because a brand-new request with a 20-token prompt inherits it. Candidates I can see from the log: the resident expert mapping in src/core/expert_source.cpp, and the prompt path's borrowed cache slots ("the prompt path borrows 7,233 CUDA0 cache slots"), with that request sitting at 6 checkpoints.

What I could NOT reproduce (each ended with a correct answer)

150k-token prompt + max_tokens 40000 on a long list task → wrote ~7k characters and stopped at end_turn
aborting mid-decode at 150k twice (~700 KB of SSE streamed, then the socket closed)
aborting during prefill at 64k three times

So the trigger looks content- or history-dependent and I cannot hand you a small reproducer. What I can hand you is a cheap observable: longest run of one repeated character ≥ 20 in the answer. We now run a canary every 5 minutes on that rule and restart when it fires.

Two asks

Which in-process state can outlive a request (or a cancelled one) and pin the output to a constant token? serve engine deadlocks in SleepConditionVariableSRW during a long streaming reply; stall watchdog never fires #481 (a silent engine) is the neighbouring failure, but nothing went silent here — the engine kept answering at full speed with the same character, so the silence watchdog would never fire.
A degeneration guard inside the engine? Re Reported expert-cache hit rate excludes --pcie-frac experts, so it rises as tok/s falls #588: I read it, and the --pcie-frac denominator artifact does not cover this — in resident mode there were 0 blob reads from the file, and the reported hit rate rose to ~100% while decode stayed at 157–183 tok/s rather than falling. My reading is that hit rate ≈ 100% + drafts accepted ≈ 100% + one character repeating is a symptom triple: a constant token stream makes every expert lookup and every suffix-lookup draft trivially right. If you agree, that triple would make an inexpensive in-engine breaker — with --suffix-draft on, a collapse is fast (158 tok/s) and, because a client that omits max_tokens gets max_new = room (66,967 here), it can run for four minutes before anyone stops it.

Side note, not the bug

serve/server.py gives a request that omits max_tokens the whole remaining context as its runway. POST /settings can cap it, but that also changes the default for every client that leaves the field out, so an engine-side "same token N times in a row → stop" would be the better knob.

Happy to test candidates on this box (5090 D 32 GB / 64 GB RAM / IQ3_S resident / 224k ctx), and to re-run this on a newer engine — we are on 0.1.35 and I see 0.1.36–0.1.38 landed since, including the #463 greedy-determinism change.

@mcoffin

Follow-up with a concrete lead on the "first non-finite value" half, from reading this engine's own kernels (0.1.35 working tree; I also checked the v0.1.38 tag — the code is unchanged there apart from being refactored into a q8_1_store helper).

Both q8_1 activation quantizers store the block sum as the fp16 of the sum of 32 raw activations:

src/kernels/cuda/iq_kernels.cu — 0.1.35 lines 979-994, 0.1.38 q8_1_store lines 1022-1034:

float amax = fabsf(xi), sum = xi;
#pragma unroll
for (int o = 16; o > 0; o >>= 1) {
    amax = fmaxf(amax, __shfl_xor_sync(0xffffffffu, amax, o));
    sum += __shfl_xor_sync(0xffffffffu, sum, o);
}
const float d = amax / 127.0f;
const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
y[ib].qs[iqs] = q;
if (iqs == 0) y[ib].ds = make_half2(d, sum);   // fp16: |sum| > 65504 -> +-inf

and the same expression in src/kernels/cuda/native_mmvq.cu:141-152 (native_quantize_q8_1_kernel).

That is the defect isolated in ggml-org/llama.cpp#23606 ("store the q8_1 block sums as sums of the quantized values. The fp16 sum of 32 raw activations overflows to -inf with large activations … a deterministic NaN source on master for models with large outlier activations"). A block only needs |sum| > 65504 — one massive-activation dimension of that magnitude inside 32 consecutive values is enough, and this architecture is known for them. Deterministic, not a race. d = amax/127 overflows the same way at amax > 8.3M.

Call sites, i.e. how much of the forward pass this covers:

src/kernels/cuda/iq_kernels.cu:1463, inside native_expert_grouped — the main MoE expert GEMM
src/kernels/cuda/shared_expert.cu:175, 180, 267, 285 — shared-expert x and gate
src/core/verify.cpp:717 — the speculative verify window
src/core/remote_experts.cpp:297 — the --pcie-frac expert path (we run pcie_frac 0.55)

On the persistence half, our engine is not the scale-by-0 case: src/core/session.cpp:139-160 (session_zero) memsets the GDN recurrent state, the QSA state, the PLE conv history and the residual, and src/program/generate.cpp calls it at 2460 / 3261 / 5016-5039, with the comment that a warm conv history "is a real contribution and not a zero". So a fresh sequence does start from zeros here — yet a 23-token request still answered !×1500 after the event, and only a process restart (re-reading 40 GiB of experts) cleared it. That points at state which outlives session_zero: the resident expert pool (39.82 GiB page-locked, 206,168 VRAM-tier exchanges, 0 blob reads from the file) or the prompt-cache checkpoints. Two built-in probes should settle it and I will run them: STRATA_STATE_HASH / STRATA_STATE_HASH_GDN (generate.cpp:5519-5567, prefill.cpp:1960) and STRATA_DBG_NAN (prefill.cpp:1814, "the first layer of a chunk whose MoE produced non-finite values").

One note on STRATA_DBG_NAN's cost: per layer per chunk it D2H-copies T*K*1280 + T*K*N + T*N + T*K*640 floats and scans them on the host, so with --prefill auto:32768 it is unusable on a live box. If that check were per chunk instead of per layer, or sampled (the last row of each chunk would have caught ours), it could stay on in production — that is the "pretty simple error detection for things like NaN floats in persistent buffers" @mcoffin asked for, and it is what turns this from a dead server into one bad response.

Our mitigation meanwhile is a server-side canary: a fixed 20-token prompt every 5 minutes, and "longest run of one repeated character ≥ 20" ⇒ dump /metrics plus the log tail, then restart. It bounds the damage at about 5 minutes; before it, one collapse ran 232 s and emitted 36,689 tokens, because a request that omits max_tokens is given max_new = room (66,967 in our case).

Hardware note for anyone else reading: our card is an RTX 5090 D (ecc.mode.current = N/A, so no ECC) with a boot-time overclock applied by a systemd service — nvidia-smi --query-gpu=clocks.max.mem reads 16001 MHz where stock is 14001, i.e. +2000 MHz on the memory (32 Gbps) and +100 on the core. Given @PippBauda's finding — 5 single-bit VRAM errors in 30 minutes at stock clock, 0 errors with the memory clock 500 MHz lower — I am treating our memory clock as a co-suspect for the first non-finite value and will report back with it lowered.

Just want to add that I encountered the same problem, although model was a Q8 I had my AI adapt to use.
Max context set to 240K, with INT8 KV. Completely fine for 157K tokens, then started repeating one word on the next turn

5090, 192GB DDR5 RAM

q8_1 activations: keep the block's fp16 scale and sum finite (#606)

serve: end a reply that repeats one token 256 times in a row (#606)

serve tests: the long fixture answer alternates two tokens (one repea…

Engine 0.1.39: the #577 file-tier regression fixed, the OpenAI Respon…

sycl: port 0.1.39 (re-migration with dpct 2025.3, 3-way merge per doc…

Thank you for the detailed write-up and for pointing at the q8_1 sums. 0.1.39 has two changes. (1) Engine: both q8_1 activation quantizers (iq_kernels.cu and native_mmvq.cu) now clamp the block's fp16 sum and scale to the largest finite half, so a block of huge activations gives a finite, approximate product instead of inf and NaN. Every block that was finite before is stored bit for bit as before, so normal answers do not change. (2) Server: a reply that repeats one token 256 times in a row is ended there with finish_reason "length" and a line in the server window, so a broken state costs one short bad reply instead of minutes. "repeat_stop_tokens": N in strata-.json sets the run length, and 0 turns it off. If it happens again on 0.1.39, a short new request that also repeats would point to state that outlives the request. Please post the engine log around it then. Your memory-clock retest would also help.

Following the #606 fix one path further: the fused decode SwiGLU quantiser stores its block summary raw, so the
clamp is missing on one of the three q8_1 activations quantisers.

src/kernels/cuda/native_mmvq.cu:166-169 (main, 2026-10-04):

const float d = amax / 127.0f;
const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
y[i / Q8K].qs[i % Q8K] = q;
if (i % Q8K == 0) y[i / Q8K].ds = make_half2(d, sum);

native_quantize_q8_1_kernel, eight lines above it in the same file, uses q8_1_finite(amax / 127.0f),
q8_1_quant and q8_1_ds — and q8_1_finite.hpp is already included at the top of that file. The IQ path's
q8_1_store (iq_kernels.cu:1340) has it too. So this looks like the one site a sweep missed, and the #606
parity check (iq_multi_parity.cpp:329) covers the other two quantisers by name, which is probably why.

Both call sites are the shared expert — decode and multi-token:

src/kernels/cuda/shared_expert.cu:203   decode
src/kernels/cuda/shared_expert.cu:310   multi-token

behind fused_swiglu_q81_enabled() (shared_expert.cu:179-185), which is on by default (STRATA_FUSED_SWIGLU_Q81
unset → on).

What the missing clamp costs, from the consumers of the block:

stored field	read by	if it is `inf`
`ds.x` (scale)	every dot: Q4_K/Q5_K (`d8[i] = __low2float(bq8i->ds)`), Q5_1, Q8_0, IQ4_NL, Q4_0	the down projection returns ±inf, and NaN where the group's integer sum is 0 — into the residual
`ds.y` (sum)	only Q4_0 and Q5_0 (`d * (sumi * ds.x - {4,8} * ds.y)`, `native_mmvq.cu:710/735`)	±inf / NaN in those two dots

Reachability, stated plainly because it is the part I cannot test here:

scale: only above amax > 8.3e6 (d = amax/127). That is the margin numerics: saturate the prompt path's FP16 SwiGLU products; log1pf in the unfused GDN gate #281 named when it kept the decode
q8_1 path and fixed the prompt path instead ("range ~8e6 stays finite") — the clamp is what turns a token past
that margin into a finite approximation rather than an inf, which is precisely what After a 36,689-token request at 155k, every later request answers one repeated token until the engine reloads #606 added to the other
two quantisers.
sum: passes 65504 far more easily (32 finite products can sum past it) — but it is inert while the
shared expert's down type is one that ignores it. It would not be inert with a Q4_0 or Q5_0 shared-expert
down: native.down_type = w_sdown->native_type — src/core/layer.cpp:438 — comes from the model file.

What I measured, and what I did not. I have no NVIDIA card, so there is no end-to-end repro and I am not
claiming one. What I do have is an independent Vulkan port of these kernels with a gate that reproduces the block
bytes: with gate = up = 3000.0 (32 products of ~9e6, so amax and the sum both past fp16 range) the unclamped
rule stores ds = inf/inf where the clamped rule stores 65504/65504, byte for byte. Happy to add that as a
regression case in the port once a fix lands, or to run a specific configuration if a repro on CUDA hardware
would help.

If you agree it is an oversight, the fix is the sibling's three lines: q8_1_finite on d and on sum before
make_half2, and the same for the else branch's two-step path (native_swiglu_kernel → native_quantize_q8_1)
already goes through the clamped quantiser, so the fused branch is the only one affected.

Merge upstream main (0.1.39) into fix/tool-call-opens-inside-thinking

@rwkeyes

@rwkeyes's finding, with a CUDA test: #838 clamps native_swiglu_quantize_q8_1_kernel and the gfx906 native_gu_fused_kernel (the same raw store, STRATA_EXP_MODE=8). iq_multi_parity now runs the fused SwiGLU quantizer through the #606 check. On an RTX 5090 it fails on v0.1.39's kernel (0 of 3 overflowed blocks clamped) and passes with the change, with every finite block unchanged.

Data point from a 0.1.39 fork, RTX 5090: a served conversation degenerated into '!' at about 99K positions, mid-thinking. Later requests carrying that history repeated it. Replaying the conversation up to that point did not reproduce it, so it is the same symptom, not a confirmed cause.

vulkan: the q8_1 activation quantiser, byte-exact - and the one arith…

vulkan: the fused SwiGLU+q8_1 quantiser - and the engine's two rules …

docs: the upstream follow-up on the missing Niko1221#606 clamp (Niko1…

q8_1 activations: keep the block's fp16 scale and sum finite (upstream …

Upstream's fixes in this branch's engine: finite q8_1 scales, the NaN…

q8_1 activations: keep the block's fp16 scale and sum finite (upstream …

Thanks for the report. To keep the tracker current we are closing issues older than three days: 0.1.44 is out with a long list of fixes, and many older reports are covered by it or by 0.1.43.1.

If this still happens on 0.1.44, please leave a comment here with your engine log (and the version line it prints), and we will reopen it. A new issue that links this one works too.

Niko1221#606: clamp the three unclamped q8_1 scale/sum emit sites via…

serve: the reasoning loop check also sees a loop written as one long …

Non-finite logits (Niko1221#879): a guard with self-healing

serve: stop_cause only with repeat_stop_period set; the repetition er…

maxwell: the i-quant gate/up dots and the batched IQ gup kernel (sm_50)

serve: stop_cause only with repeat_stop_period set; the repetition er…

At most GAP_MAX newlines are held before a call: no quadratic rescan

serve: a reply that is token 0 repeated restarts the engine before th…

hip: clamp the three q8_1 scale/sum emit sites (port of upstream Niko…

Merge pull request #1 from sub-xerox/q8_1-clamp

[Engine] clamp Q8_1 activation scales in SYCL native_mmvq (Niko1221#606)

Navigation Menu

Repository navigation

After a 36,689-token request at 155k, every later request answers one repeated token until the engine reloads #606

Summary

Environment

How it unfolded (serve log, the same process throughout)

The tell I would like your read on

What I could NOT reproduce (each ended with a correct answer)

Two asks

Side note, not the bug

49 remaining items

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

After a 36,689-token request at 155k, every later request answers one repeated token until the engine reloads #606

Description

Summary

Environment

How it unfolded (serve log, the same process throughout)

The tell I would like your read on

What I could NOT reproduce (each ended with a correct answer)

Two asks

Side note, not the bug

Activity

66419118nnn commented on Oct 3, 2026

mavyric commented on Oct 3, 2026

Niko1221 commented on Oct 4, 2026

rwkeyes commented on Oct 4, 2026

sergqwer commented on Oct 4, 2026

49 remaining items

Niko1221 commented on Oct 13, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions