qwen36_35b_dflash_repro.zip

Summary

Serving Qwen3.6-35B-A3B with the DFlash block-diffusion draft model (--speculative-config '{"method":"dflash",...}') crashes the engine with CUDA error: an illegal memory access was encountered (cudaErrorIllegalAddress) once request concurrency reaches ~16–32. The crash is not caused by OOM, KV-cache exhaustion, or long generations — GPU memory and KV-cache usage are low at the time of the crash. It is reproducible and is isolated to the DFlash + CUDA graph combination: disabling either DFlash or CUDA graphs (--enforce-eager) makes the crash disappear completely.

Environment

Models

The installed vLLM 0.22.1 already ships native DFlash support (vllm/v1/spec_decode/dflash.py, vllm/v1/spec_decode/utils.py, vllm/model_executor/models/qwen3_dflash.py).

Reproduction

Launch command

CUDA_VISIBLE_DEVICES=4,5,6,7 setsid python -m vllm.entrypoints.openai.api_server \
  --model /gfs/space/public/public_models/Qwen3.6-35B-A3B \
  --served-model-name qwen36-35b-a3b-dflash-repro \
  --port 18094 \
  --tensor-parallel-size 4 \
  --dtype bfloat16 \
  --trust-remote-code \
  --max-model-len 65536 \
  --max-num-seqs 64 \
  --gpu-memory-utilization 0.90 \
  --speculative-config '{"method": "dflash", "model": "/gfs/space/public/public_models/Qwen3.6-35B-A3B-DFlash", "num_speculative_tokens": 15}' \
  --max-num-batched-tokens 32768 \
  --gdn-prefill-backend triton \
  --api-key sk-local

Trigger: 16+ concurrent requests

import concurrent.futures, json, urllib.request
URL = "http://127.0.0.1:18094/v1/chat/completions"
HDR = {"Authorization": "Bearer sk-local", "Content-Type": "application/json"}
PROMPT = ("Prove 1^3+2^3+...+n^3=(n(n+1)/2)^2 with two methods, verify cases, "
          "discuss generalizations. Be exhaustive.")
def one(i):
    body = json.dumps({"model": "qwen36-35b-a3b-dflash-repro",
        "messages": [{"role": "user", "content": PROMPT}],
        "max_tokens": 512, "temperature": 0.7}).encode()
    with urllib.request.urlopen(urllib.request.Request(URL, data=body, headers=HDR), timeout=300) as r:
        return r.status
with concurrent.futures.ThreadPoolExecutor(max_workers=16) as ex:
    print(list(ex.map(one, range(16))))

Result: ~7/16 requests return HTTP 500, then the server shuts down (Application shutdown complete). The same run with max_workers=1 succeeds.

Control experiments (the key evidence)

All four runs use the same target model, prompt, and concurrency (16). Only the flagged variable changes:

Interpretation: the crash requires both DFlash and CUDA graphs. Removing either eliminates it. The fact that CUDA_LAUNCH_BLOCKING=1 also masks it (run #4) indicates an asynchronous out-of-bounds access whose timing window disappears under synchronous execution.

Crash signature

The engine dies with EngineDeadError propagated from the worker. The root CUDA error is reported by the NCCL watchdog:

[rank1]:[E709 12:05:21.290990045 ProcessGroupNCCL.cpp:2119] [PG ID 2 PG GUID 3 Rank 1]
Process group watchdog thread terminated with exception:
CUDA error: an illegal memory access was encountered
Search for `cudaErrorIllegalAddress' in https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html

vllm.v1.engine.exceptions.EngineDeadError: EngineCore encountered an issue.
See stack trace (above) for the root cause.

What the Python stack does NOT show

The propagated traceback only contains the EngineCore → multiproc_executor → shm_broadcast chain — no draft/target model kernel frame appears:

File ".../vllm/v1/engine/core.py", line 1238, in _process_engine_step
    outputs, model_executed = self.step_fn()
File ".../vllm/v1/executor/multiproc_executor.py", line 386, in get_response
    status, result = mq.dequeue(timeout=dequeue_timeout)
File ".../vllm/distributed/device_communicators/shm_broadcast.py", line 772, in dequeue
    with self.acquire_read(timeout, indefinite) as buf:
File ".../vllm/distributed/device_communicators/shm_broadcast.py", line 694, in acquire_read
    raise RuntimeError("cancelled")
RuntimeError: cancelled

This is consistent with a CUDA kernel going out of bounds asynchronously: the Python step_fn returns, and the illegal access surfaces later when the NCCL watchdog inspects the next all-reduce. Because execution is async, the offending kernel cannot be pinpointed from the Python traceback alone. CUDA_LAUNCH_BLOCKING=1 would localize it, but (as shown in run #4) it also prevents the crash.

State at crash time (from dump_input.py)

vLLM dumps the scheduler state on fatal error. Two independent crashes:

Crash A — 32 concurrent (first reproduction), still in prefill:

num_running_reqs        = 32
num_waiting_reqs        = 0
kv_cache_usage          = 0.199   (19.9%)
total_num_scheduled_tokens = 512
num_computed_tokens     = [117 .. 144]  (min=117, max=144)   # very short contexts
Avg prompt throughput   = 323.2 tokens/s   (still prefilling)
Avg generation throughput = 25.2 tokens/s  (just started)
SpecDecoding: mean acceptance length 8.80, acceptance rate 52.0%

Crash B — 16 concurrent (minimal repro), mid-decode:

num_running_reqs        = 7      (others already finished)
num_waiting_reqs        = 0
kv_cache_usage          = 0.0437 (4.4%)
total_num_scheduled_tokens = 112
num_computed_tokens     = [365, 518, 574, 365, 413, 432, 440]   # short contexts
num_output_tokens       = [339, 470, 526, 317, 365, 384, 392]   # short generations
per-req num_scheduled_tokens = 16  (1 bonus + 15 spec, standard DFlash)
SpecDecoding: mean acceptance length 1.86, acceptance rate 5.7%

Both crashes occur with low KV usage (4–20%), short contexts (<600 computed tokens), and short outputs (<530 tokens). This rules out OOM, KV-cache exhaustion, and long-generation accumulation as causes. The acceptance rate at crash time is not consistent (52% vs 5.7%), so a specific acceptance value is not the trigger either.

Root-cause analysis

The bug is in the DFlash proposer path under CUDA graph capture/replay. Concrete observations:

1. Buffer sizing vs. graph capture size mismatch. In vllm/v1/spec_decode/dflash.py:

   self.max_query_tokens = self.max_batch_size * (1 + self.num_speculative_tokens)
   # = max_num_seqs(64) * (1 + 15) = 1024

   while the configured CUDA graph capture sizes top out at 512 (cudagraph_capture_sizes = [..., 512], from the engine config dump). Buffers (self.input_ids, self.positions, self._slot_mapping_buffer, self._context_positions_buffer) are sized to max_num_tokens / max_query_tokens, but the graph is captured for a padded token count ≤ 512. Under concurrency, the actual number of query tokens written by the DFlash input-setup kernel can exceed the shape the graph was captured for.

2. Unmasked/under-protected loads in the DFlash input-setup triton kernel. In vllm/v1/spec_decode/utils.py, copy_and_expand_dflash_inputs_kernel (the kernel launched from DFlashProposer.set_inputs_first_pass):

   # utils.py:526 — no mask on this load
   last_pos = tl.load(target_positions_ptr + valid_ctx_end - 1)

   where valid_ctx_end = ctx_end - num_rejected (when HAS_NUM_REJECTED). last_pos then feeds query_pos = last_pos + 1 + query_off → block_num = positions // block_size → slot = block_id * block_size + (positions % block_size). A poisoned last_pos corrupts the entire slot_mapping chain, and the subsequent attention uses those slots to index the KV cache → illegal address. The reject/padding bookkeeping (num_rejected_tokens = num_draft_tokens + 1 - valid_count, range 0–16) varies per request per step, so the out-of-bounds window is data-dependent and asynchronous — matching the observed non-deterministic timing.

3. Why CUDA graphs matter. Without graphs (--enforce-eager), each step re-launches kernels with the current token counts and the triton kernel re-evaluates bounds, so the OOB window is not hit (run Performance Optimization Suggestions for DFlash Decoding #2). With graphs, the captured graph replays a fixed control flow over dynamically-sized inputs; a stale/padded dimension from capture time no longer matches the live tensor extents, exposing the unmasked access.

Note on certainty: items 1–2 are reasoned from source + crash dumps, not from a captured synchronous stack trace (which CUDA_LAUNCH_BLOCKING=1 would have provided but also masks the crash). They are the most probable location; the proven facts are the four control experiments above isolating the bug to DFlash ∩ CUDA graphs.

Workarounds

• --enforce-eager (verified): disables CUDA graphs. DFlash still works (acceptance length 4–9 tokens observed) and the server is stable under 16–32 concurrency. Cost: loses the CUDA-graph decode speedup (~20–40% higher per-step latency), but keeps DFlash's 1.4–2.9× speculative speedup.
• Reduce --max-num-seqs to ≤ 32: lowers max_query_tokens toward the graph capture boundary; not yet verified whether this fully avoids the crash while keeping graphs.
• Disable DFlash (--speculative-config removed) entirely: stable, but loses all speculative speedup.

Suggested investigation for maintainers

1. Capture a synchronous stack trace of the offending kernel without masking the crash — e.g. via compute-sanitizer --tool memcheck on a worker, or a targeted CUDA_LAUNCH_BLOCKING=1 run with a single carefully shaped request that still trips the OOB (the async timing that masks it under blocking may not apply to a single-request boundary case).
2. Audit copy_and_expand_dflash_inputs_kernel (vllm/v1/spec_decode/utils.py:459) for unmasked tl.load/tl.store against dynamic num_ctx / num_rejected / valid_ctx_end bounds, especially the target_positions_ptr + valid_ctx_end - 1 load at line 526.
3. Reconcile DFlash buffer dimensions (max_query_tokens = max_batch_size * (1 + num_spec_tokens)) against cudagraph_capture_sizes, and ensure the input-setup kernel respects the padded token count the graph was actually captured for.

Attachments

Full crash logs and reproduction scripts are available locally:

• Crash log (32-conc): …/qwen36_35b_dflash_repro/logs/vllm_port18094.log.20260709_113854.bak
• Crash log (16-conc): …/qwen36_35b_dflash_repro/logs/vllm_port18094.log.20260709_121101.bak
• Repro launcher: …/qwen36_35b_dflash_repro/launch.sh
• Repro client: …/qwen36_35b_dflash_repro/repro_crash.py

(Paths are on the reporting cluster; can provide sanitized excerpts or the raw .bak logs on request.)