Eval bug: CUDA kernel stall during model execution, killed by watchdog

Name and Version

version: 0.1.0-dev (build 10438, commit 9d57ce4)
built with GNU 16.1.1 for Linux x86_64

Operating systems

Linux

GGML backends

CUDA

Hardware

RTX Pro 6000 Blackwell MAX-Q

Models

unsloth/Qwen3.8-27B-GGUF, UD-Q8_K_XL

Problem description & steps to reproduce

During sustained token generation, llama-server crashes with a CUDA launch timeout. The NVIDIA driver's RC watchdog detects a locked GPU channel (XID 8) and terminates the process. The issue is non-deterministic — it did not occur on every run, and has so far only reproduced on a Blackwell GPU.

First Bad Commit

No response

Relevant log output

Logs

llama-server -hf unsloth/Qwen3.8-27B-GGUF:UD-Q8_K_XL --host 0.0.0.0 --jinja
--batch-size 8192 --parallel 1 -fa on --main-gpu 0 --split-mode none
--spec-type draft-mtp,ngram-mod --spec-draft-n-max 3
--spec-ngram-mod-n-match 40 --spec-ngram-mod-n-min 0 --spec-ngram-mod-n-max 16
--temp 1.0 --top-p 0.95 --top-k 20 --min-p 0.0
--reasoning on --reasoning-preserve --cache-type-k q8_0 --cache-type-v q8_0 --fit on --kv_unified --load-mode none

ggml-cuda.cu:106: CUDA error
CUDA error: the launch timed out and was terminated
current device: 1, in function ggml_backend_cuda_synchronize
at ggml-cuda.cu:2530
cudaStreamSynchronize(cuda_ctx->stream())

NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked! Notify Timeout Seconds: 7
NVRM: GPU at PCI:0000:03:00
NVRM: GPU Board Serial Number: 1794125074179
NVRM: Xid (PCI:0000:03:00): 8, pid=45087, name=llama-server, channel 0x0000000d
audit: ANOM_ABEND pid=45087 comm="llama-server" sig=6 res=1

#0 __syscall_cancel_arch (libc.so.6)
#3 __poll (libc.so.6)
#4 n/a (libcuda.so.1 + 0x35d3f7)
#5 n/a (libcuda.so.1 + 0x41f0c5)
#6 n/a (libcuda.so.1 + 0x341953)

CUDA 13.2
Drivers 610.57.04

Second independent case, same signature — and notably without speculative decoding enabled, so the stall is not specific to the MTP draft path.

Build: 0.1.0-dev (build 10452, commit 412f3c8), GCC 16.2.1, Linux x86_64 (CachyOS, kernel 7.1.8)
GPU: RTX 5090 (Blackwell, sm_120), driver 610.57.04, CUDA 13.2
Model: bartowski/Qwen3.8-27B-GGUF, Q4_K_M (imatrix), local file

Invocation:

llama-server -m Qwen3.8-27B-Q4_K_M.gguf --host 127.0.0.1 --port 8080 \
  -ngl 999 -c 131072 -fa on -ctk q8_0 -ctv q8_0 --jinja

What happened: after ~5h20m of uptime with light intermittent client traffic, the server aborted mid-generation:

ggml-cuda.cu:106: CUDA error
CUDA error: the launch timed out and was terminated

Kernel log at the same second:

NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked! Notify Timeout Seconds: 7
NVRM: Xid (PCI:0000:01:00): 8, pid=709378, name=llama-server, channel 0x00000027

Backtrace (via ggml_print_backtrace):

ggml_abort → ggml_cuda_error (libggml-cuda) → ggml_backend_sched_synchronize
→ llama_context::synchronize() → server_queue::yield_to_queue → server_queue::start_loop

SIGABRT; 6.6 GB coredump retained if useful. GPU thermals were normal (40 °C, no throttle reasons), VRAM at 25.5/32.6 GB used — no memory pressure. The request traffic shortly before the crash included repeated client-error 500s (jinja template rejects of reasoning_effort, and image-input rejects due to missing mmproj) — mentioning in case the error path is relevant, though they may be unrelated.

Also seeing model has unused tensor blk.64.attn_norm.weight -- ignoring at load (MTP head of this GGUF) — presumably unrelated since MTP/speculative decoding is not enabled here, but noting it for completeness.

Does options nvidia NVreg_RegistryDwords="RmWatchDogTimeout=60" in /etc/modprobe.d/nvidia.conf fix this for you?

Third independent report, same signature.

Build: 0.1.0-dev (build 10433, commit 364f96b), paru/AUR build, GNU 16.2.1, Linux x86_64 (CachyOS, kernel 7.1.8-1-cachyos)
GPU: RTX 5090 (sm_120), driver 610.57.04, CUDA 13.3, power-capped 450 W (default 575 W)
Model: Qwen3.8-27B Q6_K (hybrid Mamba/attention + MTP layers) + mmproj BF16

Invocation (as it was during the crashes):

llama-server -m Qwen3.8-27B-Q6_K.gguf --mmproj mmproj-BF16.gguf \
  --fit off --reasoning on --cache-type-k q8_0 --cache-type-v q8_0 \
  --image-min-tokens 1024 --image-max-tokens 4096 --jinja \
  --chat-template-kwargs '{"preserve_thinking": true, "reasoning_effort": "xhigh"}' \
  -ngl -1 -c 98304 -t 16 \
  --spec-type draft-mtp --spec-draft-n-max 2

(parallel auto → 4 slots with kv_unified; flash attention auto/on for the q8_0 V cache; CUDA_LAUNCH_BLOCKING=1 was set for debugging — the stall is still a GPU-side kernel hang, reported at cudaStreamSynchronize)

What happened: 15 crashes in ~36 h (2026-08-15 03:17 → 2026-08-16 16:29), all mid-generation (steady ~73 t/s right up to the hang) and at any context size — observed with 15k–89k evaluated tokens. Identical signature each time:

ggml-cuda.cu:106: CUDA error
CUDA error: the launch timed out and was terminated
  current device: 0, in function ggml_backend_cuda_synchronize at .../ggml-cuda.cu:2530
  cudaStreamSynchronize(cuda_ctx->stream())

NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked! Notify Timeout Seconds: 7
NVRM: Xid (PCI:0000:08:00): 8, name=llama-server

SIGABRT every time. The driver recovers cleanly (channel killed, no GPU reset); no Xid 13/31/45/48/63/64/79, no AER/PCIe errors, no thermal events; throttle reasons 0x0 and 46 °C idle.

@gg7

@gg7 Applied options nvidia NVreg_RegistryDwords="RmWatchDogTimeout=60" via /etc/modprobe.d/nvidia.conf (rebuilt initramfs, rebooted, confirmed active in /proc/driver/nvidia/params).

Then ran a sustained-generation stress test on the exact same setup that crashed yesterday (RTX 5090, driver 610.57.04 open kernel module, kernel 7.1.8, llama.cpp b10452 CUDA, Qwen3.8-27B Q4_K_M @ 128K ctx, -fa on -ctk q8_0 -ctv q8_0): 93 back-to-back completions × 3072 tokens each (~286K tokens over ~66 min), GPU pinned at ~98% / ~490W / 66°C.

Result: no crash, but also no stall events at all — every completion ran at a uniform 68–75 t/s with no outlier-slow generations, so the raised watchdog timeout was never actually exercised. Given the bug is non-deterministic (my original crash happened after ~a day of uptime, not under fresh load), I can't yet attribute anything to the option. I'll keep it applied during normal daily use and report back if the Xid 8 recurs — a recurrence with the 60s timeout would at least tell us it's a hard hang rather than a slow kernel.

@pwilkin

Does options nvidia NVreg_RegistryDwords="RmWatchDogTimeout=60" in /etc/modprobe.d/nvidia.conf fix this for you?

Keep us posted - I also observe this with RTX 5090 at random times under heavy load. Last time this was discussed, @pwilkin suggested the same fix in #20146 (comment).

Ideally, we should find a way to fix this without having to make changes to the driver settings. cc @ggml-org/ggml-cuda @ggml-org/nvidia for ideas

I am also having this exact problem with this exact model. Testing my own solutions as well as those I find in discussions here. Will update if I find a fix.

Update: I was able to avoid the problem by setting GGML_CUDA_DISABLE_GRAPHS=1 in my conf. There is a bug in CUDA graph caching with this model/family of models. I spent all day benchmarking, isolating variables, and every single one failed loudly and in the same way, except this one. I slam my model with a task that takes about 700k tokens to finish, and every time it failed as it approached the max cache headroom, except when GGML_CUDA_DISABLE_GRAPHS=1 was included.

That's your temporary fix. Costs me about ~10-15% on tok/s but worth it, since it avoids the deadlock problem entirely.

After some time, I haven't seen the bug come up in other Blackwell architectures - DGX Spark doesn't trigger it, RTX 5060Ti doesn't seem to either. I strongly suspect a subtle timing deadlock when the card is faster or a very specific architecture difference - race condition somewhere, and I suspect interaction with cache. I'm not deep enough into the code to trace it down however.

@awrgold

Correction + follow-up data on my earlier comment:

Correction: I wrote that my original crash happened "after ~a day of uptime" — that was inaccurate. The relevant metric is model usage time, not machine uptime: the original Xid 8 hit after roughly 20–30 minutes of sustained generation use. Machine uptime is irrelevant to the trigger pattern; sorry for the noise.

Status after ~2 days with RmWatchDogTimeout=60: cumulative model usage since applying the option is now well past the original ~20–30 min trigger window (regular daily use, including long agentic sessions with large prompts), and there has been zero recurrence — kernel journals across 3 boots show no Xid / NVRM errors at all. Same setup as before (RTX 5090, driver 610.57.04 open, kernel 7.1.8, Qwen3.8-27B Q4_K_M, -fa on -ctk q8_0 -ctv q8_0, llama.cpp CUDA builds b10452→current).

@awrgold — thank you for isolating this and sharing the GGML_CUDA_DISABLE_GRAPHS=1 workaround, much appreciated; the CUDA-graph-cache lead also fits the fast-Blackwell-only pattern noted above.

That said, I'm deliberately not applying it on my box for now, so the watchdog setting stays isolated as the only variable — if Xid 8 recurs with the 60s timeout active, that will cleanly indicate a hard hang rather than a slow kernel, and I'll report it here.

@ccebelenski

@ccebelenski I got it on my 2x5070Ti box.

Does this repro also on FP16 KV caches?

I also use Q8_0 KV cache on my RTX machine. Not sure if it happens with FP16.

Same issue with my 5090. It randomly crashes with this error:

1.35.381.106 E CUDA error: the launch timed out and was terminated
1.35.381.108 E   current device: 0, in function ggml_backend_cuda_synchronize at /home/mukmckenzie/Documents/GitHub/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2530
1.35.381.109 E   cudaStreamSynchronize(cuda_ctx->stream())
/home/mukmckenzie/Documents/GitHub/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:106: CUDA error

[New LWP 322963]
[New LWP 322962]
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/usr/lib/x86_64-linux-gnu/libthread_db.so.1".
0x00007eef9cb10813 in _GI__wait4 (pid=323678, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30 ../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x00007eef9cb10813 in _GI__wait4 (pid=323678, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30 in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x00007eef9d14d6c3 in ggml_print_backtrace () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libggml-base.so.0
#2  0x00007eef9d14d86b in ggml_abort () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libggml-base.so.0
#3  0x00007eef9646b8a7 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libggml-cuda.so.0
#4  0x00007eef96473488 in ggml_backend_cuda_synchronize(ggml_backend*) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libggml-cuda.so.0
#5  0x00007eef9d167c6e in ggml_backend_sched_synchronize () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libggml-base.so.0
#6  0x00007eef9c0f7cc9 in llama_context::synchronize() () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama.so.0
#7  0x00007eef9c0fc3df in llama_state_seq_set_data_ext () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama.so.0
#8  0x00007eef9c628f3d in common_prompt_checkpoint::load_tgt(llama_context*, int, unsigned int) const () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-common.so.0
#9  0x00007eef9d3b76b1 in server_context_impl::pre_decode()::{lambda(server_slot&)#6}::operator()(server_slot&) const () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#10 0x00007eef9d3abc07 in server_context_impl::iterate(std::vector<server_slot, std::allocator<server_slot> >&, std::function<void (server_slot&)>) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#11 0x00007eef9d3ac257 in server_context_impl::pre_decode() () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#12 0x00007eef9d3af1fc in server_context_impl::update_slots() () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#13 0x00007eef9d35682d in server_queue::start_loop(long) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#14 0x00007eef9d2f4603 in llama_server(common_params&, int, char**) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#15 0x00007eef9d2f5626 in llama_server(int, char**) () from /home/mukmckenzie/Documents/GitHub/llama.cpp/build/bin/libllama-server-impl.so
#16 0x00007eef9ca2a1ca in __libc_start_call_main (main=main@entry=0x5ea791f0d270 <main>, argc=argc@entry=47, argv=argv@entry=0x7ffef7a38438) at ../sysdeps/nptl/libc_start_call_main.h:58
warning: 58 ../sysdeps/nptl/libc_start_call_main.h: No such file or directory
#17 0x00007eef9ca2a28b in __libc_start_main_impl (main=0x5ea791f0d270 <main>, argc=47, argv=0x7ffef7a38438, init=<optimized out>, fini=<optimized out>, rtld_fini=<optimized out>, stack_end=0x7ffef7a38428) at ../csu/libc-start.c:360
warning: 360 ../csu/libc-start.c: No such file or directory
#18 0x00005ea791f0d2a5 in _start ()
[Inferior 1 (process 322960) detached]

Running Qwen 3.8 27B Q6_K_XL GGUF with Q8_0 KV cache

My workaround for now is using spec-type = draft-mtp,ngram-map-k instead of just spec-type = draft-mtp; this way the problem does not reproduce. As far as I understood, in this config MTP is invoked only in those cases when ngram-map-k did not return good result, so the load on GPU is lower.

FWIW RmWatchDog does not necessarily time out on long kernels, so something else may be tripping it. This one does not trip it for me https://gist.github.com/ORippler/7853d586237faf518c48b0e284160e12.
- Do you folks have displays attached to your GPUs?
- Can somebody give full dmesg/logs of the watchdog + XID 8 event?
The intitial repro of Eval bug: CUDA crash on tensor copy #20146 was reported on 30ea7bc, decreasing likelihood of PDL-involvement

Do you folks have displays attached to your GPUs?

I do for my 5070 ti, if the logs are accurate it seems to be happening for both my cards, the 5060 ti has no displays attached, my crash that led me to this report:
https://gist.github.com/Corosauce/3bd5b1dd503f6eb704748d0e23232cc3

my llama swap model run params:
https://gist.github.com/Corosauce/84122e262380da85bd98153beed202f0

Can somebody give full dmesg/logs of the watchdog + XID 8 event?

Here's my AI curated and gathered logs if at all helpful, it's been happening a lot for me past couple days:
https://gist.github.com/Corosauce/567dfbffc53aa44c365c3ec2aa4c5bfd

Thought I managed to reproduce it even with the 60 second timeout adjustment but it wasn't setup right.

Currently waiting for it to reproduce with the 60 second fix + q8 kv cache.

Do you folks have displays attached to your GPUs

I don't.

@ORippler

Re: @ORippler's two questions (displays attached / full logs of the watchdog + Xid 8 event) — answering for the repro I reported earlier in this thread.

1. Displays attached: yes. On this box the RTX 5090 is the primary desktop GPU — two displays connected (one DP, one HDMI), actively driven (KDE Plasma / Wayland) at the moment of the event. An AMD iGPU is present but has no outputs connected, so the 5090 was doing display + compute simultaneously. (Noting this as a data point alongside @gg7's headless repro, not as support for a display-only theory.)

2. Full kernel-side log of the event. The four NVRM lines below are the complete kernel record — there are no other kernel messages within ±5 minutes of the event, before or after:

aug 16 23:01:35 kernel: NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked!  Notify Timeout Seconds: 7
aug 16 23:01:35 kernel: NVRM: GPU at PCI:0000:01:00: GPU-3da9f59b-4d76-1187-950a-e2beea666068
aug 16 23:01:35 kernel: NVRM: GPU Board Serial Number: 0
aug 16 23:01:35 kernel: NVRM: Xid (PCI:0000:01:00): 8, pid=709378, name=llama-server, channel 0x00000027

Environment at the time: kernel 7.1.8-1-cachyos, driver 610.57.04 (open kernel module), llama-server serving Qwen3.8-27B with quantized KV cache — both K and V were Q8_0 at the time of this event (the V→q4_0 change I mentioned earlier happened the next day). This is the only Xid on record on this machine (journal retention reaches back 3 days before the event); RmWatchDogTimeout=60 was applied the following morning and there have been zero recurrences through Aug 22, with the setting still active (RegistryDwords: "RmWatchDogTimeout=60" in /proc/driver/nvidia/params).

Possibly relevant to "RmWatchDog does not necessarily time out on long kernels": the entity that fired here is literally krcWatchdog_IMPL (the RC watchdog), and its logged notify window — Notify Timeout Seconds: 7 — matches the default value of RmWatchDogTimeout, the exact parameter that eliminated recurrences when raised to 60. So on this box the trip condition looks like the channel failing to notify the RC watchdog within the window (whatever the underlying GPU-side stall is), rather than a per-kernel launch-length limit — which would be consistent with your long-kernel gist not tripping it.

Userspace at the same instant (systemd-coredump backtrace of pid 709378): llama-server aborted on the CUDA launch-timeout error inside a stream synchronize called from the server main loop — same shape as @void-mckenzie's trace above:

#4  ggml_cuda_error                       (libggml-cuda.so.0)
#5  <ggml_backend_cuda_synchronize>       (libggml-cuda.so.0)
#7  llama_context::synchronize()          (libllama.so.0)
#8  server_queue::yield_to_queue(...)     (libllama-server-impl.so)
#10 server_queue::start_loop(long)        (libllama-server-impl.so)

The application-side server log from that run was unfortunately rotated away, so the above is everything preserved from the event.

   * Do you folks have displays attached to your GPUs?

Not on that card, no.

Since increasing the RmWatchDogTimeout from 7 to 60 I went from multiple timeouts an hour to zero, 7 hours not a single problem so far.

EDIT: 3+ days, 100% solved.

Are all the user reports from using the quantized kv cache? is there anyone who faced this issue with a f16 kv cache?

Are all the user reports from using the quantized kv cache? is there anyone who faced this issue with a f16 kv cache?

Okay when switching from FP16 to Q8_0 it repros almost immediately for me

So far, I've receives time-outs in

FLASH_ATTN ~~(Misc. bug: CUDA FlashAttention synccheck reports divergent __syncthreads #27678) may be a root-cause~~ (error repro'd after this)
MMQ

When launched with

GGML_CUDA_PDL=0
GGML_CUDA_DISABLE_GRAPHS=1

(+ some more debug statement logs)

Small data point from my box (the third report in this thread, RTX 5090 / 610.57.04 open module): the RmWatchDogTimeout=60 workaround was not actually applying on my system. It was present in /proc/cmdline and echoed in /proc/driver/nvidia/params, but the next crash still logged:

NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked!  Notify Timeout Seconds: 7

so the RC watchdog window was still the 7 s default. One confound I should be upfront about: my dwords string also contained an entry I'd picked up from elsewhere (RMOverclockingCtrl=1), which is not a valid registry dword for this driver — so I can't fully disentangle whether that invalid entry (first in the string) broke parsing, or whether the spelling matters.

For reference, in NVIDIA's open-source header (src/nvidia/interface/nvrm_registry.h) the canonical name is RmWatchDogTimeOut (capital O in "TimeOut"), default 7, valid range 5–60:

#define NV_REG_STR_RM_WATCHDOG_TIMEOUT      "RmWatchDogTimeOut"
#define NV_REG_STR_RM_WATCHDOG_TIMEOUT_LOW  5
#define NV_REG_STR_RM_WATCHDOG_TIMEOUT_HI   60
#define NV_REG_STR_RM_WATCHDOG_TIMEOUT_DEFAULT 7

After changing my cmdline to exactly nvidia.NVreg_RegistryDwords=RmWatchDogTimeOut=60 I've been ~5 days with zero crashes, including with MTP speculative decoding enabled (the same config previously gave me 15 Xid 8s in ~36 h).

Suggestion for anyone running the 60 s workaround: it's worth verifying it actually took effect — the next time a watchdog event fires, check that dmesg says Notify Timeout Seconds: 60 and not 7. (Also, the GSP registry table in gsp_ga10x.bin only contains RmWatchDogTimeOut, RmWatchDogInterval, RmRcWatchdog — no RMOverclockingCtrl.)

One workaround would be to disable the kernel launch timeouts, but that is just masking something that shouldn't happen (7s kernels means we are at <0.2 t/s)

Navigation Menu

Repository navigation

Eval bug: CUDA kernel stall during model execution, killed by watchdog #27102

Name and Version

Operating systems

GGML backends

Hardware

Models

Problem description & steps to reproduce

First Bad Commit

Relevant log output

6 remaining items

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Eval bug: CUDA kernel stall during model execution, killed by watchdog #27102

Description

Name and Version

Operating systems

GGML backends

Hardware

Models

Problem description & steps to reproduce

First Bad Commit

Relevant log output

Activity

ccebelenski commented on Aug 15, 2026

razvandinculescu commented on Aug 16, 2026

gg7 commented on Aug 17, 2026

EdwardSierra commented on Aug 17, 2026

razvandinculescu commented on Aug 17, 2026

ggerganov commented on Aug 17, 2026

awrgold commented on Aug 17, 2026

ccebelenski commented on Aug 18, 2026

razvandinculescu commented on Aug 18, 2026

pwilkin commented on Aug 18, 2026

6 remaining items

ggerganov commented on Aug 19, 2026

void-mckenzie commented on Aug 20, 2026

portnov commented on Aug 21, 2026

ORippler commented on Aug 21, 2026

Corosauce commented on Aug 22, 2026

gg7 commented on Aug 22, 2026

razvandinculescu commented on Aug 22, 2026

ccebelenski commented on Aug 22, 2026

Corosauce commented on Aug 23, 2026

am17an commented on Aug 23, 2026

ORippler commented on Aug 24, 2026

ORippler commented on Aug 25, 2026

EdwardSierra commented on Aug 25, 2026

ORippler commented on Aug 26, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions