Repository navigation
Eval bug: 2-GPU CUDA tensor split is unstable with Qwen3.8-27B on Windows, including with MTP disabled #27750
Description
Name and Version
llama.cpp b10595
commit e8eed45
Official Windows x64 CUDA 13.3 build:
https://github.com/ggml-org/llama.cpp/releases/tag/b10595
Operating systems
Windows
GGML backends
CUDA
Hardware
Ryzen 9 9950X
96 GB DDR5
2x RTX 3090 24 GB
Windows 11
Both GPUs are used with:
--device CUDA0,CUDA1
--split-mode tensor
--tensor-split 1,1
Models
Qwen3.8-27B-UD-Q6_K_M.gguf
BF16 mmproj
https://huggingface.co/unsloth/Qwen3.8-27B-GGUF
Problem description & steps to reproduce
I have been investigating persistent instability in the 2-GPU tensor-split path on Windows with Qwen 3.8 27B.
There are already reports such as #27122 where the problem appears tied to MTP. In my case I was also able to reproduce a failure on an official upstream llama.cpp b10595 build with MTP completely disabled, so MTP is not required for at least one failure surface.
The upstream test used:
Qwen3.8-27B-UD-Q6_K_M.gguf
BF16 mmproj
Windows 11
2x RTX 3090 24 GB
--device CUDA0,CUDA1
--split-mode tensor
--tensor-split 1,1
--ctx-size 131072
--fit off
-b 4096
-ub 1024
--cache-type-k bf16
--cache-type-v bf16
--kv-unified
mmproj enabled with --no-mmproj-offload
MTP disabled
On b10595:
- Initial small canary passed.
- A large request with 24,344 prompt tokens passed.
- The following image-first request failed.
- The client connection was reset and llama-server later aborted with a CUDA error at
ggml-cuda.cu:107.
The minimum free VRAM during this run was about 8,536 MiB on CUDA0 and 8,440 MiB on CUDA1.
As a preface, I am not claiming that every failure described below has the same root cause. The upstream failure surfaced as a CUDA abort, while some of my fork reproductions surfaced as a no-progress GPU state instead. What I can establish is that the affected execution area is not limited to MTP.
I have been working on this in BeeLlama, which is a llama.cpp fork. I pushed the current investigation checkpoint here:
Branch: v0.4.4-wddm
Commit: 0c594be
This is NOT a fix or a stable candidate. It contains the useful correctness work that survived the investigation, but the remaining multi-GPU failure is still unresolved.
The investigation covered tensor-split metadata propagation, projected graph reuse/lifetime, scheduler topology reuse, CUDA resource/device ownership, and the Windows two-device no-P2P collective path.
Some useful negative results:
- Disabling CUDA graphs did not eliminate the failure.
LLAMA_GRAPH_REUSE_DISABLE=1did not eliminate the BeeLlama reproducer either. One fresh request still stopped making progress after 20,535 prompt tokens while running at about 1,368 prompt tok/s.- Reducing the local WDDM collective to one reusable generation slot did not fix it.
- Broad synchronization/retirement of prior projected graph generations moved the failure but did not eliminate it.
- The exact projected operation observed near one failure was:
Q5_K [3072,5120] x F32 [3072,1024] -> F32 [5120,1024].
The same operation completes standalone on both RTX 3090s at about 79-80 TFLOPS. - Routing that exact projected Q5_K operation through cuBLAS did not fix the real workload.
- Broadly disabling CUDA fusion/MMQ also did not establish a fix.
- Several slower synchronization/single-stream variants could make the test sequence pass, but prefill dropped from about 1,349 tok/s to roughly 795-806 tok/s. I treated those as timing/order sensitivity evidence, not fixes.
- Reproduced failures frequently occurred with roughly 8 GiB of dedicated VRAM still free per GPU and without progressive shared-memory growth.
There is also no established deterministic token boundary.
During prompt minimization an approximately 14-15k-token request initially failed, but after freezing the exact request it tokenized to 14,425 prompt tokens and passed three fresh processes. So I do not think this is a simple "Nth 4096-token batch always fails" problem. The failure appears timing/order sensitive.
Retained same-process state is also not universally required. On a clean diagnostic-free BeeLlama build I reproduced a first-request stall around 16,439 / 24,345 prompt tokens at about 1,365 tok/s, again with about 8.4 GiB free per GPU.
At the same time, retained workloads are clearly affected too. After publishing the checkpoint I tried using it normally with:
--ctx-size 262144
-b 4096
-ub 2048
--split-mode tensor
--tensor-split 1,1
--spec-type draft-mtp
mmproj enabled
It ran normally for several requests and reached 62,261 retained tokens. On the next request it generated 247 tokens and then crashed with:
CUDA error: an illegal memory access was encountered
current device: 0, in function ggml_cuda_ar_wddm_allreduce
cudaEventSynchronize(p->published[i][slot])
So the current state is:
- 2-GPU CUDA tensor split is demonstrably unstable in this workload family.
- MTP can trigger failures, but an upstream MTP-off reproduction also exists.
- VRAM exhaustion is not required.
- CUDA graph reuse alone is not sufficient to explain it.
- The exact projected Q5_K arithmetic works in isolation.
- Slower/synchronizing execution often suppresses the problem, which suggests an ordering/lifetime race, but I have not isolated the owning invariant.
I am stuck after a fairly extensive investigation and would appreciate collaboration from anyone familiar with the meta backend, tensor-split scheduler, CUDA event/resource lifetime, or multi-GPU execution.
I can run targeted builds, diagnostics, or reduced test cases.
First Bad Commit
Unknown.
Confirmed on upstream b10595 / e8eed45.
Relevant log output
Upstream b10595 result
Configuration:
Windows 11
2x RTX 3090
Qwen3.8-27B-UD-Q6_K_M
tensor split 1,1
ctx 131072
batch 4096
ubatch 1024
MTP off
mmproj on
Initial canary: PASS
Large request:
prompt tokens: 24344
result: PASS
Following image-first request:
result: connection reset / server abort
CUDA error surfaced at ggml-cuda.cu:107
Minimum free VRAM:
CUDA0: ~8536 MiB
CUDA1: ~8440 MiB
</details>
Activity
Anbeeld commentedon Aug 28, 2026
Follow-up on the issue.
A similar illegal-memory-access failure now reproduces with layer split, not tensor split.
Configuration:
6fdd0ac89--split-mode layer --tensor-split 50,50--no-mmproj-offload--ctx-size 262144-b 4096 -ub 1024--spec-type draft-dflash,ngram-modThe process first completed a 125,280-token prompt and finished that request at 131,450 total tokens.
It then completed many retained-context requests and reached 170,607 tokens without failure.
Later, a new request successfully processed a 63,449-token prompt at ~1.26k tok/s and generated another 4,597 tokens at ~29.4 tok/s before failing with:
CUDA error: an illegal memory access was encounteredcurrent device: 1ggml_backend_cuda_synchronizecudaStreamSynchronize(cuda_ctx->stream())cudaStreamSynchronize()is the point where the asynchronous error was detected, not necessarily the operation that caused it.The draft model was pinned entirely to CUDA0, while the illegal access was reported on CUDA1. This makes a direct DFlash2 draft-kernel failure less likely, although DFlash2 may still be a trigger through the target verification workload.
It seems like tensor split is not required for this failure surface.
This also does not fit a simple context-length threshold. The same process had already operated successfully above 170k tokens before later failing on a request around 68k total tokens.
I am not claiming that the layer-split and tensor-split failures have the same root cause.
Current evidence points more broadly toward a multi-GPU CUDA ordering/lifetime problem shared by both execution modes, potentially involving target graph reuse, scheduler/buffer lifetime, cross-device transfers, or another asynchronous CUDA path.
Anbeeld commentedon Aug 28, 2026
Additional tensor-split result: 3/3 fresh runs have failed.
Same configuration in all three:
--split-mode tensor --tensor-split 1,1b=4096,ub=2048The first two runs both eventually failed with the same CUDA illegal-memory-access report:
Their failure points were very different. One reached ~74k retained tokens; another completed an ~80k-token request and failed during the following prompt.
The third fresh run failed much earlier. It loaded normally, started the first request, processed:
and then the process terminated abruptly at 12,288 prompt tokens. No CUDA error or graceful shutdown was written to the log.
The two explicit CUDA failures still converge on the same AllReduce slot-reuse synchronization, which makes that path a strong suspect for the severe tensor-split failure mode. The third run provides no error location, so I cannot attribute that termination to AllReduce.
Anbeeld commentedon Sep 6, 2026
Life lesson: if your inference is unstable, but crashes appear to be random and you can't find specific reason no matter how deep you dig... then it's probably not software, it's your undervolt being unstable at certain voltages, like it was for me, apparently.