Repository navigation
CUDA error: misaligned address in ggml_cuda_op_unary<op_sigmoid> with long prompts on RTX 5090 Laptop (sm_120) — regression between d3146f2b5 and ad6c66839 #28877
Description
Environment
| Item | Value |
|---|---|
| OS | Ubuntu 24.04.5 LTS (kernel 7.0.0-31-generic) |
| GPU | NVIDIA GeForce RTX 5090 Laptop GPU (sm_120 / compute capability 12.0, 24 GB) |
| Driver | 595.99.02 |
| CUDA | 13.0.88 (V13.0.88) |
| Compiler | GCC 13.3.0 |
| Build | CMake + Ninja, -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=120 -DGGML_CUDA_FORCE_CUBLAS=OFF -DGGML_CUDA_MXFP4=OFF -DCMAKE_CUDA_FLAGS_RELEASE="-O2" |
Model
Qwen3.8-27B-IQ4_NL.gguf (~14.6 GiB, Qwen3.5 arch) + mmproj-BF16_Qwen3.8-27B.gguf
Server launch args
--mmproj ./models/mmproj/mmproj-BF16_Qwen3.8-27B.gguf --cache-type-k q4_0 --cache-type-v q4_0 -c 131072 -n -1 --reasoning-budget 8192 --flash-attn on -ngl 99 -b 2048 --ubatch-size 64 -t 16 -tb 16 --poll 0 -lm mmap+mlock --temp 0.6 --top-p 0.95 --top-k 20 --min-p 0.0 --parallel 2 --kv-unified --cache-ram 20480 --presence-penalty 0.0 --jinja
Steps to reproduce
- Serve the model with the args above (
--parallel 2). - Send long-prompt requests (≥2000 tokens) — the crash was observed while two slots were processing prompts of ~3,845 and ~3,034 tokens concurrently:
LONG=$(python3 -c "print('stress test text ' * 400)")
for i in 1 2 3; do
curl -s -m 180 http://127.0.0.1:9001/v1/chat/completions \
-H "Content-Type: application/json" \
-d "{\"model\":\"local\",\"max_tokens\":20,\"messages\":[{\"role\":\"user\",\"content\":\"$LONG\"}]}" \
-o /dev/null -w "round $i: %{http_code}\n"
doneExpected behavior
Requests complete; server stays alive (verified on the previous build).
Actual behavior
Server aborts (CUDA device fault), process dies:
ggml/src/ggml-cuda/ggml-cuda.cu:108: CUDA error
CUDA error: misaligned address
current device: 0, in function ggml_cuda_kernel_launch at .../ggml-cuda/common.cuh:1710
cudaLaunchKernelEx(&pdl_cfg.cfg, kernel, std::forward<Args>(args)... )
Backtrace (host side):
#3 ggml_cuda_error(char const*, char const*, char const*, int, char const*)
#4 void ggml_cuda_op_unary<&(op_sigmoid(float))>(ggml_backend_cuda_context&, ggml_tensor*)
#5 ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*)
#6 ggml_backend_sched_graph_compute_async(...)
#7 llama_context::graph_compute(ggml_cgraph*, bool)
#8 llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&)
#9 llama_context::decode(llama_batch const&)
#10 llama_decode(...)
The faulting path is the unary sigmoid kernel launched through cudaLaunchKernelEx (PDL). In the server log the crash happens right after:
slot print_timing: id 0 | task 36 | prompt processing, n_tokens = 3845, progress = 1.00, t = 15.60 s / 246.47 tokens per second
slot print_timing: id 1 | task 37 | prompt processing, n_tokens = 3034, progress = 0.66, t = 11.55 s / 262.71 tokens per second
Bisect result (regression window)
| Build | Commit | Result under identical load/args |
|---|---|---|
| 10919 | d3146f2b5 (ggml-webgpu: Dawn update, #28683) |
✅ Good — 4 consecutive long-prompt tests (~3,900 tokens, 2 slots busy): all HTTP 200, 0 CUDA errors |
| 10950 | ad6c66839 (ggml-cuda: BF16→F32 fallback, #28846) |
❌ Bad — crashed twice, reproducible |
30 commits between the two; the only ggml-cuda-touching ones are:
ad6c66839ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (ggml-cuda: fallback to F32 on device without BF16 hardware acceleration #28846)c8edceb06ggml-cuda: hip add specific config table for AMD GCN (ggml-cuda: hip: add missing AMD GCN MMQ config #27841)
Note: on this device fast_bf16_hardware_available(cc) should return true (sm_120 ⇒ cc = 1200 ≥ GGML_CUDA_CC_AMPERE = 800), so the newly added fallback branch should be a no-op here — yet the regression correlates with this range, and reverting to d3146f2b5 fully resolves it. This may point to an interaction rather than the branch itself.
Additional context
- A similar-sounding prior issue (CUDA sm_120 (Blackwell, RTX 5090):
flash_attn_ext_f16<256,256,32,2>raises "Warp Misaligned Register" device exception — intermittent crash;-fa offworks around it #24398, RTX 5090/sm_120,flash_attn_ext_f16"Warp Misaligned Register") was closed in July; disabling FA was the workaround there. In our case-fa onis used on both the good and the bad build, so FA config alone does not explain the regression (reverting the source is what fixes it). - Setting
-cto a smaller value is not a workaround (same crash); only reverting the commit range helps. - Host has 62 GB RAM,
-lm mmap+mlock,--cache-ram 20480.
Happy to run further tests (e.g. bisect within the window, extra env vars, different models) if that helps narrow it down.
Activity
suntao12138 commentedon Sep 14, 2026
Correction — please disregard the bisect conclusion above
The "regression between
d3146f2b5andad6c66839" claim in my original report is wrong, and I want to retract it before anyone spends time on it.Today build 10919 (
d3146f2b5, the build I previously reported as "good") crashed with the exact same backtrace — includingggml_cuda_op_unary<&(op_sigmoid(float))>andCUDA error: misaligned address. So both builds crash; it is not a regression introduced in that commit range. It is load-dependent, which is why the shorter test on the older build appeared to pass. Apologies for the noise.Observed crash conditions on the failing run (build 10919, ~112 s after start):
Under sustained use the server had processed ~13,600 tasks when it died.
New lead: PDL kernel launch path
The faulting launch is the PDL branch in
ggml_cuda_kernel_launch(common.cuh), i.e.cudaLaunchKernelExwithcudaLaunchAttributeProgrammaticStreamSerialization. PDL is on by default and can be disabled at runtime via theGGML_CUDA_PDLenv var:Current status
I have restarted the server with
GGML_CUDA_PDL=0(falls back to plainkernel<<<...>>>launches). A first stress test of 20 rounds (4 parallel long-prompt requests per round + mid-flight request cancellation) passed with 0 CUDA errors and the server alive throughout. A longer/harder run (40 rounds, 6 parallel mixed-length requests per round, ~10 min) is in progress; I will report the result.Environment: RTX 5090 Laptop (sm_120, compute 12.0), driver 595.99.02, CUDA 13.0.88, GCC 13.3.0, Ubuntu 24.04.5.
If this hypothesis is right, the
programmaticStreamSerializationlaunch path may be unstable on sm_120 / CUDA 13.0 for some kernels (at least the unary/sigmoid one) under concurrent graph compute — happy to run any additional test or provide a fuller log if useful.suntao12138 commentedon Sep 14, 2026
Follow-up: longer stress test passed with PDL disabled
Result of the longer run on build 10919 with
GGML_CUDA_PDL=0:Combined: ~8 minutes of sustained concurrent load with PDL disabled → no faults, GPU at ~98% utilization / 60 °C throughout.
For contrast, on the same machine with PDL enabled:
So
GGML_CUDA_PDL=0looks like a solid workaround here. Hypothesis to confirm: theprogrammaticStreamSerializationlaunch path (sm_120 + CUDA 13.0.88, driver 595.99.02) faults on the unary/sigmoid kernel under concurrent graph compute.Happy to run further targeted tests if useful, e.g.:
--parallel 1(to see whether concurrency is required)ggml_cuda_kernel_can_use_pdlreturning false for specific opsCUDA_LAUNCH_BLOCKING=1