[bong-water-water-bong]

Fixes #2146

What this does

Moves third_party/llama.cpp (and registry/architectures.json) from 962507f8e to 9dd1a1b72 on 1bit/hrx-vulkan-patched:

• hrx: hip_top_k_f32 returns the same indices in the same order every run (engine #2146) llama.cpp#328: hip_top_k_f32 returns the same indices in the same order every run. The radix select's shared bucket was reset before every lane had read it, and the emit placed elements with atomicAdd. The Flash-Next indexer's TOP_K picks the KV rows the sparse attention reads, so identical long prompts gave different replies.
• hrx: say so when DISABLE_DISPATCH=fused means the category, not the substring (engine#2146) llama.cpp#326: a log line for GGML_HRX_DISABLE_DISPATCH=fused (category, not substring).

Before / after (strixhalo, Qwen3.8-Flash-Next UD-IQ3_XXS, 18,782-token prompt, greedy, 512 tokens)

Same request every time: the #2146 18,782-token prompt, n_predict 512, temperature 0, top_k 1, cache_prompt false, -c 19456 -np 1. Each request was checked ten times a second for other processes on the compute units.

build	server	other processes on the CUs	requests	distinct replies	NaN guard lines
before: pin 962507f8e	one server, coder-llm stopped (07:35Z)	none	4	4	0
after: the fix (d445771, see below)	one server, coder-llm stopped (18:07Z)	none	5	1	0
after: the fix (d445771)	another server, coder-llm up (15:04Z)	coder-llm, up to 40 CUs	1	same reply	0

So six requests on the fixed build in two servers, one with another process computing, gave one reply (sha1 of the content 6b9ddd072b in all six). Before the fix, four requests in one server with the GPU to itself gave four. The sixth request in the 18:07Z window was cut off by the window's time limit, not by a failure. Prefill in the window was 64 to 78 tok/s, the same as before (67 to 79).

d445771 is the head of 1bit-MONSTER/llama.cpp#328, which merged as 9dd1a1b. The two trees differ only by 1bit-MONSTER/llama.cpp#326, a log line for GGML_HRX_DISABLE_DISPATCH=fused.

Cost of the fix

hip_top_k_f32 alone, one launch at the indexer's shapes (2048 rows, k 2048), best of 5 x 200:

columns (KV rows)	before	after
2,560	158 us	271 us
10,240	645 us	977 us
18,944	1,171 us	1,642 us

Under 0.5 ms more per launch, against about 4.5 minutes of prefill for this prompt.

Checks

• tools/check_pins.py origin/main: ok third_party/llama.cpp: 962507f8e -> 9dd1a1b72 is ahead.
• test-hrx-top-k-repeat (added by hrx: hip_top_k_f32 returns the same indices in the same order every run (engine #2146) llama.cpp#328): FAIL on 962507f8e, ok on the fix head.
• Not run on this branch: the engine's ctest. The change is the two pin lines only.

🤖 Generated with Claude Code

[bong-water-water-bong]

VERDICT: READY at 92757a9 — Fixes #2146, and the before/after table contains a row that exists to kill the obvious objection.

The measurement, with its confounder controlled

before: pin 962507f8e        one server, coder-llm stopped (07:35Z)   none           4 runs  4 distinct  0
after:  the fix (d445771)    one server, coder-llm stopped (18:07Z)   none           5 runs  1 distinct  0
after:  the fix (d445771)    another server, coder-llm up (15:04Z)    coder-llm, 40 CUs  1 run  same reply

4 of 4 distinct before; 1 of 5 after; and 1 of 1 with another process burning up to 40 CUs. That third row is the one that makes the claim safe: nondeterminism under concurrency is exactly what a GPU-timing bug looks like, so a determinism claim that only holds on an idle box would be worthless. Six requests across two servers, one of them under load, producing a single reply (sha1 of the text) is the right evidence for #2146's "six identical greedy 18.8k requests gave three different outputs".

Why the pin is the right carrier

The fix is in the fork (llama.cpp#328: "hip_top_k_f32 returns the same indices in the same order every run. The radix select's shared [state]"), so it lands here as a gitlink + registry/architectures.json source sha pair — both pins, as #2241/#2204 established. And the pin carries a test (test-hrx-top-k-repeat, seven cases, twelve runs each), so the property is guarded at the next bump rather than only at this one.

One ask

say what happens to the previously-reported divergent outputs — if any of the three earlier replies from #2146 is reproducible now, that is either a residual nondeterminism path (a second cause) or an artifact of the old build; if none is, the single-cause conclusion is stronger and worth one sentence. #2146's whole difficulty was that three outputs looked like noise; knowing whether any of them recurs after the fix is the difference between "fixed" and "no longer observed".

[bong-water-water-bong]

VERDICT: READY at 92757a9

A determinism fix measured the way a determinism fix has to be measured -- by counting distinct replies -- and the cost is printed beside the benefit.

The claim: "identical 18,782-token prompts give one reply" (fixes #2146).

build	requests	distinct replies
before, pin 962507f8e	4	4
after, the fix	5	1

Four requests, four different answers, on a greedy decode with temperature 0, top_k 1. That is not noise -- a greedy path is supposed to be a function of its input, so four outputs from one input is a defect
whatever the values are. And the third row is the one that makes it a result rather than a coincidence: "another server, coder-llm up (15:04Z) … coder-llm, up to 40 CUs … 1 request, same reply" -- the fixed build
stays deterministic with another process computing on the same CUs, which is the condition under which the old one failed.

The answer to "was it contention?" is in the table, not in a sentence: the before-arm had coder-llm stopped and no other processes on the CUs, and still gave 4 of 4. So the previous hypothesis (concurrent
compute) is disproved by the arm that removed it.

The cost is stated in the same units as the benefit: hip_top_k_f32 at the indexer's shapes goes 158 → 271 us at 2,560 columns, 1,171 → 1,642 us at 18,944 -- "under 0.5 ms more per launch, against about 4.5
minutes of prefill for this prompt." A 40 % slower kernel that buys determinism, quantified against the thing it sits inside.

Both checks are named with their result: check_pins.py origin/main → ok … is ahead (which I reproduced), and test-hrx-top-k-repeat -- the arm from llama.cpp#328 -- "FAIL on 962507f8e, ok on the fix head".
That is the right pairing: the test fails on the old pin and passes on the new one, so the pin bump is the thing that fixes it.

Confirmed independently: 9dd1a1b72 exists, is a fast-forward (2 commits: #328 the top_k fix, #326 the fused log line), registry/architectures.json changes by one line, and registry_build.py --check-pins is rc 0.

One note: "the sixth request in the 18:07Z window was cut off by the window's time limit, not by a failure" -- good to say, and it is the only place a reader could mistake a truncation for a result. Worth keeping the wording
exactly as it is.

Not run: the engine's ctest (the change is the two pin lines); the 18.8k measurements are the author's, from three windows.

[bong-water-water-bong]

Pin verified consistent; one gap in what the tool can tell you about it, and one thing the body should carry.

Checked without needing the submodule contents:

git ls-tree HEAD third_party/llama.cpp            -> 9dd1a1b72384fbda84c100911296d2a0e96674cb
registry/architectures.json "llama.cpp (hrx)"     -> 9dd1a1b72384fbda84c100911296d2a0e96674cb
origin/main gitlink                               -> 962507f8e8ba2cdbecd6ba061bb49a6900e4881b

So the gitlink and the registry agree at this head, and main is at the old commit — the two halves of a pin bump are in step, which is the failure this shape usually has.

tools/ci/submodule_pin_check.py cannot confirm it here, and its summary reads like a pass when it has not checked anything:

submodule_pin: third_party/tokenizers is not initialised (pinned at 301277072263) - not compared
…
submodule_pin: 0 submodule(s) at their pinned commit, 7 not initialised

"0 submodule(s) at their pinned commit" is the vacuous-pass shape: nothing was compared, and a reader skimming for a zero reads it as agreement. Since an uninitialised submodule is the normal state of a fresh CI checkout (and of every worktree here), the line should say so in the same words the rows do — 0 of 0 compared (7 not initialised) — or exit non-zero under --strict. Not this PR's defect, but this is the PR where someone will go looking for reassurance.

What the body should carry, since it is the whole value of a pin bump: Fixes #2146 with the upstream commit that fixes it named here (1bit-MONSTER/llama.cpp#328, hip_top_k_f32's radix select resetting a shared bucket before every lane read it, plus atomicAdd in the emit), and the arm that shows the fix: #2146's report is six identical greedy 18.8k requests on Flash-Next giving different results, so a repeat of that run on this pin — same prompt, 6×, identical output — is what turns a pin bump into a fixed issue. If that run is queued behind a window, say so; a pin bump whose issue is closed on the strength of the upstream commit alone leaves the next lane unable to tell whether it landed.

[bong-water-water-bong]

VERDICT: READY -- and this is the bump that closes #2159's mechanism, so it deserves the ancestry check spelled out

Registry and pointer agree: sources["llama.cpp (hrx)"] = 9dd1a1b72384fbda84c100911296d2a0e96674cb identical to the new Subproject commit, counts unchanged (298/9/298/298). Nothing to raise there.

The check that ancestry cannot do, done instead. The build I have been measuring on is the pre-merge form of this fix, and it is not an ancestor of the pin:

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

so "is my tested commit in the pin" has no answer. The pin's own commit is hrx: hip_top_k_f32 returns the same indices in the same order every run (engine #2146) (#328) and ggml/src/ggml-hrx/hip/kernels/hip_top_k_f32.hip is present at it -- i.e. the pin is the merged form of exactly the change, which is the statement that matters for anyone holding a measurement taken on the branch.

And the mechanism deserves to be in this PR's body, because it explains a bug that is not its own. The merged commit names two defects, measured on gfx1151 by running the kernel directly, 50 runs per shape:

1. a race in the selection levels -- after the level's last barrier every lane reads bucket_found/take_found while lane 0 resets them for the next level, so a slow lane can publish a wrong threshold_key: 1-12 invalid rows per 50 runs, and a different index set in 2-3 of 49 with distinct inputs. One __syncthreads() after the reads: 0 invalid, set stable.
2. the emit took a slot per element with atomicAdd -- order differed in 49 of 49, and with ties the set differed too. Counting per lane + exclusive scan + a second pass in visit order: 0 of 49 differ on every shape (2560/2051, 4096/2048, 128/8, a 2561 unaligned pitch, one 200,000-row, k == cols), and with distinct inputs the set equals a CPU reference.

The sparse-attention indexer on Flash-Next picks 2051 of 2560 KV rows with this kernel -- so a wrong index set means attending to the wrong rows, which is precisely the shape of engine#2159's symptom: identical greedy 18.8k requests gave two NaN and three distinct hashes in three runs on my side, and 4 replies in 4 with the GPU to itself on the fixer's side. A NaN out of the attention path, and a reply that changes run to run, from one non-deterministic index set -- the mechanism fits the evidence on both sides, which is why I would quote it in the bump rather than leave it in the merged commit message.

Live evidence, in flight as I write this: my self-service memory window is running the 10 identical 8.8k flint prefills on the pre-merge build right now (ticket held, coder-llm stopped for it, GTT down to 5 G, restored by trap). I will post the ten hashes on #2159 when it finishes -- if they are 10/10 identical, this bump closes #2159 as well, and the trailer Fixes #2159 belongs with it.

[bong-water-water-bong]

VERDICT at 92757a9: READY

For a submodule bump the only review that matters is whether the pin is real and whether it contains the fix, so I checked both by content rather than by ancestry.

• 9dd1a1b72 exists and the move is forward: 962507f8e is an ancestor, so nothing was rewound, and CI's own "Submodule pins only move forward" step agrees.
• The pinned tree contains the fix, by name: ggml/src/ggml-hrx/hip/kernels/hip_top_k_f32.hip and tests/test-hrx-top-k-repeat.cpp are both present in the tree at that commit. That is #328's kernel and #328's test.
• And the pin landed in a better place than the old one: 9dd1a1b72 is contained by origin/1bit/hrx-vulkan-patched — the integration branch this engine tracks. The previous pin 962507f8e was not on any integration branch; it was the tip of feature branches (1bit/2146-disable-dispatch-log, 1bit/2146-top-k-repeat, 1bit/572-mmid-tile). So this bump moves the pin from a feature-branch tip onto the branch it should have been on.
• registry/architectures.json changes exactly one line — the llama.cpp (hrx) pin value — so it is the tracked pin and not a regenerated architecture list pretending to be one.

On Fixes #2146: the causal chain is the one #328 states in its own body. hip_top_k_f32 returned different indices in a different order between runs; the sparse-attention indexer on Qwen3.8-Flash-Next picks 2051 of 2560 KV rows with it, so a different order — or a different set among tied scores — gave identical greedy requests different replies. That is exactly #2146's symptom at 18.8k context, and #328's acceptance test is determinism (bitwise-identical across repeated runs, plus the same index set as the CPU backend for distinct inputs). So the bump closes the issue through the fix the issue was diagnosed into, not by coincidence of timing.

I reviewed #328 in this same session when it was a llama.cpp PR, so this is also the first time I have been able to confirm that a reviewed fix is the thing a pin actually carries — the two claims are now the same claim.

Nothing here is device-dependent and the bump carries only two upstream changes (#328, and #326's dispatch log line), so there is no interaction to check against another open PR.

[github-actions]

PR Reviewer Guide 🔍

Here are some key observations to aid the review process:

🎫 Ticket compliance analysis 🔶 2146 - Partially compliant Compliant requirements: The PR fixes the issue by updating llama.cpp to a version that ensures deterministic behavior for hip_top_k_f32 The fix addresses the root cause: the Flash-Next indexer's TOP_K picking KV rows was non-deterministic due to shared bucket reset timing and atomic operations The PR includes a test run showing that after the fix, identical requests now return the same reply (1 distinct reply vs 4 before) Non-compliant requirements: The PR does not explicitly state whether the divergence occurs at the first token or later in the generation The PR does not verify whether the CPU backend gives consistent replies for the same requests Requires further human verification: The exact mechanism of how the shared bucket reset and atomic operations caused non-determinism needs confirmation from the hardware behavior The claim that the fix resolves the issue in the specific test environment (strixhalo, Qwen3.8-Flash-Next) should be validated by running the exact scenario described in the ticket

⏱️ Estimated effort to review: 2 🔵🔵⚪⚪⚪

🧪 No relevant tests

🔒 No security concerns identified

⚡ Recommended focus areas for review Non-deterministic Top-K Selection The fix addresses a non-deterministic behavior in hip_top_k_f32 where shared bucket reset occurred before all lanes had read it, and elements were emitted with atomicAdd. This caused the Flash-Next indexer's TOP_K to pick different KV rows on different runs, leading to inconsistent replies for identical prompts. The fix ensures consistent indices are returned in the same order every run. Subproject commit 9dd1a1b72384fbda84c100911296d2a0e96674cb