Repository navigation
gfx1201 (R9700, HIP): the ! degeneration is #879 — GDN state overflow (max |x| 2.26e+32) at 16K–93K; 0.1.44.1's guard catches and recovers it on this box #2519
Description
Following the maintainer's note in #1468 ("if this still happens on 0.1.44, leave a comment here … a new issue that links this one works too"). This is a single-card gfx1201/HIP data point, and it comes with a mechanism trace plus a confirmation that 0.1.44's NaN guard catches this fault on a machine where your own 1,870-request R9700 run could not fire it.
Environment
| GPU | 1x Radeon AI PRO R9700 (gfx1201, RDNA4), 32 GB |
| Host | 62 GiB RAM, Linux 6.17, NVMe; no swap in use during any run |
| Model | Qwen3.8-Flash-Next GSQ-RCO IQ3_XXS (ISTA-DASLab/…, rev ed59f920; shard 1 47,039,860,096 B, sha256 219ea929900dfa9ef091f3aa473fdba6874b65fcb36526d7d851ac9e95856d15) |
| Build | local HIP build against the ROCm 7.10 "TheRock" userspace (setup's venv wheels, not system ROCm) |
| Versions | most measurements on 0.1.42; the guard confirmation below is on 0.1.44.1 (same box, same build chain, upgraded today) |
| Args | --max-context 131072 --kv int8 --prefill auto:32768 --spec 4 --spec-min-p 0.5 --mtp rt --vram-reserve-mib 1000 |
| Env | STRATA_PLE_BATCH=0; WMMA off; STRATA_PF_FUSED unset (0.1.42 default) / explicitly =0 on 0.1.44.1 |
1. The ! degeneration is the non-finite family — traced here on 0.1.42
STRATA_DBG_NAN=1 already exists in 0.1.42, so no guard is needed to see it. Same config, two 16K requests, one degenerate and one clean:
# degenerate (reply = 40 x '!')
strata dbg: layer 0 (mmq 1, types 17/20, 0 experts): non-finite GU 0 H -1 Dm 0 bo 33226240 of T 17075
strata dbg: prompt end: last residual row 10240 non-finite (max |x| 0); PLE history 0 non-finite (max |x| 21.1);
GDN state 0 9264 non-finite (max |x| 2.26e+32)
strata dbg: verify window at position 17075, row 0: 248320 of 248320 logits non-finite (token out 0)
# clean request, same config
strata dbg: prompt end: last residual row 0 non-finite (max |x| 5.13); PLE history 0 non-finite (max |x| 12.4);
GDN state 0 0 non-finite (max |x| 3.34)
248320 of 248320 logits non-finite— the same count the reporter in HIP layer split / gfx1030: batched prompt path produces NaN logits (output "!!!!") unless STRATA_HIP_PROMPT_F16=1 and the gfx1030 card is the first stage #1208 got on gfx1030 (whole vocab).GDN state … 9264 non-finite (max |x| 2.26e+32)— the recurrent state's magnitude has run away to 1e32 before the residual goes non-finite; the clean request prints 0 for both fields. So here the chain is: GDN state overflow → residual row non-finite → all logits NaN → sampler collapses to token 0 (!).- This fires at a 16K prompt, while the same box reads 63K/76K/85K cleanly — so on this hardware it is not a context-capacity effect.
- Cost of the instrument, for anyone reusing it:
STRATA_DBG_NAN=1takes a 16K-prompt prefill from ~1,400 tok/s to ~310 tok/s (~4.5x), and it prints informational (all-zero) lines on passing requests too.
2. On 0.1.44.1 your guard catches exactly this, and the retry recovers it
Interleaved arm, 40K cold needle prompts, fresh nonce per request, engine restarted per arm, with our amplifier STRATA_PF_FUSED=1 (see §5):
non-finite logits (#879) in a prompt window at position 38621 (row 0): its tokens are not emitted,
the request ends (DONE nonfinite) and the next request reads its prompt from token 0
non-finite logits (#879): dropped the live session, 2 checkpoints, 0 parked conversations and 0 idle slots …
The three probes in that arm all returned the correct answer (via the retry), and the log line names the position. On 0.1.42 the same fault was emitted as 40 ! characters with no log line at all, so from a user's point of view the guard turns silent corruption into a visible, recoverable event. Thank you for #1892.
3. Server-side behaviour worth documenting
After a degeneration, re-sending the same prompt is served from the prefix cache (prompt 38438 tokens = 38431 reused + 7 read in 38 ms) and degenerates again; a different prompt sent immediately after, same engine, answers correctly. So the poisoned state follows the conversation/prefix rather than the engine — consistent with #2452's retry-from-0 and with #879's "session-scoped, not engine-global". Users may be wasting a full re-read by retrying the same request; a new conversation recovers without a reload.
4. Two levers that do nothing here (so other 32 GB / gfx12 users don't spend the time)
STRATA_HIP_PROMPT_F16=1(HIP layer split / gfx1030: batched prompt path produces NaN logits (output "!!!!") unless STRATA_HIP_PROMPT_F16=1 and the gfx1030 card is the first stage #1208's fix on gfx1030): interleaved 8 rounds x 2 arms, 40K,STRATA_PF_FUSED=1in both arms, on 0.1.42 — 3/8 vs 2/8, Fisher one-tailed p = 0.50. It does not transfer to single-card gfx1201.STRATA_KV_STAGE_OWN=1:take_stage()returns unless the session is KV-streaming, so on a KV-resident config the variable has no object at all (our logs never containKV streaming:). We then built the condition it needs (--kv-resident 32768) and first verified the lever acts —the prompt path borrows 8066 → 7988 CUDA0 cache slots (13.14 → 13.01 GiB), the 13.01 GiB being exactly the resident-run figure your comment inprefill.cpppredicts — and the A/B is null: 1/8 vs 1/8 (Fisher p = 0.77), 0 engine faults in the 21 requests of that configuration. Worth a line in the docs: on a 32 GB card with resident KV this arm can only produce a false negative.
5. The amplifier, for anyone trying to reproduce
STRATA_PF_FUSED=1 (0.1.42, opt-in there): at 40K it raises the degeneration rate from 0/8 (baseline) to 3/8 (37.5%) in an interleaved arm — consistent with our earlier 6/18 vs 0/8. That is why we pin STRATA_PF_FUSED=0 in production and why the fact that 0.1.43 turned it on by default for IQ3_XXS matters to us.
6. And the other signature, for the record
prefill copy_i32: an illegal memory access at ~93K tokens (91,511, frozen at 65,536 of 91,511), ≈1/13 requests on 0.1.42 with the production config; 63K/76K/85K read cleanly. We have only one 93K pass on 0.1.44.1 so far, so we cannot say anything about whether that changed.
7. What we could not test
Our engine is built against the TheRock 7.10 wheels. The follow-up comment in #1208 reports the default batched MMQ path as finite on ROCm 10 system packages and NaN with the TheRock 7.10 wheels, which points at the ROCm side rather than the engine. We have not A/B'd ROCm userspaces here (this box also serves two other engines that depend on the current setup). If that comparison is worth having, say which userspace and we will build it.
We can also bisect, or dump token ids around a failing position, or run any specific env/arm on gfx1201 — the config above is the exact one. If STRATA_DBG_NAN-style counters at the GDN update are easier than at the logits, that would localise the overflow source faster than we can from outside.
Activity
MKheru commentedon Oct 17, 2026
Same card, same
!: on our R9700 it follows the ROCm userspace (7.2.1's libraries fire, 10.0.0's do not)Thank you for the trace and for the offer. You asked which userspace is worth comparing. Here is what we measured on
one R9700 (gfx1201), posted in #879 (main results,
follow-up).
What we found. Swift 1.5 Qwen3.8-Flash-Next IQ3_XXS (GSQ-RCO native pack, MTP head),
--resident-experts --kv int8 --max-context 262144 --kv-resident 262144 --prefill 16384 --short-read 150 --spec 4,STRATA_HIP_WMMA=1. The loadwas a stress script: reads of 195-260 new tokens on resumed 12K-60K conversations. Fires were counted from the engine's
non-finite logits (#879)lines.So on our card the fault follows the runtime libraries, not the compiler. 0.1.43.1 on 7.2.1 still fired (1 in 574).
With yours (TheRock 7.10) and cjdell's in #879 (7.2.3, long prompts), that makes three single-R9700 reports of this
family on a pre-10 userspace. Vop3p's comment in #1208 (gfx1030, finite on ROCm 10.0.0) points the same way.
One difference from your trace. Our bad rows were always in the full-attention (QSA) layers, and they were already
non-finite at the MoE input, with small finite values beside them (largest |x| 10.5). None of the 36 GDN layers was
ever reported. Yours starts as a GDN state overflow (2.26e+32). The entry point differs, the end is the same (all
logits non-finite, token 0). From here I can't tell whether it is one fault or two.
The userspace to compare: ROCm 10.0.0. We used the
docker.io/rocm/dev-ubuntu-24.04:10.0.0-fullimage, with/opt/rocm/libregistered with the loader (the image does not do it). Inside: HIP 7.15.26333, hipBLASLt 1.4. Therecipe is in the follow-up linked above.
The cheapest test is the cross run: keep your current binary and swap only the libraries (
LD_LIBRARY_PATH), so theengine code is the same in both arms. That worked for us between 7.2.1 and 10.0.0 (same sonames:
libamdhip64.so.7,libhipblas.so.3,libhipblaslt.so.1). With yourSTRATA_PF_FUSED=1amplifier at 40K (3 in 8), about 16interleaved rounds per arm would show whether the rate drops to zero: 6 of 16 against 0 of 16 is p ≈ 0.009 (Fisher,
one-tailed).
What ROCm 10 does not fix: card hangs. We saw one stall under 10.0.0, on 0.1.43 (
no progress for 60 s … waiting for the GPU (attention, router) at layer 19 of the prompt chunk from token 16384); the engine restarted itself.ROCm/TheRock#9279 (formerly ROCm#6630) has a new report of resets on 2× R9700 with ROCm 10.0. Since Oct 15 our daily
engine is 0.1.44.1 on ROCm 10.0.0: 2,660 prompts, no
non-finite logits (#879)line and no stall so far. That iseveryday use with
--short-read 260, which keeps short reads off the batched path, not a stress run.@Niko1221, for setup.
docs/AMD_HIP.md(v0.1.45, lines 31-38): a system ROCm 7 is used as it is, even 7.2.x, andotherwise gfx120X gets the TheRock
7.10.0a20251120wheels, this report's userspace. A setup warning for theseversions on gfx12, or a newer ROCm line by default where the index offers one, would save the next R9700 user the
hunt.
aiworkeryuan-chuxi commentedon Oct 19, 2026
Your four-arm table was the variable we had left. We ran the cross test — with one difference: our CLI has no per-request
LD_LIBRARY_PATHswap, so each arm is its own engine start (the config'slib_dirsbecomes the loader path). Same binary (0.1.44.1, built against the TheRock 7.10 headers), same args, same prompts; only the runtime libraries change.non-finite logits (#879)/opt/rocm-7.2.4/lib)Then production, 0.1.44.1 on the 10.0.0 libraries, 571 prompts of ordinary use:
!!!!0,non-finite logits0, panic 0, engine restarts 0, 99.3% normal completions, median decode 102.2 tok/s. We have been on that pair since Oct 17, and we now pin the library path in the systemd unit as well as the config, because a setup/update run silently restored the TheRock paths once and the engine came back up on the old runtime — exactly the "logs say nothing" failure you would expect.Where our trace differs from yours. Our degenerate requests never start in a QSA layer. They start as a GDN state overflow (
GDN state 0 9264 non-finite (max |x| 2.26e+32)with the residual still finite) and end with all 248,320 logits non-finite — the same count as in #1208.STRATA_DBG_NAN=1on a passing request prints0 non-finite (max |x| 3.34)for the same field, so the state is the thing that runs away. Your rows were already non-finite at the MoE input with small finite neighbours beside them (max |x| 10.5) and never touched one of the 36 GDN layers. It may be one runtime fault surfacing in whichever kernel gets there first, but we cannot separate "one fault" from "two" either.Rate depends on an amplifier. Our 7.2.4-arm rate above is not the base rate: it comes largely from
STRATA_PF_FUSED=1at 40K prompts. Interleaved against the same build withPF_FUSED=0, the amplifier arm degenerated 3/8 (37.5%) and the baseline 0/8. With the 10.0.0 libraries andPF_FUSEDleft at the 0.1.44.1 default, 571 prompts produced no event at all — so on this card the amplifier seems to need the pre-10 runtime to do anything.0.1.45 and 0.2.0-beta on this box.
prefill copy_i32about 1 run in 13. That may be a fix — n = 1, and we have not re-run the 1-in-13 experiment.--spec 4+--mtpconfiguration.STRATA_HIP_WMMA=1andSTRATA_PF_FUSED=1on, and those are precisely the two switches that break this card (WMMA at default on 0.1.42: a deterministicprefillcrash at 28K tokens; PF_FUSED: the 3/8 rate above). For a gfx1201 reader the release notes could say it plainly — if prompts degenerate, setSTRATA_HIP_WMMA=0andSTRATA_PF_FUSED=0first — because on this hardware those are ambiguous defaults, not free speed.We can run the same 40K amplified comparison against a ROCm 10 build of the engine, or bisect the 93K crash against 0.1.42, on the exact configuration above, if either is useful to you or to @Niko1221.
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.
sdroege commentedon Oct 31, 2026
Single-card gfx1201 (R9700) confirmation, engine 0.1.54, UD-Q4_K_XL.
Non-finite logits at exactly
prompt_length - 7in every occurrence (16K: 16051, 32K: 32042, agent: 10185), independent of the verify window size:--spec 2(4-token window), 4 (6) and 8 (10) all show the same position.At -60 mV VDDGFX it happened in most 16K/32K runs; at -25 mV and 0 mV it did not occur in two full runs each. That fits your GDN state overflow mechanism and adds a voltage-sensitivity data point: the same prompt is clean at 0 mV and non-finite at -60 mV.