Skip to content

Eval bug: 2-GPU CUDA tensor split is unstable with Qwen3.8-27B on Windows, including with MTP disabled #27750

Closed

Description

@Anbeeld

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:

  1. Initial small canary passed.
  2. A large request with 24,344 prompt tokens passed.
  3. The following image-first request failed.
  4. 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=1 did 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

  1. Anbeeld commented on Aug 28, 2026

    @Anbeeld
    Author

    Follow-up on the issue.

    A similar illegal-memory-access failure now reproduces with layer split, not tensor split.

    Configuration:

    • BeeLlama v0.4.4, based on llama.cpp 6fdd0ac89
    • Windows 11
    • 2× RTX 3090 24 GB
    • Qwen3.8-27B UD-Q5_K_XL target
    • Qwen3.8-27B DFlash2 Q4_K_M draft
    • DFlash2 draft pinned to CUDA0
    • --split-mode layer --tensor-split 50,50
    • BF16 KV
    • BF16 mmproj with --no-mmproj-offload
    • --ctx-size 262144
    • -b 4096 -ub 1024
    • --spec-type draft-dflash,ngram-mod

    The 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 encountered

    current device: 1

    ggml_backend_cuda_synchronize

    cudaStreamSynchronize(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.

  2. Anbeeld commented on Aug 28, 2026

    @Anbeeld
    Author

    Additional tensor-split result: 3/3 fresh runs have failed.

    Same configuration in all three:

    • Windows 11
    • 2× RTX 3090
    • Qwen3.8-27B UD-Q6_K_M
    • MTP
    • --split-mode tensor --tensor-split 1,1
    • BF16 KV
    • b=4096, ub=2048
    • 262k context

    The first two runs both eventually failed with the same CUDA illegal-memory-access report:

    current device: 0
    ggml_cuda_ar_acquire_slot
    allreduce.cu:368
    cudaEventSynchronize(p->ev_pool[i][slot].ker)
    

    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:

    8192 tokens  @ 1406.65 tok/s
    12288 tokens @ 1412.29 tok/s
    

    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.

  3. Anbeeld commented on Sep 6, 2026

    @Anbeeld
    Author

    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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions