CUDA --ssd-streaming: prefill dies with illegal memory access at exactly 128 prompt tokens
- Lingua principale
- C
- Stelle
- 22.3k
- Fork
- 2.1k
- Merge medio
- 1g 3h
- PR unite (30g)
- 4
Descrizione
Any prompt of 128 tokens or more fails to prefill under --cuda --ssd-streaming.
127 tokens works fine, 128 dies.
Clean worktree of origin/main b030961, with DS4_CUDA_MMQ=0 set to get past the MMQ OOM in #732 .
./ds4-bench -m ./ds4flash.gguf --cuda --ssd-streaming --prompt-file \
--ctx-start 128 --ctx-max 128 --gen-tokens 0
ctx 127 prefills at 9.09 t/s. ctx 128 dies, and to see where you need CUDA_LAUNCH_BLOCKING=1:
ds4: CUDA launch dense token-tile HMMA attention failed: an illegal memory access
ds4: gpu layer 2 attention batch encode failed
Without the blocking flag you get "CUDA model range upload sync failed for attn_out_a" instead,
which sent me looking at allocations for a while.
It is just the sticky error from the kernel above surfacing at the next sync.
Tthe token-tile fast path is gated on n_tokens >= 128, so decode and short prompts never enter it.
Layer 2 is the first layer with ratio==4 and so the first to take it.
If it helps, the launch parameters at the fault are:
pos0=0 n_tokens=128 n_raw=128 raw_cap=128 raw_start=0 n_comp=32 ratio=4 n_tiles=32.
I tried to find the bad access and failed. It is not an empty compressed cache (n_comp is 32).
It is not the raw mirror buffer: I over-allocated it by 64 rows and nothing changed, and the
highest index the kernel can reach is 254 against 255 rows.
It is not an index underflow on the compressed side.
And it is not shared memory: the launch attribute is accepted, and the kernel's running offsets add up
to the same total the budget declares.
What I have left is the cp.async row issue and the MMA fragment loads.
The reason I stopped is that I have no memcheck on this box: compute-sanitizer refuses the device under WSL2,
and cuda-gdb in CUDA 13 dropped "set cuda memcheck". On native Linux this is probably a five minute memcheck run.
DS4_CUDA_NO_WINDOW_ATTENTION=1 works around it. Prefill then runs at 128, 512 and 2048 tokens (9.26 / 17.34 / 8.86 t/s).
Same machine as #732 .
Guida per i contributori
Apri la guida per i contributori
Valutazione
Questa issue non è ancora stata valutata.