CUDA illegal memory access after long prefill with q4_0 K/V + FA (prism-b10770, regression from b10709) #346

Activity

1. professorpalmer commented on Oct 10, 2026

   

   professorpalmer

   on Oct 10, 2026

   More actions

   A pointer for the bisect, from the details in the report: b10709 is good, b10770 crashes, and --ctx-checkpoints 0 avoids the crash.

   In prism-b10709-9a9394a..prism-b10770-6684606, one commit touches context checkpoints in the server:

   • bf197c43d server : reuse the buffers of evicted prompt checkpoints (server : reuse the buffers of evicted prompt checkpoints #312)

   The others in that range that touch the server or speculative code are drafter changes (#334, #304), and this report runs without a drafter. So #312 is the first candidate: a reused checkpoint buffer that is smaller than the state written into it would fit an illegal access right after a long prompt, which is when checkpoints are created and evicted.

   For comparison: our fork (bonsai-ada-surgery, base before b10709) runs with up to 32 checkpoints per slot, plus checkpoints inside long messages, on 8 GB and 12 GB cards at 64k-262k. We have not seen this crash there. A revert of bf197c43d on b10770 would confirm or clear it.
2. professorpalmer commented on Oct 10, 2026

   

   professorpalmer

   on Oct 10, 2026

   More actions

   Reproducing on an RTX 4070 (sm_89, 12 GB) with your command: prism-b10770-6684606 as is, then the same build with bf197c43d (#312) reverted. Results here within a few hours.
3. professorpalmer commented on Oct 10, 2026

   

   professorpalmer

   on Oct 10, 2026

   More actions

   First result from the RTX 4070 (sm_89, 12 GB, CUDA 13, Windows): not reproduced with your command (-c 65536 -np 1 -fa on -ctk q4_0 -ctv q4_0 --jinja --reasoning-budget 2048 ..., default -b/-ub, no drafter), a needle prompt with the code word at the start, max_tokens 3000:

   Build	~28.5k-token prompt	~42.7k-token prompt
   prism-b10770-6684606	no error, server alive	no error, server alive
   same with bf197c43d (#312) reverted	no error, server alive	no error, server alive

   Your 16k observation does reproduce here, though: with --reasoning-budget 2048 all 3000 tokens went to thinking and content was empty, on both builds and both prompts. So the budget is not applied on b10770 with this command. That is a separate problem from the crash.

   Differences from your machine: 12 GB vs 8 GB of VRAM, sm_89 vs sm_86, CUDA 13 vs 12.4. Next I hold ~4 GB of VRAM with a second process, so the server has about what an 8 GB card has, and run the same pair again. A compute-sanitizer --tool memcheck run on your machine would name the faulting kernel directly (it is slow, so the 30k prompt is enough).
4. Powerful-Plantain commented on Oct 10, 2026

   

   Powerful-Plantain

   on Oct 10, 2026

   Author

   More actions

   Thanks for the pointer. Tested it: reverting bf197c43d (#312) does not fix the crash.

   git revert --no-commit bf197c43d on prism-b10770-6684606 (applies cleanly, only tools/server/server-context.cpp), clean rebuild, same flags, CUDA_LAUNCH_BLOCKING=1:

   Prompt	b10770	b10770 − #312
   ~30K tokens	crash	crash — same illegal memory access, same location (launch_fattn, fattn-common.cuh:1113)
   ~45K tokens	crash	crash — same
   ~16K tokens	empty content, finish_reason: length	same

   So #312 is cleared. --ctx-checkpoints 0 still avoided the crash in my one run, so whatever faults seems to interact with checkpoint creation/restore rather than buffer reuse — on a hybrid model that would include the recurrent (GDN) state. Remaining candidates in range that touch this path on sm_86: #221, #318, #189, #216, #214.

   Happy to revert any of those individually, or run a git bisect across prism-b10709-9a9394a..prism-b10770-6684606 (~6 steps; each build takes ~18 min on this machine). Just say which you'd prefer.
5. professorpalmer commented on Oct 10, 2026

   

   professorpalmer

   on Oct 10, 2026

   More actions

   Second round on the RTX 4070, with 4 GiB of VRAM held by a second process (7,755 MiB free for the server, about what your 8 GB card has), and prompts sized to yours:

   Build	prompt	result
   prism-b10770-6684606	28,484 / 42,704 tokens	no error, server alive
   same, #312 reverted	28,484 / 42,704 tokens	no error, server alive
   prism-b10770-6684606	30,176 / 44,954 tokens	no error, server alive

   So on sm_89 / CUDA 13 / Windows it does not reproduce, with or without VRAM pressure, and with or without #312. What is left that only your machine has: sm_86 kernels, CUDA 12.4, Linux, a laptop GPU with the display possibly on it. Two runs on your side would locate it:

   1. compute-sanitizer --tool memcheck ./llama-server ... with the 30k prompt. It reports the faulting kernel and the address.
   2. If that is too slow: the same command with --ctx-checkpoints 0 crossed with -ub 256. Your table already shows that checkpoints 0 avoids the crash; -ub 256 changes how the last part of the prompt is split, which tells whether the crash depends on the batch shape at the end of the prompt.

   Scripts: bench/prism346_repro.sh and tools/vram-ballast/ in professorpalmer/bonsai-ada-surgery.
6. Powerful-Plantain commented on Oct 10, 2026

   

   Powerful-Plantain

   on Oct 10, 2026

   Author

   More actions

   Thanks for testing on the 4070. Results from my machine (RTX 3070 Laptop, sm_86, CUDA 12.4, Linux):

   Environment: no display on the dGPU (nvidia-smi -q: Display Active: Disabled); the only process on the GPU is llama-server; compute-sanitizer is available.

   Matrix on prism-b10770-6684606, 30,184-token prompt, CUDA_LAUNCH_BLOCKING=1:

   ctx-checkpoints	-ub	-c	result
   default	256	65536	crash (same location)
   0	512	65536	correct answer
   0	512	65536	correct answer (repeat)
   0	256	65536	correct answer
   default	512	32768	crash

   • Batch shape at the end of the prompt doesn't matter: -ub 256 crashes the same way.
   • Checkpoints are the switch: with --ctx-checkpoints 0, 4/4 runs pass (30K and 45K); with default checkpoints, every run ≥30K crashes (6/6, including the two with server : reuse the buffers of evicted prompt checkpoints #312 reverted).
   • Context size doesn't matter: it reproduces at -c 32768.

   Running compute-sanitizer --tool memcheck on the 30K prompt overnight (at -c 32768 for VRAM headroom); will post the faulting kernel and address.
7. professorpalmer commented on Oct 10, 2026

   

   professorpalmer

   on Oct 10, 2026

   More actions

   Thanks, that matrix narrows it a lot: checkpoints on/off is the switch, and #312 is cleared (it crashes with #312 reverted).

   Why checkpoints matter here: with checkpoints on, the server processes the last 4 tokens of the prompt (and the last 4 + n_ubatch) as separate small batches, so it can save a checkpoint just before the end (checkpoint_offsets[] = {4 + n_ubatch, 4} in server-context.cpp). With --ctx-checkpoints 0 the last batch is a normal one. So the first new thing after "progress = 1.00" is a 4-query attention op over ~30k q4_0 cells, plus the recurrent-state copy for the checkpoint.

   Two more variants on the RTX 4070, same command, 30,176- and 44,954-token prompts, 4 GiB of VRAM held, same prompt split as yours (… 29,660 → 30,172, then the 4-token batch):

   Build	result
   b10770 compiled for sm_86 only (your SASS)	no error, server alive
   b10770 sm_89 with the scheduler told the GPU has 40 SMs (your card's count, so the stream-k split of the attention kernels is the same as yours)	no error, server alive

   So neither the sm_86 code nor the 40-SM work split reproduces it here. What remains is specific to your machine: CUDA 12.4 / driver 595 on Linux. The sanitizer run will say which kernel it is. If it is not conclusive, a bisect on your side is short: prism-b10709..prism-b10770 is 52 first-parent commits (about 6 builds). The ones that touch this path, to try first as single reverts on b10770:

   • 0d1825501 cuda: avoid short-query MMA FA crash with quantized V (cuda: avoid short-query MMA FA crash with quantized V #189)
   • 5244ceade in-place q4_0/q8_0 K/V flash attention (cuda: Bonsai 2 27B at the full 262k window with the MTP head on 12 GB Ada: #218 + #215 hybrid PTQ1_0 dispatch, in-place q4_0/q8_0 K/V flash attention (integration checkpoint) #221)
   • 8e18ce3be cuda: keep the GDN state gather when its ids memory is reused (cuda: keep the GDN state gather when its ids memory is reused #318)
   • 56f21fb10 kv-cache: limit seq_rm to the used cell range (kv-cache : limit seq_rm to the used cell range #311)

   Separate finding, also on b10770 here: --reasoning-budget 2048 is not applied with this command (all 3000 tokens are thinking, content empty), which matches your 16k observation.
8. Powerful-Plantain commented on Oct 11, 2026

   

   Powerful-Plantain

   on Oct 11, 2026

   Author

   More actions

   memcheck found it. compute-sanitizer --tool memcheck on prism-b10770-6684606, 30,184-token prompt, -c 32768, default checkpoints:

   Invalid __global__ write of size 2 bytes
       at void dequantize_block_q4_0<__half>(const void *, T1 *, int)+0x4b0
       by thread (0,0,0) in block (43003,0,0)
       Address 0x79097465f702 is out of bounds
       and is 95,747 bytes after the nearest allocation at 0x79096a000000 of size 174,358,784 bytes
   Host backtrace at launch:
       dequantize_row_q4_0_cuda<__half>(void const*, __half*, long, CUstream_st*)
       launch_fattn<256, 8, 8>(...)
       ggml_cuda_flash_attn_ext_mma_f16_case<256, 256, 8, 8, (ggml_type)1>
       ggml_cuda_compute_forward
       ggml_cuda_graph_evaluate_and_capture
       ggml_backend_cuda_graph_compute
       ggml_backend_sched_compute_splits
   

   So the fault is the q4_0 → f16 conversion of K/V inside launch_fattn for the MMA f16 kernel (D=256, 8×8): the f16 destination is ~95 KB smaller than what's written into it. That's the earlier kernel the CUDA_LAUNCH_BLOCKING trace couldn't name.

   My guess (not verified): with checkpoints on, the end of the prompt is evaluated in a batch shape that routes this op to the convert-to-f16 path instead of #221's in-place quantized read, and the scratch reserved for that conversion (ggml_cuda_flash_attn_ext_get_alloc_size after #221) is sized for less than the ~30K-cell K/V view. With --ctx-checkpoints 0 the path isn't hit. #189's short-query routing may also be involved.

   Full memcheck log available if useful. Happy to test a fix branch on the same setup.
9. professorpalmer commented on Oct 11, 2026

   

   professorpalmer

   on Oct 11, 2026

   More actions

   Thanks, the memcheck trace pinned it. Fix in #350.

   With q4_0/q8_0 K/V and 3-4 queries, the short-query MMA route from #189 converts K and V to f16 in the op's extra buffer. But get_alloc_size (from #221) reserved no f16 buffer when the in-place quantized kernel exists. So the f16 K was written past the allocation: your dequantize_block_q4_0 from launch_fattn, mma_f16_case<256, 256, 8, 8, F16>. Checkpoints evaluate the last 4 prompt tokens as their own batch, and that is why --ctx-checkpoints 0 avoids it.

   On my 4070 the 3-query D=256 q4_0 case trips a new assert without the fix and passes with it. The whole FLASH_ATTN_EXT suite passes 4007/4007. Could you run your crashing command on branch fix-346-short-query-scratch (PR #350)?
10. Powerful-Plantain commented on Oct 12, 2026

    

    Powerful-Plantain

    on Oct 12, 2026

    Author

    More actions

    Tested on the machine from #346 (RTX 3070 Laptop, sm_86, CUDA 12.4, Linux), e48c865d1 on top of prism-b10772-e8fd209, same flags, default context checkpoints, CUDA_LAUNCH_BLOCKING=1. Fixed.

    Test	b10770	b10772 + #350
    ~30K needle	crash	correct, stop
    ~45K needle	crash	correct, stop
    ~16K needle	empty content, length	correct, stop
    short prompt	ok	ok, decode 42.4 tok/s
    tool call	ok	ok (tool_calls)

    No illegal access, no assert, no FA fallback warning in any log; 7083 MiB loaded. Against b10709 at 45K: prefill 271 → 437 t/s, decode 15.6 → 21.1 tok/s.

    The 16K runaway answer is gone too. I can't tell whether that was this fix or #340/#347 in the base.

    Thanks for the quick turnaround. Happy to run compute-sanitizer on this build if you want a 0-error confirmation.

    Now running this build in production; a live 35K-context Hermes session is clean.
11. professorpalmer commented on Oct 13, 2026

    

    professorpalmer

    on Oct 13, 2026

    More actions

    Thanks for the full test on the 3070. All five cases pass, with no assert and no fallback warning. That closes the crash for #350.

    • Yes, please run compute-sanitizer --tool memcheck on this build with the ~30K prompt, if it is easy for you. A 0-error result is a good record on cuda: reserve the f16 K/V scratch for the short-query MMA route (#346) #350.
    • The 16K runaway answer may be this fix too, but I have not proven it. At 16K the checkpoint also evaluates the last 4 tokens as their own batch, so the same route runs. A write past the f16 buffer can land in another live allocation without a crash. That gives a wrong answer instead of an illegal access.
12. Powerful-Plantain commented on Oct 14, 2026

    

    Powerful-Plantain

    on Oct 14, 2026

    Author

    More actions

    memcheck on the same 3070 (sm_86, CUDA 12.4, Linux): 0 errors.

    Binary: the installed production build, b10773 = e48c865d1 on prism-b10772-e8fd209 (sha256 identical to the build that passed the earlier tests). compute-sanitizer --tool memcheck, 30,184-token prompt, default context checkpoints, -c 32768, q4_0 K/V, FA on — the same setup that reported Invalid __global__ write in dequantize_block_q4_0<__half> on b10770.

    	b10770	b10773 (#350)
    memcheck	invalid write, 95,747 B past the f16 buffer	ERROR SUMMARY: 0 errors
    answer	crash	correct, stop

    (The log ends with a CUDA error after "Received second interrupt" — my test script sent a second stop signal before the slow sanitized shutdown finished. That's after the request completed and isn't related to the fix.)

    That's consistent with your explanation of the 16K runaway as the same overflow landing in a live allocation. I've been running this build in production since the 12th, including Hermes sessions past 35K context — no issues. Thanks again for the quick fix.
13. professorpalmer commented on Oct 14, 2026

    

    professorpalmer

    on Oct 14, 2026

    More actions

    Thank you. 0 errors on the same setup that showed the invalid write (95,747 B past the f16 buffer) on b10770 is the clear record we wanted for #350.

    #350 now has a merge conflict with the current base. I am rebasing it today and will post the new commit there. The fix itself does not change.
14. professorpalmer commented on Oct 14, 2026

    

    professorpalmer

    on Oct 14, 2026

    More actions

    Update: the fix for this crash is in the PrismML base now, through #355 (it keeps the quantized K/V type in the short-query route, so the conversion that wrote past the buffer does not happen). I closed #350 as superseded. On the current base, all 4335 FLASH_ATTN_EXT op tests pass on an RTX 4070, including the 3-4 query cases with q4_0 / q8_0. A build from the current prism branch includes the fix, so you no longer need the #350 patch.