Eval bug: GPU ring buffer crash in llm_build_dflash_kv_update at long context prefill (~158K tokens)

[ppsx]

Name and Version

version: 9343 (233c99e)
built with GNU 13.3.0 for Linux x86_64

Operating systems

Linux

GGML backends

CUDA

Hardware

Intel(R) Core(TM) i7-14700K + NVIDIA RTX 5090 32GB VRAM

Models

unsloth/Qwen3.6-27B-GGUF:Q5_K_S + spiritbuun/Qwen3.6-27B-DFlash-GGUF:Q4_K_M

Problem description & steps to reproduce

Setup:

• BeeLlama v0.1.0 (commit b9343-233c99ed7), built CUDA 13.2, RTX 5090 Blackwell
• Target: unsloth/Qwen3.6-27B-GGUF Q5_K_S
• Drafter: spiritbuun/Qwen3.6-27B-DFlash-GGUF Q4_K_M
• Cache: --cache-type-k turbo4 --cache-type-v turbo3_tcq
• Config: --ctx-size 262144 --kv-unified --spec-draft-n-max 8
  --no-spec-dm-adaptive --spec-dflash-cross-ctx 1024 (also tried 512)
  -ub 1024 -b 1024

Symptom:
NIAH-style prefill at ~200K tokens consistently aborts during
update_drafter_kv_cache at progress 158720/200010 (deterministic checkpoint).

Stack trace:
ggml-impl.h:318: fatal error
ggml_backend_sched_set_tensor_backend
llm_build_dflash_kv_update
llama_context::dflash_kv_cache_update_gpu
common_speculative_state_dflash::update_drafter_kv_cache
common_speculative_state_dflash::flush_prefill

Tried:

• --spec-dflash-cross-ctx 512: crashes same way
• 128K NIAH (same config): 40/40 passes
• GGML_DFLASH_GPU_RING=0: WORKS at 200K (much slower, but completes)
• DFlash disabled, only TCQ cache on Q5_K_S target: works at 200K

Diagnosis: appears localized to GPU ring path in DFlash long-context prefill.

Repro: 200K-token prompt to /v1/chat/completions, prefill in -ub 1024 chunks.
NIAH JSONL test set available on request.

First Bad Commit

No response

Relevant log output

Stacktrace:

slot update_slots: id  0 | task 0 | n_tokens = 157696, memory_seq_rm [157696, end)
slot update_slots: id  0 | task 0 | prompt processing progress, n_tokens = 158720, batch.n_tokens = 1024, progress = 0.793560
srv  update_slots:   decode ubatch: 1024 tok, 1196.6ms (1.17ms/tok)
slot update_slots: id  0 | task 0 | n_tokens = 158720, memory_seq_rm [158720, end)
slot update_slots: id  0 | task 0 | prompt processing progress, n_tokens = 159744, batch.n_tokens = 1024, progress = 0.798680
srv  update_slots:   decode ubatch: 1024 tok, 1203.2ms (1.17ms/tok)
slot update_slots: id  0 | task 0 | n_tokens = 159744, memory_seq_rm [159744, end)
slot update_slots: id  0 | task 0 | prompt processing progress, n_tokens = 160768, batch.n_tokens = 1024, progress = 0.803800
srv  update_slots:   decode ubatch: 1024 tok, 1210.6ms (1.18ms/tok)
/data/ai/beellama.cpp/ggml/src/ggml-impl.h:318: fatal error
[New LWP 9970]
[New LWP 9969]
[New LWP 9968]
[New LWP 9967]
[New LWP 9966]
[New LWP 9965]
[New LWP 9964]
[New LWP 9963]
[New LWP 9962]
[New LWP 9961]
[New LWP 9960]
[New LWP 9959]
[New LWP 9958]
[New LWP 9957]
[New LWP 9956]
[New LWP 9955]
[New LWP 9954]
[New LWP 9953]
[New LWP 9952]
[New LWP 9951]
[New LWP 9950]
[New LWP 9949]
[New LWP 9948]
[New LWP 9947]
[New LWP 9946]
[New LWP 9945]
[New LWP 9944]
[New LWP 9943]
[New LWP 9942]
[New LWP 9941]
[New LWP 9940]
[New LWP 9936]
[New LWP 9931]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007e547e110813 in __GI___wait4 (pid=10532, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007e547e110813 in __GI___wait4 (pid=10532, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007e547ec2c503 in ggml_print_backtrace () from /data/ai/ai-beellama/bin/../lib/libggml-base.so.0
#2  0x00007e547ec2c6ab in ggml_abort () from /data/ai/ai-beellama/bin/../lib/libggml-base.so.0
#3  0x00007e547ec469e3 in ggml_backend_sched_set_tensor_backend () from /data/ai/ai-beellama/bin/../lib/libggml-base.so.0
#4  0x00007e547e8bb34f in std::_Function_handler<void (llama_ubatch const&, ggml_tensor*, char const*, int), llama_context::graph_get_cb() const::{lambda(llama_ubatch const&, ggml_tensor*, char const*, int)#1}>::_M_invoke(std::_Any_data const&, llama_ubatch const&, ggml_tensor*&&, char const*&&, int&&) () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#5  0x00007e547e90835e in llm_graph_context::cb(ggml_tensor*, char const*, int) const () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#6  0x00007e547e9086c1 in llm_graph_context::build_norm(ggml_tensor*, ggml_tensor*, ggml_tensor*, llm_norm_type, int) const () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#7  0x00007e547ea3713a in llm_build_dflash_kv_update::llm_build_dflash_kv_update(llama_model const&, llm_graph_params const&) () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#8  0x00007e547e960e58 in llama_model::build_graph(llm_graph_params const&) const () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#9  0x00007e547e8c6c07 in llama_context::dflash_kv_cache_update(int) () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#10 0x00007e547e8c74e9 in llama_context::dflash_kv_cache_update_gpu(void const*, int, int, int, void (*)(void*, void const*, unsigned long, unsigned long)) () from /data/ai/ai-beellama/bin/../lib/libllama.so.0
#11 0x00007e547f0de749 in common_speculative_state_dflash::update_drafter_kv_cache(int) () from /data/ai/ai-beellama/bin/../lib/libllama-common.so.0
#12 0x00007e547f0df0a7 in common_speculative_state_dflash::flush_prefill() () from /data/ai/ai-beellama/bin/../lib/libllama-common.so.0
#13 0x00007e547f0d6acd in common_speculative_flush_prefill(common_speculative*) () from /data/ai/ai-beellama/bin/../lib/libllama-common.so.0
#14 0x000062354a2a1e1b in server_context_impl::update_slots() ()
#15 0x000062354a340d2c in server_queue::start_loop(long) ()
#16 0x000062354a1fc617 in main ()
[Inferior 1 (process 9930) detached]

[ppsx]

It seems that the following fixes that issue. Still testing thou.

diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index ce9fb1909..04c32c31d 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -3307,6 +3307,10 @@ bool llama_context::dflash_kv_cache_update(int n_tokens) {
         return false;
     }
 
+    // graph_get_cb() inserts "norm" tensors into sched->hash_set per build; without
+    // reset the table fills across calls and ggml_hash_find_or_insert aborts.
+    ggml_backend_sched_reset(sched.get());
+
     const size_t max_nodes = graph_max_nodes(n_tokens);
     auto res = std::make_unique<llm_graph_result>(max_nodes);

[ppsx]

The one-line fix ggml_backend_sched_reset(sched.get()) eliminates the
crash but introduces a critical performance regression:

• short context: 213 t/s -> 46 t/s (-78%)
• medium: 129 t/s -> 43 t/s (-67%)
• long: 97 t/s -> 35 t/s (-64%)
• xlong: 117 t/s -> 25 t/s (-78%)

For comparison, the same config WITHOUT DFlash at all (only Q5_K_S target

• TCQ cache) yields ~67 t/s on short. So 46 t/s is BELOW the no-DFlash
  baseline, indicating two compounding regressions:

1. DFlash speculative speedup is gone (fell below the no-spec baseline)
2. Per-call overhead from the reset itself

The reset is too coarse. It almost certainly clears scheduler state that
DFlash relies on - cached graph topology, tensor backend assignments,
or speculative state that survives across token decodes.

A precise fix needs to scope to only the offending state. Likely options:

• Clear only the hash_set entries inserted by graph_get_cb()::lambda
  for "norm" tensors, leave other sched state untouched.
• Detect when hash_set is approaching capacity and rebuild selectively,
  not on every graph build.
• Fix the root cause: graph_get_cb()::lambda should check if the
  tensor is already in the hash_set before inserting (find-or-insert
  becomes find-only-if-present), rather than always inserting.
• Investigate whether the graph build should be using a per-build
  hash_set rather than the parent sched's.

Original repro and verification path unchanged from previous task.
The fix must preserve short-context throughput AT or NEAR 213 t/s
while still resolving the long-ctx crash.

[Anbeeld]

I do have what seems to be more correct fix for this issue. I'm rebuilding right now and will soon push it to main for you to test, after smoke testing from my side.

[Anbeeld]

It's in. Seems ok from short smoke testing. Please try it with your crash case.

[ppsx]

That's perfect!
I can confirm both - not crashing and excellent performance.
Thank you!

[aamsellem]

Confirming the same crash signature reproduces on a different setup — and confirming the fix in 75ae2a6 lands at the right place.

Hit on: RTX 5090M Laptop (24 GB sm_120 consumer Blackwell mobile), CUDA 13, BeeLlama v0.1.2 (633cd34fb), custom amd64+CUDA13+sm_120 build (aamsellem/beellama-cpp:0.1.2).

Config differences from the original report:

• Target: unsloth/Qwen3.6-27B-GGUF UD-Q3_K_XL (smaller quant, more headroom)
• Drafter: spiritbuun/Qwen3.6-27B-DFlash-GGUF dflash-draft-3.6-q8_0 (q8_0 not q4_k_m)
• mmproj F16 also loaded (vision app variant)
• KV: --cache-type-k turbo3 --cache-type-v turbo3
• ctx-size 200K (max we can fit at 24 GB with mmproj added)
• batch 2048 / ubatch 2048 (mandatory ≥ image_max_tokens for the multimodal path)

Crash signature observed: same cublasSgemm_v2 abort during DFlash spec cycles after sustained generation. Importantly it doesn't manifest as a clean ggml abort here — instead it surfaces as a HAMi-core (Olares' GPU sharing layer) OOM with the unsigned-underflow tell:

[HAMI-core ERROR (pid:1 thread=130863170578240 allocator.c:56)]:
  Device 0 OOM 18446744073709500456 / 25480396800
CUDA error: an internal operation failed
/app/ggml/src/ggml-cuda/ggml-cuda.cu:98: CUDA error
  current device: 0, in function ggml_cuda_op_mul_mat_cublas at /app/ggml/src/ggml-cuda/ggml-cuda.cu:1631
  cublasSgemm_v2(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, row_diff, src1_ncols, ne10, &alpha, ...)

The 18446744073709500456 request size is 2^64 - 51160 — characteristic underflow in HAMi-core's internal counter after many DFlash spec micro-allocations. The signal looks like an OOM but the underlying issue is the same GPU ring path the issue documents.

Token count at crash: consistently around 50-100K active session tokens, observed from a Hermes Agent workload that accumulates browser tool results + reasoning. Far earlier than your 158720 checkpoint, likely because mmproj + smaller VRAM headroom + parallel slot reservations push the ring further into stress sooner.

Tried GGML_DFLASH_GPU_RING=0: haven't been able to test yet — the deployment chart didn't expose the env var, just patched it locally. Will report back once verified.

Working around it for now: bumped CUDA_DEVICE_MEMORY_LIMIT_0 from 24300m to 30000m (above physical) so HAMi-core stops strict accounting before the underflow can corrupt the request size. Buys runway but doesn't fix the underlying ring path.

Plan: rebuilding the sm_120 image from main (incl. 75ae2a6) — once the qemu cross-compile finishes I'll re-test with the same Hermes workload that reproduces today, and report whether the crash is gone with the actual fix vs the workaround. Expect to confirm within the next 24h.

Thanks for the fast turnaround on this — the symptom was eating ~hourly sessions on Olares One.