Eval bug: Qwen3.5-27B CUDA illegal memory access regression in prompt cache (single GPU, agentic tool-call pattern)

[cutlerbenjamin1-cmd]

Name and Version

Last known working: 149b2493c (March 19, 2026)
First observed broken: f49e91787 (April 3, 2026, current master)

Both built from source with identical configuration.

Operating systems

Windows 11 (10.0.26200)

GGML backends

CUDA

Hardware

• GPU: NVIDIA GeForce RTX 3090 (24 GB VRAM, compute capability 8.6) — single GPU
• CPU: AMD Ryzen 9 5900XT (16c/32t)
• RAM: 32 GB
• CUDA: 12.8
• Compiler: MSVC 19.44.35222.0

Models

mradermacher/Huihui-Qwen3.5-27B-abliterated-GGUF (i1-Q4_K_M)

Architecture: qwen35 (hybrid attention + Mamba2/Gated Delta Net, 64 layers, full_attention_interval=4)

Problem description & steps to reproduce

Summary

Qwen3.5-27B crashes with CUDA error: an illegal memory access was encountered when used with agentic frameworks that send large, dynamically-changing prompts with tool call/result patterns. The crash occurs in the prompt cache save path after a successful response. This is a regression — commit 149b2493c (March 19) is stable; current master crashes.

Trigger pattern (agentic tool calls)

This specifically reproduces with frameworks like Claude Code CLI that use the OpenAI-compatible /v1/messages endpoint with:

1. Large prompts (~29K tokens) containing full system prompts, tool schemas, conversation history, and dynamic context injections
2. Rapidly changing prefixes — each turn injects new <system-reminder> blocks, tool call results, and context that shifts the early prompt tokens
3. Multi-turn tool-call loops — model generates tool call JSON → tool executes → new prompt with result appended → repeat

The crash does not reproduce with simple chat interfaces (e.g., OpenWebUI) using the same model and server configuration, because those send smaller prompts with stable prefixes.

Steps to reproduce

1. Start llama-server with Qwen3.5-27B:

llama-server -m Qwen3.5-27B.Q4_K_M.gguf -c 120064 -ngl 99 --flash-attn --port 8090 -np 1

2. Send a first small request (~233 tokens) — succeeds fine, prompt cache checkpoint saved.

3. Send a second large request (~29K tokens, typical of an agentic framework with tool definitions) — processes successfully, returns HTTP 200, then:

CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_synchronize at ggml/src/ggml-cuda/ggml-cuda.cu:2924
  cudaStreamSynchronize(cuda_ctx->stream())

The crash occurs after the response is sent (during prompt cache save/update), not during inference.

Observed behavior from logs

srv  log_server_r: done request: POST /v1/messages 127.0.0.1 200
reasoning-budget: deactivated (natural end)
CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_synchronize at ggml/src/ggml-cuda/ggml-cuda.cu:2924

Key log entries showing the hybrid model's prompt cache behavior:

• cache_reuse is not supported by this context (hybrid model limitation)
• forcing full prompt re-processing due to lack of cache data
• Multiple checkpoints created during processing (created context checkpoint 1-5 of 32)
• Fused Gated Delta Net enabled (autoregressive + chunked)

Workaround

Adding --cache-ram 0 prevents the checkpoint save crash, but a second crash still occurs during subsequent tool-call turns (likely in memory_seq_rm partial truncation fallback for the recurrent state). Rolling back to 149b2493c resolves both crashes.

Bisection

Commit	Date	Status
149b2493c	March 19	✅ Stable with agentic tool calls
f49e91787	April 3	❌ Crashes on prompt cache save

Not yet narrowed further. The 93+ commits between these contain several prompt cache and checkpoint changes (#20087, #19924, #19877, #19849, #20955).

Key differences from existing reports

Aspect	This issue	#20176	#21140	#20225
Backend	CUDA	HIP/ROCm	HIP/ROCm	HIP
GPUs	Single	Single	Multi (2x)	Single
Trigger	Agentic tool calls (~29K dynamic prompts)	Multi-turn chat	Sequential agents	Multi-turn chat
Crash point	Post-response prompt cache save	Checkpoint load	Checkpoint restore	No crash (perf)

First Bad Commit

Not yet bisected to a single commit. Regression window: 149b2493c..f49e91787

Relevant log output

Full server log (working first request, crashing second request)

ggml_cuda_init: found 1 CUDA devices (Total VRAM: 24575 MiB):
  Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6, VMM: yes, VRAM: 24575 MiB
system_info: n_threads = 16 (n_threads_batch = 16) / 32 | CUDA : ARCHS = 860 | USE_GRAPHS = 1

srv    load_model: prompt cache is enabled, size limit: 8192 MiB

# First request (233 tokens) - OK
slot update_slots: id  0 | task 0 | processing task
slot update_slots: id  0 | task 0 | prompt processing done, n_tokens = 233
srv  log_server_r: done request: POST /v1/messages 127.0.0.1 200
prompt eval time = 689.16 ms / 233 tokens (338.09 tokens per second)
eval time = 3760.64 ms / 144 tokens (38.29 tokens per second)

# Second request (29306 tokens) - processes OK then crashes
slot update_slots: id  0 | task 2 | new prompt, n_ctx_slot = 120064, n_keep = 4096, task.n_tokens = 29306
slot update_slots: id  0 | task 2 | forcing full prompt re-processing due to lack of cache data
slot update_slots: id  0 | task 2 | created context checkpoint 1 of 32 (pos_min = 8191)
slot update_slots: id  0 | task 2 | created context checkpoint 2 of 32 (pos_min = 16383)
slot update_slots: id  0 | task 2 | created context checkpoint 3 of 32 (pos_min = 24575)
slot update_slots: id  0 | task 2 | created context checkpoint 4 of 32 (pos_min = 28789)
slot update_slots: id  0 | task 2 | prompt processing done, n_tokens = 29306
slot update_slots: id  0 | task 2 | created context checkpoint 5 of 32 (pos_min = 29301)
srv  log_server_r: done request: POST /v1/messages 127.0.0.1 200
reasoning-budget: deactivated (natural end)

CUDA error: an illegal memory access was encountered
  current device: 0, in function ggml_backend_cuda_synchronize at ggml/src/ggml-cuda/ggml-cuda.cu:2924
  cudaStreamSynchronize(cuda_ctx->stream())

Build configuration

cmake -B build     -DGGML_CUDA=ON     -DCMAKE_CUDA_ARCHITECTURES="86"     -DGGML_NATIVE=ON     -DGGML_LTO=ON     -DGGML_CUDA_FA_ALL_QUANTS=ON     -DLLAMA_CURL=OFF     -DCMAKE_BUILD_TYPE=Release     -G "Ninja"

[cutlerbenjamin1-cmd]

Correction on --cache-ram 0 workaround: Further testing shows this did not actually prevent the crash — both crashes appear to be the same underlying CUDA illegal memory access in the recurrent (SSM/GDN) state handling, just hit from different entry points. The first run with --cache-ram 0 happened to succeed, but the crash returned on subsequent tool-call turns.

Rolling back to 149b2493c (March 19) is the only reliable fix. Both crash paths (prompt cache save and multi-turn seq_rm/truncation) are eliminated by the rollback, suggesting a single regression in the recurrent memory CUDA codepath between March 19 and current HEAD.

[cutlerbenjamin1-cmd]

Root Cause Analysis: 744c0c731 — "rotate activations for better quantization" (#21038)

After extensive code analysis of all ~170 commits in the regression window (149b2493c..f49e91787), three independent reviewers converge on the same commit as the culprit:

Commit: 744c0c731 (April 1) — PR #21038
Author: @ggerganov

Why this commit

This commit adds Walsh-Hadamard rotation matrices to Q/K/V activations before KV cache storage when the cache is quantized. On Qwen3.5-27B:

• Rotation is active — confirmed by server logs: attn_rot_k = 1, attn_rot_v = 1 (KV cache is q8_0, head dims are multiples of 64)
• 16 of 64 layers are attention (full_attention_interval=4), each now gets 4 extra ggml_mul_mat + reshape operations for rotation
• The rotation was merged April 1 — right in the regression window (March 19 working → April 3 broken)

Proposed crash mechanism (K-shift graph buffer overflow)

The most likely crash path is through build_graph_shift in llama-kv-cache.cpp, which runs after the response is sent during KV cache maintenance:

1. After the 29K-token response completes (HTTP 200 sent), the server runs KV cache maintenance
2. build_graph_shift → build_rope_shift now generates ggml_mul_mat_aux ops on views of the quantized KV cache spanning the entire context length
3. For Qwen3.5-27B (n_heads_kv=4, n_embd_head_k=256), ggml_mul_mat_aux reshapes the dequantized K view to [256, n_heads_kv × ctx_size] F32 — a massive intermediate tensor
4. Each rotation mul_mat produces ~250MB F32 intermediates per attention layer, applied twice (rotate back before RoPE, rotate forward after)
5. The backend scheduler's scratch buffers were reserved for the inference graph topology during sched_reserve(), not for the shift graph's new massive rotation intermediates
6. ggml_backend_sched_alloc_graph places these intermediates into scratch space that was sized for a different topology — tensors overlap or land out-of-bounds
7. CUDA kernels write to invalid memory asynchronously; cudaStreamSynchronize catches the error after the response is already sent

This explains the crash timing perfectly: the error occurs at cudaStreamSynchronize AFTER "done request: POST /v1/messages 127.0.0.1 200".

Pre-rotation build_rope_shift (3 nodes per layer):

cast(K, F32) → rope_ext → cpy(back to quantized)

Post-rotation build_rope_shift (9 nodes per layer):

cast(K, F32) → reshape_2d → mul_mat(rot) → reshape_4d   ← rotate BACK
             → rope_ext
             → reshape_2d → mul_mat(rot) → reshape_4d   ← rotate FWD
             → cpy(back to quantized)

That's +96 compute nodes across 16 attention layers, with intermediate F32 tensors sized proportional to head_dim × n_heads_kv × ctx_length.

Alternative mechanism (inference graph memory pressure)

Even if the K-shift path isn't the specific trigger, the rotation adds 64 extra ggml_mul_mat operations to the inference graph across 16 attention layers. On a hybrid model where the graph already includes 48 recurrent/GDN layers with their own state tensors, plus 16 attention layers, the extra rotation intermediates for a 29K-token prompt may push compute buffer usage past what sched_reserve allocated — especially on a 24GB card where ~20GB is already consumed by model weights + KV cache + recurrent state.

Commits evaluated and dismissed

Commit	Why not the cause
3fc6f1aed — graph reuse + pipeline parallel	Pipeline parallelism not active on single GPU
db9d8aa42 — bf16 FA vec kernel	RTX 3090 (sm_86) doesn't have native bf16 compute
86221cf6d — FA kernel selection fix	K and V head dims are equal (256) for this model; the guard is a safety check that prevents crashes
f93c09e26 — recurrent state_read_meta bounds	Tightens bounds check (prevents bad state), doesn't cause crashes
17193cce3 — don't quantize SWA cache	Only touches ISWA path; Qwen3.5-27B uses non-ISWA hybrid memory
ec16a072f — MOE GEMV optimization	Qwen3.5-27B is not MoE
88458164c — FA head dim 512	Qwen3.5-27B has head dim 256
64ac9ab66 — CUB argsort fix	Not in the attention/KV path

Verification test

The commit includes a kill switch. To confirm, run the current HEAD build with:

# Windows
set LLAMA_ATTN_ROT_DISABLE=1
llama-server.exe [normal flags]

# Linux
LLAMA_ATTN_ROT_DISABLE=1 llama-server [normal flags]

If the crash stops, 744c0c731 is confirmed. Will test later today and report back.

Environment reminder

• Single GPU: RTX 3090 24GB
• CUDA 12.8, MSVC, Windows 11
• Model: Qwen3.5-27B Q4_K_M (hybrid attention + Mamba2/GDN, 64 layers)
• Trigger: Agentic tool-call pattern (Claude Code CLI) with ~29K-token dynamic prompts

[cutlerbenjamin1-cmd]

Confirmed: 744c0c731 (#21038) is the culprit

Setting LLAMA_ATTN_ROT_DISABLE=1 on the current HEAD build (f49e91787) completely eliminates the crash. The server survives multiple agentic tool-call rounds with 29K+ token prompts, checkpoint save/restore, and slot reuse — all of which previously crashed.

Before (rotation enabled — default on current HEAD):

llama_kv_cache: attn_rot_k = 1
llama_kv_cache: attn_rot_v = 1
→ CUDA error: an illegal memory access was encountered (after 2nd request)

After (LLAMA_ATTN_ROT_DISABLE=1):

llama_kv_cache: attention rotation force disabled (LLAMA_ATTN_ROT_DISABLE)
llama_kv_cache: attn_rot_k = 0
llama_kv_cache: attn_rot_v = 0
→ Multiple tool-call rounds, checkpoint restore from pos 16383, all slots release cleanly

Key log evidence from the working run:

• Checkpoint restore succeeded: restored context checkpoint (pos_min = 16383, n_tokens = 16384)
• Second prompt reused ~16K cached tokens, only processed remaining ~13K
• Both slots completed and released without any CUDA errors
• all slots are idle — clean shutdown

Summary

The Hadamard activation rotation introduced in 744c0c731 (#21038) causes CUDA illegal memory access on hybrid attention+SSM models (Qwen3.5-27B) with quantized KV cache (q8_0). The crash is triggered by large prompts (~29K tokens) and manifests during post-response cache maintenance. The rotation adds extra ggml_mul_mat operations to both the inference graph and the K-shift graph, with intermediate F32 tensors sized proportional to head_dim × n_heads_kv × ctx_length that appear to exceed or conflict with the scheduler's pre-reserved scratch buffers for hybrid model compute graphs.

Affected: All hybrid/recurrent architectures (Qwen3.5, Falcon-H1, Jamba, etc.) with quantized KV cache on CUDA.

Workaround: LLAMA_ATTN_ROT_DISABLE=1 environment variable before launching llama-server.

cc @ggerganov

[github-actions]

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