Eval bug: CUDA SOFT_MAX "invalid argument" on qwen4_exp / Qwen3.8-Flash-Next (Blackwell, compute_120a)

[celalergun]

Name and Version

build: b10808-64a155d24
llama-cli --version
version: 0.3.0-dev (build 10808, commit 64a155d)
built with GNU 15.2.0 for Linux x86_64

Operating systems

Linux
lsb_release -a
No LSB modules are available.
Distributor ID: Ubuntu
Description: Ubuntu 26.04 LTS
Release: 26.04
Codename: resolute

Which llama.cpp modules do you know to be affected?

llama-server, llama-cli

GGML backends

CUDA

Hardware

NVIDIA RTX PRO 4500 Blackwell (Workstation Edition), 32 GB GDDR7, compute capability 12.0 (compute_120a)
Driver 595.71.05, CUDA Version (nvidia-smi) 13.2
nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2026 NVIDIA Corporation
Built on Tue_Jun_09_02:43:40_PM_PDT_2026
Cuda compilation tools, release 13.3, V13.3.73
Build cuda_13.3.r13.3/compiler.38244171_0

Models

unsloth/Qwen3.8-Flash-Next-GGUF, quant UD-Q4_K_XL (4-shard GGUF: Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf)
Architecture: qwen4_exp (Qwen3.8-Flash-Next), support merged via #27742 on 2026-08-27

Operating systems

Linux

GGML backends

CUDA

Hardware

Intel(R) Xeon(R) Silver 4410Y + NVIDIA RTX PRO 4500 Blackwell Server Edition

Models

Qwen3.8-Flash-Next-GGUF

Problem description & steps to reproduce

Every attempt to actually run inference on this model with the CUDA backend crashes with a CUDA "invalid argument" error in the SOFT_MAX op, always at the same source location (ggml-cuda.cu:2417). This is fully reproducible and I've isolated it to the CUDA execution path specifically — the identical model file runs correctly on CPU only.

Tested configurations, in order:

1. llama-server with -ngl 99 --cpu-moe (all experts offloaded to CPU, dense/attention layers on GPU) — crashed on the first real client request (via curl to /v1/chat/completions), after the model finished loading successfully.
2. llama-server with -ngl 99 --cpu-moe --parallel 1 --flash-attn off — crashed earlier, during the server's own internal startup warmup pass (before "model loaded" is printed), so this combination fails even without any client request.
3. llama-cli with -ngl 0 (no GPU layers) — still crashed on the first prompt. This turned out to be because --no-op-offload is required to actually prevent CUDA from being used at all; -ngl 0 alone still routes CPU-resident tensor ops through the GPU via the automatic "op offload" optimization.
4. llama-cli with -ngl 0 --no-op-offload (fully CPU, no CUDA involvement) — succeeds, produces coherent output (~3.8–4.1 tok/s, as expected for CPU-only on this model size).

So: any configuration that lets CUDA execute the SOFT_MAX op for this architecture crashes; the only working configuration is the one that avoids CUDA entirely.

First Bad Commit

Not bisected. Architecture support (qwen4_exp) was merged via #27742 on 2026-08-27; this may be present since that merge or introduced by one of the follow-up PRs (#27880, #27941, #27977, #27836). Happy to bisect if a maintainer can point me at likely candidate commits given this is CUDA/SOFT_MAX-specific.

Command line

Crash (config 1):
llama-server
-m unsloth/Qwen3.8-Flash-Next-GGUF/UD-Q4_K_XL/Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf
-ngl 99 --cpu-moe
-c 32768 --jinja
--temp 1.0 --top-p 0.95 --top-k 20 --min-p 0.0
--host 0.0.0.0 --port 8080
--api-key "..."

Crash (config 2):
llama-server
-m unsloth/Qwen3.8-Flash-Next-GGUF/UD-Q4_K_XL/Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf
-ngl 99 --cpu-moe
-c 32768 --jinja
--temp 1.0 --top-p 0.95 --top-k 20 --min-p 0.0
--host 0.0.0.0 --port 8080
--api-key "..."
--parallel 1 --flash-attn off

Crash (config 3):
llama-cli
-m unsloth/Qwen3.8-Flash-Next-GGUF/UD-Q4_K_XL/Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf
-ngl 0
-c 4096
-p "Write a haiku about GPUs running out of memory."
-n 64

Works (config 4):
llama-cli
-m unsloth/Qwen3.8-Flash-Next-GGUF/UD-Q4_K_XL/Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf
-ngl 0 --no-op-offload
-c 4096
-p "Write a haiku about GPUs running out of memory."
-n 64

Relevant log output

Crash trace (identical failure point across configs 1–3)

/home/bartererpatio/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:108: CUDA error E ggml_cuda_compute_forward: SOFT_MAX failed E CUDA error: invalid argument E current device: 0, in function ggml_cuda_compute_forward at /home/bartererpatio/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2417 E err /home/bartererpatio/llama.cpp/build/bin/libggml-base.so.0(+0x1b236) [...] /home/bartererpatio/llama.cpp/build/bin/libggml-base.so.0(ggml_print_backtrace+0x20d) [...] /home/bartererpatio/llama.cpp/build/bin/libggml-base.so.0(ggml_abort+0x166) [...] /home/bartererpatio/llama.cpp/build/bin/libggml-cuda.so.0(ggml_cuda_error+0xb5) [...] /home/bartererpatio/llama.cpp/build/bin/libggml-base.so.0(ggml_backend_sched_graph_compute_async+0x89f) [...] /home/bartererpatio/llama.cpp/build/bin/libllama.so.0(llama_context::graph_compute+0xa0) [...] /home/bartererpatio/llama.cpp/build/bin/libllama.so.0(llama_context::process_ubatch+0x119) [...] /home/bartererpatio/llama.cpp/build/bin/libllama.so.0(llama_context::decode+0x3e6) [...] /home/bartererpatio/llama.cpp/build/bin/libllama.so.0(llama_decode+0x12) [...] /home/bartererpatio/llama.cpp/build/bin/libllama-server-impl.so(server_context_impl::decode+...) [...] /home/bartererpatio/llama.cpp/build/bin/libllama-server-impl.so(server_queue::start_loop+0x138) [...] Aborted (core dumped)

(full untruncated trace available on request — captured identically across configs 1, 2, and 3)

Successful CPU-only run (config 4) for comparison

build : b10808-64a155d24 model : unsloth/Qwen3.8-Flash-Next-GGUF/UD-Q4_K_XL/Qwen3.8-Flash-Next-UD-Q4_K_XL-00001-of-00004.gguf ftype : Q4_K - Medium

Write a haiku about GPUs running out of memory.
[Start thinking]

[ Prompt: 19.7 t/s | Generation: 3.8 t/s ]

Coherent output produced, no crash, confirming the model/quant/build are otherwise sound.

[murattopcu73]

Resolved on my box — root cause was a CUDA 12 runtime linked against a CUDA 13 driver, not qwen4_exp

Following up on my own report: this turned out not to be architecture-specific. The crash is gone after a rebuild, with no source changes.

Diagnosis

nvcc was 13.3, but the built library was linked against the system CUDA 12 runtime:

$ ldd build/bin/libggml-cuda.so | grep -iE 'cudart|cublas'
    libcudart.so.12 => /usr/lib/x86_64-linux-gnu/libcudart.so.12
    libcublas.so.12 => /usr/lib/x86_64-linux-gnu/libcublas.so.12
    libcublasLt.so.12 => /usr/lib/x86_64-linux-gnu/libcublasLt.so.12

while the driver reports CUDA 13.2. CMakeCache.txt had only:

CMAKE_CUDA_COMPILER:FILEPATH=/usr/local/cuda/bin/nvcc

with no CUDAToolkit_ROOT — so CMake compiled with the 13.3 toolkit for sm_120 but linked whatever libcudart it found on the system path. This is exactly the mismatch @mangrove-one isolated in #25060 by re-linking alone, and the CMake resolution problem @github-pmunoz described in the same thread.

Fix

cmake -B build-cuda13 -DGGML_CUDA=ON \
  -DCUDAToolkit_ROOT=/usr/local/cuda-13.3 \
  -DCMAKE_CUDA_COMPILER=/usr/local/cuda-13.3/bin/nvcc \
  -DCMAKE_CUDA_ARCHITECTURES=120 \
  -DCMAKE_BUILD_TYPE=Release

Resulting link:

libcudart.so.13 => /usr/local/cuda/targets/x86_64-linux/lib/libcudart.so.13
libcublas.so.13 => /usr/local/cuda/targets/x86_64-linux/lib/libcublas.so.13
libcublasLt.so.13 => /usr/local/cuda/targets/x86_64-linux/lib/libcublasLt.so.13

Every configuration that previously aborted now runs clean. Same commit (64a155d24, b10808), same model file, same driver — only the CUDA runtime changed.

Results

unsloth/Qwen3.8-Flash-Next-GGUF UD-Q4_K_XL (103.68 GiB, 176.94B params, qwen4exp A3B), RTX PRO 4500 Blackwell 32 GB + Xeon Silver 4410Y, 123 GB DDR5:

config	pp512 (t/s)	tg128 (t/s)
-ngl 0 --no-op-offload (CPU only, the previous workaround)	19.7	3.8
-ngl 99 -ncmoe 99 -fa 1	97.79 ± 8.66	18.08 ± 0.70
-ngl 99 -ncmoe 48 -fa 1	93.06 ± 3.48	18.00 ± 0.79
-ngl 99 -ncmoe 40 -fa 1	109.41 ± 2.75	20.54 ± 0.56
-ngl 99 -ncmoe 36 -fa 1	166.11 ± 12.33	22.38 ± 1.70
-ngl 99 -ncmoe 32 -fa 1	152.03 ± 4.92	22.74 ± 1.12
-ngl 99 -ncmoe 28 -fa 1	fails to load (out of VRAM)	—

Also verified at -c 32768 -ncmoe 36 -fa on via llama-cli: 18.3 t/s generation, coherent output, no aborts.

So ~6x generation throughput over the CPU-only workaround.

Caveat — this probably does not cover every report

This explains my case and @mangrove-one's, but not @pmsudhi's or @Emilien-Etadam's in #25060, both of whom reported a correct sharedMemPerBlockOptin and a version-consistent toolchain and still hit the identical abort. So the per-device (rather than per-kernel-instantiation) CUDA_SET_SHARED_MEMORY_LIMIT cache — "Bug 2" in #25060 — still looks like a real, separate defect worth fixing.

Two suggestions that would have saved me a day:

1. Warn at CUDA backend init when the linked libcudart major version does not match the driver's CUDA major version. This is silent today and produces a failure that looks nothing like its cause.
2. Validate sharedMemPerBlockOptin before passing it to cudaFuncSetAttribute (the Bug 1 clamp proposed in Misc. bug: Blackwell GGML-CUDA SOFT_MAX Crash #25060) — cheap, generic, and turns this class of failure into something diagnosable.

Happy to close this one if maintainers agree it is a duplicate of the #25060 cluster; leaving it open for now since the qwen4_exp reproduction may still be useful for whoever tackles Bug 2.