Repository navigation
[Engine] dsv41 on gfx1100 (HIP, self-built): switching back to a conversation after another one ran fails with "illegal memory access" #4350
Description
Area
[Engine] (speed, kernels, models, the server and API)
Desktop app, Terminal or source?
Source (START-HERE.bat, setup.sh or a build)
What happened
strata-dsv41 on a self-built HIP engine dies with an illegal memory access as soon as a second conversation takes the model: continuing conversation A works, switching to another conversation works, but coming back to A kills the engine. After that every request returns HTTP 400 until the engine is restarted.
Response body of the failing request:
{"error": {"type": "invalid_request_error", "message": "cudaStreamSynchronize(s): an illegal memory access was encountered"}}ds4: cudaStreamSynchronize(s): an illegal memory access was encountered
ds4: cudaMemcpyAsync(d.emb_rows, h_emb, (size_t) T * D * 2, cudaMemcpyHostToDevice, d.st): an illegal memory access was encountered
ds4: cudaEventCreateWithFlags(&e, cudaEventDisableTiming): an illegal memory access was encountered
One teardown also printed rocBLAS error during freeing of allocated memory in handle destructor: rocblas_status_internal_error.
Smallest steps or command that reproduce it
Five requests with curl/any client, all max_tokens ~96, temperature 0.7:
- Start the engine (config below), wait for ready.
- Conversation A, request 1:
system= "you are a physics teacher", user = "why is the sky blue?" → keep the answer, append it plususer: "summarize in one sentence", send again (A#2), repeat once (A#3). - Conversation B, one request:
system= "you are a Japanese chef", user = "why must miso soup not boil long?" - Conversation A again: resend the A transcript (A#1..A#3 plus the answers) → the engine faults here.
Reproduced 2 times out of 2 engine starts with that script, no variation in where it dies. A much longer run of one-conversation-only traffic (40 requests, random-prefix prompts, max_tokens 64) passed without a single error, so the trigger is the conversation switch, not the request count.
Reproduction status
Reproduced locally
Strata version
Strata 0.2.7 (commit 5e8a4cca, tag v0.2.7) · engine 0.1.52 · self-built build-hip/strata-dsv41 (setup.sh --update, then cmake --build build-hip --target strata-dsv41)
Graphics card, memory and system
RX 7900 XTX 24 GB (gfx1100, Card Model 0x744c) · 64 GB DDR5 (4×16 GB) · Ubuntu, kernel 7.0.0-38-generic, ROCm 7.2.0. No display is attached to the 7900 XTX: the desktop runs on the Raphael iGPU, and the engine is pinned with HIP_VISIBLE_DEVICES=0, ROCR_VISIBLE_DEVICES=0.
Model, size and settings
DeepSeek-V4.1-Flash Uncensored GSQ-RCO IQ2_XXS, maya-ds4 pack · 65K context · pictures off.
--gpus 0 --max-context 65532 --ssd-cache 45 \
--expert-profile /home/lyy/Strata/data/dsv41/profile-2.25.bin --experts 5
env: HIP_VISIBLE_DEVICES=0 ROCR_VISIBLE_DEVICES=0 DS4_CPU_SMT=1 DS4_AHEAD_EXTRA=2 STRATA_DS4_HIPK=1
Startup lines (identical on good and bad starts):
ds4: experts per word: 5 (as asked; full quality: --experts 6)
ds4: RAM store: up to 4380 slots of 10.3 MB (45.0 GB), policy 0, 8 reader threads, unbuffered reads
ds4: RAM store: 4380 of the profile's hottest experts read in (34.0 GB) in 5.5 s
ds4: prompt path: the chunked pass from 96 tokens, windows of 16 below
ds4: CUDA0 layers [0, 39] + head: weights 6.20 GB, experts 14.54 GB (1647 of 15360 in VRAM), caches 0.06 GB, 574 MiB free
ds4: 13713 routed experts on the CPU lane (118.3 GB, in place in the model's files), 8 threads
Switches I tried, and they do not help:
STRATA_DS4_SLOT_KEEP=0— still faults at the same step (A#4). So it is not the conversation slot save/load (deepseek_slot_env()inserve/server.py, on by default in 0.2.7).STRATA_DS4_REUSE=0— faults earlier (A#2, the first request that has to read its prompt in full again), so the reuse/checkpoint path is involved rather than the slot folder.--experts 5vs the default 6 — both fault.--ssd-cache 45vs50— both fault.reasoning_effort: none/high/ unset — not the trigger; these requests carry no reasoning control at all. A long run with default thinking (max_tokens4096, 3493 tokens generated, 7530-char reasoning) completed normally while it was the only conversation.- The Qwen path (
sc117, qwen4exp) on the same build is stable in the same session (a 208,723-token request read at 2211 tok/s), so this looks specific tostrata-dsv41.
Logs
One of the two runs (strata-dsv41.log; line 103 is A#1, line 109 is the other conversation, line 110 is the traffic back in A, then the fault):
103: maya ds4 serve: prompt 50 tokens = 0 reused + 50 read in 3098 ms (16.1 tok/s), 48 generated in 4077 ms (11.77 tok/s); CPU lane 67.9% of the experts
104: maya ds4 serve: prompt 115 tokens = 98 reused + 17 read in 777 ms (21.9 tok/s), 31 generated in 2012 ms (15.40 tok/s); CPU lane 61.7% of the experts
105: maya ds4 serve: prompt 91 tokens = 49 reused + 42 read in 1044 ms (40.2 tok/s), 48 generated in 2605 ms (18.43 tok/s); CPU lane 51.2% of the experts
106: maya ds4 serve: prompt 156 tokens = 139 reused + 17 read in 450 ms (37.8 tok/s), 30 generated in 1200 ms (25.00 tok/s); CPU lane 47.6% of the experts
107: maya ds4 serve: prompt 130 tokens = 90 reused + 40 read in 692 ms (57.8 tok/s), 48 generated in 2404 ms (19.97 tok/s); CPU lane 42.0% of the experts
108: maya ds4 serve: prompt 195 tokens = 178 reused + 17 read in 409 ms (41.5 tok/s), 27 generated in 1002 ms (26.94 tok/s); CPU lane 40.2% of the experts
109: maya ds4 serve: prompt 53 tokens = 0 reused + 53 read in 2578 ms (20.6 tok/s), 49 generated in 2602 ms (18.83 tok/s); CPU lane 41.0% of the experts
110: maya ds4 serve: prompt 119 tokens = 102 reused + 17 read in 457 ms (37.2 tok/s), 30 generated in 1273 ms (23.57 tok/s); CPU lane 40.0% of the experts
111: ds4: cudaStreamSynchronize(s): an illegal memory access was encountered
The extra lines per conversation come from the server's own side requests (for example the chat title); my script sends exactly the five described above.
Related reports I found before filing: #3431 (Windows HIP, same gfx1100, "~1 of 8 engine starts" silently misreads prompts — also per-start) and #4090 (dsv41 on Linux wedges in the prompt path; different symptom). I could not find one for the conversation switch.
Activity
github-actions commentedon Oct 27, 2026
Thanks for the issue. To keep the tracker readable, every item needs an area prefix in the title and the parts of the issue form filled in. This one is missing:
Edit the issue (title and description) and this note goes away on its own. If it is still missing parts in 5 days it will be closed, which is not a rejection: fix it and comment
/reopen.Details: CONTRIBUTING.md
github-actions commentedon Oct 27, 2026
Thanks, this follows the format now.
udqe commentedon Oct 28, 2026
Note
AI-assisted static analysis - model: Qwen3.8-Flash-Next (Unsloth UD-Q4_K_XL) served by this repo's own engine 0.1.52, driven by a DeepSeek-based harness. The maintainer re-read the cited paths line by line in the v0.2.7 (engine 0.1.52) checkout; the failure below has NOT been triggered on my machine (DeepSeek runs in a single conversation here). Posting it in this thread rather than as a new issue because it is the same device-current shape as this report, on an adjacent path.
slot_load()sets the current device once from the lead layer and then writes every tensor-parallel member under it (ds4_model.cu ~5430-5452, second loop ~5466-5473):lm.winis allocated on the member's OWN device (load_dense ~3855-3919 allocates per memberl.dev), and every other per-member loop in the file re-selects the device (restore ~5227-5237, reset ~5142-5150). With the defaulttp = 2and two identical cards (tp_alltrue; ds4_model.hpp:55), a 2-wide group's members sit on different GPUs, so the second member's H2D runs with the wrong current device. The copy fails (cudaErrorInvalidValue class),slot_loadreturns false, the FailClosed path resets the model, andslot_take'sslot_drop(best)(ds4_main.cpp ~394-403, called at ~552-557) deletes the slot's.binAND.meta- the saved conversation is destroyed with no retry and every later switch re-reads the full prompt.Separate observation on the same handler: ANY failed take-back - including a transient file-read error during the ifstream at ~538-550 (backup/indexer mid-file) - takes the same
slot_droppath.What would confirm: two identical GPUs, defaults, take a conversation back, log the cudaMemcpy error string on the first take-back.
Workaround adopted on this machine until a fix lands:
--tp 1on the DeepSeek profile (the bug is latent when nm==1, i.e. single-device or--tp 1).Niko1221 commentedon Oct 29, 2026
Two separate things. 1) The
slot_load()finding is right: it selected the lead layer's device once and copied into every member of a tensor-parallel group, whose buffers sit on their own GPUs. Fixed on our side (it ships in 0.2.8); it only applies with tp above 1. On CUDA that copy works silently; we cannot say what it did on HIP.STRATA_DS4_REUSE=0faults at the first request with 115 new tokens, above the 96-token line where the chunked pass starts, points at that pass withSTRATA_DS4_HIPK=1. This needs a gfx1100 to settle.giostark commentedon Oct 29, 2026
Hi Niko ,
First , I have to express to you how grateful I am for the piece of software you bring to us. In local , for personal use, this is helping me a lot , really , so THANKS!!! Then , I'm not skilled at all so be patience , I'll try to be useful at the conversation. I cannot past a single line of the error so I tell my experience.
I have been trying Strata from the 0.1.45 release. I used a Qwen 27B to help me to compile Strata and now I'm using a Qwen3.8 FLASH NEXT on your software.
This is my config: 1) server with webui , embedding , voices stuffs, search stuffs . Those run on a 9950x , 32Gbram , 2x 2070 8GB , several M2. 2) The ws has just strata . And run a 5950x with 128GB ram and 2x7900xtx , 4TB m2 990pro.
I use Arch Linux 7.2.11-arch1-1 , ROCm 10.0 . Never a crash with any model till the follow.
From the first release this is the progression where the bug occur :
v0.1.45 - first custom install (stock engine, ROCm/gfx1100). No crashes documented during this version; setup, context 512K, k8v4 KV, MTP/spec and vision validated cleanly.
v0.2.1.3 - base already present in /home/ws/Strata-custom before the next step. No crashes recorded.
v0.2.2 - engine 0.1.47. Built and running; no crash reports tied to this version.
v0.2.4.2 - engine 0.1.49.2. This is where the crashes start being reported. On this live version I opened the next round with "fix this, all these crashes may not be our fault" (to the AI) and went looking for a newer release, which led to v0.2.6. Note the upstream fixes used as the reason to upgrade are engine-level, e.g. v0.2.7 ships "free memory of the second card shrinking with every long cold prompt on a layer split - a leaked copy stream (#3470, #3350)". So the crashes observed from around v0.2.4.2 onward are engine/Strata userspace crashes, not host reboots: Strata exits with code 1 and systemd (Restart=on-failure) brings it back on the next request.
v0.2.6 - engine 0.1.51. Updated hoping the stress-test fixes would help; crashes still present.
v0.2.7.1 - engine 0.1.53. Current live version. The error below is reproduced on THIS version, so it is not fixed by the leaked-copy-stream fix (#3470/#3350).
At first I thought was limited to multichat but I was wrong.
It is NOT limited to multi-chat. It reproduces with a single active user via the WebUI front end. The trigger is the query itself and the concurrency of simultaneous connections, not the number of open chats - one user going through the WebUI can already hit it. Two related error strings were observed; the second one comes from the restore path ("restoring parked conversation"), not from saving a snapshot. The AI checked this , comparing the moments of the queries and the crashes.
Exact engine error strings seen:
Log excerpts (latest version):
From the open-webui proxy that fronts the engine (journal on the front-end host, what I read this session):
I can revert to use the a old version to check it twice. Or If I can , tell me what I can do to help.
cuzlin226 commentedon Oct 29, 2026
Thanks for digging into this. Coming back with three measurements from the same single 7900 XTX (gfx1100, ROCm 7.2.0, Ubuntu, self-built
build-hip/strata-dsv41, Strata 0.2.7 / engine 0.1.52,HIP_VISIBLE_DEVICES=0 ROCR_VISIBLE_DEVICES=0).1. Reproduced again: 3/3 engine starts, always at the same step.
Same five-request script as in the report (A#1–A#3, one request in another conversation B, then A resent in full),
max_tokens96,temperature0.7. A#4 returns{"error": {"type": "invalid_request_error", "message": "cudaStreamSynchronize(s): an illegal memory access was encountered"}}and the engine log ends on the same line. Same place every time, no variation.
2. With
STRATA_DS4_HIPK=0it faults identically.I copied the config and changed only that one variable — verified in
/proc/<engine pid>/environthat the engine really gotSTRATA_DS4_HIPK=0, everything else untouched (--ssd-cache 45,--experts 5,--max-context 65532, same expert profile). Same script:Two things follow. The fault does not need the HIP kernel path on this device. And the last successful pass before the fault read only 17 new tokens, far below the 96-token line where the chunked pass starts — so this reproduction does not exercise that pass at all.
Side note, unrelated to the fault: with HIPK off the same requests decoded at 20.2–21.8 tok/s versus 15.1–19.6 tok/s with it on. On
strata-dsv41here, HIPK costs speed and buys nothing.3.
ds4-slots/is empty after the fault.The folder contains the saved conversation before the crash and is an empty directory afterwards. That matches
slot_take→slot_drop(best)removing both.binand.metawhen a take-back fails: the conversation we just failed to restore is destroyed, and the retry re-reads the whole prompt.Where that leaves it (tp = 1): the remaining suspect looks like the state restore path, not a kernel. Consistent with the earlier note that
STRATA_DS4_REUSE=0faults earlier — at A#2, the first full re-read, 115 new tokens — the failure follows the path that reconstructs state, not the path that computes it. Combined with (2), the chunked pass is out of the picture on this device as well.I have the box and the script here, so if you want a specific A/B (a build with extra assertions in the restore path,
STRATA_DS4_REUSE=0with HIPK off, a longer interleaving, anything) just say which and I will run it and report the logs.DeepSeek engine: slot_load() writes each tensor-parallel member's win…
giostark commentedon Oct 29, 2026
I can say after 10h of intense use that the v0.2.2 , on my setup with 2x7900xtx , is not affected. Multi-chat , multi-query with one chat that execute commands and other chats with heavy thinking and huge tokens. Not a issue.
Instead newer version present the crash very frequently.
Niko1221 commentedon Oct 29, 2026
0.2.8 fixes the tensor-parallel half:
slot_load()now writes each tensor-parallel member's state under its own device. We have not reproduced the single-card fault on gfx1100, so this stays open. If you can, runstrata updateand send the engine log from a single-card run that faults.giostark commentedon Oct 30, 2026
Hi ,
Tested 0.2.8 / engine 0.1.54 on the gfx1100 box; the illegal-memory-access fault is still here. Build: pipeline split, --layer-split 22 (not tensor-parallel), Qwen path, so per your note the tp>1 slot_load fix doesn't cover this.
Attached is the engine log around 5 faults today plus three controlled repro attempts. Short version: the fault does not depend on traffic shape. Cold restart with an immediate conversation switch = 0/10. Cache forced to parked=4 with many evictions on small chats = 0 faults. Under real mixed load it faults within tens-to-hundreds of requests, then often stays clean for hundreds more.
The one thing common to every real crash and absent from all my scripted repros: the restore/park happens on a very large conversation (70503 / 48745 / 13247 tokens). Every fault also has parked=4 with evictions, and "verify: layer N never rang" lands exactly on the split boundary (layer 22 / layer 0).
report_4350_with_log.txt
engine_faults.txt
engine_meta.txt
catch-26 commentedon Oct 30, 2026
also ran into this issue today. can provide logs if needed.
Niko1221 commentedon Oct 31, 2026
Two things. We could not find the cause of the fault; it did not reproduce on our cards (one crash in about 40 runs on another AMD card). But what happened after it is fixed in code: when the GPU reports a fatal error the engine now ends and the server starts it again on the next request, instead of failing every request until you restart it. The request that was running gets a 503 you can send again; three such errors within ten minutes stop the restarts and say so. We tested this with an injected fault on DeepSeek and Qwen, not a real one. It goes into the next update. If you can, send the engine log of a run that faults.
dturini12 commentedon Nov 1, 2026
Note
AI-assisted report. The investigation and this text were prepared with Claude (Anthropic, Claude Opus 5.5 in Claude Code), working read-only on my machines. I reviewed it before posting. The log excerpts are copied from my system and trimmed; nothing else in them is edited.
A data point on the post-fault behavior, plus one server detail your pending fix may or may not cover.
Setup:
strata-dsv41v0.2.9c62d5e3a1(engine 0.1.55), Linux gfx1151 (Ryzen AI Max+ 395), single GPU, tp 1, ISTA-DAS 2.00 pack. The fault itself wasmoe_sort_kernel<false>; I posted the details in #4825. After it, the engine kept failing every request, as reported here, while/healthstill answered"status": "ok". The restart-on-fatal-error change you described for the next update would have covered that part.The extra detail: a leftover
DONEturns each error into an empty"length"reply for the next request.ds4_main.cppprintsERR <msg>and thenDONE 0 <n> 0.0 0.0 cancel …(lines 919-920; the picture checks at 764-765 do the same).Engine.generate()inserve/server.py(2442-2444) raises on theERRwithdone = True, so theDONEis never read. That follows serve: a request the engine rejects with ERR waits the full 300 s before failing (STOP + _drain_control("DONE") after an ERR), holding the control lock #1059/serve: release batch control after terminal native ERR #1183: "the engine refused it and is idle: no DONE follows". This holds for the main engine but not for these two dsv41 paths.DONEat once. It ends with 0 tokens and keeps the defaultfinish = "length"(server.py:5512), and the log prints "the reply reached max tokens while still thinking".ERRorDONE, so replies alternate:error, then emptylength.In OpenCode this looked like a real output-limit problem. After three empty
lengthreplies it gave up with "Response still truncated after two output token limit continuations". The actual cause was the GPU fault.Server log: the last good request, then the alternation (progress lines removed)
/metrics requests[] for the same window (times UTC); each error/length pair has the same prompt
What is observed vs. not:
ds4_main.cpp:764-765path), which is not fatal, should leave the same leftoverDONE, so the next request would return an emptylength. I have not triggered this.If the fix in the next update doesn't already do it, reading the
DONEthat follows a dsv41ERR(with a short timeout) would close it.cuzlin226 commentedon Nov 2, 2026
Still reproduces on v0.2.9 / engine 0.1.55 (
c62d5e3a1) — and the post-fault empty-reply pattern shows up on dsv41 tooThe last data point in this thread was 0.1.52, so I re-ran the same minimal script on the engine released 2026-10-31. It still faults, first try.
Setup
strata-dsv41v0.2.9, engine 0.1.55 build c62d5e3, self-built for gfx1100 (STRATA_BUILD_DSV41=ON)HIP_VISIBLE_DEVICES=0 ROCR_VISIBLE_DEVICES=0--max-context 65532,--ssd-cache 45,--experts 5,--gpus 0; RAM store 34.0 GB read in 5.5 sResult — 1/1, same signature as 0.1.52
cudaStreamSynchronize(s): an illegal memory access was encounteredPost-fault state
serve.serverandstrata-dsv41), port staysLISTEN, VRAM stays at 25.45 / 25.75 GiBds4 slots: ... 0 conversation(s) kept from earlier runs (0.00 GB)— the failed take-back drops the slot filesso the leftover-
DONEbehaviour described above is present on this path too, and the two symptoms are one failure mode: the engine is dead and every following request reads the previous request's leftoverERR/DONE.One observation that may or may not be related to the
slot_loaddevice pick. At startup on this single-card box the engine prints:ds4_model.cu(thetpreduction before the groups are formed) starts ato.tp= 2 and lowers it because the GGUF's 256-value blocks cannot splitn_ff 2304; only one GPU is visible, so the runtime really is tp 1. I mention it only in case a slot/device path is still reading the pre-reduction value — with--tpunset every run on this card passes through that2 -> 1transition, which would explain why this reproduces here and not on your cards.On the pending fix: as I read the timeline, the restart-on-fatal-error change was posted 2026-10-31 19:12 and 0.2.9 shipped 2026-10-31 08:02, so this build does not carry it — which matches what we saw (no restart, no 503, only zombie replies). Glad to re-test as soon as a build with it lands.
This box reproduces deterministically with the 5-request script, so a single-variable experiment here is cheap. Say the word and I will run whatever you want to isolate — with the engine log of any run.
Engine log, start to fault (this run, trimmed to 36 lines)
cuzlin226 commentedon Nov 2, 2026
Note
AI-assisted report. The investigation was run with a coding agent on my machine; every number below comes from
output I saw in the engine log, the kernel log, or the two commands listed at the end. No log is edited.
Root cause found on gfx1100 — it is not the expert cache or the slot path. The gfx11 default dense dequantizer
(
dq_rows_kernel<T>, the row-per-lane one thatdq_use_old_kernel()selects ongfx11*) overruns its shared stagingbuffer for Q6_K dense tensors, so
gg::rows()produces garbage weights there, and the garbage overflows FP16 intoinf. From there:inf→ GEMM outputNaN→ hidden stateNaN→ gate logitsNaN→route_kernelnever finds amaximum (
bs[i] > -FLT_MAXis false for NaN) and writes0x7FFFFFFFas the expert id →moe_sort_kernel<false>readsvs[0x7FFFFFFF]≈ +8.6 GB off the table → GPU page fault.Reproduction status: reproduced locally, 5/5. Environment: engine 0.1.55 (
c62d5e3a1), Linux, RX 7900 XTX(gfx1100), ROCm 7.2, single GPU, tp 1, a maya-ds4 pack (DeepSeek-V4.1 weights: dense tensors Q6_K, experts IQ2_XXS),
--experts 5 --ssd-cache 45.1. The fault is in a kernel, not in the copy that reports it
With
HIP_LAUNCH_BLOCKING=1the error moves fromcudaStreamSynchronizetocudaMemcpyAsync(c.bounds_pin, c.bounds, (n_vs + 1) * 4, D2H), but the kernel log shows what actually faulted:Synchronising between the three launches inside
moe_sort()names it:moe_sort_kernel<false>fails first — thesame kernel dturini12 reported on gfx1151 in #4825.
2. NaN appears at one dense GEMM, and the weights it was fed already had
infStage probes on layer 0 of the restored chunk (a NaN/Inf/absmax read of each buffer, D2H + sync):
The
w16_zeroline is the decisive one: the buffer was zeroed, so the dequant did run — it wrote theinfs. Thetensor it was asked for:
fmt=4 K=14 (Q6_K) k=8192 n=5120 ld=6720.3. Single variable: the dequantizer
dq_rows_kernel)w16has 5 inf-laden column blocksDS4_DQ_OLD=1(force the old one)DS4_DQ_OLD=0(forcedq_rows_kernel2)w16clean, two full A→B→A cycles pass4. Why the old kernel breaks Q6_K
dq_rows_kernel<T>stages a warp's 32 sub-blocks in shared memory through contiguousuint4loads:For the 256-value K-quants
kSpan = 4 * B:4*BSo for Q6_K each warp writes 52
uint4into a 44-uint4slot (into the next warp's slot, and past the array for thelast warp), and
kSpan / 16= 52 truncates the span, so the last 8 bytes of the fourth superblock are never stagedeither. The decoded values are garbage, and some of the garbage overflows FP16.
Ty<kQ6_K>andTy<kQ5_K>are marked"GEMV only" but the chunked prompt's dense GEMM (
gemm()→gg::rows()) does callrows()for them.5. Why only after switching conversations
That dense GEMM is not reached on a plain prefill. Same request, same model, only the middle conversation removed: the
gemm(k=8192, n=5120)descriptor is never printed and nothing faults. So the trigger is "the restored chunk takes thechunked dense path", and the bug is the dequantizer it meets there.
6. Why CUDA does not see it
The old kernel is HIP-only (
gfx11*selects it), so a CUDA build always uses the other one; and an out-of-range readthat lands in a mapped VA on CUDA silently returns garbage instead of raising a protection fault — which is also why
this stayed invisible in scripted replays.
Workaround
export DS4_DQ_OLD=0Verified here: two full A→B→A cycles, five requests, no
illegal, replies normal.Suggested fix
Size the staging slot from the type instead of the Q5_K case, and round the span up:
Separately — and I can file it on its own if you prefer —
route_kernel(ds4_kernels.cu:508) turns any upstream NaNinto a machine-wide failure:
bixstays0x7FFFFFFF, that value is written as an expert id, andbs[ix] = -FLT_MAXthen indexes out of
__shared__ float bs[1024]. A guard (fall back tor, clampixto[0, n)) would have kept thisa bad reply instead of a fault plus a stuck server. That is a hardening change, not this fix.
I can prepare a PR with the
kSpanfix (and the Q6_K/Q5_K dense-GEMM case it should be tested with) if you want it.Reproduce
After a fatal GPU error the engine ends itself and the server starts …