diff --git a/.agents/issue-index.md b/.agents/issue-index.md index 1ffc08641..7864093ad 100644 --- a/.agents/issue-index.md +++ b/.agents/issue-index.md @@ -491,6 +491,7 @@ rather than merged. `scripts/check-agent-record.py` gates both. | [#1447](https://github.com/mudler/vllm.cpp/issues/1447) | `ENG-EXPERT-STREAM` | **`docs/USAGE.md` said EVERY Qwen3.8 decode figure came from the W0e C ABI harness; the 66.7 s/token streaming-off row of 16 August 2026 came from `vllm-server`, the same binary the section tells the reader to run.** Introduced by #1211 and fixed in the same flow: the sentence is scoped to the W0e and W0f runs and the exception is named, in `docs/USAGE.md` and in the spec paragraph that mirrors it. Provenance read at the source, `.agents/specs/expert-streaming.md:837` (server entry point) and `:905`. | record | | [#1403](https://github.com/mudler/vllm.cpp/issues/1403) | `ENG-CUDAGRAPH-BREAK` | **`test_qwen3_5_decode_graph_seam` SIGSEGVs — a DUPLICATE of [#1394](https://github.com/mudler/vllm.cpp/issues/1394) that was already fixed when it was filed, so no code change was owed.** Filed against `96ed8346f`, which PREDATES the fix `7dec1d990` (PR #1393, merged 2026-08-20 00:49); `git merge-base --is-ancestor` puts the fix outside the reported tree. Re-measured at `b537a5344`, CPU-only Release x86_64: the test passes 5 runs of 5, `8 cases`/`8 passed`, `assertions: 138`, `Status: SUCCESS!`, exit 0. The cause was re-derived by single-variable mutation rather than inherited from #1394's report, and NAMED by a sanitizer rather than reasoned. Debug + `VLLM_CPP_SANITIZE=address,undefined` (NDEBUG OFF) with BOTH halves of `7dec1d990` reverted to their `96ed8346f` bytes: `AddressSanitizer: SEGV on unknown address 0x5045f5f84900`, `The signal is caused by a READ memory access`, at `src/vt/cpu/cpu_paged_attn.cpp:59` in `KvElem` reached from `:224`, on threadpool worker T1, inside case `W6: two spec shapes of EQUAL S and different q get two graphs` (`tests/vllm/models/test_qwen3_5_decode_graph_seam.cpp:800`). **It is NOT a use-after-free and there is no free site**, and `AddressSanitizer can not provide additional info` is the tell: the out-of-bounds block-table read is itself IN-BOUNDS for ASan because it lands inside the neighbouring pooled allocation, so this is not a `heap-buffer-overflow` either — the VALUE read becomes a wild block index that `KvElem` multiplies by the KV block stride and dereferences, which is why the fault address is unmapped rather than a redzone and why the crash is order-dependent instead of deterministic. Both halves are load-bearing: reverting ONLY the `SpecAttnMeta` fixture leaves the kernel bound to refuse by name at `src/vt/cpu/cpu_paged_attn.cpp:152` (exit 1, the case THROWS); reverting the bound as well restores the crash (Release exit 139, 3 runs of 3; ASan exit 1); restoring both gives exit 0, 8/8, zero sanitizer findings. The issue's `5f68e60df` pointer was right about the commit and wrong about the half — `git log -S'SpecAttnMeta'` and `git log -S'W6: two spec shapes of EQUAL S'` each return that commit ALONE, and both hits sit in the TEST fixture, not in the graph eligibility it widened. [#1405](https://github.com/mudler/vllm.cpp/issues/1405)'s truncation reproduced (`assertions: 135` crashed vs `138` complete) and is SHARPER under a threadpool fault: doctest printed the summary TWICE with different totals, `135` then `141`, because the main thread kept running after a worker died. That a red reached `main` unreported is [#1285](https://github.com/mudler/vllm.cpp/issues/1285)/[#1376](https://github.com/mudler/vllm.cpp/issues/1376), not this row | bug | | [#1458](https://github.com/mudler/vllm.cpp/issues/1458) | `VT-ACT-ROUND-POLARITY` | **`4712dac40` reds FOUR suites on `main` — `test_ltx2_text_encoder`, `test_muse_glimmer_text`, `test_muse_glimmer_text_fallback`, `test_minimax_music3_ar` — by exceeding bf16 error floors none of them had re-derived.** Found while gating [#1403](https://github.com/mudler/vllm.cpp/issues/1403) (PR #1457) and filed in flow, not that row's defect. Deterministic, not a load artifact: first seen under `-j4` and re-measured SERIALLY, 101 s across the four. `test_ltx2_text_encoder.cpp:2407` reads `CHECK( 0.1323 <= 0.109394 )` and `:2409` reads `CHECK( 0.0752773 <= 0.0573374 )`, both over by 20-30%, with the file at `26 passed / 1 failed`, `4118 assertions`, `Status: FAILURE!`. ATTRIBUTED BY MUTATION rather than inferred, both directions, in one build directory changing only `src/vt/cpu/cpu_ops.cpp`, compile rc 0 on every arm and `sha256sum` taken before and after: at `0adeb8b0e` four fail; reverted to `4712dac40^` four pass and `test_ltx2_text_encoder` is 27/27 `SUCCESS!`; restored to a re-matched `sha256` four fail again. They also passed at `b537a5344`, three commits earlier, in a full 567-test `ctest` run at 3-6 s each. `4712dac40` is +42/-3 in `src/vt/cpu/cpu_ops.cpp` plus a new `tests/vt/test_ops_activation.cpp`, and touches none of the four. NOT FIXED IN FLOW, deliberately: the question is a numerics decision, not a defect with an obvious repair — either the floors were calibrated against the rounding polarity that commit corrected and need re-deriving against the oracle, or the narrowing is wider than upstream's — and editing the four floors to pick the first branch is the scope-widening `AGENTS.md` prohibits | bug | +| [#1401](https://github.com/mudler/vllm.cpp/issues/1401) | `BACKEND-ROCM` | `GdnPostConvK` decomposes by `(token, head)` into `t*(hk+1)` items, and the ONE item with `head == hk` copies the entire `value_dim` row while the other `hk` items do about `4*dk` element-ops each. On Qwen3.5/3.6 geometry (`Hk=16, Hv=32, Dk=Dv=128`) at decode (`t=1`) that is 17 items in a single workgroup — 17 of 256 threads — with one lane doing 4096 element-ops against sixteen doing 512, an 8x imbalance on top of using one CU of 32. The kernel's duration is set by that lane: 413-479 us per call, 30 calls per token, **11.24 ms/token = 19% of ROCm GPU decode time**, the second largest kernel behind the activation quantizer ([#1400](https://github.com/mudler/vllm.cpp/issues/1400), [#1294](https://github.com/mudler/vllm.cpp/issues/1294)). Unlike the quantizer there is nothing to hoist — `Ld`/`St` are compile-time overloads and the kernel is already templated on all three dtypes — so this is purely a decomposition defect. FIXED IN FLOW by giving a token `hk` q/k slots, then `value_dim/dk` copy chunks, then one gate slot, which takes max per-thread element-ops from 4096 to 512 and `items` from 17 to 49. Every element operation is identical and independent (the v copy is elementwise, the gates are per-`h`), so this changes WHICH thread does an element and never the arithmetic. Same-binary A/B behind `VT_ROCM_GDN_POSTCONV_CHUNK`, measured on RX 9060 XT (gfx1200), ROCm 7.2.3, `rocprofv3 --kernel-trace --stats` with arms interleaved on a near-idle host: Qwen3.6-14B-A3B `qwen35moe` 468.6 -> 84.1 us per call (5.6x) and 12.560 -> 14.170 tok/s; Qwen3.6-35B-A3B with 7 GiB host-resident 479.1 -> 83.0 us (5.8x) and 7.846 -> 8.579 tok/s; Ornith-1.5-9B **DENSE** `qwen35` 413.2 -> 76.5 us (5.4x) and 16.275 -> 18.451 tok/s. **Output byte-identical on all three, both arms.** The dense 9B reaches the kernel through `Qwen3_5ForCausalLM` with no MoE block and no offload, so nothing else in the path accounts for the win; the 35B's smaller end-to-end ratio is expected because its decode is PCIe-bound from offloaded experts. NOT addressed and stated rather than implied: occupancy is unchanged at one workgroup (`items` 17 -> 49), and full element-parallelism (~8192 items) needs a wave-per-head reduction for the q/k sums, which is a separate change; the effect is decode-shaped and will not appear in a throughput benchmark because larger `t` spreads the original across more items | bug | | [#1294](https://github.com/mudler/vllm.cpp/issues/1294) | `BACKEND-ROCM` | ROCm decode on gfx1200 is kernel-bound, not transfer-bound, and the profile names which kernels. `rocprofv3` (rocprofiler-sdk 7.2.3) on RX 9060 XT with `Qwen3.6-14B-A3B-VibeForged-v2-Q4_K_M.gguf`, decode isolated by differencing `--max-tokens 4` against `--max-tokens 36` and dividing by 32, idle host: wall 88.7 ms/token (11.3 tok/s, matching the unprofiled 11.06-11.18 so the profiler is not the observer), GPU kernel dispatch 59.99 ms/token over 1738 dispatches, and **real H2D+D2H transfer 0.885 ms/token over 160 copies**. `hipMemcpyAsync` reports 59.31 ms/token of API time across 689 calls, which is the CPU blocking on unfinished GPU work rather than PCIe cost — so the host round-trips in the MoE reference path (`qwen3_5.cpp:6736`) are NOT the cost, which refutes the obvious hypothesis. By class: `vt::` kernels 44.92 ms/token (74.9%), hipBLASLt GEMM 13.81 (23.0%), copy/fill plumbing 1.26 (2.1%). Two leads, both measured: `QuantizeQ8KK` at 281 calls/token, 21.25 ms/token, 75.6 us per call, **35% of GPU time**; and `GdnPostConvK` (`rocm_gdn_postconv.hip:57`) at 30 calls/token, 11.24 ms/token, 375 us per call, 19%. Lead one has a named cause ON main and backend-agnostic: `MoeBlock` hands the SAME activation buffer to `KqGrouped` twice (`qwen3_5.cpp:6832-6833`, gate then up), so the identical activation is quantized twice per layer per token, and CPU and CUDA pay it too. NOT FIXED IN FLOW: the quantizer itself is `rocm_grouped_gemm.hip`, which is not on main and arrives with [#523](https://github.com/mudler/vllm.cpp/pull/523), so the 35% figure is feedback on that pull request rather than a main defect, and `GdnPostConvK` has no hypothesis attached yet — both want their own red-first evidence rather than a drive-by. Dispatch counts were byte-identical (281/30/60/20/21) across a contended and an idle run, so the structure is deterministic and only durations moved. Measured on `main` @ `4ee5f4a6` plus #523 plus the #559/#570 `AttnQkNormRopeGate` fix; no pinned-oracle comparison, one model, one prompt, batch 1 | record | | [#1292](https://github.com/mudler/vllm.cpp/issues/1292) | `FIX-GCC15-ARRAY-BOUNDS` | `cmake/CompilerWarnings.cmake:39-42` demoted `-Warray-bounds` to non-fatal for GCC >= 16 only, and its comment asserted "Everything <= 15 is unchanged and still fails the build on a real out-of-bounds." gcc 15.2.0 reproduces the same false-positive class from the OTHER libstdc++ container: `std::vector`'s inlined copy-construct plus `_M_allocate`, reported at `ltx2_samplers.cpp:161,163` as "array subscript -1 is outside array bounds of `float [2305843009213693951]`". That bound is `SIZE_MAX/4`, the allocator's unconstrained size range rather than a real object, which is the same inlining-attribution mechanism the guard's own comment documents for `_Sp_counted_base::_M_release()`. The subscripts are `sigmas.back()` guarded three lines above by `VT_CHECK(sigmas_in.size() >= 2)` on the container it copies, so the code is correct and GCC does not propagate the throwing check through the inlined copy. `main` therefore did not build on gcc 15.x at all, stopping at 146/1552. No CI lane covers the release: the matrix is gcc 13 (ubuntu-latest distro `g++`, every ordinary lane) and gcc 16 (`build-newest-gcc`, `container: gcc:16`, already exempt via the guard), so the failure falls exactly between the two points, and gcc 15 is the current toolchain on Arch, NixOS unstable and Fedora 42. Introduced by `4d774864` (`LTX25-RES2S-LOOP`, 2026-08-17), which added the loop; the guard predates it and was written for the `shared_ptr`/json case. FIXED IN FLOW by widening the existing guard to `VERSION_GREATER_EQUAL 15`, which keeps the diagnostic visible as a warning. Red-before/green-after measured on `4ee5f4a6` with gcc 15.2.0: the TU is `exit 1` with two `-Werror=array-bounds=` errors, and `exit 0` with the same two as warnings once the guard's own `-Wno-error=array-bounds` applies; `scripts/agent-preflight.sh` fails the identical 10 gates before and after, so the change adds no failure. Whether CI should also cover gcc 14/15 is a SEPARATE question and is deliberately NOT bundled: the widened guard already prevents this defect recurring on 15, so a lane would guard a different class (the transitive-include portability the `build-newest-gcc` comment describes) and deserves its own issue rather than riding a build fix | bug | | [#1404](https://github.com/mudler/vllm.cpp/issues/1404) | `SPEC-DFLASH2` | The upstream head this row mirrors MOVED. `SPEC-DFLASH2` W1 and W2 were written against [vllm-project/vllm#52816](https://github.com/vllm-project/vllm/pull/52816) head `19c9351904df4c63042671bc67a866ca48dc7d6f`; the PR head is now `66e5414c6d75a8529473d977f7458c140bbab8a0`, which also FOLDED IN the previously stacked LM-head guard fix [vllm#52883](https://github.com/vllm-project/vllm/pull/52883). `## Risks/decisions` D1 accepted exactly this cost when it chose to mirror an unmerged pull request: "the anchors can move under review, and the port reconciles if they do". Reconciled by W3 on 2026-08-20, which re-read every anchor it cites at the new head and recorded the delta rather than rewriting the old citations in silence. Diffing the two heads changes five files; of the two this row ports, `qwen3_dflash2.py` is +24/-11 and IS exactly two things, both infrastructure rather than math, while `vllm/v1/worker/gpu/spec_decode/dflash2/speculator.py` is +37/-34 and is neither -- it is math and state layout in W4's own territory, so a W4 implementer who reads only this row would under-scope that wave. Both deltas measured with `git diff --no-index --numstat` over the blobs at each head; the per-file table and the enumeration live in the spec's `## Upstream chain`. `set_model_tag("dflash2_candidate_selector")` around the selector's construction is a DELIBERATE NON-PORT (spec D11): upstream needs it because `CandidateSelector` carries its own `@support_torch_compile` and is built under the draft's model tag, so the two would share one compile-cache namespace and the selector would load the draft's graph -- and this engine has no torch.compile and no compile cache. The widened LM-head guard IS ported (spec D12), in its WIDE form: `UnquantizedLinearMethod` beside `UnquantizedEmbeddingMethod`, because a `ParallelLMHead` returns the LINEAR method whenever a quant config leaves the head unquantized. `_score_edges`, `CandidateSelector`, `hidden_projection`, the two codebooks, `_topk`, `output_multiplier` and `final_logit_softcapping` are BYTE-IDENTICAL at the two heads, so the SELECTOR's math is unaffected -- "the anchors moved" and "the port must change" are different statements, and which one is true depends on which file. This does NOT advance the parity pin `555967922`, which does not carry the architecture at all | record | diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index 1120e1ddf..8069a0aab 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -167,6 +167,7 @@ VT_QWEN3_ROPE_CACHE VT_RELEASE_HOST_WEIGHTS VT_RMSNORM_DECODE_FAST VT_RMSNORM_GATED_FAST +VT_ROCM_GDN_POSTCONV_CHUNK VT_ROCM_GEMM_COMPUTE VT_ROCM_GEMV VT_ROCM_HIPBLASLT diff --git a/src/vt/rocm/rocm_gdn_postconv.hip b/src/vt/rocm/rocm_gdn_postconv.hip index 041eb4d77..7fc23cd13 100644 --- a/src/vt/rocm/rocm_gdn_postconv.hip +++ b/src/vt/rocm/rocm_gdn_postconv.hip @@ -17,6 +17,7 @@ #include #include +#include #include #include "vt/ops.h" @@ -104,6 +105,86 @@ __global__ void GdnPostConvK(Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, float* g_out } } +// Same work as GdnPostConvK, decomposed so no single thread carries the whole +// v copy. The original splits by (token, head) into t*(hk+1) items, and the ONE +// item with head==hk copies the entire value_dim row -- 4096 elements at +// Hv=32/Dv=128 -- while the other hk items do ~4*dk each. At decode (t=1) that +// is 17 items in a single workgroup with an 8x imbalance, so the kernel's +// duration is set by that one lane. +// +// Here a token's slots are: hk q/k heads, then v_chunks copy chunks of v_chunk +// elements each, then one slot for the gates. Every element operation is +// identical and independent (the v copy is elementwise, the gates are per-h), so +// this reorders WHICH thread does a given element and never the arithmetic -- +// output is bit-exact. +template +__global__ void GdnPostConvChunkedK(Tqkv* q_out, Tqkv* k_out, Tqkv* v_out, float* g_out, + float* beta_out, const Tconv* conv, const Tgate* araw, + const Tgate* braw, const float* a_log, + const float* dt_bias, int64_t t, int64_t hk, int64_t dk, + int64_t hv, int64_t dv, int64_t a_row_stride, + int64_t b_row_stride, float eps, int64_t v_chunk, + int64_t v_chunks) { + const int64_t key_dim = hk * dk, value_dim = hv * dv; + const int64_t conv_dim = 2 * key_dim + value_dim; + const int64_t slots = hk + v_chunks + 1; + const int64_t items = t * slots; + const int64_t step = static_cast(gridDim.x) * blockDim.x; + for (int64_t item = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + item < items; item += step) { + const int64_t tok = item / slots; + const int64_t slot = item % slots; + const int64_t row = tok * conv_dim; + if (slot < hk) { + const int64_t head = slot; + const Tconv* qin = conv + row + head * dk; + const Tconv* kin = conv + row + key_dim + head * dk; + float qss = 0.0f, kss = 0.0f; + for (int64_t j = 0; j < dk; ++j) { + const float qv = Ld(qin, j); + const float kv = Ld(kin, j); + qss += qv * qv; + kss += kv * kv; + } + const float qinv = 1.0f / sqrtf(qss + eps); + const float kinv = 1.0f / sqrtf(kss + eps); + Tqkv* qo = q_out + (tok * hk + head) * dk; + Tqkv* ko = k_out + (tok * hk + head) * dk; + for (int64_t j = 0; j < dk; ++j) { + St(qo, j, Ld(qin, j) * qinv); + St(ko, j, Ld(kin, j) * kinv); + } + } else if (slot < hk + v_chunks) { + const int64_t c = slot - hk; + const Tconv* vin = conv + row + 2 * key_dim; + Tqkv* vo = v_out + tok * value_dim; + const int64_t j0 = c * v_chunk; + const int64_t j1 = (j0 + v_chunk) < value_dim ? (j0 + v_chunk) : value_dim; + for (int64_t j = j0; j < j1; ++j) St(vo, j, Ld(vin, j)); + } else { + for (int64_t h = 0; h < hv; ++h) { + const int64_t idx = tok * hv + h; + const float av = Ld(araw, tok * a_row_stride + h); + const float bv = Ld(braw, tok * b_row_stride + h); + const float x = av + dt_bias[h]; + const float sp = x > 20.0f ? x : log1pf(expf(x)); + g_out[idx] = -expf(a_log[h]) * sp; + beta_out[idx] = 1.0f / (1.0f + expf(-bv)); + } + } + } +} + +// Same-binary A/B. Default ON; VT_ROCM_GDN_POSTCONV_CHUNK=0 restores the +// original single-thread-per-v-row kernel. Output must be byte-identical. +inline bool GdnPostConvChunkedEnabled() { + static const bool on = [] { + const char* e = std::getenv("VT_ROCM_GDN_POSTCONV_CHUNK"); + return e == nullptr || std::string(e) != "0"; + }(); + return on; +} + } // namespace void GdnPostConvKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out, @@ -121,11 +202,26 @@ void GdnPostConvKernelRocm(Queue& q, Tensor& q_out, Tensor& k_out, Tensor& v_out VT_CHECK(conv.dtype == DType::kF32 || conv.dtype == DType::kBF16, "rocm gdn_post_conv: conv must be f32 or bf16"); const int64_t items = t * (hk + 1); + // Chunked decomposition: split the value_dim copy into dk-sized pieces so it + // is not one thread's serial loop. dk is the natural chunk -- it is what the + // q/k slots already do, so every slot carries comparable work. + const int64_t value_dim = hv * dv; + const int64_t v_chunk = dk > 0 ? dk : value_dim; + const int64_t v_chunks = v_chunk > 0 ? (value_dim + v_chunk - 1) / v_chunk : 1; + const int64_t items_chunked = t * (hk + v_chunks + 1); hipStream_t s = AsStream(q); auto launch = [&](auto qkv_tag, auto conv_tag, auto gate_tag) { using Tqkv = decltype(qkv_tag); using Tconv = decltype(conv_tag); using Tgate = decltype(gate_tag); + if (GdnPostConvChunkedEnabled()) { + GdnPostConvChunkedK<<>>( + q_out.Ptr(), k_out.Ptr(), v_out.Ptr(), g_out.Ptr(), + beta_out.Ptr(), conv.Ptr(), araw.Ptr(), braw.Ptr(), + a_log.Ptr(), dt_bias.Ptr(), t, hk, dk, hv, dv, araw.stride[0], + braw.stride[0], args.eps, v_chunk, v_chunks); + return; + } GdnPostConvK<<>>( q_out.Ptr(), k_out.Ptr(), v_out.Ptr(), g_out.Ptr(), beta_out.Ptr(), conv.Ptr(), araw.Ptr(), braw.Ptr(),