Eval bug: multi-GPU setups not working

[Anbeeld]

Name and Version

0.1.0

Operating systems

Other? (Please let us know in description)

GGML backends

CUDA

Hardware

2 identical cards

Models

No response

Problem description & steps to reproduce

Per NickCanCode in Reddit comment https://www.reddit.com/r/LocalLLaMA/comments/1t88zvv/comment/okxsz31/

Doesn't work for me. It gives [log] whenever I make a request.

P.S. Using 2 identical cards.

First Bad Commit

No response

Relevant log output

beellama.cpp-main\ggml\src\ggml-cuda\ggml-cuda.cu:98: CUDA error 
CUDA error: an illegal memory access was encountered

[a-51-50]

I'm getting the same error, when i remove DFlash settings it works fine.
Specs:
Dell Precision T7610
CPU: 2xXeon E5-2680 (AVX)
RAM: 512GB DDR3
GPU: 3x3060
GPU (for display): 1050ti (I've tried removing 1050ti, still the same error)

[louis-szeto]

same here

[walmis]

Description

CUDA Illegal Memory Access occurs during the prefill stage when using dflash speculative decoding. The crash happens specifically during the KV cache synchronization step in a multi-GPU setup.

---

Environment

• OS: Ubuntu 26.04
• CUDA: 13.1
• GPU: Dual NVIDIA RTX 3090 (2x24GB)
• Model: Qwen3.6-27B-UD-Q4_K_XL.gguf
• Drafter: dflash-draft-3.6-q4_k_m.gguf

---

Build Configuration

Built with the following flags:

cmake -B build -DGGML_CUDA=ON \
  -DGGML_CUDA_NCCL=ON \
  -DLLAMA_BUILD_TESTS=OFF \
  -DGGML_CUDA_FA=ON \
  -DGGML_CUDA_FA_ALL_QUANTS=ON \
  -DGGML_NATIVE=ON \
  -DCMAKE_BUILD_TYPE=Release

---

Reproduction Command

The server is started with these arguments:

MODEL=Qwen3.6-27B-UD-Q4_K_XL.gguf

ARGS=(
  -m "$MODEL"
  --no-mmproj-offload
  --spec-draft-model "dflash-draft-3.6-q4_k_m.gguf"
  --spec-type dflash
  --spec-dflash-cross-ctx 1024
  --port 8082
  -np 1
  --kv-unified
  -ngl all
  --spec-draft-ngl all
  -b 2048
  -ub 256
  --ctx-size 65536
  --cache-type-k turbo4
  --cache-type-v turbo3_tcq
  --flash-attn on
  --jinja
  --no-mmap
  --mlock
  --no-host
  --metrics
  --log-timestamps
  --log-prefix
  --log-colors off
  --reasoning on
  --chat-template-kwargs '{"preserve_thinking":true}'
  --temp 0.6
  --top-k 20
  --min-p 0.0
  --host 0.0.0.0
  --direct-io
)

./build/bin/llama-server "${ARGS[@]}"

---

Stack Trace

The crash is triggered during common_speculative_state_dflash::flush_prefill(), leading to an illegal memory access in the CUDA backend.

[PROJECT_ROOT]/ggml/src/ggml-cuda/ggml-cuda.cu:98: CUDA error
[LOG_ID] E CUDA error: an illegal memory access was encountered
[LOG_ID] E   current device: 0, in function ggml_backend_cuda_synchronize at [PROJECT_ROOT]/ggml/src/ggml-cuda/ggml-cuda.cu:3099
[LOG_ID] E   cudaStreamSynchronize(cuda_ctx->stream())

#7  0x000071bdd0aabf68 in ggml_backend_cuda_synchronize(ggml_backend*) () from [PROJECT_BIN]/libggml-cuda.so.0
#8  0x000071bdd56c4f38 in llama_context::dflash_kv_cache_update(int) () from [PROJECT_BIN]/libllama.so.0
#9  0x000071bdd56c5288 in llama_context::dflash_kv_cache_update_gpu(void const*, int, int, int, void (*)(void*, void const*, unsigned long, unsigned long)) () from [PROJECT_BIN]/libllama.so.0
#10 0x000071bdd5cf6e58 in common_speculative_state_dflash::flush_prefill() () from [PROJECT_BIN]/libllama-common.so.0
#11 0x000071bdd5ced9ad in common_speculative_flush_prefill(common_speculative*) () from [PROJECT_BIN]/libllama-common.so.0
#12 0x0000580cb8df2d6b in server_context_impl::update_slots() ()
#13 0x0000580cb8e9ebc2 in server_queue::start_loop(long) ()
#14 0x0000580cb8d45e65 in main ()

---

Additional Context

The error consistently appears during the initial prefill on a dual RTX 3090 system. The issue occurs specifically when the speculative dflash state attempts to synchronize the KV cache on the GPU. This suggests a potential race condition or pointer invalidation in the dflash prefill logic when interacting with multi-GPU memory spaces, unified KV caches, or turbo cache types. Disabling speculative decoding resolves the crash.

[Anbeeld]

Commited some changes that might fix the issue. Can't verify myself as I don't have multi-GPU setup. Please try it out everyone and report back, will follow up with more fixes if needed.