#4 fix(cuda): sync device before cuMem pool extension to fix VMM allocator race

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 can race with previously-queued async work on the default stream and return CUDA_ERROR_NOT_READY ("device not ready"). The CU_CHECK macro then fires GGML_ABORT, but in some process configurations the abort doesn't actually terminate (the gdb-fork backtrace path interferes, the SIGABRT gets swallowed somewhere, etc.) and execution continues. The mapping is registered but access permissions were never set, so every subsequent read/write into that region silently misbehaves — typically returning zeroed or stale data without surfacing any error.

Visible failure mode: workloads that interleave compute with ggml_backend_alloc_ctx_tensors on the same backend. Reproduced deterministically by dflash's prefix-cache path (Luce-Org/lucebox-hub#59), where snapshot copies fire during prefill, then a follow-up alloc triggers pool extension, and subsequent cached-state restores return zeroed KV → decode emits 0 tokens with finish_reason=stop.

With this sync added, the same workload runs cleanly with VMM enabled and delivers 2.5× TTFT / 1.5× wall speedup on long agentic prefixes (RTX 3090 Ti, CUDA 13.2, Qwen 3.6 27B).

Sync placement

Just before the cuMem* calls 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(...));
    ...
}

The sync only fires when the pool actually needs to grow, so steady-state hot-path allocations (where pool_size already covers the request) are unaffected. Added cost is one cudaDeviceSynchronize per pool extension event — negligible on any real workload (pool extensions are infrequent by design — the pool is sized in granularity-rounded chunks and rarely shrinks).

Validation

Bench in dflash with VMM enabled, prefix cache, slots=2, max-ctx=24576:

session	turns	TTFT cold	TTFT warm	TTFT ×	wall ×	empties
short (~10K chars)	10	40.0s	39.9s	1.00×	0.77×	0/10
medium (~25K)	6	55.9s	44.9s	1.24×	0.75×	0/6
long (~38K)	10	133.8s	51.7s	2.59×	1.27×	0/10
huge (~70K)	10	242.1s	97.3s	2.49×	1.52×	0/10

Pre-fix: workloads at multi-slot prefix cache produce empty bodies on every cache hit after call 2 (deterministic). Post-fix: 0/36 empties.

Test plan

• dflash prefix-cache repro reproduces empty bodies pre-fix; produces real content post-fix
• No regressions in standard llama.cpp tests (caveat: only ran the dflash workload — full ggml-cuda test suite hasn't been run against this patch)
• Test on Blackwell / sm_120 — only validated on RTX 3090 Ti / sm_86 so far. The race is platform-agnostic, but the timing window may differ.

🤖 Generated with Claude Code