Activity

1. ByungHyun21

   added

   bug-unconfirmed

   on Aug 22, 2026
2. thc1006 commented on Aug 24, 2026

   

   thc1006

   on Aug 24, 2026

   More actions

   Reading through the second corruption path you flagged, the one you traced to output_reorder() and then set aside because excluding embd_nextn from the swaps broke single-stream. I think the gate belongs on the mode rather than on the tensor, and the file already makes that distinction in two other places.

   embd_nextn is indexed differently depending on cparams.embeddings_nextn_masked. The allocation says so directly, with the reason in the comment:

   embd_nextn.size = has_embd_nextn ? n_embd_out*n_outputs_max  : 0;
   
   if (has_embd_nextn && !cparams.embeddings_nextn_masked) {
       // unmasked: nextn row exists for every token in the batch, not just
       // those flagged via batch.logits[i] -> size by token count instead.
       embd_nextn.size = (size_t) n_embd_out * n_batch;
   }

   The extraction site switches on the same flag:

   const int64_t offset = masked ? n_outputs_prev : n_tokens_prev;

   output_reorder() does not:

   if (embd_nextn.size > 0) {
       for (uint64_t k = 0; k < n_embd_out; k++) {
           std::swap(embd_nextn.data[i0*n_embd_out + k], embd_nextn.data[i1*n_embd_out + k]);
       }
   }

   and the swaps it applies are output-row indices, taken from the selection sort over out_ids:

   for (uint32_t i = 0; i < n_outputs - 1; ++i) {
       ...
       output_swaps.push_back({ i, j_min });
   }

   So in masked mode the rows are output rows and the swap is right. In unmasked mode they are token rows, n_batch of them rather than n_outputs, and the same swap moves rows that have nothing to do with the sort. That is cross-sequence row shuffling of exactly the kind you suspected, and it only happens when more than one sequence is present in a batch, since out_ids is already ordered otherwise.

   It would also explain the result that made you stop. Removing embd_nextn from the swaps unconditionally takes the reorder away from the masked path too, and that path needs it, which is consistent with acceptance collapsing single-stream at around 5 t/s. The narrower change is to keep the swap and gate it:

   if (embd_nextn.size > 0 && cparams.embeddings_nextn_masked) {

   To be clear about what this is: source reading against c060ca9, not a measurement. I have not built it, I have not run your reproducer, and I am not claiming this is the whole of the second path. What I can say is that the three sites disagree about how embd_nextn is indexed, and that the disagreement is only reachable in the configuration where you saw the collapse.

   On the first path, the async copy: your instrumentation makes the ordering argument on HIP, and the call is unchanged at src/llama-context.cpp:1954 in c060ca9, so a CUDA build goes through the same path. I have an RTX 3090 with the same unsloth/Qwen3.8-27B-GGUF UD-Q4_K_XL target and can run your reproducer at -np 4 with long prompts to say whether this is HIP ordering or backend-independent lifetime, which seems worth settling before anyone commits to a fix shape. The card is tied up on an unrelated measurement for the next day or so; I will post the result either way, including if CUDA turns out not to reproduce it.
3. ByungHyun21 commented on Aug 24, 2026

   

   ByungHyun21

   on Aug 24, 2026

   Last edited by ByungHyun21

   Author

   More actions

   Follow-up from the reporter's side — I kept digging after filing, built and measured several fix shapes, and the picture changed enough to be worth posting before anyone commits to a fix. Everything below is measured on the same rig as the issue (gfx1151 / ROCm 7.2.2, UD-Q4_K_XL + mmproj-F16, -np 4 -c 1048576 --spec-type draft-mtp --spec-draft-n-max 4, the repro script from the issue), except where noted.

   1. The output_reorder() mode gate works — confirmed by build+run

   @thc1006's reading is correct. I implemented exactly that gate on b10176:

   // output_reorder()
   if (embd_nextn.size > 0 && cparams.embeddings_nextn_masked) {
       for (uint64_t k = 0; k < n_embd; k++) {
           std::swap(embd_nextn.data[i0*n_embd + k], embd_nextn.data[i1*n_embd + k]);
       }
   }

   and their explanation of my earlier failed attempt also matches: my first try removed the swap unconditionally, which breaks the masked (draft-context) path — that's why single-stream collapsed to ~5 t/s back then.

   But the gate alone does not fix the concurrent collapse. With the gate in and the sched-wide synchronize removed, I still got 2–6 of 8 empty replies per repro run.

   2. The real first-path mechanism: graph-copy alternation, not buffer reuse

   The issue text says the async copy "races with later graphs that reuse the extra buffer". I now believe that's not the mechanism. What I actually found:

   ggml_backend_sched double-buffers graphs (cur_copy = next_copy; next_copy = (next_copy+1) % n_copies in alloc_graph). The async extractions in llama_context::decode() read through the original graph tensors. When the copy selection alternates, an extraction can observe the copy that did not run this round — stale/uninitialized device memory — and land NaN rows on the host even though the stream itself is perfectly ordered.

   The evidence chain (all on b10176 + the mode gate):

   • End-of-decode backend sync (no synchronize inside the ubatch loop, one ggml_backend_synchronize after the loop): repro → 2–6/8 empty replies. A NaN scan of embd_nextn rows right after the sync fired 2,442 times in one run — the copies landed with garbage content, so this is not a host-visibility problem.
   • Allocator-lifetime theory eliminated: I instrumented cudaFree and every extraction copy (≥32 MiB threshold for the free counter, all copies logged). One repro run: 7,742 extraction copies, 1 device free, 0 address overlaps. Nothing is being freed under pending copies.
   • No background threads in ggml_backend_sched graph dispatch; graph launch and the memcpy are issued from the same thread on the same stream. Stream-ordering arguments say this can't corrupt — and empirically it doesn't, as long as the same graph copy is used.
   • The sched-wide synchronize was doing something else: it resets sched->next_copy = 0 ("this ensures that during generation the same copy is used every time" — the comment in ggml_backend_sched_synchronize). My original workaround was accidentally pinning the copy selection.

   3. Minimal host-side fix: pin the graph copy

   I added a small API and call it after each graph compute; with this alone the NaN count goes to zero — no GPU synchronize anywhere in the decode loop:

   --- a/ggml/src/ggml-backend.cpp
   +++ b/ggml/src/ggml-backend.cpp
   @@ void ggml_backend_sched_synchronize(ggml_backend_sched_t sched) { ... }
   +
   +// pin the graph-copy selection to copy 0: async output extractions must always
   +// read the copy that actually computed. Host-only, no GPU synchronization.
   +void ggml_backend_sched_pin_graph_copy(ggml_backend_sched_t sched) {
   +    GGML_ASSERT(sched);
   +    if (!sched->is_alloc) {
   +        sched->next_copy = 0;
   +    }
   +}

   --- a/src/llama-context.cpp  (graph_compute)
        auto status = ggml_backend_sched_graph_compute_async(sched.get(), gf);
   +    ggml_backend_sched_pin_graph_copy(sched.get());

   (plus the one-line declaration in ggml/include/ggml-backend.h)

   Result: NaN 2,442 → 0 on the same repro, with the decode loop completely free of synchronizes, and h rows verified clean after every decode.

   Remaining wrinkle: with pin + end-of-decode sync + the mode gate, the NaN corruption is gone but ~2/8 early-EOS empty replies remain per run. Not root-caused yet (h is clean; suspicion on the backend-sampler output copies — copy_tensor_async_* for t_sampled* — being consumed mid-flight, or another copy-selection-sensitive consumer).

   Variant matrix, same repro, same conditions (empty replies out of 8):

   build	h NaN	empty/8
   no sync at all (stock b10176 behavior)	yes	collapse (the issue)
   per-ubatch sched sync (issue workaround)	0	0
   per-ubatch backend-only sync	0	0 (but tg ~12% slower than sched sync)
   end-of-decode backend sync, no pin	2,442 events	2–6
   pin + end-of-decode sync	0	2
   pin + per-ubatch sched sync + mode gate	0	0

   The "backend-only sync is slower than sched sync" row is worth noting for anyone shrinking the workaround: the next_copy = 0 reset also keeps CUDA/HIP graph reuse stable; with a backend-only synchronize the copies alternate and the log shows constant graph-warmup resets, costing ~12% tg (15.9 → 14.0 t/s) in my measurements.
4. thc1006 commented on Aug 24, 2026

   

   thc1006

   on Aug 24, 2026

   More actions

   Thanks for building the gate and reporting back, and for pinning down why the first attempt failed.

   One question on the copy-alternation mechanism before more gets built on it: can you print sched->n_copies on your rig?

   There is a single assignment to it,

   sched->n_copies = parallel ? GGML_SCHED_MAX_COPIES : 1;

   and parallel is cparams.pipeline_parallel, whose first condition is model.n_devices() > 1. On one device that leaves n_copies == 1, which makes next_copy = (next_copy + 1) % 1 a constant zero, and both the alternation and the next_copy = 0 inside ggml_backend_sched_synchronize no-ops.

   Smaller second question: ggml_backend_sched_pin_graph_copy assigns only when !sched->is_alloc, and is_alloc is set true at the end of ggml_backend_sched_alloc_graph and cleared only on reset. Called right after graph_compute_async, does the assignment actually run for you, or is it skipped?

   Asking because 2,442 to 0 is not noise. If n_copies is 1 on your build then something other than copy selection changed between those two runs, and that seems worth knowing before the fix shape is settled.

   Set against your measurements this is only source reading, and gfx1151 with --mmproj may enumerate devices differently from what I am assuming here.

   My reproduction is a single CUDA device, so n_copies == 1 by construction. That makes it a control for this specific question: if the collapse appears there too, copy alternation cannot be the mechanism. I will log n_copies and the is_alloc value at that call site alongside the acceptance numbers. The card is on an unrelated run for another couple of hours.
5. ByungHyun21 commented on Aug 25, 2026

   

   ByungHyun21

   on Aug 25, 2026

   Author

   More actions

   I traced it further and I think I found the cause. Not 100% proven, but the mechanism is consistent with everything I measured, and the fix works reliably on my Strix Halo. If your CUDA run turns out clean, that would support this being the right track.

   What I found

   set_input() writes input tensors via ggml_backend_tensor_set, which on the CUDA backend does:

   // ggml/src/ggml-cuda/ggml-cuda.cu:782
   static void ggml_backend_cuda_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
       ggml_backend_cuda_buffer_context * ctx = (ggml_backend_cuda_buffer_context *) buffer->context;
       ggml_cuda_set_device(ctx->device);
       cudaMemcpyAsync((char *) tensor->data + offset, data, size, cudaMemcpyHostToDevice, cudaStreamPerThread);
       cudaStreamSynchronize(cudaStreamPerThread);
   }

   On CUDA, the legacy default stream implicitly synchronizes with all blocking streams, so this is safe even when the compute graph is running on a different stream. On HIP, hipStreamPerThread has no implicit synchronization with other streams.

   So under continuous batching with multiple ubatches per decode: ubatch k's graph is still executing on the compute stream while the host is already writing ubatch k+1's input tensors on hipStreamPerThread. The H2D write lands in the middle of the previous graph's input reads, and the graph computes from a mix of old and new input rows. The result is NaN in the output tensors — not just t_h_nextn, the device tensor itself is NaN.

   There is already a comment in llama_context::process_ubatch that describes this exact race:

   // with pipeline parallelism, the previous graph_compute_async may still be running
   // on the GPU. we must synchronize before set_inputs to avoid overwriting input tensors
   // that the previous compute is still reading.
   if (cparams.pipeline_parallel) {
       ggml_backend_sched_synchronize(sched.get());
   }

   But it only fires when pipeline_parallel is true, which requires multiple devices. On a single GPU (which is my setup — one gfx1151), pipeline_parallel is false and the guard is skipped. The CUDA default stream saves it on NVIDIA. On HIP it doesn't.

   How I got there

   The thing that narrowed it down was checking the device tensor, not just the host buffer. I added a probe that reads both sides at the spec-hook entry:

   probe code

   // src/llama-context.h — expose the last t_h_nextn for probing
   struct llama_context {
       ggml_tensor * t_h_nextn_dbg = nullptr;  // stashed at extraction
   
       const void * debug_nextn_probe(float * dev_rows, float * host_rows, int n, uint32_t n_embd) {
           if (t_h_nextn_dbg == nullptr || t_h_nextn_dbg->data == nullptr) return nullptr;
           ggml_backend_t b = ggml_backend_sched_get_tensor_backend(sched.get(), t_h_nextn_dbg);
           if (b) {
               ggml_backend_synchronize(b);
               ggml_backend_tensor_get(t_h_nextn_dbg, dev_rows, 0, (size_t) n * n_embd * sizeof(float));
           }
           if (host_rows && embd_nextn.data) {
               memcpy(host_rows, embd_nextn.data, (size_t) n * n_embd * sizeof(float));
           }
           return t_h_nextn_dbg->data;
       }
   };

   // common/speculative.cpp — at draft_mtp::process entry, before any drafts:
   if (getenv("HOOK_PROBE")) {
       static float dev_rows[4 * 5120], host_rows[4 * 5120];
       llama_debug_nextn_probe(ctx_tgt, dev_rows, host_rows, 4, 5120);
       // count NaN rows in each side
       // log: "HPROBE tokens=%d host_nan=%d/4 dev_nan=%d/4"
   }

   Result without the fix:

   HPROBE tokens=1024 host_nan=4/4 dev_nan=4/4   ← 647 out of 652 probes
   

   The device tensor itself was NaN at the spec-hook entry, meaning the graph compute was producing NaN, not the async D2H copy delivering garbage. That ruled out everything on the copy/extract side (which is what the original issue title assumed) and pointed at the graph inputs.

   Along the way I also eliminated:

   • graph-copy alternation (n_copies is 1 on my setup, so next_copy = (next_copy+1) % 1 is always 0 — the sched_synchronize next_copy = 0 reset is a no-op, answering your earlier question with a measurement: n_copies=1, is_alloc=1, next_copy=0 on every call)
   • upstream 8497981 (cherry-picked, no effect)
   • GGML_HIP_GRAPHS=OFF (no effect)
   • draft context on CPU via -ngld 0 (no effect)
   • stream mismatch between the memcpy and the graph compute (logged both, always identical)
   • CPU splits (n_splits=1 throughout)
   • buffer frees under pending copies (instrumented cudaFree: 7,742 extraction copies, 1 device free, 0 address overlaps)

   Fix

   Full context in llama_context::process_ubatch (around line 1370 on my tree):

       // set the input data for the input tensors
       {
           // LOCAL FIX (HIP): hipStreamPerThread (used by ggml_backend_tensor_set
           // for input writes) has no implicit sync with the compute stream,
           // unlike the CUDA legacy default stream. Without this sync, the H2D
           // write for ubatch k+1 races ubatch k's still-running graph reading
           // the same input tensors, producing NaN in the graph outputs.
           ggml_backend_sched_synchronize(sched.get());
   
           // FIXME this call causes a crash if any model inputs were not used
           // in the graph and were therefore not allocated
           res->set_inputs(&ubatch);
       }

   This replaces the per-ubatch sched-wide synchronize from the original workaround (which was placed after the t_h_nextn extraction). Same cost, but now at the point where the race actually happens. The extraction-point sync is removed.

   A more targeted fix would be making ggml_backend_cuda_buffer_set_tensor use the compute stream (or at minimum also synchronize it) on HIP, but I haven't attempted that. The sched-wide sync before set_inputs is sufficient and doesn't add cost beyond what the workaround already paid.

   Test results (all on gfx1151 / ROCm 7.2.2, same config as the issue)

   test	result
   4×19k concurrent repro, 4 runs (3 at temp 1.0, 1 at temp 0.6)	32/32 healthy, 0 empty
   device tensor NaN probe at hook entry	0 NaN events (was 647/652 before fix)
   5-min mixed soak (100 requests, 3 concurrent)	0 errors
   tg @7.4k tok	17.2 t/s (workaround: 16.6)
   tg @19.1k tok	14.9 t/s (workaround: 15.0)

   For your CUDA run: if the default stream does implicitly synchronize on NVIDIA as expected, this race shouldn't reproduce on CUDA even without the fix. If you see the collapse anyway, then something else is going on and I'd want to know. Either way, your n_copies logging will confirm the measurement on a different setup.

   I'm not going to chase this further unless your CUDA results point somewhere unexpected. The fix is deployed and stable on my Strix Halo, and the mechanism fits everything I've been able to measure.

   @thc1006
6. thc1006 commented on Aug 25, 2026

   

   thc1006

   on Aug 25, 2026

   Last edited by thc1006

   More actions

   Dear @ByungHyun21 ,

   Thanks for the detailed follow-up. The device-side NaN probe is very useful, and your measurement that n_copies == 1 confirms that graph-copy alternation was not the cause on this single-GPU setup.

   I agree that the original D2H t_h_nextn extraction theory is now unlikely, because the device tensor was already NaN before the host copy was consumed. An input-tensor lifetime or write-after-read race between consecutive ubatches is a reasonable direction.

   There is one CUDA stream-semantics detail that I think needs correction, though.

   cudaStreamPerThread is the per-thread default stream, or PTDS. It is not the legacy default stream, and it does not implicitly synchronize with arbitrary other streams. llama.cpp also creates its compute streams with cudaStreamNonBlocking, so legacy-default-stream synchronization would not provide the required ordering against those streams either.

   Therefore, I do not think CUDA is automatically protected here while HIP is not. For both backends, the important question is whether there is an explicit happens-before relationship between:

   1. graph k finishing its reads from the reused input tensor; and
   2. the H2D write for ubatch k+1 overwriting that tensor.

   A clean CUDA result would still be useful, but it would be a bounded non-reproduction on that hardware and software configuration. It would not prove that CUDA has an implicit synchronization guarantee that prevents the same race.

   I also want to clarify one line from your measurements:

   stream mismatch between the memcpy and the graph compute — logged both, always identical

   Does “the memcpy” refer to the D2H t_h_nextn extraction copy, or to the H2D copy performed by set_inputs()?

   If it refers to the H2D input copy, then the identical stream handles would be important, because operations submitted to the same stream should already be ordered. If it refers to the D2H extraction copy, it would be helpful to state that explicitly so it is not confused with the new input-overwrite hypothesis.

   For reference, I completed the CUDA run on an RTX 3090 with CUDA 12.0 and llama.cpp c060ca9:

   • -np 4, prompts from 256 through 16,384 tokens: no zero acceptance;
   • -np 4, approximately 19,000 tokens: 4/4 requests healthy, acceptance 0.34 to 0.44;
   • -np 4, 16,384 tokens, 12 repeated concurrent rounds: 48/48 requests had non-zero acceptance;
   • -np 8, approximately 4,500 tokens: 8/8 requests healthy;
   • no successful response had empty final content.

   The 24,576- and 32,768-token cases were not valid tests because they exceeded the per-slot context limit. -np 16 was also not tested successfully because the 24 GB card could not allocate the compute buffers.

   My environment was not identical to yours: it used a later commit, no mmproj, a smaller total context, and shorter generations. I would therefore describe the result as:

   The visible failure was not reproduced within the tested RTX 3090 CUDA envelope. This does not prove that CUDA is free from the same missing dependency or lifetime bug.

   To test the input-overwrite hypothesis more directly, an event-based dependency would be stronger than a scheduler-wide synchronization:

   1. record an event after graph k has finished using the reused input buffer;
   2. make the input-write stream wait for that event;
   3. then write the inputs for ubatch k+1.

   If that removes the device-side NaNs without ggml_backend_sched_synchronize(), it would provide much stronger evidence that the missing input-buffer ordering is the actual cause.

   Thanks again for continuing the investigation. The device-side probe has significantly narrowed the problem, even though I think the CUDA-versus-HIP explanation should be revised.
7. ByungHyun21 commented on Aug 25, 2026

   

   ByungHyun21

   on Aug 25, 2026

   Last edited by ByungHyun21

   Author

   More actions

   It was the D2H extraction copy (ggml_backend_tensor_get_async on t_h_nextn) compared against the compute stream on the same context. The H2D input write from set_inputs was not part of that comparison — you're right that this is the pair that matters for the hypothesis. Regardless of which exact mechanism turns out to be the cause, the fix is deployed and stable on my Strix Halo and I believe I'm at the theoretical bandwidth limit for single-stream decode on this hardware, so I'll leave further investigation to anyone who needs it.

   @thc1006

   👍React with 👍1Reacted by 秀吉
8. nabe2030 commented on Aug 25, 2026

   

   nabe2030

   on Aug 25, 2026

   More actions

   Follow-up, and first — thank you for tracking this down. Your finding (the H2D input write in
   process_ubatch() racing a still-running graph, with the pipeline_parallel guard skipping
   single-GPU setups) turned out to explain more than the acceptance-rate collapse you originally
   reported.

   We tested the fix you posted here against two symptoms on our end: our own multi-slot output
   corruption (--parallel 8/16/32, MTP off — clean at all three after your fix, reliably
   corrupted before) and the single-slot deep-context tool-array corruption reported in #27579
   (16/16 correct after your fix vs. 0/8 unpatched). Both resolved. Details:
   #27579 (comment) and
   #27579 (comment).

   We're withdrawing our original t_h_nextn async-copy theory in favor of your input-write-race
   finding, same as you did. We haven't filed a PR for the sync fix — if you're planning to (sounds
   like you already have a working patch), we'd rather defer to you as the one who found it; happy
   to add our multi-slot + tool-array test results as supporting data on a PR if that's useful,
   otherwise we're glad to help however's useful.

   👍React with 👍1Reacted by ByungHyun
9. ByungHyun21 commented on Aug 25, 2026

   

   ByungHyun21

   on Aug 25, 2026

   Author

   More actions

   Thanks for the independent verification — that it also resolves #27579 and the multi-slot corruption without MTP is more than I expected. Please go ahead and submit the PR. The fix is a single ggml_backend_sched_synchronize(sched.get()) in llama_context::process_ubatch immediately before res->set_inputs(&ubatch), replacing the pipeline_parallel-only guard that was already there for the same race. If you can add me as a co-author on the commit (Co-authored-by), that would be appreciated.

   @nabe2030
10. nabe2030 commented on Aug 26, 2026

    

    nabe2030

    on Aug 26, 2026

    More actions

    Thanks — and I owe you a straight answer rather than a PR.

    I went looking before filing, and #27311 (pwilkin, "Scheduler UMA ring buffer",
    open since Aug 18) already modifies the exact call site your fix targets. From its
    commit 9bdc6ea:

    Add ggml_backend_sched_prepare_inputs(), which llama.cpp calls before writing
    the inputs of a reused graph. (…) It replaces an unconditional
    ggml_backend_sched_synchronize() that was already at that call site for pipeline
    parallelism, so multi GPU setups now rotate instead of stalling on every reused graph.

    Same diagnosis as yours, reached independently: the host writes ubatch k+1's inputs
    while ubatch k's graph is still reading them, and the pipeline_parallel guard is the
    only thing standing there. The difference is the remedy — instead of making the
    synchronize unconditional, it gives graph inputs a ring of 2 and rotates, so the wait
    disappears rather than being extended to single-GPU setups. Reported cost on gfx1151 is
    inside measurement noise (prefill 193.1 vs 196.9 t/s at 104k tokens), and it ships a
    sanitizer showing write-after-read conflicts on inp_tokens, K/V indices, the attention
    mask and positions with the ring off, and zero with it on.

    So filing your patch separately would put a competing implementation on the same lines,
    which I don't think helps either of us. What I'd rather do:

    1. Build Scheduler UMA ring buffer (+ sanitizer and fixes) #27311 on our Strix Halo and run your -np 4 × 19k reproducer against it,
       ring on vs off, and post the acceptance numbers to that PR. Nobody there has tested
       MTP at all — if the ring fixes your collapse, that's a fourth issue it closes.
    2. Credit the diagnosis to you explicitly there, linking back here.

    The PR needs two approving reviews and has four pending, so evidence that it fixes more
    than perplexity is probably worth more right now than another patch.

    If you'd rather still have the sync version upstream as a backportable stopgap for people
    on release builds, say so and I'll file it that way — explicitly marked as a stopgap,
    cross-linked to #27311, with you as Co-authored-by. Your call; you found it.

    One thing I could use either way: is the acceptance check in your reproducer scriptable
    as a pass/fail, so what I post is a number rather than my reading of a log?
11. ByungHyun21 commented on Aug 26, 2026

    

    ByungHyun21

    on Aug 26, 2026

    Author

    More actions

    To answer @nabe2030's question — yes, the acceptance check is scriptable as pass/fail, and the repro already had the first half (any EMPTY reply = fail). The server-side half parses the --metrics draft acceptance lines; the one subtlety is that a collapsed slot dies fast, so its report shows zero acceptance on a small generated count, while healthy short tasks can hit zero by variance on one or two tokens. Gating on acc < 0.01 && generated >= 8 separates them cleanly across all three builds I tested.

    #!/usr/bin/env python3
    """Pass/fail check for the draft-mtp np-N race (#27572).
    
    Gates:
      1. Content gate: any EMPTY reply in the concurrent repro = FAIL
         (the repro script already exits non-zero on this)
      2. Acceptance gate: any slot reporting acceptance < 0.01 with
         generated >= 8 tokens = FAIL
    
    Usage: python3 check_mtp_race.py <server_log>
    Server must run with --metrics. Run the concurrent repro
    (4 x ~19k prompts) while the log is written, then invoke this.
    Exit 0 = PASS, exit 1 = FAIL.
    """
    import re, sys
    
    LOG = sys.argv[1] if len(sys.argv) > 1 else "/tmp/server_test.log"
    
    pat = re.compile(
        r'slot print_timing: id\s+(\d+)\s*\|\s*task\s+\d+\s*\|\s*'
        r'draft acceptance = ([0-9.]+)\s*\(\s*(\d+) accepted /\s*(\d+) generated')
    
    rows = []
    for line in open(LOG, errors="replace"):
        m = pat.search(line)
        if m:
            rows.append((int(m.group(1)), float(m.group(2)),
                         int(m.group(3)), int(m.group(4))))
    
    if not rows:
        print("FAIL: no draft acceptance lines found (is --metrics on?)")
        sys.exit(1)
    
    n_reports = len(rows)
    min_acc = min(r[1] for r in rows)
    zero_reports = sum(1 for r in rows if r[1] < 0.01)
    # A slot that collapsed generates garbage and hits EOS within a few
    # tokens, so its report shows (acc=0, generated small). Healthy short
    # tasks can also show low acceptance by variance, but never exactly
    # zero with more than a handful of generated tokens.
    collapsed = sum(1 for r in rows if r[1] < 0.01 and r[3] >= 8)
    total_acc = sum(r[2] for r in rows)
    total_gen = sum(r[3] for r in rows)
    overall = total_acc / total_gen if total_gen else 0.0
    
    print(f"reports={n_reports}  overall={overall:.5f}  "
          f"min={min_acc:.5f}  zero_reports={zero_reports}  "
          f"collapsed(acc=0,generated>=8)={collapsed}")
    
    if collapsed > 0:
        print("RESULT: FAIL — at least one slot collapsed to zero acceptance")
        sys.exit(1)
    
    print("RESULT: PASS — no zero-acceptance collapse detected")
    sys.exit(0)

    Verified against all three builds from the comparison I posted on #27311:

    A (master d222767c7 + sync fix):  reports=21 overall=0.51358 min=0.36000 collapsed=0 → PASS (exit 0)
    B (PR #27311 b3823d816):         reports=21 overall=0.51309 min=0.42188 collapsed=0 → PASS (exit 0)
    C (stock master d222767c7):      reports=9  overall=0.45647 min=0.00000 collapsed=1 → FAIL (exit 1)
    

    So for your run against #27311: python3 check_mtp_race.py <your server log> after the repro is all you need — a single exit code plus four numbers. The repro script's own output (8/8 OK vs any EMPTY) remains the primary gate; this adds the numeric confirmation that acceptance did not collapse even when content happens to survive.

    Full comparison data and methodology are on #27311. Keeping this issue open until that PR merges.
12. nabe2030 commented on Aug 26, 2026

    

    nabe2030

    on Aug 26, 2026

    More actions

    Thank you — that's more than I asked for, and the generated >= 8 gate is the part I
    would have gotten wrong. A collapsed slot and an unlucky short task both showing zero is
    exactly the ambiguity that made me ask for a script rather than trust my own reading.

    One data point toward the gate's robustness at larger N: our run on #27311 was
    --parallel 8 over 200 tasks rather than 4 × 19k, and the separation was just as clean —
    ring off gave acceptance 0.0000 with 191/200 outputs flagged, ring on gave 0.5083 with
    0/200. Nothing sat near the boundary in either direction, so the threshold doesn't look
    fragile as the task count grows.

    Agreed on keeping this open until #27311 merges.
13. iceatao commented on Aug 29, 2026

    

    iceatao

    on Aug 29, 2026

    More actions

    Independent data point (Windows / CUDA, dual RTX 4090) — matches this race:

    • llama.cpp: buun-llama-cpp fork (latest as of 2026-08-27, HEAD 2174ad6)
    • Model: DavidAU/Qwen3.8-27B-Cold-Fusion-GAIN-V1.1-NM-DAU-NEO-MAX-NEO-MTP-Q4_K_M.gguf (qwen35, built-in MTP head, --mmproj mmproj-Q8_0.gguf)
    • Server: llama-server, n_slots=2 (-np 2), n_ctx_slot=262144, kv_unified=false, --chat-template-file with the froggeric v22.4 fixed template

    Observed per-slot verify/rollback histograms after swapping to the fixed template:

    • slot 0: healthy — (2/0:75, 2/1:17, 2/2:39) → ~70% of cycles with at least one accepted draft token, mean len ≈ 1.89
    • slot 1: fully collapsed — (2/2:486) → 486/486 cycles rejected both draft tokens, acceptance = 0.0 on every client, regardless of who sends the request

    Controls I ran:

    • Restored the original stock chat template: both slots collapsed (slot 1 acceptance 2.8%, mean len 1.06) → the stock template (xhigh reasoning injection) alone kills MTP acceptance for this model.
    • Fixed template + -np 1: acceptance returns to a healthy range → the per-slot divergence is specific to parallel serving.

    This pattern (one slot healthy, one slot at exactly 0.0, server-side and client-independent) looks exactly like the input-write race described in this issue. Happy to provide more logs or run your repro script against this setup if useful.
14. nabe2030 commented on Sep 3, 2026

    

    nabe2030

    on Sep 3, 2026

    More actions

    @iceatao — this is the first report in these threads on CUDA rather than HIP, which makes
    it worth more than another gfx1151 confirmation. The signature matches closely: one slot
    healthy, one slot at exactly 0.0 across 486/486 cycles, server-side and independent of
    which client sends the request, and healthy again at -np 1. I measured the same
    discontinuity on gfx1151 — 55.41% acceptance at -np 1 collapsing to 0.00000 (0 of
    136,713 draft tokens) at -np 2. Not a gradual degradation, a cliff at the transition
    from one slot to two.

    That it reproduces on CUDA matters for the fix. The HIP-specific theories in these threads
    route through props.integrated and the host-buffer path, which your dual-4090 setup does
    not take. #27311 addresses this in the scheduler layer instead — a write-after-read on
    graph input tensors, backend-independent — so a CUDA reproduction is evidence for that
    shape of fix rather than the HIP-local ones.

    One caveat about your data, offered so nobody over-reads it. You have two variables moving:
    the stock template alone collapses acceptance on both slots by your own control, so that
    run cannot separate template effects from the race. The clean pair is fixed template with
    -np 2 (slot 1 at 0.0) versus fixed template with -np 1 (healthy) — everything else
    held constant. If you post again, leading with just those two would make the case harder
    to dismiss.

    Taking you up on the offer: yes, please. @ByungHyun21's check_mtp_race.py earlier in this
    thread is the right tool — it parses the draft acceptance lines and fails when any slot
    shows acc < 0.01 with generated >= 8, so the result is a pass/fail rather than someone's
    reading of a histogram. Running it against #27311 with GGML_SCHED_UMA_RING=2 versus =1
    on your dual-4090 box would be the first CUDA A/B anyone has, and #27311 currently has no
    CUDA multi-GPU data at all.
15. mitchmindtree

    added a commit that references this issue on Sep 24, 2026

    llama : wait for token-row readbacks before the next ubatch runs

    ...

    b5a74ad
16. github-actions

    added

    stale

    on Oct 4, 2026
17. github-actions commented on Oct 19, 2026

    

    github-actionsbot

    on Oct 19, 2026 – with GitHub Actions

    Contributor

    More actions

    This issue was closed because it has been inactive for 14 days since being marked as stale.
18. github-actions

    closed this as not plannedon Oct 19, 2026