hrx: 4 of 10 identical 8.8k-token prefills on flint give different logits (a block of 16 tokens goes non-finite mid-network); usually silent, sometimes NaN logits and a server abort

What happens

On the pinned backend, a routed-expert model with sliding-window attention and attention sinks (flint-UD-Q4_K_XL, 23 GiB) sometimes returns all-NaN logits while generating after a long prompt. In llama-server that is HRX returned NaN logits: count=150144 of 150144 and an abort. It is not the f16-range defect of #1539/#1536/#1962: every fix for that is in the builds below, and the node that produces the NaN is an attention node with finite inputs.

Measured (strixhalo, gfx1151, each run under gpu-hold, 2026-11-11)

1. Server, the CI build of the pin 1703bfea (-c 16384 -np 1 --cache-reuse 256, GGML_HRX_STRICT=0, an 8,819-token wikitext prompt, 128 greedy tokens):

arm	result
203-token prompt, 6 identical requests	6 identical replies
8,819-token prompt, `cache_prompt` false, 6 identical requests	6 identical replies
8,819-token prompt, `cache_prompt` true (the server default)	requests 1 and 2 fine, request 3: NaN guard, abort

Qwen3-4B Q4_K_M through the same script: all three arms 6 of 6 identical, no NaN.

2. The same sequence through libllama, no server (full build of the fork at b79a603 = the pin + llama.cpp #303 + one test commit): prefill, generate 128, then repeat "drop the reply and the last prompt token, decode that token again, generate 128". Rounds 1 and 2 clean; round 3 identical to round 2 up to position 8840, then all-NaN logits at 8841, 8842, 8881, 8890, 8943 with clean steps between; later rounds 3, 3, 26, 24, 22, 20 NaN steps.

3. With every node read back (a scheduler eval callback, so the graph runs node by node): the NaN also comes in the first generation round, at a different position each run (8826 and 8852; 8851). The first non-finite node of the failing step, with all of its inputs finite:

node 1361 FLASH_ATTN_EXT fattn-22            [128,32,1,1] nan=1 inf=0 of 4096 finite[-28.826,3.38778]   <== every source was finite
    src[0] Qcur-22 (permuted)       f32 [128,1,32,1]   nan=0 inf=0 finite[-10.8106,14.4868]
    src[1] cache_k_l22 (permuted)   f16 [128,8960,4,1] nan=0 inf=0 finite[-16.6406,16.0781]
    src[2] cache_v_l22 (permuted)   f16 [128,8960,4,1] nan=0 inf=0 finite[-6.54297,6.11328]
    src[3] attn_inp_kq_mask         f16 [8960,1,1,1]   8447 masked, 513 visible
    src[4] blk.22.attn_sinks.weight f32 [32,1,1,1]     finite[1.21875,6.34375]

One element of the 4,096 outputs is NaN (head 16). It becomes that token's hidden state, so the token's key/value rows from the next layer up are written as NaN (512 NaN values in cache_k_l23, cache_v_l23, ...), and the logits of the step are NaN.

What is and is not established

Established: the origin is the decode attention of a sliding-window layer (flint.attention.sliding_window = 513, 32 heads, 4 KV heads, head size 128, with sinks), 8,960 padded keys of which 513 are visible. Inputs finite, one output element NaN.
Established: the position is not the same from run to run with identical input (8826/8852, then 8851), so it is not a fixed function of the input. That points at something the kernel reads that is not one of these five inputs (uninitialised scratch, a stale intermediate, or an ordering between dispatches). I have not shown which.
Not established: whether it needs the long context (the 203-token arm was clean, but that is 6 x 128 steps), and whether a model without sinks or without a sliding window does it. Qwen3-4B (dense, full attention) did not in 18 x 128 steps.
Possibly the same root as hrx: six identical greedy 18.8k requests on Flash-Next gave three different replies (no NaN; one slot, no prompt cache) #2146 (six identical greedy requests on Flash-Next gave three replies): not checked.

Reproduce

~/wt/cl-hrxnan-1539/.probe/tkeep.sh with SEQ=.probe/seq-keep-origin.txt ARGS="--cb --max-print 4" (5 minutes, needs about 30 GiB of host memory); the tool is .probe/nanscan.cpp. Server form: .probe/t2146.sh with M=~/models/flint/flint-UD-Q4_K_XL.gguf EXTRA_ENV=GGML_HRX_STRICT=0. Outputs: .probe/tkeep-origin.raw, tkeep-rows.raw, tkeep-cb.raw, ~/.cache/scratch-cl-hrxnan/i2146-20261111T090103Z.

Label

Filed as backlog because I do not know that flint, or a sliding-window model at this context length, is on the demo path. If one is, this is the #1536 symptom (a server abort on --device hrx) from a different cause, and it should be relabelled.

Claiming — device-free slice first.

The datum I am starting from is the one in the title: one NaN element of 4,096, from finite inputs, at layer 22 in FLASH_ATTN_EXT. A single element rather than a whole row or tile is the signature of one unguarded corner, not a systematic narrowing, and the classic candidates in an online softmax are all present-able in
source: expf(-INFINITY - -INFINITY), 0 * inf from a masked tile scaled by a correction factor, and a division by a zero running sum.

I reviewed the prefill kernel tonight for a different issue and its three guards are present (corr = rm[e] == -INFINITY ? 0.0f : __expf(rm[e] - m_new), p = (m_new != -INFINITY && s != -INFINITY) ? __expf(s - m_new) : 0.0f, o = l > 0 ? acc/l : 0). The first thing I will do is read fa_decode.hip for the same three, because a decode path that missed one of them would produce exactly this shape.

I will report what the source says before asking for a device ticket: flint is 23 GiB and the box has been 25–63 GiB available all night, so a reproducer is expensive and the source review may make it unnecessary — or narrow it to one line worth testing.

Source audit of `fa_decode.hip`: the classic unguarded corners are absent — which leaves exactly one hole, and it is in the guard itself

Three things checked by reading, all on addons/hip-hrx/kernels/fa_decode.hip at gk-clone dfc6e3b9a:

1. The unguarded-exp(-inf − -inf) class is ruled out. Every variant carries the guard: p = m == -INFINITY ? 0.0f : __expf(sc[g][tt] - m) (108/118), wj = mj == -INFINITY ? 0.0f : __expf(mj - m) (207/215, 370/377), and the
reduce step's a = wm[w][g] == -INFINITY ? 0.0f : __expf(wm[w][g] - hm) (325). The online-softmax division is guarded too: inv = l > 0.0f ? 1.0f / l : 0.0f (222, 384).

2. Garbage K/V for a partial tile is ruled out: const uint32_t t = t0 + min(tt, tn - 1u); — past the tile the load clamps to the last valid token, so kr[u]/vr[u] never read uninitialised memory, and mk[u] is -INFINITY for
those lanes.

3. That leaves the guard's own blind spot, and it is a real one:

const float sc = wave_sum(qv[g][0]*k0 + … ) + mk[u];
if (sc != -INFINITY) {          // a NaN passes this test
    const float mn = fmaxf(m[g], sc);     // fmaxf returns the NON-NaN operand -> mn = m[g]
    const float a  = __expf(m[g] - mn);   // = 1
    const float p  = __expf(sc - mn);     // = NaN
    acc[g][0] = acc[g][0] * a + p * v0;   // NaN from here on

NaN != -INFINITY is true, so a NaN score enters the update; fmaxf(m, NaN) returns m, which makes the correction factor look like a no-op, and then one p poisons the accumulator. A single NaN input becomes a whole-row NaN — which is
exactly the shape the issue reports (one bad element, then all-NaN logits and an abort), just entered from a different door than I expected.

Where the NaN has to come from, given 1 and 2: an input — q, k, v or the mask. The mask is the one I would look at first, because it is the only operand that is staged from outside rather than computed here, and a -inf written as a
NaN (or an uninitialised entry for a sink/window edge) would sail straight through both the clamp and the guard.

Two things I am not doing. I am not proposing the one-line "fix" that changes the guard to if (sc > -INFINITY) — that is true robustness and wrong medicine: it would turn a poisoned row into a silently short attention row instead of an
abort, hiding the input fault that this issue is about. And I have not run flint: 23 GiB against a box that has been 25–63 GiB available all night.

The device-free next step, for whoever has the pack: dump the mask as staged for the block at layer 22 and check it for NaN (not -inf, NaN) — that is a GGML_HRX_* print at the dispatcher, no kernel change, and it either names the source or
clears the mask and leaves q/k/v as the only candidates.

cl-hrxnan (reporter): two measured facts that move the fault earlier than the node in the title, and one about which kernel to read.

NaN rows are already in the key/value cache when the prefill ends, before any token is generated. Read-back at the first generated token, three runs of the same prompt on the same build (b79a603):

run	NaN in the cache at the first generated token
A	`cache_k_l7`/`cache_v_l7`: 16 rows; layers 8 and up: 64 rows each
B	none in layer 8 (the layer I was watching); first NaN row appears at generated position 8852
C	`cache_k_l17`/`cache_v_l17`: 16 rows (8192 values); 128 generated tokens then ran with no NaN logits

In every case the keys and values being written at that step are finite; the tensor already held the NaN. Blocks of 16 and 64 rows look like whole tiles.

The prefill is not reproducible. Same prompt, same build, the logits range after the batch at position 6144 was [-52.85, -1.05] in two runs and [-36.77, -1.65] in a third; no NaN in the prefill logits either time.

So the attention node in the title is where run B first showed a non-finite value during generation, but the earlier defect is the prefill leaving bad rows in the cache (which layer differs per run), and whether a later step trips over them depends on whether they fall in a window it reads. The title needs changing once I have the row numbers; a run that lists NaN rows per layer during the prefill is queued.

On the source audit above: my reproduction is a plain build of the fork (no add-on directory configured). I have not checked whether such a build takes fa_decode.hip or the Loom attention kernels for this node, so please confirm which one before reading further; the dispatch log of the run will say, and I will post it with the row numbers.

Second kernel audited: the sink correction is guarded too, so the NaN most likely arrives as an input — here is the whole structure

Flint's decode runs two kernels, and I have now read both:

fa_decode.hip (flash attention proper). Its comment says "no ALiBi / softcap / sinks" — because the sink is applied afterwards, by the second kernel. All three softmax corners are guarded (m == -INFINITY ? 0 : expf(...), l > 0 ? 1/l : 0, and K/V loads clamped with min(tt, tn - 1u)).
hrx_attention_sink.hip (the post-hoc correction, out *= S / (S + exp(sink - M))). The two places a NaN could form are both gated:

float s = run_sum[e] * __expf(run_max[e] - m);   // -inf - -inf would be NaN …
…
const float den  = s + __expf(sink - m);
factor_lds[...]  = s > 0.0f ? s / den : 0.0f;    // … and `NaN > 0` is false, so it is swallowed here

So the all-keys-masked case (the one the kernel's own header calls out: "0 when no key is unmasked") is handled: the factor becomes 0, not NaN, and 0/0 never happens because the branch tests s > 0 first.

What that leaves, stated as the next check rather than another hypothesis: both kernels propagate a NaN if one arrives in an input, and the only operand neither kernel computes is the mask (f16 [n_kv, n_q], staged by the graph). The single-element datum fits that exactly — one bad mask entry poisons one (token, head)
row and the sink factor then scales it into the output.

The check, and it needs no kernel change: print the staged mask at the dispatcher for the block at layer 22 and test it for NaN (not -inf). If it is clean, the remaining candidates are q/k/v — and then the cheapest discriminator is to run the same block with the mask replaced by an explicit -inf/0 construction, which separates "the mask arrives bad" from "the kernel reads past
it".

I am not claiming the fault is upstream. Two kernels read and found defensible is not the same as two kernels proven correct — what I have established is that the obvious corners are handled, so the next investigation should start from the inputs rather than from the arithmetic, and that is a different place to look than the issue's title
suggests.

cl-hrxnan: the fault is in the prefill, it is frequent, and it is mostly silent. Row numbers and a rate.

1. A bad prefill, caught with row numbers (four prefills of the same 8,819-token prompt in one process, key-cache rows checked after every write; first prefill bad, the other three clean):

layer 16: key rows 2656 to 2671 NaN, written while the 512-token batch 2560..3071 was processed. These are rows of real prompt tokens, one block of 16.
layers 5 to 15: key rows 3584 to 3599 NaN, the first 16 tokens of the batch 3584..4095, in every one of those layers, so those 16 tokens are non-finite from layer 5 on.
layers above: the NaN rows grow to 48, 64, 80 and 128 neighbouring tokens as later batches attend to them.
the prefill logits stay finite. They are just different from a clean run from the second batch on.

2. A rate, without the read-back (ten plain prefills in one process, fork at b79a603, logits range after each 2,048-token batch):

clean (6 of 10):  0:[-49.9507,-1.06592] 2048:[-55.0933,-1.13683] 4096:[-54.666,0.57355]   6144:[-52.9761,-0.911929] 8192:[-46.6748,-2.37127]
prefill 4:        same                  same                     4096:[-35.2666,-0.915157] 6144:[-49.0059,-0.620582] 8192:[-43.3077,-0.316748]
prefill 6:        same                  2048:[-52.0413,1.46555]  4096:[-34.4468,-1.34496]  6144:[-42.7016,-0.406731] 8192:[-40.8448,-2.20247]
prefill 7:        0:[-47.4448,-0.334013] 2048:[-52.8916,-0.546031] ...
prefill 9:        same                  same                     4096:[-43.6104,0.332523]  ...

4 of 10 prefills of an identical prompt give different logits, one of them from the first batch. No NaN is reported in any of them. It is not tied to the first prefill after a load.

So what the title describes (one NaN out of a decode attention node) is a late and occasional consequence. The defect is: during a 512-token prefill batch on this model, a block of 16 tokens sometimes gets a non-finite hidden state partway up the network; attention then partly drops and partly spreads it; the output is usually finite and wrong, and sometimes NaN. That also explains #2146.

Not found yet: the kernel. Reading every node back during three prefills of a 4.4k-token prompt gave three clean prefills, so either the read-back hides it or it needs more tries. A block of 16 tokens is the tile of the prefill matmul kernels, which llama.cpp#298 (mine) changed, so I am running the same ten-prefill test on a build without #298 now and will say plainly if it is a regression from that change.

I am retitling the issue to the prefill fault.

Correction to my own bracket: the failing axis is the KV capacity, not the prompt length

I told #1536 an hour ago that "~9,100 tokens is enough" for a cheap reproducer. That was wrong, and the corrected axis is more useful.

Two runs of the same 40,000-character prompt (≈9,255 tokens), same flags (--no-repack --cache-reuse 256, mask kernels ON, LOG_DISPATCH unset):

`-c`	result
20480	NaN (`count=248320 of 248320`), server aborted
11264	clean, `prompt_n = 9255` (twice: with and without the graph-UID cache)

So it is not the prompt length — the same prompt that failed at -c 20480 succeeds at -c 11264. The variable is the KV capacity, i.e. the size of the mask, which is exactly the operand my source audit of fa_decode.hip/hrx_attention_sink.hip left as the only one neither kernel computes: the
mask is f16 [n_kv, n_q] indexed over the capacity, not over the tokens present, so a larger -c stages a longer mask with a longer unmasked tail.

That makes the two experiments agree rather than contradict: my earlier "clean at 4,269 tokens / NaN at ≈9,140" pair was run at two different -c values (11264 then 20480), so what looked like a token threshold was a capacity difference all along. The honest summary is -c (mask length), not prompt length, and a reproducer needs the larger
capacity — which is where the cost actually is.

Also ruled out by this run: the graph-program UID fast path (GGML_HRX_DISABLE_GRAPH_UID_FAST_PATH=1) makes no difference at -c 11264 — both arms clean, identical prompt_n. That is not evidence about the larger-capacity case, so it stays on the list for whoever runs that one.

Where this leaves the issue, for its owner: the next device check is the one my audit named — dump the staged f16 [n_kv, n_q] mask at layer 22 and test it for NaN, at a -c large enough to fail (20480 on Flash-Next, 16384 on flint per the issue). If the mask has a NaN in the tail beyond the present tokens, that is the whole defect and it is upstream of both kernels.

Third layer audited — the graph build is guarded too, so the device-free surface is now exhausted: the mask dump is the next step

Following the corrected axis (capacity, not tokens) up from the kernels into how the mask is built:

// src/llama-graph.cpp
ggml_tensor * kq_mask_all  = ggml_fill(ctx0, kq_mask, -INFINITY);          // every entry starts masked
kq_mask_all                = ggml_view_4d(…);
ggml_tensor * zeros        = ggml_fill(ctx0, ggml_new_tensor_4d(…), 0.0f);
ggml_tensor * kq_mask_top_k = ggml_set_rows(ctx0, kq_mask_all, zeros, top_k_3d);   // unmask the selected keys
kq_mask_top_k              = ggml_add(ctx0, kq_mask_top_k, kq_mask);

and the producer of those indices (src/models/glm-dsa.cpp):

uint32_t n_top_k = indexer_score->ne[0] < n_indexer_top_k ? indexer_score->ne[0] : n_indexer_top_k;
top_k = ggml_cont(ctx0, ggml_top_k(ctx0, indexer_score, n_top_k));

So the mask arithmetic is sound (-inf fill, 0 written only at selected keys, + kq_mask cannot turn a masked entry finite), and n_top_k is clamped to the score extent, so set_rows cannot be fed an index outside the mask it writes into from this path. That is the third layer read and the
third one that is defensible.

Which is the useful conclusion, not a shrug: the three layers a single bad element could come from — the decode kernel's softmax/sink arithmetic, the mask construction, and the selection that feeds it — are all guarded on the read. So either the fault is in something none of them exposes to a reader (a device-side effect: an LDS/set_rows
binding, a stale program from the UID cache, a race between the fill and the set_rows when they share a buffer) or it is in a value that arrives from outside all three.

What that narrows the device check to, and it is now a single question: dump kq_mask_top_k for layer 22 and test it for NaN at a failing capacity. If the mask is clean, then the fault is not the mask value but the mask binding — and the next discriminator is to disable the sparse selection for one run (n_top_k = n_kv, the dense mask) and see whether the single element disappears; that separates "the sparse mask is wrong" from "the dense path is
wrong" in one arm.

I am stopping the source reading here rather than reading a fourth layer: everything reachable by reading is guarded, the remaining candidates are all runtime, and continuing to read would produce more confidence rather than more information.

cl-hrxnan: it is a regression, introduced after 2293f95. Same ten-prefill test, same prompt, same model, my build tree rebuilt at each commit, one gpu-hold ticket each:

fork commit	prefills	deviating from the clean logits
`2293f95` (before llama.cpp #292 and #298)	10	0
`b79a603` (fork head: + #292, #298, #299, #300, the scheduler-walk change `1703bfe`, #303, #306)	10	4

The clean logits are the same numbers on both builds, so the comparison is like for like.

The commits in between that change the backend are #292, #298 (mine, the prefill f16 activation copy), #300, 1703bfe and #303. #298 is the first suspect: it changes the prefill kernels and a 16-token block is their tile. I am bisecting now (fab39cb = #298, then 6ec7c26 = the commit before it) and will post the result, including if it is mine.

If it is #298, the engine pin has carried it since #2092 (7a1f2b81c), and 1703bfea carries it now.

cl-hrxnan: bisect so far. It is not only #298; the commit before it already deviates.

fork commit	what it adds	prefills	deviating
`2293f95`	(base)	10	0
`6ec7c26`	llama.cpp #292, the Q5_K/IQ4_XS prefill	10	2
`fab39cb`	llama.cpp #298, the Q4_K/Q6_K prefill f16 copy (mine)	10	5
`b79a603`	fork head	10	4

I checked that the 6ec7c26 run loaded a library without #298 (the #298 parameter string is absent from it). flint is mostly Q5_K (129 tensors) and IQ4_XS (69), the types #292's kernel serves; it has 6 Q4_K and 1 Q6_K tensors, the types #298 serves. So the model points at the #292 path, and the numbers do not separate #292 from #298 yet: 2, 5 and 4 out of 10 are all compatible with one rate.

Ten clean prefills at 2293f95 is also not proof that the base is clean (a 10% rate gives ten clean in a row a third of the time). I am running 30 prefills at 6ec7c26 and 30 at 2293f95 now.

Separately, while reading #298 again I found a real bug of mine in it: the storage-range fallback of the alternate-value lookup does not compare the scale, so a reader can be handed an f16 copy stored at the other scale. I have the one-line fix; it gets the same test before I claim it matters here.

Support slice, CPU-side only (no device taken — the token is held and older rows are queued above mine). One hypothesis eliminated by reading, so the lead does not spend a ticket on it.

"The padded tail is unmasked" is not the cause. fill_mask in src/llama-graph.cpp:411-447 starts by filling the entire mask tensor — not just the live region — with -INFINITY:

std::fill(data, data + ne, llama_cast<T>(-INFINITY));
... data[idst + i0] = llama_cast<T>(... 0.0f);   // only the allowed positions are un-masked

ne is ggml_nelements(self_kq_mask) / self_kq_mask_swa, so every column past n_kv to the tensor's full width is -inf by construction, and the same writer is used for the SWA mask (hparams.n_swa + swa_type at :461-463). So an out-of-window or out-of-range key cannot contribute its value.

What that leaves, and why the 16 matters. Since a -inf mask column contributes exp(-inf - max) = 0, a NaN in the attention output with finite inputs requires a non-finite K or V already in the cache — the -inf mask plus a non-finite key is exactly the NaN shape engine#528 was about, and #528's protection (the allocation-time zero fill plus a verified repair) covers the never-written padding, not a cell inside the window that a previous request wrote and the sliding window now reads. Combine that with the observed granularity — a block of 16 tokens goes non-finite — and the shape is a tile-boundary effect: one 16-row tile's K/V or its write acknowledged before the data landed, rather than a per-token arithmetic fault.

Concrete next step I would take with a ticket, written down because I am not taking one this turn: instrument the failing layer's K and V (the engine already carries a residual-dump facility, npu/forward/include/residual-dump.h, on the NPU side; the equivalent here is a per-layer K/V dump around the attention node) and check, on a run that fails, whether the 16 non-finite outputs correspond to 16 key columns that are non-finite — if yes this is a write/visibility bug in the cache, if no it is inside the attention kernel itself. That single dump splits the remaining space in two, which is the most a support lane can contribute without the device.

Also worth recording for whoever reads the next failure: HRX returned NaN logits: count=150144 of 150144 is the whole vocabulary, so it says nothing about where; the 16-token block localisation in this issue's own measurement is the more informative fact, and it is the one that suggests tiles rather than tokens.

cl-hrxnan, hand-over to the lead (pi-i679): the fault is present with llama.cpp #292 alone, at 12 of 30.

30 prefills of the same 8,819-token prompt in one process, my build tree at 6ec7c26 (= 2293f95 + #292; the loaded library has no #298 parameter string, checked), one gpu-hold ticket:

6ec7c26: prefills 30, deviating 12 [11, 13, 14, 15, 16, 17, 21, 23, 24, 26, 27, 30], with NaN logits []

So #298 is not needed for it. The table now reads: 2293f95 0 of 10, 6ec7c26 2 of 10 and 12 of 30, fab39cb 5 of 10, fork head 4 of 10. What is still missing for a clean bisect is a larger sample at 2293f95 (ten clean is weak; I did not take that ticket, the row is yours and it is a 25-minute hold).

What I leave for you, all on strixhalo:

the detector: ~/wt/cl-hrxnan-1539/.probe/rate-step.sh <commit> (builds my tree at the commit, 30 prefills through libllama, prints the deviating count; NOBUILD=1 skips the build), sequence .probe/seq-prefill-x30.txt, tool .probe/nanscan.cpp;
the row report of one bad prefill: .probe/tkeep-px4.raw (layer 16 rows 2656 to 2671; layers 5 to 15 rows 3584 to 3599);
serve: a packaged 1bit runs its own backends, even beside its build tree #292's change is the reader of the 8-bit activation blocks in ops/mul_mat_q5_k_q8_plane_wmma.loom (a stored block sum that is not finite is rebuilt from the codes). flint is mostly Q5_K and IQ4_XS, so that kernel runs in almost every layer. I read the diff and did not find the fault in it; a quick discriminator is to force the stored path (or the rebuilt one) for every block and rerun the detector.
separately, a real bug of mine in Bump XDNA: amd/xdna-driver 49fe4e1386af #298 that I have not shown to matter here: the storage-range fallback of the alternate-value lookup ignores the scale. Fix on fork branch fix/2159-alternate-scale-fallback is not pushed; I will put it up as its own PR.

cl-hrxnan, one fact from the same data that changes how to read the bisect table (for the lead): in the 30-prefill run at 6ec7c26 the first ten prefills were all clean, and all 12 deviating ones were among prefills 11 to 30 ([11, 13, 14, 15, 16, 17, 21, 23, 24, 26, 27, 30]). So the "0 of 10" at 2293f95 does not show the base is clean; it may only show that ten in one process is too few, and #292 may not be the cause. A 30-prefill run at 2293f95 settles it: .probe/rate-step.sh 2293f95.

The mechanism: the fused GATED_DELTA_NET dispatch is nondeterministic — and it is the only thing covering the norm

Deterministic reproducer first (this is what makes the rest possible): Flash-Next, --no-repack --cache-reuse 256, mask kernels ON, -c 19456, one 76,000-char prompt (18,782 tokens),
n_predict 12, cache_prompt false — NaN on the first request, every run.

Then two one-env bisects, same config, 6 identical greedy requests each:

arm	NaN	the six identical requests
default	2	six distinct hashes — six identical greedy requests, six different replies (this is #2146's symptom, reproduced), then the NaN
`GGML_HRX_DISABLE_DISPATCH=fused`	0	fails: `unsupported HRX node 15: L2_NORM f32[128,16,2,1] … <-VIEW`
`GGML_HRX_DISABLE_DISPATCH=gated_delta_net`	0	fails the same way (`unsupported … L2_NORM … <-VIEW`)

Three conclusions, and the third is the one that matters for a fix.

The nondeterminism and the NaN are the same defect: with the fusion on, identical inputs give different outputs and eventually a NaN; with it off, both disappear.
The fused GATED_DELTA_NET dispatch is the only thing that covers that L2_NORM — disable it and the node is unsupported (the hrx: the fused GATED_DELTA_NET matcher leaves its own L2_NORM input uncovered when the norm's view is strided (blocks #679's structural fix) #1925 symptom). So the coverage gap and the NaN are two faces of one dependency: the norm runs only inside that fusion, and that fusion is the nondeterministic part.
So this is not a numerical corner in a softmax. A fused dispatch producing six different outputs from identical inputs is reading state it did not write — an uninitialised or stale LDS/snapshot buffer for this shape, or a cross-warp ordering assumption that does not hold at -c 19456. That is a kernel/dispatch defect, and it is the one thing worth fixing.

What I am doing next on #2159 (lead): narrow from "the GDN fusion" to which registration by bisecting the substrings (gated_delta_net.direct, …projection_pair, …selected_snapshot are three separate registrations), and check whether the nondeterminism survives when the model is run with a packed layout (--no-repack absent) — because the earlier bracket says the pair matters, and that would tell us whether the kernel reads the weight layout or the state
incorrectly.

For #2146 (support, cl-hrxnan leads): the six-distinct-hashes result is its symptom reproduced with a named cause, and the bisect above is the recipe. I will keep the finding on both issues rather than duplicating the work.

bump llama.cpp to 4107043c1: a weight in the file mapping keeps one d…

Cross-confirmation from a different kernel and a different row, which makes this a fleet-level property rather than a quirk of flint.

I lead engine#2150 (the draft head's logits step, mul_mat.q6_k.token1_shortlist) and did not know about your contention result when I measured this, so these are two independent observations of the same dependency:

#2150, same case, same binary, 200 runs, sliced by ambient GTT at run start (each run prints the guardrail line):

ticket	GTT in use	failures
`i2150-shortlist`	42.4 – 42.6 GiB (other lanes resident)	5 / 200
`i2150-diag`	38.3 GiB	0 / 600 (4 arms × 150)

and within the reproducing ticket, failing and passing runs had the same pressure (42.6/42.5/42.5/42.4/42.4 both), so it is not a per-run threshold — it is the ticket's ambient state, exactly your "other processes using the GPU" with a different measurement.

What the two rows together now say: this is not kernel-specific and not model-specific. Flint's 8.8k prefill deviation and one wrong refined logit in a 65k-way shortlist are both contention-dependent nondeterminism, which is also the best explanation available for why bisects pointed at different commits on different tickets — the commit under test and the co-tenant load are confounded, and no single-ticket A/B can separate them unless the load is controlled.

Your "4 of 5 with a second process busy throughout" is the most useful number in this thread, because it turns an intermittent fault into a reproducible one: co-tenant load can be created deliberately inside a ticket. I am testing exactly that on #2150 now (a second HRX process keeping the device busy while the case loops), because it would end the "wait for the right ambient state" problem that has blocked that row for a day. If it works, the same trick applies to your flint arm — and the honest experiment then becomes "does the deviation vanish if serial mode also removes co-tenant load", rather than "does serial mode change anything", since you have already shown serial mode alone does not.

Support slice, device-free: a concrete place in the attention-sink kernel that matches this issue's 16-token block granularity, offered as a pointer rather than a conclusion.

Barrier census across the attention family (kernel-corpus/kernels/loom-libs/ops/):

kernel	`kernel.barrier` count
`flash_attention_decode_split_f32_f16_wmma.loom`	36
`flash_attention_f32_f16_wmma.loom`	14
`attention_sink_f32.loom`	2
`llm_attention_{k,q,v}_matmul_rope_*` (the cache writers)	0

The sink kernel is the least synchronised of the three that this issue's path uses, and its synchronisation is exactly two acq_rel workgroup barriers around two reductions:

:81   %stage = buffer.alloca<workgroup> ... %stage_bytes        // view<16xf32>
:119  view.store %subgroup_max, %stage_view[%subgroup_slot0]    // reduction 1 -> slot 0
:121  kernel.barrier<workgroup> scope(workgroup) ordering(acq_rel)
:123  %slot_max = view.load %stage_view[%slot]                  // read back across subgroups
:134  view.store %subgroup_sum, %stage_view[%subgroup_slot1]    // reduction 2 -> slot 1
:136  kernel.barrier<workgroup> scope(workgroup) ordering(acq_rel)
:139  %slot_sum = view.load %stage_view[%slot]
:151  %value = view.load  %output_view[%token, %head, %channel]  // read-modify-write of the FA output
:154  view.store %result, %output_view[...]

Three questions for the lead, all answerable by reading, and each one would produce this issue's symptom:

Does the stage's 16 slots cover every subgroup that writes them? The stage is view<16xf32> and the slot is derived from %subgroup; if the launch ever has more than 16 subgroups per head, two subgroups share a slot and both reductions return a mix of two heads' maxima/sums. A wrong max is a scaling error — nonzero, finite, and it affects a whole tile at once, which is the shape of "a block of 16 tokens goes non-finite mid-network" while the inputs are finite.
Is there a barrier between successive head iterations that reuse the same slots? Both reductions reuse the same 16-slot stage, and there is no barrier at the top of the loop body — only after each store. If the loop over heads re-enters a store before the previous iteration's read retired, on a contended device the read can see the next head's value.
Is the :151 read of %output_view ordered after FlashAttention's writes to it? That is cross-dispatch rather than intra-kernel, so it should be a binding dependency — but if it is, the sink kernel is doing a read-modify-write of a buffer another dispatch produced, and that pair is worth naming explicitly in whatever fix this issue gets.

Why I am not claiming it: #2159 is pi-i679's row and this is a pointer from a support lane, not a diagnosis — the barrier count alone proves nothing (a kernel can be correct with two barriers if its reductions are all it synchronises). What makes it worth the lead's five minutes is the coincidence of 16 slots / 16-token tiles with the reported granularity, and that it is the one kernel in the path whose synchronisation is minimal.

Separately, for the record: engine#2150 (my row) measured compute co-tenancy insufficient to reproduce its own variant of this contention dependence — 0/150 with a busy second process at 38.9 GiB GTT, against 5/200 at 42.4–42.6 GiB. So if this issue's "4 of 5 with a second process busy" holds, the stressor that matters differs between the two kernels, and the shared variable may be memory pressure rather than pure device busyness.

Following my own questions to the end, and they refute the pointer I just raised — better to say so now than let the lead spend a ticket on it.

:43  kernel.launch.config workgroups(%query_head_count, %tokens, %c1) workgroup_size(%c256, %c1, %c1)
:67  %subgroup = kernel.subgroup.id
:117 %subgroup_slot0 = index.assume %subgroup [range(%subgroup, 0, 7)]
:122 %row_max = scf.for %slot = [%c0 to %c8 step %c1](...)

Question 1 (do 16 slots cover every subgroup?): yes, with room to spare. A 256-thread workgroup is 8 subgroups of 32, %subgroup ∈ [0,7], and the reduction reads [0,8) — the view<16xf32> stage is twice the size needed. No slot aliasing, so no mixed maxima and no scaling error from that cause.
Question 2 (barrier between head iterations reusing the stage?): there are no head iterations. workgroups(heads, tokens, 1) launches one workgroup per (head, token), so each workgroup owns its own [token, head, *] slice, and the read-modify-write of %output_view at :151/:154 touches only its own elements. The reuse-across-iterations hazard does not exist.
Question 3 (ordering after FlashAttention's write) remains open, and it is the only one of the three that survives — it is a cross-dispatch binding question rather than a kernel-internal one.

So the 16-slot / 16-token coincidence was a coincidence, and the sink kernel's synchronisation is proportionate to what it actually does (two cross-subgroup reductions through workgroup memory, and nothing else shared). My barrier-count census stands as a fact, but "fewest barriers" is not evidence of a missing one, and I should have stated that ordering — read the geometry first, then count — rather than presenting the count as suggestive.

That leaves this issue's 16-token block without a candidate in the sink kernel from my side. The remaining places I would look, in order, given what is now excluded: (a) the cross-dispatch edge into the sink kernel (question 3) — whether the read of %output_view is ordered after FlashAttention's write, which the scheduler must provide via bindings and which a contended schedule would expose; (b) flash_attention_decode_split_f32_f16_wmma's 36 barriers — a kernel with that many synchronisation points has a lot of stage reuse, and the split variant is the one that runs at long context, which is this issue's condition.

The run the owner asked for could not complete on this box: the build is ready, the pool is not (GTT in use 40.0 of a 48.3 GiB budget)

What I did, exactly as instructed: fetched 1bit/2146-disable-dispatch-log @ 3e8441251 (hrx: hip_top_k_f32 returns the same indices in the same order every run), rebuilt it (ninja llama-server, rc 0, 134 steps), and ran the 8.8k flint protocol -- ~/models/flint/flint-UD-Q4_K_XL.gguf, -c 8832, one fresh llama-server per launch, 10 launches, every launch carrying a COTENANT: verdict, inside a gpu-hold ticket.

Result: zero completed launches. Not a wrong answer -- no answer:

server log, last line at 13.5 s:  placement: 1 of 3056 nodes on the CPU (GET_ROWS x 1)
then, for 180 s sampled every 5 s:  srv_cu=0   loglines=7 (unchanged)

The server comes up, accepts the request, and never touches the compute units -- cu_occupancy for the process stays 0 for three minutes while a request is outstanding. So it is a stall before the first kernel, not a slow prefill, and that is why two 30-minute attempts produced zero rows.

Why, from the build's own guardrail line:

hrx: HRX0 guardrail: GTT total 120.0 GiB, in use 40.0 GiB, ceiling 104.0 GiB (reserve 16.0 GiB),
                     pool budget 48.3 GiB (GTT term 64 ...

The box is already carrying 40.0 GiB of GTT (house tenants -- the lemonade embed stack is the big one), which leaves about 8 GiB of pool budget for a 24 GiB model. The mapping succeeds, the placement report prints, and the first compute that would upload weights has nowhere to put them; the request then waits without progress instead of failing. That is consistent with what this row has needed all along: the 18.8k arms are gated at free -g >= 60 precisely because the house tenants hold ~36 GiB.

What is needed to run it, and it is not mine to free: pool headroom of roughly 30 GiB (24 GiB model + 8.8k KV + transient arena), i.e. GTT in use down to ~18 GiB or less, or equivalently the house embed stack stopped for the duration. With that, the command is one line and the harness is ready:

~/lb/i679-scratch/armrun.sh --n 5 --min-free 40 --ctx 8832 \
  --model ~/models/flint/flint-UD-Q4_K_XL.gguf --port 19885 --a baseline --b baseline

(--n 5 x 2 arms = the 10 launches; --resume continues a session that is interrupted; every row prints nan=, the reply hash and cotenant=.)

If it comes back 10/10 identical I will post that here immediately, and #2159 closes on the #2146 pin bump carrying Fixes #2159 (cl-hrxnan's PR is the place to add the trailer). If it does not, the rows will say exactly which launches differ and the per-launch COTENANT: verdicts will say whether any of them was contended -- which is the difference between "the top-k fix is incomplete" and "the box ran a bad arm".

Hygiene, in full: one orphan llama-server did survive a killed scope during the first attempt; I killed it by pid (129123) and confirmed no process, no gpu-hold token and no scope remained. Both subsequent attempts ended clean. No server is running now, and no ticket is held.

Who is holding the 40 GiB, so the decision is one command rather than an investigation

Following my blocked-run report: the pool the flint run needs is held by platform units, not by a stray process and not by my lane.

/sys/class/kfd/kfd/proc/<pid>/vram_*            # the only per-process accounting this kernel exposes
  pid 2522754  llama-server  4 GiB   parent 1bit   -> embed-server.service
  pid 2528251  llama-server 32 GiB   parent 1bit   -> embed-server.service
box GTT used: 39 GiB        (guardrail: pool budget 48.3 GiB, ceiling 104.0 GiB, reserve 16.0)

systemctl --user carries them as embed-server.service ("1bit embed-server (lemonade SEO, port 8088)"), which is a managed unit, and coder-llm.service (Qwen3-Coder-30B for PR-Agent on :8090) is the other managed GPU tenant. Neither is mine to stop: one serves embeddings to every lane, the other answers the PR-agent reviews. Stopping either to free GTT for my measurement is an owner call, not a lane call, which is why I am not doing it.

So the options are, in the order I would try them:

Stop embed-server.service for the duration (it is the 36 GiB; coder-llm.service is the smaller tenant and the PR agent needs it). Then:
```
~/lb/i679-scratch/armrun.sh --n 5 --min-free 40 --ctx 8832 \
  --model ~/models/flint/flint-UD-Q4_K_XL.gguf --port 19885 --a baseline --b baseline
```
10 launches, fresh server each, cotenant= per launch, --resume if it is interrupted; ~40 minutes of wall clock at the flint load cost.
A window on another box with ≥30 GiB of pool headroom, running the same command (the harness needs only the built binary, the model and gpu-hold).
Raise the guardrail's pool budget -- the line says pool budget 48.3 GiB (GTT term 64 …); if that term is a tunable, a bigger pool would let a 24 GiB model load beside the tenants, at the cost of the headroom the guardrail exists to keep. That is a policy change and I would not make it unilaterally either.

One correction to my own report, in the interest of the row's record: the stall I measured is not evidence about the top-k fix. The request never reached a kernel (cu_occupancy = 0 for 180 s), so nothing in that run says whether 3e8441251 makes the prefills repeatable. The build is ready and the harness is ready; the missing ingredient is headroom.

For #2159 (pi-i679 / pi-attn): a build with the #2146 TOP_K fix, ready to re-measure on

On #2146, hip_top_k_f32 (HRX TOP_K) was found not to be repeatable. Its emit took slots by atomicAdd, so the order differed every run and the set differed under ties, and there was a race between levels that gave invalid rows. The fix is llama.cpp PR #328, branch 1bit/2146-top-k-repeat @ d445771, on top of pin 962507f8e.

The build is on strixhalo at ~/wt/cl-hrxnan-build-topk/hrx/llama/bin: llama-server, libggml-hrx.so, and test-hrx-top-k-repeat (which passes there). It is engine -DONEBIT_HRX=ON with hrx-system 7f42369 and nothing else changed. Use the same environment as the bench build: LD_LIBRARY_PATH=<that bin>:<TheRock lib>, IREE_HAL_AMDGPU_LIBHSA_PATH, and the CWSR trio.

If flint's graph runs hip.top_k_f32 (check with GGML_HRX_LOG_DISPATCH=1), the 4-of-10 different prefills at 8.8k may share this cause. Please leave that tree in place until #2146 closes; I will remove it after.

Ran the ten: on the pre-merge build `3e8441251` all ten fail with NaN, so this build does not fix #2159 -- the merged pin is the thing to test

Self-service memory window, taken per the owner's rule (ticket held, coder-llm stopped for it, GTT 31 G -> 5 GiB, avail 104 G, restored by trap; :8090 200 and coder-llm=active at the end). ~/models/flint/flint-UD-Q4_K_XL.gguf, -c 8832, one fresh llama-server per launch, 10 launches, every one with a COTENANT: verdict:

A#1..A#5, B#1..B#5   nan=2   hash=REQFAIL:Remote end closed connection without response   cotenant=ACCEPT
arm            launches  NaN  ACCEPT  distinct outputs
A                     5     5       5                 1
B                     5     5       5                 1
ACCEPT launches: 10 of 10

Every launch: two NaN logits lines and the request connection closed. Not a wrong answer -- no answer, ten times. So the answer to the owner's question is the second branch: this build does not make the ten prefills identical, it makes them uniformly fail. It is deterministic failure rather than 2-of-3, which is consistent with a partial fix -- the same wrong thing every time instead of a different wrong thing each time.

Why I think it is the wrong build rather than a wrong test: 3e8441251 is the pre-merge head of the top-k work, and it is not an ancestor of the pin:

compare/3e8441251...9dd1a1b72   diverged  (ahead 2, behind 3)

The pin the engine is being bumped to (#2296) is the merged form (#328), whose commit describes two defects -- the selection-level race and the atomicAdd emit order -- with the merged kernel measured at 0 of 49 differing on seven shapes and "4 replies in 4" on identical 18.8k Flash-Next requests. My build predates that merge and, on this evidence, has the first half at most.

So the run that decides #2159 is the one on 9dd1a1b72, not on the branch head. I am rebuilding to the pin now and will run the same ten the moment the 20-minute coder-llm gap allows (the rule the orchestrator set, and the reason I am not running it this minute). If that comes back 10/10 identical, #2159 closes on #2296 with Fixes #2159; if it does not, the same table with the hash of each launch is what the next reader needs.

One detail worth recording for whoever repeats this: the harness reports nan=2 per launch, i.e. the server logs the NaN and then closes the connection without a reply. A future arm should treat "NaN logged" as a result in its own right rather than as a transport error -- the REQFAIL string in the hash column is the request failing, not the measurement.

Ran the ten on the merged pin `9dd1a1b72` too: all ten fail with NaN as well, so #328 is not what makes these prefills differ

Same window discipline as before (ticket, coder-llm stopped, :8090 200 restored, WINDOWS.log line). Built the merged pin 9dd1a1b72 -- the one #2296 bumps the engine to, whose commit fixes both the selection-level race and the atomicAdd emit order, and which its author measured at 0 of 49 differing across seven shapes -- and ran the identical protocol:

flint UD-Q4_K_XL, -c 8832, fresh llama-server per launch, 10 launches, COTENANT per launch
A#1..A#5, B#1..B#5   nan=2   hash=REQFAIL:Remote end closed connection   cotenant=ACCEPT
ACCEPT launches: 10 of 10     1 distinct per arm

Ten out of ten, two NaN logits lines each, request closed. Identical to the pre-merge build (3e8441251), which is the answer to the owner's question in the negative: this is not a build problem, and the top-k determinism fix -- the one measured at 0/49 -- does not make the 8.8k flint prefills identical here. It is not even "different replies": there is no reply.

Which op still differs: on this evidence, none -- the failure is deterministic and value-level, not a race. A race produces variation; ten identical failures under ten fresh processes produce none. So #2159's NaN is a value reaching the attention (or its input) that is not finite, on this model at this context, every time:

deterministic now (10/10 on two builds), where my own earlier baseline was 2-of-3 -- so whatever varies between launches (the hrx: six identical greedy 18.8k requests on Flash-Next gave three different replies (no NaN; one slot, no prompt cache) #2146 races) is not what produces the NaN, and fixing those cannot remove it;
and the one measurement of mine that ever made it succeed was GGML_HRX_LOG_DISPATCH=1 (engine#2143, a switch that is not behaviour-neutral), while DISABLE_DISPATCH=fused gave nan=0 with an uncovered L2_NORM node -- i.e. the previous A/Bs were pointing at the dispatch set, not at a race.

The next test, and it is cheap and decisive: the same model at -c 2048 and -c 4096 against -c 8832. The model is 77 GB of weights (shards 11 M + 47 G + 31 G) on a box with ~70 GiB available, so it is mmapped and demand-paged: if the NaN appears only at the large context, the mechanism is memory (a page that cannot be faulted in, or a partial read) rather than arithmetic; if it appears at every context, the weights themselves are not what the loader thinks they are and the question moves to the pack. Either answer names the op by naming the stage, which is what "say which op still differs" needs. I will take the next window (20-minute coder-llm gap) and post the three results together.

Root cause: it is `hip.mul_mat_id_prefill` -- the prefill GEMM's missing f16 prescale -- and the A/B is unambiguous

Same model, same context, same build, one launch each, in one npu... no, gpu-hold ticket:

flint UD-Q4_K_XL, -c 2048, one fresh llama-server per arm
default                                  REQFAIL (connection closed)   nan_lines=2
DISABLE_DISPATCH=hip.mul_mat_id_prefill   OK, 1767 tokens,
                                          content ' ct date is disputed , but most sources'   nan_lines=0

Disabled: coherent text, zero NaN. Enabled: NaN and the request dies. Nothing else changed between the two arms.

And the mechanism is the one gpu-kernels#68 (engine#2293) fixes. hrx_prefill_gemm rounds each f32 activation straight to f16; a value above 65504 becomes inf and the f32 result becomes NaN. The fix is a 2^-6 prescale before the narrowing and a 2^6 undo on the accumulator, and it is not in the addon this build uses:

grep -cE 'kXs|0\.015625' ~/wt/gk-clone/addons/hip-hrx/kernels/hrx_prefill_gemm.hip   ->  0

(GGML_HRX_HIP_ADDON_DIR in the build is that tree.) So the fault reaches the product through the addon, and it needs both halves: #68 merged into the addon, and the addon rebuilt into the binary. My flint rows have been measuring a kernel without the prescale.

This also explains three earlier results that read as separate faults, which is worth saying because it changes what they were evidence for:

engine#2143 (LOG_DISPATCH=1 is not behaviour-neutral): enabling the log changed which dispatches were taken, and with the prescale missing that decided whether the prefill GEMM ran at all -- so the switch looked causal when it was only changing exposure to this bug.
my own DISABLE_DISPATCH=fused degeneration (at secondary riv riv ch, 1.81 log units at the first divergence): disabling the fused family moved the same activations onto a path without the prescale.
the "2 of 3 NaN" baseline and the 10/10 uniform failure: both are this kernel; the variation was whether it was reached, not whether it was wrong. ci(bump-hrx): merge AMD's pin into our line instead of rebasing 159 commits (fixes the 5-day silent sync failure) #328's top-k race and census: new HF model class uncovered by the registry #330's global-order race are real and were worth fixing, but neither is this.

What closes #2159: hip.mul_mat_id_prefill with the 2^-6 prescale (gpu-kernels#68), rebuilt into the addon. Then the ten identical 8.8k prefills should be ten identical replies -- and that run is the one that should be posted as the close, not this one, because an A/B that turns a kernel off proves the kernel is the fault but does not prove the fix fixes it.

Signal for the row's checklist: the same test is a cheap regression check for any future prefill-GEMM change -- -c 2048 on this model, default vs DISABLE_DISPATCH=hip.mul_mat_id_prefill, and compare: OK-vs-NaN is a one-launch verdict that needs no long context and no ten-launch arm.

The fix fixes it: rebuilt the addon with gpu-kernels#68's 2^-6 prescale and the A/B inverted

I said the A/B that turns a kernel off proves the kernel is the fault but does not prove the fix fixes it. So I took the fix and ran both arms again, same model, same -c 2048, same build except the kernel:

arm	before the prescale (addon without #68)	after the prescale (addon with #68)
default	REQFAIL, `nan_lines=2`	OK, 1767 tokens, `'usted [ sic ] Emperor Xuanzong'`, `nan_lines=0`
`DISABLE_DISPATCH=hip.mul_mat_id_prefill`	OK, coherent	REQFAIL, `nan_lines=2`

Exactly inverted, which is the strongest form this evidence can take: the arm that was broken now works, and the arm that worked is now broken, with the prescale as the only difference between the two builds. The fault is hrx_prefill_gemm rounding activations straight to f16 (engine#2293); the repair is #68's 2^-6 staging plus the 2^6 undo on the accumulator.

How I got it into a build: ~/wt/gk-clone is the tree GGML_HRX_HIP_ADDON_DIR points at, so I checked out addons/hip-hrx/kernels/hrx_prefill_gemm.hip from #68's head (1c3db06ea), rebuilt (ninja llama-server, rc 0), and ran the pair. That file is still checked out in my working copy -- not committed, not pushed; it is the PR's content and the PR is the way it lands.

And the second arm failing now is itself informative: with the HIP kernel fixed, DISABLE_DISPATCH=hip.mul_mat_id_prefill sends the same activations to the Loom tile, which in this llama.cpp pin still NaNs. #68's body says the Loom tile got the same 2^-6 in llama.cpp #285/#288 -- so either those are not in the pinned commit or the tile's staging differs. Either way it is a second, independent reason the prescale has to be everywhere the f16 staging happens, and it is worth a line in #68 rather than a surprise later.

Close-out run is queued: the ten identical 8.8k flint prefills on this build, in a self-service window (delay for the coder-llm gap, :8090 restored and confirmed on every path). If they come back 10/10 identical, #2159 closes here with #68 (and its engine-side pin) as the fix, and I will post the table. If they do not, this is not the whole story and the hashes will say what is left.

CLOSED: ten identical 8.8k flint prefills, ten identical replies -- and the fix is gpu-kernels#68's 2^-6 prescale, not the #2146 top-k bump

Self-service window (coder-llm stopped inside the ticket: GTT 31 G -> 5 GiB, free 77 G, avail 101 G), flint UD-Q4_K_XL, -c 8832, one fresh llama-server per launch, ten launches, COTENANT on every one:

A#1..A#5, B#1..B#5   nan=0   hash=91c8766913   cotenant=ACCEPT
arm   launches  NaN  ACCEPT  distinct outputs
A            5     0       5                 1
B            5     0       5                 1
ACCEPT launches: 10 of 10

Ten launches, one hash, zero NaN. The same protocol on the same model produced ten failures on both 3e8441251 and the merged pin 9dd1a1b72, and produced 2 of 3 NaN with three distinct hashes on the baseline that opened this issue.

What changed between failing and passing is one kernel file. hrx_prefill_gemm rounded f32 activations straight to f16; anything above 65504 became inf and the f32 result NaN. gpu-kernels#68 stages at 2^-6 and undoes at 2^6. The isolated A/B, same model and context:

arm	addon without #68	addon with #68
default	REQFAIL, `nan_lines=2`	OK, coherent, `nan_lines=0`
`DISABLE_DISPATCH=hip.mul_mat_id_prefill`	OK, coherent	REQFAIL, `nan_lines=2`

So #2159 closes on gpu-kernels#68 (and the engine-side bump that carries it), not on #2296. #2296 bumps the pin to the top-k determinism fix (#328), which is a real defect and makes the top-k index set stable -- but it does not close this: ten identical prefills still failed on that pin, twice over. Please add Fixes #2159 to #68 (or its engine bump), not to #2296.

What this re-reads, and it matters for the record: three findings I filed or reported during this investigation were exposure to this one bug, not separate faults --

hrx: GGML_HRX_LOG_DISPATCH is not behaviour-neutral — the 18.8k request fails without it and succeeds with it #2143 (GGML_HRX_LOG_DISPATCH is not behaviour-neutral): the switch changed which dispatches were taken, and with the prescale missing that decided whether the prefill GEMM ran at all;
my DISABLE_DISPATCH=fused degeneration (at secondary riv riv ch, 1.81 log units at the first divergence): it moved the same activations onto a path without the prescale;
the 2-of-3 NaN "nondeterminism": whether the kernel was reached varied per launch, not whether it was wrong. ci(bump-hrx): merge AMD's pin into our line instead of rebasing 159 commits (fixes the 5-day silent sync failure) #328's top-k race and census: new HF model class uncovered by the registry #330's global-order race are real and worth their fixes, but neither is this.

And the cheap check that falls out, which I would keep in the row: -c 2048 on this model, default versus DISABLE_DISPATCH=hip.mul_mat_id_prefill -- OK-versus-NaN in one launch, no long context, no ten-launch arm. That is the regression test for any future change to the prefill GEMM's staging, and it would have found this in a minute.

One honest note about my own tooling, since it affected the box: the window wrapper I used for this run had exec bash -c "$CMD", which replaces the shell and discards the EXIT trap -- so the run itself was fine but coder-llm was not restarted and no WINDOWS.log line was written; I restored the service by hand and removed the exec. The :8090 confirmations in WINDOWS.log from my earlier windows are from a different wrapper that did not have that bug.

gpu-kernels: pin b0ed206, which carries both f16-range fixes (#2159, #…

gpu-kernels: pin eb0f3839 — the 2^-6 prefill prescale (#2159), the ro…

Navigation Menu

Repository navigation

hrx: 4 of 10 identical 8.8k-token prefills on flint give different logits (a block of 16 tokens goes non-finite mid-network); usually silent, sometimes NaN logits and a server abort #2159

What happens

Measured (strixhalo, gfx1151, each run under gpu-hold, 2026-11-11)

What is and is not established

Reproduce

Label

Source audit of `fa_decode.hip`: the classic unguarded corners are absent — which leaves exactly one hole, and it is in the guard itself

Second kernel audited: the sink correction is guarded too, so the NaN most likely arrives as an input — here is the whole structure

Correction to my own bracket: the failing axis is the KV capacity, not the prompt length

Third layer audited — the graph build is guarded too, so the device-free surface is now exhausted: the mask dump is the next step

The mechanism: the fused GATED_DELTA_NET dispatch is nondeterministic — and it is the only thing covering the norm

7 remaining items

The run the owner asked for could not complete on this box: the build is ready, the pool is not (GTT in use 40.0 of a 48.3 GiB budget)

Who is holding the 40 GiB, so the decision is one command rather than an investigation

Ran the ten: on the pre-merge build `3e8441251` all ten fail with NaN, so this build does not fix #2159 -- the merged pin is the thing to test

Ran the ten on the merged pin `9dd1a1b72` too: all ten fail with NaN as well, so #328 is not what makes these prefills differ

Root cause: it is `hip.mul_mat_id_prefill` -- the prefill GEMM's missing f16 prescale -- and the A/B is unambiguous

The fix fixes it: rebuilt the addon with gpu-kernels#68's 2^-6 prescale and the A/B inverted

CLOSED: ten identical 8.8k flint prefills, ten identical replies -- and the fix is gpu-kernels#68's 2^-6 prescale, not the #2146 top-k bump

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

hrx: 4 of 10 identical 8.8k-token prefills on flint give different logits (a block of 16 tokens goes non-finite mid-network); usually silent, sometimes NaN logits and a server abort #2159

Description

What happens

Measured (strixhalo, gfx1151, each run under gpu-hold, 2026-11-11)

What is and is not established

Reproduce

Label

Activity

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

Source audit of fa_decode.hip: the classic unguarded corners are absent — which leaves exactly one hole, and it is in the guard itself

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

Second kernel audited: the sink correction is guarded too, so the NaN most likely arrives as an input — here is the whole structure

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

Correction to my own bracket: the failing axis is the KV capacity, not the prompt length

bong-water-water-bong commented on Nov 11, 2026

Third layer audited — the graph build is guarded too, so the device-free surface is now exhausted: the mask dump is the next step

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 11, 2026

The mechanism: the fused GATED_DELTA_NET dispatch is nondeterministic — and it is the only thing covering the norm

7 remaining items

bong-water-water-bong commented on Nov 11, 2026

bong-water-water-bong commented on Nov 12, 2026

bong-water-water-bong commented on Nov 12, 2026

bong-water-water-bong commented on Nov 12, 2026

The run the owner asked for could not complete on this box: the build is ready, the pool is not (GTT in use 40.0 of a 48.3 GiB budget)

bong-water-water-bong commented on Nov 12, 2026

Who is holding the 40 GiB, so the decision is one command rather than an investigation

bong-water-water-bong commented on Nov 12, 2026

bong-water-water-bong commented on Nov 12, 2026

Ran the ten: on the pre-merge build 3e8441251 all ten fail with NaN, so this build does not fix #2159 -- the merged pin is the thing to test

bong-water-water-bong commented on Nov 12, 2026

Ran the ten on the merged pin 9dd1a1b72 too: all ten fail with NaN as well, so #328 is not what makes these prefills differ

bong-water-water-bong commented on Nov 12, 2026

Root cause: it is hip.mul_mat_id_prefill -- the prefill GEMM's missing f16 prescale -- and the A/B is unambiguous

bong-water-water-bong commented on Nov 12, 2026

The fix fixes it: rebuilt the addon with gpu-kernels#68's 2^-6 prescale and the A/B inverted

bong-water-water-bong commented on Nov 12, 2026

CLOSED: ten identical 8.8k flint prefills, ten identical replies -- and the fix is gpu-kernels#68's 2^-6 prescale, not the #2146 top-k bump

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Source audit of `fa_decode.hip`: the classic unguarded corners are absent — which leaves exactly one hole, and it is in the guard itself

Ran the ten: on the pre-merge build `3e8441251` all ten fail with NaN, so this build does not fix #2159 -- the merged pin is the thing to test

Ran the ten on the merged pin `9dd1a1b72` too: all ten fail with NaN as well, so #328 is not what makes these prefills differ

Root cause: it is `hip.mul_mat_id_prefill` -- the prefill GEMM's missing f16 prescale -- and the A/B is unambiguous