# M3.9 — Split-plan 10× tax investigation + fix

> Sequel to [`67f9bcb`](../experience/wins/2026-05-07-prefill-batching-confirmed.md)
> + [`4a3612b`](../experience/wins/2026-05-07-split-plan-10x-tax-confirmed.md).
> Long-ctx TTFT bottleneck is split-plan (mixed prefill+decode)
> per-token cost 10× pure prefill. M3.9 investigates the tax
> source and lands the smallest correct fix.

## P0 finding (source-code survey)

`decode_batch_with_prefill` at `batch_decode.rs:890` IS a **unified
mixed-batch path**:
- Combines decode tokens + prefill tokens into one `combined_tokens`
  array (line 955-959)
- Single H2D upload (line 960-963)
- Single `mixed_buffers` setup with shared metadata
- Single `forward_mixed_batch` call from scheduler

So the "tax" is NOT "two kernel chains launched separately within
one step". The architecture is already unified.

But **the path can fall back**:
`decode_batch_with_prefill` returns `Ok(false)` when:
- LoRA enabled
- KV format not BF16/FP8/INT8
- token count mismatch
- prefill slot already in decode_slot_indices
- prefill slot already in prefill_slot_indices
- `paged_kv_pool.seq_len(prefill.slot_idx) != start_pos`

When it returns `Ok(false)`, scheduler `step_mixed_launch`
(decode.rs:570-583) falls back to `step_prefill_batch(&launch_candidates)`
+ `launch_decode_batch_from_tokens(...)` — **TWO separate launches**.
This is the actual "two-kernel-chain" tax case.

## Three candidate hypotheses for 10× per-token tax

### H1 — Fallback path firing (true two-launch tax)

The unified `decode_batch_with_prefill` returns `Ok(false)` at
long-ctx steady state because of the `seq_len != start_pos` check
or some other condition. Scheduler falls back to two sequential
launches.

**Test**: Add a single counter/log line at each return site of
`decode_batch_with_prefill`. Re-bench longctx 4k/c=4. If the
counter shows mostly Ok(false), fix the falling-back condition.

### H2 — Attention KV scaling

At split steady state, each row's attention reads its accumulated
KV (2-4k tokens). At pure prefill steady state, KV starts empty.
The attention kernel's per-token cost is O(KV_len) — so once KV
fills up, the per-token cost is fundamentally larger.

**Test**: Compute expected per-token cost from pure-prefill data
+ KV scaling. If 154 µs/tok at split is in the ballpark of
"15.4 µs/tok × (1 + KV_len/chunk_size)", H2 is the explanation.
Cannot be "fixed" — it's inherent attention scaling.

### H3 — Varlen kernel inefficiency (mixed qlen tile shapes)

The HD128 prefill kernel handles uniform qlen=N tiles efficiently.
When mixed (qlen=1 for decode rows + qlen=N for prefill chunks),
the kernel must handle irregular tile sizes — possibly with
serialization or larger overhead per tile boundary.

**Test**: Compare kernel time with `nsys stats --report
cuda_gpu_kern_sum` filtered to the mixed-batch attention kernel.
If the kernel itself is much slower at mixed shape vs uniform,
H3 is supported.

## Phase 0 — instrument + classify (1-2 hr)

Tasks:

| # | Task | File | LOC |
|---|---|---|---|
| 0.1 | Add per-call counters: `decode_batch_with_prefill_ok_true_count`, `..._ok_false_count` exposed via /v1/stats | `infer/src/model/qwen3/batch_decode.rs` + metrics path | ~30 |
| 0.2 | Re-run longctx 4k/c=4 bench; check counter ratio | bench | 0 |
| 0.3 | Compute expected H2 attention scaling from pure-prefill data | analysis | 0 |
| 0.4 | nsys trace longctx 4k/c=4 with --delay 30 (works at this concurrency per c=8 trace experience); inspect per-kernel breakdown | trace + nsys stats | 0 |

After Phase 0, ONE of H1/H2/H3 is confirmed. Pick fix path:

### If H1 confirmed (fallback firing) — Phase 1A

Find the `Ok(false)` condition that's firing in steady state and
fix it. Likely candidate: the `seq_len != start_pos` check is too
strict for chunked prefill where multiple in-flight chunks may
have raced. ~30-60 LOC fix in `decode_batch_with_prefill` + e2e
test for byte-exact equivalence.

Expected gain: 10× per-token at split → ~1.5-2× (matching pure
prefill modulo H2 attention scaling).

### If H2 confirmed (attention scaling) — Phase 1B (deferred)

H2 is fundamental and not fixable. The "fix" is to delay split
entry by staying in pure-prefill mode until the prefill queue
empties OR to use chunked prefill with smaller chunks (so split
fires later when KV is still small per row).

This is a scheduler policy change in `step_mixed_launch` admission
heuristic. ~50 LOC plus extensive e2e testing.

Expected gain: TTFT improves because pure prefill is 10× faster;
ITL gets slightly worse (decode delayed by 1 step or so).

### If H3 confirmed (varlen kernel inefficiency) — Phase 1C (kernel)

The HD128 prefill kernel needs better varlen handling. TileLang
HD128 prefill template (`batch_prefill_paged_hd128.py`) already
supports varlen via `q_indptr`. The issue may be the kernel grid
configuration (BLOCK_M=64, BLOCK_N=16) is suboptimal for mixed
qlen=1+qlen=2048 batches.

Tunable via `tilelang.tune` autotune. ~50 LOC change in
.py template + bench validation.

Expected gain: 10× → 3-5× (kernel improvement, but H2 attention
scaling remains).

## Acceptance

- Phase 0 produces a definitive diagnosis (one of H1/H2/H3 selected).
- Phase 1 implementation lands; bench longctx 4k/c=4 shows TTFT
  reduction.
  - If H1: target ≥ 50% TTFT improvement (3403 → ≤ 1700 ms).
  - If H2: target ≥ 30% TTFT improvement via delayed split (3403 → ≤ 2400 ms).
  - If H3: target ≥ 30% TTFT improvement via kernel tuning.
- e2e + greedy_consistency tests pass.

## Out of scope

- Changing chunk size (currently 2048). Defer to a separate
  workload-sweep plan.
- F4-Big (already past parity for high-conc; not relevant to TTFT
  axis).
- M_b.2 (FP8 prefill TileLang) — kernel speed at pure prefill
  already 19× vLLM per-token (`67f9bcb`); no kernel work justified
  for prefill axis until split tax is closed.

## Cross-references

- Discovery + Phase 1 trace: `2026-05-07-m3.6-phase1-nsys-arle-s48-highconc.md`
- Cross-request batching confirmed: `2026-05-07-prefill-batching-confirmed.md`
- Split-plan 10× tax confirmation: `2026-05-07-split-plan-10x-tax-confirmed.md`
- Source: `infer/src/model/qwen3/batch_decode.rs:890` (decode_batch_with_prefill)
- Fallback site: `infer/src/scheduler/cuda/decode.rs:570-583`
- vLLM v1 prefill priority reference:
  https://github.com/vllm-project/vllm/blob/main/vllm/v1/core/scheduler.py

## Rule

- **Architecture survey before instrumentation, instrumentation before fix**.
  M3.9 followed the rule (see this plan's P0). M3.8 v1 violated the rule
  (planned a fix without first verifying the architecture's actual state).
- **One bench experiment can disprove a 156-line plan**. Embrace
  cheap experiments over expensive plans when feasible.
