Misc. bug: [draft-mtp] Draft acceptance collapses to 0.0 under `-np N` with multi-ubatch batches — async `t_h_nextn` device→host copy race

[ByungHyun21]

[draft-mtp] Draft acceptance collapses to 0.0 under -np N with multi-ubatch batches — async t_h_nextn device→host copy race

Summary

Self-speculative decoding (--spec-type draft-mtp) silently breaks under parallel slots (-np 4): as soon as concurrent requests carry long prompts (decode batches spanning multiple ubatches), draft acceptance drops from a healthy 0.48–0.85 to exactly 0.00000 on all slots, generation falls back below no-MTP speed, and consumers of the endpoint see empty message.content completions (finish_reason=stop after a handful of reasoning_content tokens).

The trigger is a race between the async device→host copy of t_h_nextn (the MTP draft's input hidden states) and any later graph that reuses the t_h_nextn extra buffer — the next ubatch within the same decode(), or the next decode() call under -np N continuous batching. When the copy loses, the MTP draft conditions on NaN h rows, its logits go NaN, and every draft is rejected.

Verified at f5b9bd39b (build 10176), HIP backend. The same binary is correct at -np 1 and for single-ubatch batches.

Environment

llama.cpp	f5b9bd39b (b10176), HIP build, AMDGPU_TARGETS=gfx1151
ROCm	7.2.2 (pinned; gfx1151 native)
GPU	AMD Radeon 8060S (Strix Halo, gfx1151), 256 GB unified
Model	unsloth/Qwen3.8-27B-GGUF UD-Q4_K_XL + mmproj-F16 (qwen35 hybrid: Gated DeltaNet + Gated Attention, embedded nextn/MTP layer)
OS	Ubuntu 24.04, kernel 6.17

Server flags:

llama-server -m Qwen3.8-27B-UD-Q4_K_XL.gguf --mmproj mmproj-F16.gguf --jinja \
  --host 0.0.0.0 --port 9999 -ngl all -fit off -np 4 \
  -c 1048576 -n 131072 -b 1024 -ub 256 -fa on -ctk q4_0 -ctv q4_0 \
  --load-mode dio --spec-type draft-mtp --spec-draft-n-max 4 -rea on \
  --reasoning-format deepseek --temp 1.0 --top-p 0.95 --top-k 20 --min-p 0.0 --metrics

Reproduction

Self-contained script (stdlib only, no client framework needed) — it fires 4 concurrent chat completions, each with a ~19k-token system prompt, and prints per-request OK/EMPTY plus finish reasons.

repro_mtp_np4.py (click to expand)

#!/usr/bin/env python3
"""Reproduce the draft-mtp acceptance collapse under -np N concurrent load."""
import json, sys, threading, time, urllib.request

URL = "http://127.0.0.1:9999/v1/chat/completions"
N_CONC = int(sys.argv[1]) if len(sys.argv) > 1 else 4
TARGET_TOKENS = int(sys.argv[2]) if len(sys.argv) > 2 else 19000

unit = ("### Skill {i}: skill_{i} handles topic {i}. Steps: analyze, plan, "
        "execute, report. Constraints: verify outputs, cite sources.\n")
SYS = "You are an agent. Available skills:\n" + "".join(
    unit.format(i=i) for i in range(TARGET_TOKENS // 9))

def one(idx, results):
    body = {"model": "local_llm", "max_tokens": 1200, "messages": [
        {"role": "system", "content": SYS},
        {"role": "user", "content": f"[{idx}] Say hello and list what you can help with, one line."}]}
    req = urllib.request.Request(URL, json.dumps(body).encode(),
                                 {"Content-Type": "application/json"})
    t0 = time.time()
    try:
        r = json.load(urllib.request.urlopen(req, timeout=1200))
        c = r["choices"][0]
        content = (c["message"].get("content") or "").strip()
        results.append((idx, "OK" if content else "EMPTY", c["finish_reason"],
                        r["usage"]["completion_tokens"], time.time() - t0))
    except Exception as e:
        results.append((idx, "ERROR", type(e).__name__, 0, time.time() - t0))

def round_test(round_n):
    results = []
    ts = [threading.Thread(target=one, args=(i, results)) for i in range(N_CONC)]
    for t in ts: t.start()
    for t in ts: t.join()
    empty = sum(1 for r in results if r[1] != "OK")
    for r in results:
        print(f"  round {round_n} req{r[0]}: {r[1]:5s} finish={r[2]:6s} ct={r[3]:5d} {r[4]:7.1f}s")
    print(f"  round {round_n}: empty {empty}/{N_CONC}")
    return empty

if __name__ == "__main__":
    total = sum(round_test(n + 1) for n in range(2))
    print(f"TOTAL empty: {total}/{2 * N_CONC}")

Expected output on an affected build

Server log (--metrics, per slot, at task end):

slot print_timing: id 0 | draft acceptance = 0.00000 (    0 accepted /   4786 generated), mean len = 1.00
slot print_timing: id 1 | draft acceptance = 0.00000 (    0 accepted /   4786 generated), mean len = 1.00
slot print_timing: id 2 | draft acceptance = 0.00000 (    0 accepted /   4786 generated), mean len = 1.00
slot print_timing: id 3 | draft acceptance = 0.00000 (    0 accepted /   4786 generated), mean len = 1.00

The same server, sequential single request on a cold slot:

slot print_timing: id 3 | draft acceptance = 0.57865 (  206 accepted /   356 generated), mean len = 3.31

Client side, affected requests return finish_reason=stop with empty content after a few reasoning tokens, or run away in reasoning until the token cap — both surface as "empty completion" to an agent/chat client.

What does not trigger it

• -np 4 with short prompts (single-ubatch batches): 0/12 failures over 3 rounds
• Long prompts processed sequentially (one slot active): healthy acceptance
• Warm prefix cache (prefill mostly cached): healthy
• -np 1 with MTP: always healthy

→ the failure requires concurrent requests producing decode calls whose batches split across ubatches (long prompts), under multi-slot continuous batching.

Root-cause analysis

Instrumentation adds two probes (patch inlined at the end of this section):

1. In llama_context::decode()'s per-ubatch t_h_nextn extraction: after the async copy, ggml_backend_sched_synchronize() then compare the device tensor contents (via synchronous ggml_backend_tensor_get) against the host buffer rows.
2. At common_speculative_impl_draft_mtp::process() entry: NaN-probe of the target-context embd_nextn row 0 (the rows the draft is about to consume).

Findings during a failing concurrent round:

ctx DBG nextn extract: ub_tokens=15 n_seqs_unq=3 n_rows=15 off=0 tensor[ne0=5120 ne1=15] dev[nanrows=0/15] host[nanrows=0/15]   ← extraction time: clean
spec DBG ENTRY: n_tokens=10 n_act=2 tgt_row0[sq8=0.000e+00 nan=8]                                                                ← hook entry: NaN
spec DBG   seq=2 beg=0 end=4 h[sq8=0.0000e+00 nan=40] pend[sq8=0.0000e+00 nan=8]

i.e. the graph computed t_h_nextn correctly (device rows have sane norms), the extraction-time copy landed cleanly for the batch it ran on, but by the time the speculative hook reads embd_nextn, the rows it needs are NaN. With -lv 5 the draft candidates during failure are junk tokens with nan probabilities:

D spec draft:  - seq_id 3, draft candidate 0, pos 1: 8 (nan) ')'
D spec draft:  - seq_id 3, draft candidate 1, pos 1: 9 (nan) '*'

The gap is ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn_out, 0, …) in the unmasked nextn-embeddings path of src/llama-context.cpp. t_h_nextn lives in a scheduler extra buffer that subsequent graph executions reuse:

• the next ubatch of the same decode() call (long prompts: ~75 ubatches at -ub 256 for a 19k prompt), and
• the next decode() call (another slot's prefill/verify under -np N continuous batching).

The async copy is not ordered against those later graphs. When it executes after the buffer has been overwritten, NaN/garbage lands in embd_nextn, the MTP head produces NaN logits, and every draft is rejected — acceptance drops to exactly zero. Sequential and single-ubatch cases complete the copy before any overwrite (probabilistically), which explains the selective trigger conditions above.

instrumentation.patch (applies to b10176, click to expand)

diff --git a/common/speculative.cpp b/common/speculative.cpp
index 5653a90b8..2cfc13763 100644
--- a/common/speculative.cpp
+++ b/common/speculative.cpp
@@ -16,6 +16,7 @@
 #include <cstring>
 #include <iomanip>
 #include <map>
+#include <cmath>
 #include <cinttypes>

 #define SPC_DBG(fmt, ...) LOG_DBG("spec %12.*s: " fmt, 12, __func__, __VA_ARGS__)
@@ -1404,6 +1405,27 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
             return true;
         }

+        // [instrumentation] probe target h_nextn rows for NaN at hook entry
+        if (batch_in.n_tokens <= 40) {
+            int n_act = 0;
+            for (llama_seq_id s = 0; s < (llama_seq_id) n_seq; ++s) {
+                for (int k = 0; k < batch_in.n_tokens; ++k) {
+                    if (batch_in.seq_id[k][0] == s) { n_act++; break; }
+                }
+            }
+            if (n_act > 1) {
+                const float * h0 = llama_get_embeddings_nextn_ith(this->params.ctx_tgt, 0);
+                if (h0) {
+                    int nn = 0; double sq = 0;
+                    for (int e = 0; e < 8; ++e) {
+                        if (std::isnan(h0[e])) nn++; else sq += (double) h0[e]*h0[e];
+                    }
+                    LOG_WRN("spec DBG ENTRY: n_tokens=%d n_act=%d tgt_row0[sq8=%.3e nan=%d]\n",
+                            batch_in.n_tokens, n_act, sq, nn);
+                }
+            }
+        }
+
         // TODO: how to make it work with vision tokens?
         if (batch_in.token == nullptr || batch_in.embd != nullptr) {
             return true;
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index 9b399d609..d418a261a 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -1992,6 +1992,35 @@ int llama_context::decode(const llama_batch & batch_inp) {

                 GGML_ASSERT((offset + n_rows)*n_embd <= (int64_t) embd_nextn.size);
                 ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn_out, 0, n_rows*n_embd*sizeof(float));
+
+                // [instrumentation] compare device tensor vs host buffer after full sync
+                {
+                    static int64_t s_calls = 0;
+                    if (s_calls < 20000) {
+                        s_calls++;
+                        ggml_backend_sched_synchronize(sched.get());
+                        int nchk = n_rows < 15 ? (int) n_rows : 15;
+                        std::vector<float> dbg_rows((size_t) nchk * n_embd, 0.0f);
+                        ggml_backend_tensor_get(t_h_nextn, dbg_rows.data(), 0,
+                                (size_t) nchk * n_embd * sizeof(float));
+                        auto nan_rows = [&](const float * base) {
+                            int nr = 0;
+                            for (int r = 0; r < nchk; ++r) {
+                                for (int e = 0; e < 8; ++e) {
+                                    if (std::isnan(base[(size_t) r * n_embd + e])) { nr++; break; }
+                                }
+                            }
+                            return nr;
+                        };
+                        LLAMA_LOG_WARN("ctx DBG nextn extract: this=%p ub_tokens=%d n_seqs_unq=%d "
+                                "n_rows=%lld off=%lld tensor[ne0=%lld ne1=%lld] "
+                                "dev[nanrows=%d/%d] host[nanrows=%d/%d]\n",
+                                (void *) this, ubatch.n_tokens, ubatch.n_seqs_unq,
+                                (long long) n_rows, (long long) offset,
+                                (long long) t_h_nextn->ne[0], (long long) t_h_nextn->ne[1],
+                                nan_rows(dbg_rows.data()), nchk, nan_rows(embd_nextn_out), nchk);
+                    }
+                }
             }
         }

Workaround (local patch, verified)

--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@
                 ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn_out, 0, n_rows*n_embd*sizeof(float));
+                ggml_backend_sched_synchronize(sched.get());

With this, repeated cold-start runs of the repro (4 × 19k concurrent, multiple rounds) complete 16/16 healthy, acceptance 0.26–0.55, no empty completions.

Known cost (A/B, single-stream, temp 0, identical server flags):

prompt tokens	pp t/s before → after	tg t/s before → after
620	291.6 → 290.2	16.3 → 15.5
4603	339.9 → 338.4	17.7 → 16.6
9154	552.1 → 547.1	25.3 → ~15
18256	507.6 → 503.9	24.9 → ~14

pp is unaffected; long-prompt single-stream tg loses ~40% because the sched-wide synchronize disturbs the scheduler's graph-reuse state (acceptance flattens at ~0.40 instead of rising to 0.85–0.94 with prompt length). A proper fix should make the copy safe without a host-side synchronize — e.g. event-ordering the copy against subsequent graph launches, or rotating the t_h_nextn extra buffer — rather than this workaround.

Alternative approaches tried (all rejected)

1. Synchronize only full ubatches (ubatch.n_tokens >= n_ubatch): insufficient — the verify tokens' rows sit in the last, partial ubatch, whose copy still races with the next decode() call under concurrent load. Reproduced acceptance 0.0.
2. ggml_backend_tensor_get (backend sync) on every extraction: tg regressed to 11–12 t/s — the masked (draft-context) extraction runs several times per draft round on tiny decodes and is latency-critical.
3. Backend-sync only the unmasked (target-context) extraction: single-stream fully restored (25.2 / 24.7 t/s at 9k/18k) but concurrency still collapsed to acceptance 0.0 → there appears to be a second corruption path under -np N. Prime suspect: output_reorder() applies output_swaps (indices recorded for logits output rows) unconditionally to embd_nextn, whose rows are token rows in unmasked mode — cross-sequence row shuffling. However, simply excluding embd_nextn from the swaps collapsed acceptance even single-stream (~5 t/s), so the interaction is subtler than a one-line gate and we did not chase it further.

Full per-variant benchmark tables and logs available on request.

Related regressions observed at current master (2115b73d8, b10581)

While bisecting we hit two further draft-mtp issues on the same hardware/model — possibly related, filed here for completeness:

1. HIP: draft acceptance is 0.00000 from the very first request, even sequentially (0 accepted / 4786 generated, mean len 1.00). Something between b10176 and b10581 regressed embedded-nextn drafting outright for qwen35. Candidate commits in the range (not yet bisected):
   • 2c6b141ef — common: fix draft-mtp with embeddings (#27400)
   • 1692f9e50 — ggml: recurrent state rollback for ggml_ssm_scan (#26623) — qwen35 is a hybrid with Gated DeltaNet (ssm_scan) layers
   • 1d2869c6e / f65e568fd — spec type auto-detection changes
2. Vulkan (RADV, mainline build b10581): --spec-type draft-mtp aborts during draft-context creation, with and without --mmproj:

common_speculative_init_result: creating MTP draft context against the target model '…'
ggml-backend.cpp:348: GGML_ASSERT(tensor->data != NULL && "tensor not allocated") failed

Happy to provide full logs, the instrumentation patch output, per-variant benchmarks, and anything else useful.

[thc1006]

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.

[ByungHyun21]

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.

[thc1006]

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.

[ByungHyun21]

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

[thc1006]

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.

[ByungHyun21]

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

[nabe2030]

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.

[ByungHyun21]

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

[nabe2030]

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?