PagedAttention: CUDA illegal memory access on a full-context request when max_batch_size>=2 and max_scheduled_tokens=2048

@tianleiwu is already working on this.

Since Sep 17, 2026.

• #2587 by @copilot-swe-agent — open

### Describe the bug

With PagedAttention, a request at (or near) the model's full context aborts the process with `CUDA failure 700: an illegal memory access was encountered` when **all three** of the following hold:

1. `engine.dynamic_batching.max_batch_size >= 2`
2. `engine.dynamic_batching.max_scheduled_tokens` / `search.chunk_size` = **2048** (1024 and 512 are fine)
3. The long request is issued **after** an earlier, shorter request in the same `Engine`

Any one of the three removed makes it pass. It is not an out-of-memory condition — peak usage is ~30 GB on a 143 GB H200, and the same configuration completes when the long request is the *only* request.

Note that in the repro the requests are issued **strictly sequentially** — one request is created, run to completion, and closed before the next begins — so `max_batch_size=2` never actually batches two requests. It only changes how the engine sizes its buffers. That, plus the fact that draft depth is irrelevant (below), suggests a capacity/stride computation that is sized from `max_batch_size` and `max_scheduled_tokens` but indexed from the live sequence length.

### Condition matrix

All rows: `num_blocks=1152`, `paged_block_size=256`, prompts issued sequentially as 32,768 then 262,112, `search.max_length=262144`.

| `max_batch_size` | `max_scheduled_tokens` = `chunk_size` | `max_draft_tokens` | prompts | result |
| ---: | ---: | ---: | --- | --- |
| 1 | 2048 | 7 | 32,768 then 262,112 | OK (peak 29,100 MiB) |
| 2 | 512 | 7 | 32,768 then 262,112 | OK (peak 27,884 MiB) |
| 2 | 1024 | 7 | 32,768 then 262,112 | OK (peak 27,948 MiB) |
| 2 | 2048 | 7 | **262,112 only** | OK (peak 30,122 MiB) |
| 2 | 2048 | 7 | 32,768 then 262,112 | **CUDA 700** |
| 2 | 2048 | **1** | 32,768 then 262,112 | **CUDA 700** |

So:

- **Draft depth is not the driver** — it fails identically at `max_draft_tokens=1` and `7`. (`0` is rejected by config validation, so a drafter-free comparison was not possible with this model.)
- **The chunk threshold sits between 1024 and 2048.**
- **A prior request is required** — the identical engine serves the 262,112-token request on its own.

### Error

```
[E:onnxruntime:, cuda_call.cc:148 CudaCall] CUDA failure 700: an illegal memory access was encountered ;
GPU=0 ; file=onnxruntime/core/providers/cuda/cuda_execution_provider.cc ; line=534 ;
expr=cudaStreamSynchronize(static_cast(stream_));

[E:onnxruntime:, cuda_call.cc:148 CudaCall] CUDA failure 700: an illegal memory access was encountered ;
GPU=0 ; file=onnxruntime/core/providers/cuda/cuda_stream_handle.cc ; line=182 ;
expr=cudaStreamSynchronize(static_cast(GetHandle()));

terminate called after throwing an instance of 'onnxruntime::OnnxRuntimeException'
what(): ... CUDA failure 700: an illegal memory access was encountered ;
file=onnxruntime/core/providers/cuda/cuda_allocator.cc ; line=98 ; expr=cudaFreeHost(p);
Aborted (core dumped)
```

The first failing sync is the giveaway that the fault is in a kernel launched during the long request, not in allocation; `cudaFreeHost` merely inherits the poisoned context during teardown.

### To reproduce

The model is Qwen3.8-27B exported by Model Builder with PagedAttention and a DFlash 2 block drafter:

```
-p int4 -e cuda --extra_options
block_size=32 op_types_to_quantize=MatMul/Gather matmulnbits_weights_prepacked=1
use_paged_attention=true paged_block_size=256 enable_cuda_graph=true
use_device_allocator_for_initializers=true
kv_cache_quant_scheme=int4_per_channel
dflash2_path= dflash2_num_draft_tokens=7 max_draft_tokens=7 dflash2_precision=int4
aux_hidden_state_layers=6,20,34,48,62 state_update_capacity=7
```

`search.max_length` / `context_length` are 262,144.

```python
import json
import numpy as np
import onnxruntime_genai as og

MODEL = ""
NUM_BLOCKS = 1152
CHUNK = 2048 # 1024 or 512 -> no crash
MAX_BATCH_SIZE = 2 # 1 -> no crash
PROMPTS = [32768, 262112] # drop the first entry -> no crash

config = og.Config(MODEL)
config.clear_providers()
config.append_provider("cuda")
config.overlay(
json.dumps(
{
"search": {"chunk_size": CHUNK},
"engine": {
"dynamic_batching": {
"max_scheduled_tokens": CHUNK,
"max_batch_size": MAX_BATCH_SIZE,
"num_blocks": NUM_BLOCKS,
}
},
"speculative": {"max_draft_tokens": 7},
}
)
)
model = og.Model(config)
tokenizer = og.Tokenizer(model)
engine = og.Engine(model)

unit = list(tokenizer.encode("The quick brown fox jumps over the lazy dog. "))
for n in PROMPTS:
tokens = (unit * (n // len(unit) + 2))[:n]
options = og.RequestOptions()
options.set_max_session_tokens(min(n + 40, 262144))
request = engine.create_request(options=options)
turn = og.TurnOptions(request)
turn.set_do_sample(False)
turn.set_max_generated_tokens(32)
request.begin_turn(np.asarray(tokens, dtype=np.int32), turn)
buffer = engine.create_event_buffer(16)
generated = 0
while engine.has_pending_requests():
for event in engine.run(buffer):
if event.flags & og.EngineEventFlags.TOKEN:
generated += 1
request.close()
print(f"prompt={n} generated={generated}")
```

Expected: both prompts complete. Actual: `prompt=32768 generated=32` prints, then the process aborts during the 262,112-token request.

Unrelated aside that cost some time while narrowing this: a prompt of exactly `search.max_length` is always rejected because the session must leave room for at least one generated token, which is why the repro uses 262,112 rather than 262,144.

### Urgency

Not blocking us — we ship `max_batch_size=1` for this configuration and take the single-stream limitation. Filing because it is a hard process abort rather than a graceful error, and the three-way condition makes it easy to hit accidentally when tuning throughput knobs.

### Platform

Linux

### OS Version

Ubuntu 24.04

### ONNX Runtime GenAI Version

0.16.0-dev (built from source)

### ONNX Runtime Version

1.31.0 (built from source)

### Execution Provider

CUDA

### Hardware

NVIDIA H200 (SM90), CUDA 13.0, driver 580.159.04