opencoti patch 0102 — streaming-FA bulk-H2D tail lift + persistent staging scratch ring (bug-2116, bug-2121) Two changes to the rolling-KV streaming-FA lift in fattn.cu: 1. bug-2116 (perf): for a HOST quant tail source at decode (n_b==1), bulk-copy the RAW QUANT bytes H2D once at full DMA rate into device scratch, then run the nc-converter from device memory. The converter reading the pinned-host tail via UVA at word granularity measured 10-16 GB/s vs 57-65 GB/s bulk on bs2 — the UVA cliff behind the S0 spill-decode curve. Measured after-curve (14B-1M q8q4 window 131k, decode tps): 800MiB 9.05->12.23, 1600 3.96->10.28, 6400 ~1->5.53. 2. bug-2121 (correctness): both bulk paths (the new quant one and the #586 f16 one) staged through ggml_cuda_pool_alloc scratch FREED on scope exit while the async memcpy+convert on the copy stream were still pending. Decode (bug-1841 single stream) is stream-ordered and safe; PREFILL runs two streams and the FA compute allocates from the SAME pool concurrently (MMA stream-k fixup buffers, every tile) -> freed-but-pending scratch re-handed cross-stream -> corrupted K/V tiles. Manifested with >=2 host-tail tiles (>32768 exercised tail cells): garbage decode / needle-LOST + racy host GPFs; <=1 tail tile clean. Fix: persistent per-slot scratch ring (g_scrK/g_scrV[4], grow-only, worst-case tile_kv_full capacity so no cudaMalloc inside captured decode graphs) — no pool sharing across streams. Gate (bs2, 14B-1M q8q4 window 131k): tails 1600/3200/6400 all needle=PRESENT, coherent text, zero crashes across three full 106k prefills (3200 previously crashed 2x + garbage decode; 6400 previously needle-LOST). --- a/llama.cpp/ggml/src/ggml-cuda/fattn.cu +++ b/llama.cpp/ggml/src/ggml-cuda/fattn.cu @@ -1350,6 +1350,30 @@ char * const slotK_base = g_slotK; char * const slotV_base = g_slotV; + // opencoti-hook: f5-rolling-kv bug-2121 — PERSISTENT per-slot H2D staging scratch. + // The bulk-H2D lifts below (bug-2116 quant path + #586 f16 path) used + // ggml_cuda_pool_alloc scratch FREED on lambda/loop exit while the + // cudaMemcpyAsync + converter on the copy stream were still pending. On decode + // (single stream, bug-1841) reuse is stream-ordered and safe; on PREFILL the FA + // compute on ctx.stream() allocates from the SAME pool concurrently (MMA + // stream-k fixup buffers, every tile), so a freed-but-pending scratch gets + // re-handed to the other stream → cross-stream aliasing → corrupted K/V tiles. + // Manifested only with ≥2 host-tail tiles interleaving with compute (>32768 + // exercised tail cells): 14B-1M q8q4 window-mode tail=3200/6400 MiB → garbage + // decode / needle-LOST (+ racy host GPFs); tail≤1600 (1 tail tile) clean. A + // persistent ring keyed by slot removes pool sharing entirely; within-ring WAR + // is stream-ordered (memcpy → convert → next lift's memcpy, all on cs). + static char * g_scrK[4] = {}; static size_t g_scrK_cap[4] = {}; + static char * g_scrV[4] = {}; static size_t g_scrV_cap[4] = {}; + auto scr_get = [](char ** ring, size_t * cap, int slot, size_t need) -> char * { + if (cap[slot] < need) { + if (ring[slot]) CUDA_CHECK(cudaFree(ring[slot])); // cudaFree syncs pending users + CUDA_CHECK(cudaMalloc((void **)&ring[slot], need)); + cap[slot] = need; + } + return ring[slot]; + }; + // R2-b S3b (#312): with >=2 slots the op runs a double-buffer ping-pong — the // per-tile lift is issued on a dedicated copy_stream while FA/lse for the // current tile run on ctx.stream(), so tile t+1's DMA overlaps tile t's @@ -1445,10 +1469,34 @@ // below are unchanged. Unifies device + pinned-host (S3c) sources // (the kernel reads host via UVA). Slot is f16 → lse half-cast OK. // 3c-5: tK/tV + loc select the window (device) or tail (host) source. - to_fp16_k((const char *)tK->data + (size_t)loc*tK->nb[1], (half *)sK, + // opencoti-hook: f5-rolling-kv bug-2116 — for a HOST quant source, + // bulk-H2D the RAW QUANT bytes (cell-major block, 2.5× fewer bytes + // than f16) into device scratch at full DMA rate, then convert from + // device memory. Calling the converter on the host pointer made it + // read the tail via UVA at word granularity (~10–16 GB/s measured + // vs 57–65 GB/s bulk on bs2). n_b>1 keeps the UVA path (nb[3] + // batch blocks are not spanned by the cell-major copy; unused in + // the spill regime where ne[3]==1). + const bool bulkq = tr.host && n_b == 1; + const char * qsrcK = (const char *)tK->data + (size_t)loc*tK->nb[1]; + const char * qsrcV = (const char *)tV->data + (size_t)loc*tV->nb[1]; + if (bulkq) { + // bug-2121: persistent per-slot scratch — NOT pool memory (see scr_get above). + // Capacity = worst-case tile (tile_kv_full), not this_kv: this_kv grows per + // decode token, and a realloc inside a captured decode graph would fail. + const size_t kqblk = (size_t)this_kv * tK->nb[1]; + const size_t vqblk = (size_t)this_kv * tV->nb[1]; + char * qscrK = scr_get(g_scrK, g_scrK_cap, slot, (size_t)tile_kv_full * tK->nb[1]); + char * qscrV = scr_get(g_scrV, g_scrV_cap, slot, (size_t)tile_kv_full * tV->nb[1]); + CUDA_CHECK(cudaMemcpyAsync(qscrK, qsrcK, kqblk, cudaMemcpyDefault, cs)); + CUDA_CHECK(cudaMemcpyAsync(qscrV, qsrcV, vqblk, cudaMemcpyDefault, cs)); + qsrcK = qscrK; + qsrcV = qscrV; + } + to_fp16_k(qsrcK, (half *)sK, head_dim, this_kv, n_head_kv, n_b, (int64_t)(tK->nb[1]/ts_k), (int64_t)(tK->nb[2]/ts_k), (int64_t)(tK->nb[3]/ts_k), cs); - to_fp16_v((const char *)tV->data + (size_t)loc*tV->nb[1], (half *)sV, + to_fp16_v(qsrcV, (half *)sV, dv, this_kv, n_head_kv, n_b, (int64_t)(tV->nb[1]/ts_v), (int64_t)(tV->nb[2]/ts_v), (int64_t)(tV->nb[3]/ts_v), cs); } else @@ -1468,17 +1516,20 @@ const bool bulk = use_2d && tr.host; const size_t kblk = (size_t)this_kv * tK->nb[1]; // contiguous cell-major span const size_t vblk = (size_t)this_kv * tV->nb[1]; - ggml_cuda_pool_alloc scrK(pool, bulk ? kblk : 1); - ggml_cuda_pool_alloc scrV(pool, bulk ? vblk : 1); const char * baseK; const char * baseV; if (bulk) { + // bug-2121: persistent per-slot scratch — NOT pool memory (see scr_get above). + // Sequential bb iterations reuse it WAR-safely (copy/repack both on cs). + // Worst-case tile capacity: see the quant path above (capture-safe growth). + char * scrK = scr_get(g_scrK, g_scrK_cap, slot, (size_t)tile_kv_full * tK->nb[1]); + char * scrV = scr_get(g_scrV, g_scrV_cap, slot, (size_t)tile_kv_full * tV->nb[1]); const char * hK = (const char *)tK->data + (size_t)bb*tK->nb[3] + (size_t)loc*tK->nb[1]; const char * hV = (const char *)tV->data + (size_t)bb*tV->nb[3] + (size_t)loc*tV->nb[1]; - CUDA_CHECK(cudaMemcpyAsync(scrK.ptr, hK, kblk, cudaMemcpyDefault, cs)); // one full-BW H2D - CUDA_CHECK(cudaMemcpyAsync(scrV.ptr, hV, vblk, cudaMemcpyDefault, cs)); - baseK = scrK.ptr; // scratch is cell-major, cell `loc` at offset 0 - baseV = scrV.ptr; + CUDA_CHECK(cudaMemcpyAsync(scrK, hK, kblk, cudaMemcpyDefault, cs)); // one full-BW H2D + CUDA_CHECK(cudaMemcpyAsync(scrV, hV, vblk, cudaMemcpyDefault, cs)); + baseK = scrK; // scratch is cell-major, cell `loc` at offset 0 + baseV = scrV; } else { baseK = (const char *)tK->data + (size_t)bb*tK->nb[3] + (size_t)loc*tK->nb[1]; baseV = (const char *)tV->data + (size_t)bb*tV->nb[3] + (size_t)loc*tV->nb[1]; @@ -1677,10 +1728,26 @@ if (!all_f16) { // S3d: one nc-converter launch per tile → packed f16 slot (see // the overlap path for the layout proof). Source strides s=nb/ts. - to_fp16_k((const char *)K->data + (size_t)kv0*K->nb[1], (half *)sK, + // bug-2116: host quant source → bulk-H2D the raw quant block at + // DMA rate + convert from device scratch; direct converter-on-host + // reads via UVA at word granularity (see the overlap-path note). + const bool bulkq = host_src && n_b == 1; + const size_t kqblk = bulkq ? (size_t)this_kv * K->nb[1] : 1; + const size_t vqblk = bulkq ? (size_t)this_kv * V->nb[1] : 1; + ggml_cuda_pool_alloc qscrK(pool, kqblk); + ggml_cuda_pool_alloc qscrV(pool, vqblk); + const char * qsrcK = (const char *)K->data + (size_t)kv0*K->nb[1]; + const char * qsrcV = (const char *)V->data + (size_t)kv0*V->nb[1]; + if (bulkq) { + CUDA_CHECK(cudaMemcpyAsync(qscrK.ptr, qsrcK, kqblk, cudaMemcpyDefault, stream)); + CUDA_CHECK(cudaMemcpyAsync(qscrV.ptr, qsrcV, vqblk, cudaMemcpyDefault, stream)); + qsrcK = qscrK.ptr; + qsrcV = qscrV.ptr; + } + to_fp16_k(qsrcK, (half *)sK, head_dim, this_kv, n_head_kv, n_b, (int64_t)(K->nb[1]/ts_k), (int64_t)(K->nb[2]/ts_k), (int64_t)(K->nb[3]/ts_k), stream); - to_fp16_v((const char *)V->data + (size_t)kv0*V->nb[1], (half *)sV, + to_fp16_v(qsrcV, (half *)sV, dv, this_kv, n_head_kv, n_b, (int64_t)(V->nb[1]/ts_v), (int64_t)(V->nb[2]/ts_v), (int64_t)(V->nb[3]/ts_v), stream); } else