PR #1 - dusterbloom/lucebox-hub

Targets feat/prefix-cache (PR Luce-Org/lucebox-hub#59). Drafted for review.

Five commits, fairly cleanly split: bench loader → bench tooling fixes → repro script → upstream cuda fix.

1. Faithful transcript replay (`bench_agent_loop.py`)

extract_user_turns reads only type=user records where content is a str. Real Claude Code transcripts have user records of tool_result blocks with content as a list — silently skipped. The bench replays only typed human prompts with bench-synthesised assistant replies, with no tool I/O.

The replacement loader (_load_transcript + _to_openai_messages) coalesces consecutive same-role records into single turns, converts Anthropic blocks → OpenAI structured messages (tool_use → assistant.tool_calls, tool_result → role=tool, thinking dropped), and at each assistant index sends the exact prefix that was sent at that point. Validated against a 65-record session: old loader extracts 7 typed-user turns; new loader walks all 32 call points with prefix at the last call = 95,536 chars. Targets dflash/scripts/server_tools.py (the production agent endpoint).

Also: streaming SSE for separate TTFT measurement, and a token-count fallback that counts content/reasoning/tool deltas when usage.completion_tokens isn't emitted (PR #59's server doesn't honour stream_options.include_usage).

2. Standalone empty-response reproducer (`repro_empty_response.py`)

Self-contained, no transcript dependency. Generates 6 growing-prefix calls (130 → 5K → 10K → 15K → 20K → 26K chars) with synthetic pylint-style tool results, runs against slots=0 (control) then slots=2 (repro), prints a side-by-side table.

Pre-fix (slots=2): warm call 2 returns content=0 / comp_tok=0 / finish=stop after 9.7s of real prefill; calls 3-6 return empty in <50ms (cache "hits" on a corrupted slot). Reproduces deterministically.

3. Root cause + upstream fix (`dflash/deps/llama.cpp`)

The bug is in upstream ggml-cuda's VMM allocator, not in PR #59 itself. PR #59 just reliably triggers it because its inline-snap pattern interleaves compute with allocations on the same backend.

Mechanism:

The VMM pool's alloc() path extends the pool via cuMemCreate + cuMemMap + cuMemSetAccess when an allocation can't fit in the existing mapped region.
cuMemSetAccess races with previously-queued async work and returns CUDA_ERROR_NOT_READY.
CU_CHECK fires GGML_ABORT, but in some process configurations the abort doesn't terminate (gdb-fork backtrace path interferes, SIGABRT swallowed somewhere — exact mechanism unclear, but reproducible).
The mapping is registered but access permissions are never set.
Subsequent reads/writes silently return zeroed/garbled data — for the prefix cache, snapshots get stored into the broken region and restore as garbage, making the model emit 0 tokens.

Fix is one cudaDeviceSynchronize() before the cuMem* sequence inside the pool extension branch:

if (size > avail) {
    // ... compute reserve_size ...
    GGML_ASSERT(pool_size + reserve_size <= CUDA_POOL_VMM_MAX_SIZE);

    CUDA_CHECK(cudaDeviceSynchronize());   // ← new

    // allocate more physical memory
    CU_CHECK(cuMemCreate(...));
    ...
}

Only fires when the pool actually needs to grow, so steady-state hot-path allocations are unaffected.

llama.cpp PR: Luce-Org/llama.cpp-dflash-ggml#4

This PR's submodule URL is temporarily pointed at easel's fork (fix/cuda-vmm-pool-extension-race branch) until that PR merges; revert to Luce-Org/llama.cpp.git@luce-dflash after merge.

Also bumps prefix_cache.startup_sync's await_reply timeout 10s → 60s for daemons with multi-slot snap pools at large --max-ctx.

Validation

repro_empty_response.py exits 0 (repro confirmed) pre-fix, exits 1 ("DID NOT REPRO") post-fix.

4-session sweep on RTX 3090 Ti, CUDA 13.2, Qwen 3.6 27B, --max-ctx=24576 --warm-slots=2:

session	calls	TTFT cold	TTFT warm	TTFT ×	wall ×	empties
lucebox-hub (≤11K char prefix)	10	40.0s	39.9s	1.00×	0.77×	0/10
nexiq-small (≤25K)	6	55.9s	44.9s	1.24×	0.75×	0/6
axon-med (≤39K)	10	133.8s	51.7s	2.59×	1.27×	0/10
helix-large (≤71K)	10	242.1s	97.3s	2.49×	1.52×	0/10

Zero empty responses across 36 warm calls. Cache wins scale with prefix size — 2.49–2.59× TTFT and 1.27–1.52× wall on real long-prefix agentic sessions is the headline win this PR (and PR #59) set out to validate. Short sessions barely break even, as expected (cache-creation overhead vs. limited reuse).

Test plan

repro_empty_response.py — pre-fix repros, post-fix doesn't
4-session sweep — 36/36 warm calls real
Build: clean nvcc, no warnings
Cross-check on Blackwell / sm_120 — only validated on RTX 3090 Ti / sm_86
Confirm test_multi_turn_prefix_cache.py still passes — should, since it's checking turn-2-5 ratios on a small fixture that probably didn't trigger the bug anyway

Status

Draft. Two PRs need review:

This PR (lucebox-hub): bench tooling + repro + submodule bump.
llama.cpp PR (Luce-Org/llama.cpp-dflash-ggml#4): the actual fix.

Happy to split this PR if reviewer prefers (bench-only vs submodule-bump), but the bench changes don't make sense without the underlying fix being available.

🤖 Generated with Claude Code

#1 bench: faithful Claude Code transcript replay + repro & build-flag fix for empty-response bug

1. Faithful transcript replay (bench_agent_loop.py)

2. Standalone empty-response reproducer (repro_empty_response.py)