diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index 2e657bd68..eabbff259 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -55,6 +55,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#170](https://github.com/mudler/vllm.cpp/issues/170) | `ENG-RELEASE-BINARIES` | Publish container images to GHCR (cuda, vulkan, cpu) | feature | | [#322](https://github.com/mudler/vllm.cpp/issues/322) | `ENG-RELEASE-BINARIES` | Release handoff collides with tracked checkout `assets` directory | bug | | [#406](https://github.com/mudler/vllm.cpp/issues/406) | `ENG-TRAILER-MERGE-ARTIFACTS` | The trailer gate fails on how commits LAND: GitHub's Co-authored-by displaces the trailer block | bug | +| [#382](https://github.com/mudler/vllm.cpp/issues/382) | `KERNEL-ATTN-PAGED` | decode-opt attention kernel is head_dim-256 only; head_dim 128 (Qwen3-dense, Llama, Mistral) falls to the block kernel | perf | | [#206](https://github.com/mudler/vllm.cpp/issues/206) | `KERNEL-SSM-MAMBA` | RTX 5070 Ti: close Qwen3.5-4B TTFT, TPOT and VRAM gaps vs vLLM — owns the sm_120 post-conv token tile and the K=4 causal-conv arm (PR #155) | feature | | [#305](https://github.com/mudler/vllm.cpp/issues/305) | `KERNEL-SSM-MAMBA` | GDN causal-conv: the `conv_state` initial-state read races the final-state write across blocks (`VT_CONV_REG` + exact chunks, both default ON) | bug | | [#396](https://github.com/mudler/vllm.cpp/issues/396) | `KV-EXTERNAL-CACHE` | `test_lmcache_connector` data race under TSan: `MockLmcacheServer` writes non-atomic `listen_fd_` before joining its accept thread | bug | @@ -98,8 +99,8 @@ issue is not yet placed. Keyed record: update in place, never append. | [#242](https://github.com/mudler/vllm.cpp/issues/242) | — | `docs/FEATURES.md` drift: arch counts say 30 (registry has 35), multimodal-over-HTTP marked ☐ though W1-W3 landed | bug | | [#243](https://github.com/mudler/vllm.cpp/issues/243) | — | `vllm-feature-gap-analysis.md` is a stale 2026-07-28 snapshot: 9 of 16 HIGH/MED gaps have since landed | bug | | [#250](https://github.com/mudler/vllm.cpp/issues/250) | — | `a5b52047` reached main without a task branch, and `check-role-discipline` cannot be waived | bug | -| [#285](https://github.com/mudler/vllm.cpp/issues/285) | — | The operator lock refuses a second coordinator; it should only RECORD who is working where (spec `specs/operator-record.md`) | bug | | [#274](https://github.com/mudler/vllm.cpp/issues/274) | — | `main` is not verified by its own CI: the per-job `github.ref` concurrency groups cancel every long job on the next push, so the suite never completes (spec `specs/main-verifiability.md`) | bug | +| [#285](https://github.com/mudler/vllm.cpp/issues/285) | — | The operator lock refuses a second coordinator; it should only RECORD who is working where (spec `specs/operator-record.md`) | bug | | [#296](https://github.com/mudler/vllm.cpp/issues/296) | — | Two limitations recorded when #285 landed: a stale TTL comment, and a publish-NAME pin `os.rename` escapes (spec `specs/operator-record.md`, "Follow-up") | bug | | [#408](https://github.com/mudler/vllm.cpp/issues/408) | — | 12 of 54 `tests/scripts` suites are executed by nothing, and `check-test-registration.py`'s fixed `REQUIRED_TESTS` cannot see the class (found while repairing #274) | bug | diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index a83e0e8b5..e66695e24 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -17,6 +17,7 @@ VT_ROCM_GEMV VT_ROCM_HIPBLASLT VLLM_MM_TOWER_PROFILE VT_ARCH_TACTIC_STATS +VT_ATTN_DECODE_D128 VT_ATTN_DECODE_GQA VT_ATTN_DECODE_OPT VT_ATTN_FLASH2 diff --git a/src/vt/cuda/cuda_paged_attn.cu b/src/vt/cuda/cuda_paged_attn.cu index aaba9b45f..87da049ee 100644 --- a/src/vt/cuda/cuda_paged_attn.cu +++ b/src/vt/cuda/cuda_paged_attn.cu @@ -299,7 +299,41 @@ __device__ inline void LoadRow8(const float* p, int64_t base, int lane, float r[ } } -template +// EPL-generic row loaders, so the decode-opt kernel can serve a head_dim other +// than 256. EPL == 8 forwards to LoadRow8, so the d256 path stays BYTE-FOR-BYTE +// what it was (same transactions, same order, same accumulation). EPL == 4 adds +// head_dim 128 -- the width Qwen3-dense / Llama / Mistral actually use -- as one +// 64-bit transaction per lane for bf16 and one 128-bit for f32. 32 lanes * 4 +// elems == 128, so a warp still covers exactly one head-dim row. +template +__device__ inline void LoadRowN(const T* p, int64_t base, int lane, float* r); + +template <> +__device__ inline void LoadRowN<8, __nv_bfloat16>(const __nv_bfloat16* p, int64_t base, int lane, + float* r) { + LoadRow8(p, base, lane, r); +} +template <> +__device__ inline void LoadRowN<8, float>(const float* p, int64_t base, int lane, float* r) { + LoadRow8(p, base, lane, r); +} +template <> +__device__ inline void LoadRowN<4, __nv_bfloat16>(const __nv_bfloat16* p, int64_t base, int lane, + float* r) { + const uint2 w = reinterpret_cast(p + base)[lane]; + const __nv_bfloat16* h = reinterpret_cast(&w); +#pragma unroll + for (int i = 0; i < 4; ++i) r[i] = __bfloat162float(h[i]); +} +template <> +__device__ inline void LoadRowN<4, float>(const float* p, int64_t base, int lane, float* r) { + const uint4 a = reinterpret_cast(p + base)[lane]; + const float* fa = reinterpret_cast(&a); +#pragma unroll + for (int i = 0; i < 4; ++i) r[i] = fa[i]; +} + +template __global__ void PagedAttentionDecodeOptKernel(Tout* out, const TQ* query, const TKV* k_cache, const TKV* v_cache, const int32_t* block_table, const int32_t* seq_lens, @@ -338,23 +372,23 @@ __global__ void PagedAttentionDecodeOptKernel(Tout* out, const TQ* query, const const int64_t g = h / (hq / num_kv_heads); const int64_t qoff = (t * hq + h) * d; - float q_reg[kDecEpl]; - LoadRow8(query, qoff, lane, q_reg); // this lane's query slice, loaded once + float q_reg[EPL]; + LoadRowN(query, qoff, lane, q_reg); // this lane's query slice, loaded once float m = -CUDART_INF_F, l = 0.0f; - float o_reg[kDecEpl]; + float o_reg[EPL]; #pragma unroll - for (int i = 0; i < kDecEpl; ++i) o_reg[i] = 0.0f; + for (int i = 0; i < EPL; ++i) o_reg[i] = 0.0f; // Warp-strided scan over keys: warp-shuffle q.k, register online softmax. for (int64_t j = jmin + warp; j <= jmax; j += kDecWarps) { const int64_t blk = block_table[r * bt_row + (j / block_size) * bt_col]; const int64_t off = j % block_size; - float k_reg[kDecEpl]; - LoadRow8(k_cache, blk * kc_blk + off * kc_pg + g * kc_hd, lane, k_reg); + float k_reg[EPL]; + LoadRowN(k_cache, blk * kc_blk + off * kc_pg + g * kc_hd, lane, k_reg); float dot = 0.0f; #pragma unroll - for (int i = 0; i < kDecEpl; ++i) dot += q_reg[i] * k_reg[i]; + for (int i = 0; i < EPL; ++i) dot += q_reg[i] * k_reg[i]; #pragma unroll for (int o = 16; o > 0; o >>= 1) dot += __shfl_down_sync(0xffffffffu, dot, o); dot = __shfl_sync(0xffffffffu, dot, 0); // broadcast full-head score to lanes @@ -363,10 +397,10 @@ __global__ void PagedAttentionDecodeOptKernel(Tout* out, const TQ* query, const const float m_new = fmaxf(m, s); const float corr = expf(m - m_new); // 0 on the first key (m == -inf) const float pw = expf(s - m_new); - float v_reg[kDecEpl]; - LoadRow8(v_cache, blk * vc_blk + off * vc_pg + g * vc_hd, lane, v_reg); + float v_reg[EPL]; + LoadRowN(v_cache, blk * vc_blk + off * vc_pg + g * vc_hd, lane, v_reg); #pragma unroll - for (int i = 0; i < kDecEpl; ++i) o_reg[i] = o_reg[i] * corr + pw * v_reg[i]; + for (int i = 0; i < EPL; ++i) o_reg[i] = o_reg[i] * corr + pw * v_reg[i]; l = l * corr + pw; m = m_new; } @@ -377,7 +411,7 @@ __global__ void PagedAttentionDecodeOptKernel(Tout* out, const TQ* query, const float* m_sh = o_sh + kDecWarps * d; // [kDecWarps] float* l_sh = m_sh + kDecWarps; // [kDecWarps] #pragma unroll - for (int i = 0; i < kDecEpl; ++i) o_sh[warp * d + lane * kDecEpl + i] = o_reg[i]; + for (int i = 0; i < EPL; ++i) o_sh[warp * d + lane * EPL + i] = o_reg[i]; if (lane == 0) { m_sh[warp] = m; l_sh[warp] = l; @@ -389,19 +423,19 @@ __global__ void PagedAttentionDecodeOptKernel(Tout* out, const TQ* query, const #pragma unroll for (int w = 0; w < kDecWarps; ++w) gm = fmaxf(gm, m_sh[w]); float gl = 0.0f; - float acc[kDecEpl]; + float acc[EPL]; #pragma unroll - for (int i = 0; i < kDecEpl; ++i) acc[i] = 0.0f; + for (int i = 0; i < EPL; ++i) acc[i] = 0.0f; #pragma unroll for (int w = 0; w < kDecWarps; ++w) { const float sc = expf(m_sh[w] - gm); // empty warp: exp(-inf) == 0 gl += l_sh[w] * sc; #pragma unroll - for (int i = 0; i < kDecEpl; ++i) acc[i] += sc * o_sh[w * d + lane * kDecEpl + i]; + for (int i = 0; i < EPL; ++i) acc[i] += sc * o_sh[w * d + lane * EPL + i]; } const float inv = 1.0f / gl; #pragma unroll - for (int i = 0; i < kDecEpl; ++i) Store(out, qoff + lane * kDecEpl + i, acc[i] * inv); + for (int i = 0; i < EPL; ++i) Store(out, qoff + lane * EPL + i, acc[i] * inv); } } @@ -1966,6 +2000,27 @@ bool DecodeGqaEnabled() { return enabled; } +// head_dim-128 arm of the warp-split-KV decode kernel (the EPL=4 instantiation). +// Without it, `d == 128` -- the width Qwen3-dense / Llama / Mistral use -- misses +// the decode-opt kernel entirely on the `d == 32 * kDecEpl` (256) gate and falls +// to the generic block kernel. +// +// DEFAULT OFF, opt in with VT_ATTN_DECODE_D128=1. It is correctness-complete but +// NOT byte-exact against the block kernel it replaces: the two reduce the KV +// sequence in a different ORDER (warp-strided online softmax vs the block +// kernel's per-tile loop), so a greedy anchor can move at an exact bf16 tie. +// Shipping it OFF keeps every existing golden byte-identical, mirroring how the +// FA2 decode GQA group-swap landed (gated OFF, correctness-complete) before it +// was flipped ON against the full gate. The flip is a separate change and owes +// the near-tie razor + distributional gate + regen under the ratified-tie rule. +bool DecodeD128Enabled() { + static const bool enabled = [] { + const char* e = std::getenv("VT_ATTN_DECODE_D128"); + return e != nullptr && e[0] == '1'; + }(); + return enabled; +} + // --- Launchers ------------------------------------------------------------- template @@ -1975,6 +2030,27 @@ void LaunchDecode(cudaStream_t s, Tensor& out, const Tensor& query, const Tensor int64_t hq, int64_t d, int64_t num_reqs, int64_t num_kv_heads, int64_t block_size) { const dim3 grid(static_cast(num_tokens), static_cast(hq)); + // head_dim 128 (EPL=4), opt-in. No GQA carve-out: PagedAttentionDecodeGqaKernel + // sits inside the `d == 32 * kDecEpl` (256) branch below, so at head_dim 128 it + // can never run, and excluding qpk == kDecGqaQG here would strand exactly those + // models (e.g. Qwen3-32B, hq/num_kv_heads == 8) on the generic block kernel this + // arm exists to replace. + if (DecodeOptEnabled() && DecodeD128Enabled() && d == 32 * 4) { + const int block = kDecWarps * 32; + const size_t shmem = (static_cast(kDecWarps) * d + 2 * kDecWarps) * sizeof(float); + auto* k4 = args.window_size.has_value() + ? PagedAttentionDecodeOptKernel + : PagedAttentionDecodeOptKernel; + k4<<>>( + out.Ptr(), query.Ptr(), k_cache.Ptr(), v_cache.Ptr(), + block_table.Ptr(), seq_lens.Ptr(), query_start_loc.Ptr(), + num_reqs, hq, num_kv_heads, d, block_size, block_table.stride[0], block_table.stride[1], + k_cache.stride[0], k_cache.stride[1], k_cache.stride[2], v_cache.stride[0], + v_cache.stride[1], v_cache.stride[2], args.scale, args.logits_soft_cap, args.causal, + WindowLeft(args), WindowRight(args)); + Check(cudaGetLastError(), "paged_attention decode-opt-d128 launch"); + return; + } if (DecodeOptEnabled() && d == 32 * kDecEpl) { // GQA-fused decode (grid.y = num_kv_heads): stage each KV head's K/V ONCE and // attend all QG q-heads of that head, killing the opt kernel's 8x redundant