Eval bug: draft-mtp with Qwen3.6-35B-A3B on 2x RTX 5060 Ti (layer split): "CUDA error: the launch timed out" in MTP path at ~88K context #30396

Name and Version

version: 0.6.0-dev (build 11572, commit efe3626)
built with GNU 15.2.0 for Linux x86_64

Operating systems

Linux

GGML backends

CUDA

Hardware

Ryzen 7 5800X, B550, 2x RTX 5060 Ti 16 GB (Blackwell, sm_120).
GPU0 in PCIe 4.0 x16 slot (no display), GPU1 in PCIe 3.0 x4 slot (drives the desktop, KDE Plasma Wayland).
Driver: 595.91.07
Cuda compilation tools, release 13.1, V13.1.115
Build cuda_13.1.r13.1/compiler.37061995_0

Models

Qwen3.6-35B-A3B Q4_K_M (bartowski/Qwen_Qwen3.6-35B-A3B-GGUF), built-in MTP head.

Problem description & steps to reproduce

--spec-type draft-mtp the server aborts after a while of agentic use (OpenCode, single slot, prompts growing to ~90K tokens). It happened twice, both times inside the MTP code path; without --spec-type draft-mtp the same setup has been stable.

Crash 1 (older build): Xid 8 + "RC watchdog: GPU is probably locked" in dmesg, backtrace in common_speculative_impl_draft_mtp::draft.

Crash 2 (this report): right after a request finished at n_tokens = 87943, the next request (prefix reuse, f_sim_best = 0.987) aborted within ~2.5 s:

CUDA error: the launch timed out and was terminated
current device: 1, in function ggml_backend_cuda_synchronize
cudaStreamSynchronize(cuda_ctx->stream())

#7 ggml_backend_cuda_synchronize
#8 ggml_backend_sched_synchronize
#9 llama_context::synchronize()
#10 llama_get_embeddings_nextn(llama_context*)
#11 common_speculative_impl_draft_mtp::process(llama_batch const&)
#12 common_speculative_process
#15 server_context_impl::decode
#16 server_context_impl::update_slots

Device 1 is the GPU attached to the display, so it has kernelExecTimeoutEnabled=<1/0, из скрипта>. NVreg_RegistryDwords="RmDisableGpuWatchdog=1" is set and active, but does not prevent it. Normal kernels at this batch size take milliseconds, so this looks like a kernel hanging in the MTP/NextN path rather than a legitimately long kernel.

Also at startup in this run the MTP draft context failed to allocate its compute buffer on device 1 and the scheduler fell back:

creating MTP draft context against the target model ...
ggml_backend_cuda_buffer_type_alloc_buffer: allocating 641.06 MiB on device 1: cudaMalloc failed: out of memory
graph_reserve: failed to allocate compute buffers
sched_reserve: compute buffer allocation failed, retrying without pipeline parallelism

Before the crash generation was normal (~104 t/s, draft acceptance 0.86).

Crash 3 (latest master, <номер сборки>): same error on device 1, this time during generation (n_gen = 276) at n_tokens = 73737. The backtrace no longer shows MTP frames, only the generic sync after decode:

#7 ggml_backend_cuda_synchronize
#8 ggml_backend_sched_synchronize
#9 llama_context::synchronize()
#10 server_queue::yield_to_queue
#11 server_context_impl::decode(int&, int)
#12 server_context_impl::update_slots()

dmesg at the same moment:

NVRM: krcWatchdog_IMPL: RC watchdog: GPU is probably locked! Notify Timeout Seconds: 7
NVRM: GPU at PCI:0000:2b:00: GPU-...
NVRM: Xid (PCI:0000:2b:00): 8, pid=..., name=llama-server, channel 0x0000002b

So the GPU1 channel made no progress for 7+ seconds, i.e. a hang, not a slow kernel. GPU1 is the display GPU, so it is the only one with the RC watchdog; the hang might originate in a cross-device wait on GPU0's work.

Isolation tests (each run until context > 90K in an agentic session):

without --spec-type draft-mtp: <результат>
with draft-mtp and GGML_CUDA_DISABLE_GRAPHS=1: <результат>

Command line:

CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=0,1
llama-server -m Qwen_Qwen3.6-35B-A3B-Q4_K_M.gguf
--ctx-size 131072 --n-gpu-layers 999 --main-gpu 0 --tensor-split 12,11
--cache-type-k q8_0 --cache-type-v q8_0 --flash-attn on
--batch-size 2048 --ubatch-size 512 --parallel 1 --cache-ram 8192
--threads 4 --threads-batch 16
--temp 1.0 --top-p 0.95 --top-k 20 --min-p 0
--spec-type draft-mtp --reasoning on --reasoning-preserve --jinja

Full environment info and log are attached (mtp-crash-info.txt).

First Bad Commit

No response

Relevant log output

This GDB supports auto-downloading debuginfo from the following URLs:
https://debuginfod.ubuntu.com
Enable debuginfod for this session? (y or [n]) [answered N; input not from terminal]
Debuginfod has been disabled.
To make this setting permanent, add 'set debuginfod enabled off' to .gdbinit.
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/usr/lib/x86_64-linux-gnu/libthread_db.so.1".
__syscall_cancel_arch () at ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S:56
⚠️ warning: 56 ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S: No such file or directory
#0 __syscall_cancel_arch () at ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S:56
56 in ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S
#1 0x00007c96ee0a039c in __internal_syscall_cancel (a1=, a2=, a3=, a4=, a5=0, a6=0, nr=61) at ./nptl/cancellation.c:49
⚠️ warning: 49 ./nptl/cancellation.c: No such file or directory
#2 __syscall_cancel (a1=, a2=, a3=, a4=, a5=a5@entry=0, a6=a6@entry=0, nr=61) at ./nptl/cancellation.c:75
75 in ./nptl/cancellation.c
#3 0x00007c96ee11cacf in __GI___wait4 (pid=, stat_loc=, options=, usage=) at ../sysdeps/unix/sysv/linux/wait4.c:30
⚠️ warning: 30 ../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#4 0x00007c96ee741f23 in ggml_print_backtrace () from ~/llama.cpp/build/bin/libggml-base.so.0
#5 0x00007c96ee7420d6 in ggml_abort () from ~/llama.cpp/build/bin/libggml-base.so.0
#6 0x00007c96ea906125 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from ~/llama.cpp/build/bin/libggml-cuda.so.0
#7 0x00007c96ea90f088 in ggml_backend_cuda_synchronize(ggml_backend*) () from ~/llama.cpp/build/bin/libggml-cuda.so.0
#8 0x00007c96ee75d22e in ggml_backend_sched_synchronize () from ~/llama.cpp/build/bin/libggml-base.so.0
#9 0x00007c96ed6212ac in llama_context::synchronize() () from ~/llama.cpp/build/bin/libllama.so.0
#10 0x00007c96ee9230f9 in server_queue::yield_to_queue(std::function<void ()>&&) () from ~/llama.cpp/build/bin/libllama-server-impl.so
#11 0x00007c96ee98355f in server_context_impl::decode(int&, int) () from ~/llama.cpp/build/bin/libllama-server-impl.so
#12 0x00007c96ee9852f8 in server_context_impl::update_slots() () from ~/llama.cpp/build/bin/libllama-server-impl.so
#13 0x00007c96ee926678 in server_queue::start_loop(long) () from ~/llama.cpp/build/bin/libllama-server-impl.so
#14 0x00007c96ee8d503e in llama_server(common_params&, int, char**, server_child&) () from ~/llama.cpp/build/bin/libllama-server-impl.so
#15 0x00007c96ee8d6c47 in llama_server(int, char**) () from ~/llama.cpp/build/bin/libllama-server-impl.so
#16 0x00007c96ee02a601 in __libc_start_call_main (main=main@entry=0x60dda7f2e2f0

, argc=argc@entry=52, argv=argv@entry=0x7ffd910822a8) at ../sysdeps/nptl/libc_start_call_main.h:59
⚠️ warning: 59 ../sysdeps/nptl/libc_start_call_main.h: No such file or directory
#17 0x00007c96ee02a718 in __libc_start_main_impl (main=0x60dda7f2e2f0 , argc=52, argv=0x7ffd910822a8, init=, fini=, rtld_fini=, stack_end=0x7ffd91082298) at ../csu/libc-start.c:360
⚠️ warning: 360 ../csu/libc-start.c: No such file or directory
#18 0x000060dda7f2e325 in _start ()
[Inferior 1 (process 4042372) detached]

mtp-crash-clean.log