Eval bug: Flash attention crash (MUL_MAT failed / cudaStreamSynchronize) on Pascal GPUs with MiniMax-M2.7 at >24K context

[david-bue]

Eval bug: Flash attention crash (MUL_MAT failed / cudaStreamSynchronize) on Pascal GPUs with MiniMax-M2.7 at >24K context

Name and Version

ggml_cuda_init: found 10 CUDA devices:
  Device 0: Tesla P40, compute capability 6.1, VMM: yes, VRAM: 22905 MiB
  [... ×10 identical]
version: 8823 (4fbdabdc6)
built with GNU 11.4.0 for Linux x86_64

Operating systems

Linux (Pop!_OS, kernel 6.17.9)

GGML backends

CUDA

Hardware

• Server: Gigabyte G431-MM0 — 4U GPU server, 10× double-width GPU slots via CPBG8A0 PCIe backplane
• CPU: AMD EPYC 3151 (4C/8T)
• RAM: 32 GB
• GPUs: 10× NVIDIA Tesla P40 (GP102GL, compute capability 6.1, 24 GB each = 240 GB total)
• Backplane: CPBG8A0 with ASM2824 PCIe switch, running at Gen1 speeds (2.5 GT/s)

Models

• Model: unsloth/MiniMax-M2.7-GGUF (UD-Q5_K_XL, 5 shards)
• Architecture: minimax_m2 — 230B total params, 10B active (MoE)

Problem description

llama-server crashes with MUL_MAT failed or cudaStreamSynchronize failure during the first decode (generation) step after prompt evaluation completes when context exceeds approximately 24,000 tokens. The crash occurs consistently regardless of KV cache quantization type. Prompt evaluation itself completes successfully — the crash triggers on the transition to token generation.

The crash rotates between different GPUs across runs (GPU 3, GPU 9, GPU 2 observed), confirming it is not a hardware fault on any specific card.

Pascal (CC 6.1) uses the vec fallback path for flash attention (no MMA/WMMA support). This fallback kernel path appears to have a bug at high KV sequence lengths with this MoE architecture.

Reproduction

Launch configuration

llama-server \
    --model MiniMax-M2.7-UD-Q5_K_XL-00001-of-00005.gguf \
    --n-gpu-layers 999 \
    --tensor-split 1,1,1,1,1,1,1,1,1,1 \
    --main-gpu 0 \
    --split-mode layer \
    --ctx-size 65536 \
    --cache-type-k q8_0 \
    --cache-type-v q8_0 \
    --batch-size 256 \
    --ubatch-size 256 \
    --parallel 1 \
    --flash-attn auto \
    --jinja \
    --cont-batching

Steps to reproduce

1. Start llama-server with the above configuration
2. Open the built-in web UI (or use curl)
3. Have a multi-turn conversation, accumulating context past ~24,000 tokens
4. Server crashes during the first generation token after prompt eval completes on the request that crosses ~24K
   Alternatively, send a single large prompt (~24K+ tokens) to /v1/chat/completions.

Crash is reproducible with both KV quantization types

KV cache type	Context size	Crash?	Crash point
q4_0	98304	Yes	~26,491 tokens
q8_0	65536	Yes	~24,544 tokens
f16	—	Cannot test	--flash-attn off required, but flash-attn off rejects quantized KV; f16 KV at large context exceeds available VRAM

Cannot reproduce below ~24K tokens

Extended conversations staying under ~22K tokens run indefinitely without issues across all tested configurations.

Crash logs

Crash 1: q8_0 KV at 65536 context

Prompt eval completes normally, crash on first decode step:

493.33.954.903 I slot update_slots: id  0 | task 3528 | prompt processing done, n_tokens = 24018, batch.n_tokens = 86
[... request completes, next request arrives ...]
494.03.390.646 I slot update_slots: id  0 | task 3647 | new prompt, n_ctx_slot = 65536, n_keep = 0, task.n_tokens = 24544
494.03.581.837 I slot update_slots: id  0 | task 3647 | prompt processing done, n_tokens = 24544, batch.n_tokens = 236
494.06.416.338 E CUDA error: unspecified launch failure
494.06.416.343 E   current device: 9, in function ggml_backend_cuda_synchronize at /home/am/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:3005
494.06.416.344 E   cudaStreamSynchronize(cuda_ctx->stream())

Crash 2: q4_0 KV at 98304 context — MUL_MAT failure

33.57.946.541 I slot update_slots: id  0 | task 481 | prompt processing done, n_tokens = 23310, batch.n_tokens = 77
34.05.579.787 E CUDA error: unspecified launch failure
34.05.579.792 E   current device: 9, in function ggml_backend_cuda_synchronize at /home/am/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:3005

Crash 3: q4_0 KV at 98304 context — explicit MUL_MAT failure

82.04.257.564 I slot update_slots: id  0 | task 30770 | prompt processing progress, n_tokens = 26491, batch.n_tokens = 256, progress = 0.907475
82.06.690.099 E ggml_cuda_compute_forward: MUL_MAT failed
82.06.690.154 E CUDA error: unspecified launch failure
82.06.690.158 E   current device: 3, in function ggml_cuda_compute_forward at /home/am/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2884

dmesg output (PCIe completion timeouts + Xid)

The CUDA crash cascades into PCIe-level errors across multiple GPUs:

[39573.533422] nvidia 0000:0a:00.0:    [14] CmpltTO                (First)
[39573.533513] nvidia 0000:0b:00.0:    [14] CmpltTO                (First)
[39573.534634] nvidia 0000:08:00.0:    [14] CmpltTO                (First)
[39573.534732] nvidia 0000:09:00.0:    [14] CmpltTO                (First)
[39573.622827] NVRM: GPU at PCI:0000:0a:00: GPU-4bb48fa5-76de-c949-65b9-969270e2a751
[39573.622835] NVRM: Xid (PCI:0000:0a:00): 32, pid=25938, name=llama-server, channel 0x0000000a intr 00008000

Analysis

• The crash consistently occurs at the boundary between prompt evaluation and token generation at >24K KV depth
• Prompt eval batches complete fine even at 26K+ tokens — it is the first decode step that fails
• Pascal GPUs (CC 6.1) lack MMA and WMMA instructions, so flash attention falls back to the vec kernel path (fattn-vec-f16.cuh / fattn-vec-f32.cuh)
• The vec fallback kernels may have a shared memory, register, or tiling issue at high KV sequence lengths with MiniMax's MoE attention head configuration
• The crash GPU rotates between runs (GPU 3, GPU 9, etc.), ruling out faulty hardware
• The model runs stable at any context length through the native llama-server web UI as long as context stays below ~24K — this is not a frontend or API client issue
• All 10 GPUs pass sustained compute workloads with other models in general.

llama-server should handle MiniMax-M2.7 inference at the full configured context length (65K–98K) on Pascal GPUs without crashing, or flash attention should gracefully fall back to a safe path when the vec kernels cannot handle the requested KV depth.

[david-bue]

Local application of PR #22137 (drop the Q->ne[1] >= 1024 || Q->ne[3] > 1 clause from the FA-vec KV_max mask scan gate in fattn-common.cuh) on llama.cpp master @ fa8feae successfully prevents the Pascal vec-kernel OOB crash.

Verified on Qwen3.5-35B-A3B-UD-Q8_K_XL at 65K ctx with -fa on -ctk q8_0 -ctv q8_0 running stable past 30K filled context on 4× Tesla P40 (CC 6.1).

I am acknowledging Johannes' point that this is a workaround masking a kernel-level bounds-check bug; carrying locally pending the proper kernel fix.

Cheers ssam18 for the heads up.

[github-actions]

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