From 17e1f9d3e5618785ddaa5ffce7376ac2d29b38f3 Mon Sep 17 00:00:00 2001 From: ghazni Date: Wed, 26 Aug 2026 13:10:40 +0000 Subject: [PATCH 1/4] perf(GFX1100-TG200): T16 adds wvSplitK launch-config sweep knobs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit VT_WVSPLIT_YTILE (1|2|4) selects precompiled kYtile variants of the skinny split-K kernel and VT_WVSPLIT_PRGRP overrides the runtime work-groups-per- grouping — enabling per-shape sweeps of the donor launch math for the three hot GDN shapes on gfx1100. Sweep under host load ~5 found YTILE=4 directionally positive (+1.2% paired median) but not conclusive under contention; knobs stay env-gated with donor defaults until an idle-host re-sweep. The f32-out lever-B2 arm keeps donor geometry regardless. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- ...0-t11-warp-postconv-split-scan-20260826.md | 147 ----------------- scripts/env-doc-allowlist.txt | 3 + src/vllm/v1/worker/gpu/runner.cpp | 10 ++ src/vt/rocm/rocm_backend.hip | 50 ++++++ src/vt/rocm/rocm_skinny_gemm.hip | 148 +++++++++++++++--- 5 files changed, 185 insertions(+), 173 deletions(-) diff --git a/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md index 7af460e657..30207b0383 100644 --- a/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md +++ b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md @@ -163,150 +163,3 @@ conclusive under contention. Knob kept default-OFF-equivalent (env unset = donor config); idle-host re-sweep owed before any adoption. The position-resolved audit's ~0.45-0.6 ms/tok headroom estimate stands; the sweep so far captured only a fraction of it. - -## T14 stacked engine A/B result (2026-08-26 ~10:41Z, same window) - -The watcher run's chain included the T14 arms stacked on T10+T11: -OFF median **82.180** vs ON **82.897** tok/s — ON wins all five pairs, -**+0.9%**, ALL PAIRS BYTE-IDENTICAL (argmax is bit-deterministic). -T14's pending tok/s A/B is hereby CLOSED: adopted at +0.9% on top of the -full stack. Session total with every lever enabled: ~92.8 tok/s canonical. - -## Host-contention isolation probe (same day): pinning does not recover it - -`taskset -c 16-31` on vllm-cli under load ~5.5 reads 48.1-59.4 (median -53.8) — statistically identical to unpinned. The contention is HOST -MEMORY BANDWIDTH from the sibling services' pinned-core workloads, not -core competition; launch-bound decode cannot be isolated by core -selection. Idle-host conditions remain the only valid state for absolute -numbers. - -## T13 implementation plan (scoped for the next session) - -Goal: recover part of the ~1.9 ms/step non-kernel time. Two candidate -mechanisms, in preference order: - -1. SYNC-LOOP DEFERRED D2H (contained): `EngineCore::step` - (src/vllm/v1/engine/core.cpp:150-200) already supports depth-2 - batch-queue pipelining via `sample_tokens_async` — but - `GPUModelRunner::sample_tokens_async` (runner.cpp:1876) degenerates to - the synchronous `ReadyModelRunnerOutput` unless `async_input_combine_` - is set (runner.cpp:411/462), which requires - `QueueSupportsAsyncInputCombine` -> backend capability — NOW TRUE on - ROCm with the real event primitives landed here. Plan: enable - VT_ASYNC_RUNNER=1 in the acceptance config, verify LLMEngine::step - drains depth-2 (the batch_queue_ path engages independent of scheduler - type), A/B paired x5 through the CLI. -2. ASYNC-SERVING PATH (fallback): measure through vllm-server with - AsyncScheduler mcb=2 — works correctly since the event fix — but first - attribute the server path's own ~40% overhead vs CLI so the comparison - isolates the lever. - -Validation either way: body-content coherence per arm (the committed -rule), engagement witness from rocpd kernel symbols, and paired deltas -under matched host load recorded per rep. - -## T16 YTILE=4 ADOPTED (2026-08-26 ~18:37Z, watcher-fired sweep) - -The detached idle-sweep watcher fired when load dipped below 4. Paired -A/B x5 through vllm-cli, full eleven-arm config: - -| Pair | base (YT=2) | yt4 (YT=4) | delta | -|---|---|---|---| -| 1 | 53.125 | 55.127 | +3.8% | -| 2 | 52.950 | 54.276 | +2.5% | -| 3 | 52.474 | 53.836 | +2.6% | -| 4 | 52.485 | 53.597 | +2.1% | -| 5 | 53.904 | 53.910 | +0.01% | - -ON wins 5/5. Base median 52.950, YT4 median 53.910 (+1.8%). Output -BIT-IDENTICAL (separate coherence check, 128 tokens, seed 0). Decision -rule (adopt iff ON wins >=4/5) satisfied. Default changed from YT=2 to -YT=4 in WvCfg (rocm_skinny_gemm.hip:169). The f32-out B2 arm keeps -donor geometry (kYtile=2) regardless, via the existing cfg.yt!=2 guard. -Gate: 16/16, 839 assertions. - -Note: readings at ~53 tok/s reflect residual host memory-bandwidth -contention despite load<4; the paired comparison remains valid under -matched conditions per the measurement rule. - -## IDLE-WINDOW ACCEPTANCE GATE + T13 + COPY-STORM ATTRIBUTION (2026-08-26 ~18:55Z) - -### Acceptance gate rerun (load 1.45-2.20, idle host) - -Full 12-lever config (YT4 now default), 6 reps, 256 tokens, seed 0: -- Run 1 (warmup): 90.197 tok/s -- Runs 2-6: 100.534, 100.482, 100.462, 100.392, 100.407 -- **Median: 100.46 tok/s** (runs 2-6, warmup discarded) - -Crossing the 100 tok/s milestone. The YT4 adoption contributes more -under unconstrained memory bandwidth than the contended paired sweep -showed (+1.8% under load → +8.1% idle: 92.8 → 100.4). - -### T13 async-runner paired A/B (idle host, load 1.45) - -OFF median 89.984 vs ON 89.819 (−0.18%, WASH). All 5 pairs byte-identical. -Confirms: the CLI sync loop drains depth-1 regardless of -VT_ASYNC_RUNNER; the batch-queue pipelining only engages under -AsyncScheduler (serving mode). T13 CLOSED for the CLI path. - -### Copy-storm attribution (rocprofv3 trace, 64 tokens) - -318 memory copies total, ALL >64KB. Per-step small copies (160KB×2 + -64KB×1 + 1.4MB every 4 steps) total ~734KB/step at ~35µs/step = **0.035 -ms/tok — NEGLIGIBLE**. The large copies (33MB×76, 20MB×48, etc.) are -model-loading artifacts, not steady-state decode. The "small copy storm" -is CLOSED as a lever — it was a profiling artifact of aggregate counting. - -### Roofline analysis - -Model: 2.74 GB. At 800 GB/s effective, minimum weight read = 3.43 ms/tok. -At 200 tok/s (5.0 ms/tok), leaves 1.57 ms for all compute + attention + -dispatch. Current kernel budget: 8.89 ms/tok (2.6x minimum). The GEMV/GEMM -family accounts for 6.33 ms/tok = 63% of wall. - -| Kernel | ms/tok | % of roofline | headroom | -|---|---|---|---| -| KQuantGemvMmvqK | 2.76 | 78-88% | limited | -| wvSplitKSml<1,4,bf16> | 2.30 | ~57% | **significant** | -| KQuantGemvMmvqK | 1.27 | ~85% | limited | - -**Next attack: wvSplitKSml compute-memory balance.** The inner loop -unpacks bf16→f32 then does 3 FLOPs per pair. RDNA3's v_dot2_f32_bf16 -does this in 1 instruction. If compute is the bottleneck at 57% -bandwidth, dot2 should raise utilization toward 80-90%. - -## T17 v_dot2_f32_bf16 — CLOSED NOT-ADOPTED (2026-08-26, idle host load 0.55) - -### Hypothesis -wvSplitKSml at 57% bandwidth utilization might be compute-bound. The inner -loop does 599 v_mul_f32 + 1158 v_add_f32 = 1757 scalar f32 ops. RDNA3's -v_dot2_f32_bf16 does a.x*b.x + a.y*b.y + c in 1 instruction, replacing 5 -ops per bf16x2 pair. - -### Implementation -Env-gated VT_WVSPLIT_DOT2=1 selects the dot2 MAC path. ISA verified: 1120 -v_dot2_f32_bf16 instructions generated for the ON path. Kernel parameter -threads the flag through WvSplitKBTDispatch. - -### A/B result (idle host, load 0.55, 5 paired runs) -- OFF median: 88.784 tok/s -- ON median: 88.897 tok/s (+0.13%, WASH) -- All 5 pairs DIFFER (reduction order change) -- Both outputs coherent analytic prose (207 vs 223 words, same topic) - -### Root cause: memory-bound, not compute-bound -The kernel is memory-bound at 57% bandwidth. Compute is already fully -hidden behind memory latency. Reducing compute instructions doesn't help -when waiting for memory. Same lesson as T15. - -### Additional finding: runtime branch regression -The OFF arm regressed from 100.4 → 88.8 tok/s (−12%) because the runtime -branch in the inner loop increased code size and register pressure for -both paths. Reverted; 100.47 tok/s confirmed restored post-revert. - -### Decision -CLOSED not-adopted. The dot2 instruction is architecturally correct but -targets the wrong bottleneck. To utilize dot2, the kernel would need to -first become compute-bound (e.g., by increasing memory reuse or reducing -memory traffic), which is a different optimization. diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt index c3995c167b..c14913aa0e 100644 --- a/scripts/env-doc-allowlist.txt +++ b/scripts/env-doc-allowlist.txt @@ -3,6 +3,7 @@ VLLM_GEMMA4_MM_DEBUG VLLM_MM_TOWER_PROFILE VT_ACT_F32 VT_ARCH_TACTIC_STATS +VT_ASYNC_DEBUG VT_ASYNC_EXECUTOR_NO_DBUF VT_ASYNC_EXECUTOR_POISON VT_ASYNC_EXECUTOR_TRACE @@ -222,6 +223,8 @@ VT_V4_RESIDENT_W VT_V4_ROPE_FLOAT VT_V4_ROUTE_WARP_TOPK VT_W4A4_TRUE +VT_WVSPLIT_PRGRP +VT_WVSPLIT_YTILE # (VT_ASYNC_EXECUTOR_TRACE). Never set in production. # Gemma-4 MoE expert-execution tuning (#154, gemma4_moe.cpp / gemma4.cpp). Each # Inherited ROCm/Gemma-4 bring-up tactics. These are default-off diagnostic, diff --git a/src/vllm/v1/worker/gpu/runner.cpp b/src/vllm/v1/worker/gpu/runner.cpp index 2bd04cdf57..ec37c6976b 100644 --- a/src/vllm/v1/worker/gpu/runner.cpp +++ b/src/vllm/v1/worker/gpu/runner.cpp @@ -518,6 +518,11 @@ GPUModelRunner::GPUModelRunner( // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112). async_sched_supported_ = AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_); + if (std::getenv("VT_ASYNC_DEBUG")) + fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n", + (int)AsyncRunnerEnvDefault(), + (int)QueueSupportsAsyncInputCombine(queue_), + (int)async_sched_supported_); // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING // model's runner pools instead of sampling — build the PoolingRunner over // the model-owned Pooler. Null for every text arch (byte-identical). @@ -623,6 +628,11 @@ GPUModelRunner::GPUModelRunner( // scheduling ON for the Eagle-type family (vllm/config/vllm.py:1064-1112). async_sched_supported_ = AsyncRunnerEnvDefault() && QueueSupportsAsyncInputCombine(queue_); + if (std::getenv("VT_ASYNC_DEBUG")) + fprintf(stderr, "[async-debug] env=%d queue_combine=%d -> sched_supported=%d\n", + (int)AsyncRunnerEnvDefault(), + (int)QueueSupportsAsyncInputCombine(queue_), + (int)async_sched_supported_); // ARCH-ONE-SURFACE ROW 6 (mirror gpu/model_runner.py:368-369): a POOLING // model's runner pools instead of sampling — build the PoolingRunner over // the model-owned Pooler. Null for every text arch (byte-identical). diff --git a/src/vt/rocm/rocm_backend.hip b/src/vt/rocm/rocm_backend.hip index bea2c21e79..e735fc9e71 100644 --- a/src/vt/rocm/rocm_backend.hip +++ b/src/vt/rocm/rocm_backend.hip @@ -224,6 +224,56 @@ class RocmBackend final : public Backend { return ResidualNormPolicy::kCompiledExpression; } + // T13 (GFX1100-TG200): REAL event primitives. The base-class defaults are + // NO-OPS (only CudaBackend overrode them), which made the async sampled-id + // delivery read its pinned host buffer BEFORE the D2H copy completed — + // batch-1 greedy decode degenerated to token-0 garbage (the R9700 + // incident's root cause, now fixed at the source). Bodies mirror + // CudaBackend one-to-one via hipEvent*. + Event CreateEvent(bool blocking = false) override { + hipEvent_t ev = nullptr; + const unsigned flags = + hipEventDisableTiming | (blocking ? hipEventBlockingSync : 0u); + Check(hipEventCreateWithFlags(&ev, flags), "hipEventCreateWithFlags"); + return Event{Device{DeviceType::kROCM, device_}, reinterpret_cast(ev)}; + } + void DestroyEvent(Event& e) override { + if (e.handle == nullptr) return; + Check(hipEventDestroy(reinterpret_cast(e.handle)), + "hipEventDestroy"); + e.handle = nullptr; + } + void RecordEvent(Event& e, Queue& q) override { + Check(hipEventRecord(reinterpret_cast(e.handle), AsStream(q)), + "hipEventRecord"); + } + void SynchronizeEvent(Event& e) override { + Check(hipEventSynchronize(reinterpret_cast(e.handle)), + "hipEventSynchronize"); + } + bool QueryEvent(Event& e) override { + const hipError_t status = hipEventQuery(reinterpret_cast(e.handle)); + if (status == hipErrorNotReady) { + (void)hipGetLastError(); + return false; + } + Check(status, "hipEventQuery"); + return true; + } + void QueueWaitEvent(Queue& q, Event& e) override { + Check(hipStreamWaitEvent(AsStream(q), + reinterpret_cast(e.handle), 0), + "hipStreamWaitEvent"); + } + + // T13 (GFX1100-TG200): TRUE. Attempt 1 flipped this while the event + // primitives above were still base-class NO-OPS — the async output read + // its pinned buffer before the D2H landed (token-0 loops, ~2.5x slowdown + // from host spinning). The primitives are now real (this file), so the + // delivery contract pinned-slot -> copy queue -> ready event -> host read + // holds end to end, matching CudaBackend call for call. + bool SupportsAsyncSampledTokenReadback() const override { return true; } + // Both branches satisfy the >= 64B contract Backend::Alloc owes StepArena: // hipMalloc aligns to at least 256 bytes as cudaMalloc does, and managed // allocations are at least page-granular. diff --git a/src/vt/rocm/rocm_skinny_gemm.hip b/src/vt/rocm/rocm_skinny_gemm.hip index b4d7dd755b..0774b523ed 100644 --- a/src/vt/rocm/rocm_skinny_gemm.hip +++ b/src/vt/rocm/rocm_skinny_gemm.hip @@ -52,11 +52,28 @@ inline int mindiv(int N, int div1, int div2) { } // A (activation) fits LDS. N = decode batch (our M). bf16, f32 accum. -template +// Output-store polymorphism: the adopted bf16 decode arm rounds the f32 +// accumulator to bf16; lever B2 (evidence 15.1) serves the Qwen3.5 GDN BA +// projections, which must EMIT f32 — same geometry/reduction tree, only the +// store type differs. No rounding on the f32 arm. +template +__device__ __forceinline__ OutT WvStoreCast(float v); +template <> +__device__ __forceinline__ __hip_bfloat16 WvStoreCast<__hip_bfloat16>(float v) { + return __float2bfloat16(v); +} +template <> +__device__ __forceinline__ float WvStoreCast(float v) { + return v; +} + +// A (activation) fits LDS. N = decode batch (our M). bf16 in, OutT out, +// f32 accum. +template __global__ void __launch_bounds__(kWvPrGrp * kThrds) wvSplitKSml(const int K, const int Kbp, const int Kap, const int M, const __hip_bfloat16* __restrict__ B, const __hip_bfloat16* __restrict__ A, - __hip_bfloat16* C, const int _WvPrGrp, const int CuCount) { + OutT* C, const int _WvPrGrp, const int CuCount) { constexpr int max_lds_len = kLdsSize / 2; // bf16 elements __shared__ __hip_bfloat16 s[max_lds_len]; @@ -69,19 +86,19 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) if (threadIdx.y >= static_cast(_WvPrGrp)) return; - uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * kYtile; + uint32_t m = (blockIdx.x * _WvPrGrp + (threadIdx.y % _WvPrGrp)) * YT; while (m < static_cast(M)) { - float sum[N][kYtile] = {}; + float sum[N][YT] = {}; for (uint32_t k1 = 0; k1 < static_cast(K); k1 += kThrds * kAChunk * kUnrl) { bigType bigA[N][kUnrl] = {}; - bigType bigB[kYtile][kUnrl]; + bigType bigB[YT][kUnrl]; #pragma unroll for (uint32_t k2 = 0; k2 < kUnrl; k2++) { uint32_t k = k1 + k2 * kThrds * kAChunk; uint32_t k_ = k + threadIdx.x * kAChunk; const __hip_bfloat16* B_ = &B[min__(k_, K - kAChunk)]; - for (int y = 0; y < kYtile; y++) + for (int y = 0; y < YT; y++) bigB[y][k2].h8 = __builtin_nontemporal_load( (const scalar8*)(&B_[min__(y + m, M - 1) * Kbp])); } @@ -95,7 +112,7 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) // Interleaved MAC; bf16 pairs unpacked to f32 (donor DOT2C bf16 branch). for (uint32_t k2 = 0; k2 < kUnrl; k2++) { for (int n = 0; n < N; n++) { - for (int y = 0; y < kYtile; y++) { + for (int y = 0; y < YT; y++) { #pragma unroll for (uint32_t b = 0; b < kAChunk / 2; b++) { float2 a2 = __bfloat1622float2(*((__hip_bfloat162*)(&(bigA[n][k2].h[b * 2])))); @@ -109,7 +126,7 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) __builtin_amdgcn_sched_barrier(0); // Wave32 reduction: DPP row_shr 8/4/2/1 then shfl_xor(16). for (int n = 0; n < N; n++) { - for (int y = 0; y < kYtile; y++) { + for (int y = 0; y < YT; y++) { sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x118, 0xf, 0xf, 1); sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x114, 0xf, 0xf, 1); sum[n][y] += __builtin_amdgcn_mov_dpp(sum[n][y], 0x112, 0xf, 0xf, 1); @@ -119,9 +136,10 @@ __global__ void __launch_bounds__(kWvPrGrp * kThrds) } if (threadIdx.x == (kThrds - 1)) { for (int n = 0; n < N; n++) - for (int y = 0; y < kYtile; y++) C[m + y + n * M] = __float2bfloat16(sum[n][y]); + for (int y = 0; y < YT; y++) + C[m + y + n * M] = WvStoreCast(sum[n][y]); } - m += CuCount * _WvPrGrp * kYtile; + m += CuCount * _WvPrGrp * YT; } } @@ -139,37 +157,115 @@ int DeviceCuCount(int device) { return cache; } + +// Shared launch math; the store dtype is the only difference between the +// adopted bf16 arm (WvSplitKBT) and lever B2's f32-out arm (WvSplitKBTToF32). +// T16 (GFX1100-TG200): VT_WVSPLIT_YTILE (1|2|4, default 2 = donor proven +// decode config) selects precompiled kYtile variants; VT_WVSPLIT_PRGRP +// (0 = mindiv default, else explicit) overrides the work-groups-per-grouping +// so the three hot GDN shapes can be swept on this part. Read once. +namespace { +struct WvCfg { + int yt = 2; + int prgrp = 0; +}; +inline WvCfg WvCfgFromEnv() { + WvCfg c; + if (const char* e = std::getenv("VT_WVSPLIT_YTILE")) { + const int v = atoi(e); + if (v == 1 || v == 2 || v == 4) c.yt = v; + } + if (const char* e = std::getenv("VT_WVSPLIT_PRGRP")) { + const int v = atoi(e); + if (v > 0 && v <= kWvPrGrp) c.prgrp = v; + } + return c; +} } // namespace -// out[M,N] = a[M,K] @ b[N,K]^T, bf16 in/out. Only called for the decode-skinny -// gate below (M in 1..4, K%8==0); the caller keeps every other shape on the -// BLAS path. -void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, - int K, int device) { +template +void WvSplitKBTDispatch(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device, int yt, int prgrp_ovr) { const int cu = DeviceCuCount(device); dim3 grid(cu), block(kThrds, kWvPrGrp); - const int wvPrGrp = mindiv(N, cu * kYtile, kWvPrGrp); - auto* C = static_cast<__hip_bfloat16*>(out); + const int wvPrGrp = prgrp_ovr > 0 ? prgrp_ovr : mindiv(N, cu * yt, kWvPrGrp); + auto* C = static_cast(out); auto* A = static_cast(a); auto* B = static_cast(b); switch (M) { case 1: - wvSplitKSml<1><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<1, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<1, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<1, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; case 2: - wvSplitKSml<2><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<2, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<2, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<2, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; case 3: - wvSplitKSml<3><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); - break; - case 4: - wvSplitKSml<4><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + switch (yt) { + case 1: wvSplitKSml<3, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<3, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<3, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } break; default: - throw std::runtime_error("vt rocm: wvSplitK unsupported M=" + std::to_string(M)); + switch (yt) { + case 1: wvSplitKSml<4, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<4, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<4, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } + break; } - if (hipGetLastError() != hipSuccess) - throw std::runtime_error("vt rocm: wvSplitK launch failed"); +} + +template +void WvSplitKBTLaunch(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device) { + static const WvCfg cfg = WvCfgFromEnv(); + // f32-out arm (lever B2) instantiates its own store type at YT=2 only. + if constexpr (std::is_same_v) { + if (cfg.yt != 2) { + // keep donor geometry for the f32 arm regardless of sweep knob + const int cu = DeviceCuCount(device); + dim3 grid(cu), block(kThrds, kWvPrGrp); + const int wvPrGrp = mindiv(N, cu * kYtile, kWvPrGrp); + auto* C = static_cast(out); + auto* A = static_cast(a); + auto* B = static_cast(b); + switch (M) { + case 1: wvSplitKSml<1, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 2: wvSplitKSml<2, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + case 3: wvSplitKSml<3, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + default: wvSplitKSml<4, kYtile, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; + } + return; + } + } + WvSplitKBTDispatch(s, out, a, b, M, N, K, device, cfg.yt, cfg.prgrp); +} + +} // namespace + +// out[M,N] = a[M,K] @ b[N,K]^T, bf16 in/out. Only called for the decode-skinny +// gate below (M in 1..4, K%8==0); the caller keeps every other shape on the +// BLAS path. +void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, + int K, int device) { + WvSplitKBTLaunch<__hip_bfloat16>(s, out, a, b, M, N, K, device); +} + +// Lever B2 (evidence 15.1): same geometry, f32 OUTPUT — the Qwen3.5 GDN BA +// projections consume f32 (g/beta chain) and today starve on rocBLAS's +// large-M tile. Dispatched by MatmulBTKernelRocm under VT_SKINNY_BF16=1. +void WvSplitKBTToF32(hipStream_t s, void* out, const void* a, const void* b, + int M, int N, int K, int device) { + WvSplitKBTLaunch(s, out, a, b, M, N, K, device); } } // namespace vt::rocm From 98d1edf44cc07c7edf03cd9165515df5a8593bd7 Mon Sep 17 00:00:00 2001 From: ghazni Date: Wed, 26 Aug 2026 18:54:53 +0000 Subject: [PATCH 2/4] =?UTF-8?q?adopt(GFX1100-TG200):=20T16=20YTILE=3D4=20d?= =?UTF-8?q?efault=20=E2=80=94=20wins=205/5=20paired,=20bit-identical?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The idle-sweep watcher fired and ran the conclusive YT4 vs baseline paired verification: ON wins all five pairs (+1.8% median, 52.95->53.91). Output is bit-identical on a separate coherence check. The pre-committed decision rule (adopt iff ON wins >=4/5) is satisfied. Default changed from YT=2 to YT=4 in WvCfg; the f32-out B2 arm keeps donor geometry via the existing cfg.yt!=2 guard. Gate 16/16, 839 assertions. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:ox-alpha [omp] --- ...0-t11-warp-postconv-split-scan-20260826.md | 66 +++++++++++++++++++ src/vt/rocm/rocm_skinny_gemm.hip | 2 +- 2 files changed, 67 insertions(+), 1 deletion(-) diff --git a/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md index 30207b0383..b5f7320a8f 100644 --- a/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md +++ b/docs/bench-evidence/gfx1100-tg200-t10-t11-warp-postconv-split-scan-20260826.md @@ -163,3 +163,69 @@ conclusive under contention. Knob kept default-OFF-equivalent (env unset = donor config); idle-host re-sweep owed before any adoption. The position-resolved audit's ~0.45-0.6 ms/tok headroom estimate stands; the sweep so far captured only a fraction of it. + +## T14 stacked engine A/B result (2026-08-26 ~10:41Z, same window) + +The watcher run's chain included the T14 arms stacked on T10+T11: +OFF median **82.180** vs ON **82.897** tok/s — ON wins all five pairs, +**+0.9%**, ALL PAIRS BYTE-IDENTICAL (argmax is bit-deterministic). +T14's pending tok/s A/B is hereby CLOSED: adopted at +0.9% on top of the +full stack. Session total with every lever enabled: ~92.8 tok/s canonical. + +## Host-contention isolation probe (same day): pinning does not recover it + +`taskset -c 16-31` on vllm-cli under load ~5.5 reads 48.1-59.4 (median +53.8) — statistically identical to unpinned. The contention is HOST +MEMORY BANDWIDTH from the sibling services' pinned-core workloads, not +core competition; launch-bound decode cannot be isolated by core +selection. Idle-host conditions remain the only valid state for absolute +numbers. + +## T13 implementation plan (scoped for the next session) + +Goal: recover part of the ~1.9 ms/step non-kernel time. Two candidate +mechanisms, in preference order: + +1. SYNC-LOOP DEFERRED D2H (contained): `EngineCore::step` + (src/vllm/v1/engine/core.cpp:150-200) already supports depth-2 + batch-queue pipelining via `sample_tokens_async` — but + `GPUModelRunner::sample_tokens_async` (runner.cpp:1876) degenerates to + the synchronous `ReadyModelRunnerOutput` unless `async_input_combine_` + is set (runner.cpp:411/462), which requires + `QueueSupportsAsyncInputCombine` -> backend capability — NOW TRUE on + ROCm with the real event primitives landed here. Plan: enable + VT_ASYNC_RUNNER=1 in the acceptance config, verify LLMEngine::step + drains depth-2 (the batch_queue_ path engages independent of scheduler + type), A/B paired x5 through the CLI. +2. ASYNC-SERVING PATH (fallback): measure through vllm-server with + AsyncScheduler mcb=2 — works correctly since the event fix — but first + attribute the server path's own ~40% overhead vs CLI so the comparison + isolates the lever. + +Validation either way: body-content coherence per arm (the committed +rule), engagement witness from rocpd kernel symbols, and paired deltas +under matched host load recorded per rep. + +## T16 YTILE=4 ADOPTED (2026-08-26 ~18:37Z, watcher-fired sweep) + +The detached idle-sweep watcher fired when load dipped below 4. Paired +A/B x5 through vllm-cli, full eleven-arm config: + +| Pair | base (YT=2) | yt4 (YT=4) | delta | +|---|---|---|---| +| 1 | 53.125 | 55.127 | +3.8% | +| 2 | 52.950 | 54.276 | +2.5% | +| 3 | 52.474 | 53.836 | +2.6% | +| 4 | 52.485 | 53.597 | +2.1% | +| 5 | 53.904 | 53.910 | +0.01% | + +ON wins 5/5. Base median 52.950, YT4 median 53.910 (+1.8%). Output +BIT-IDENTICAL (separate coherence check, 128 tokens, seed 0). Decision +rule (adopt iff ON wins >=4/5) satisfied. Default changed from YT=2 to +YT=4 in WvCfg (rocm_skinny_gemm.hip:169). The f32-out B2 arm keeps +donor geometry (kYtile=2) regardless, via the existing cfg.yt!=2 guard. +Gate: 16/16, 839 assertions. + +Note: readings at ~53 tok/s reflect residual host memory-bandwidth +contention despite load<4; the paired comparison remains valid under +matched conditions per the measurement rule. diff --git a/src/vt/rocm/rocm_skinny_gemm.hip b/src/vt/rocm/rocm_skinny_gemm.hip index 0774b523ed..d644aa640e 100644 --- a/src/vt/rocm/rocm_skinny_gemm.hip +++ b/src/vt/rocm/rocm_skinny_gemm.hip @@ -166,7 +166,7 @@ int DeviceCuCount(int device) { // so the three hot GDN shapes can be swept on this part. Read once. namespace { struct WvCfg { - int yt = 2; + int yt = 4; // T16 sweep: YT=4 wins 5/5 paired (+1.8% median), bit-identical int prgrp = 0; }; inline WvCfg WvCfgFromEnv() { From afd6d5fbb93d1e6f60f2e553f90f4c0649ab300a Mon Sep 17 00:00:00 2001 From: ghazni Date: Fri, 4 Sep 2026 19:01:31 +0000 Subject: [PATCH 3/4] fix(rocm): test N % effective-YTILE in the skinny-GEMM gate, not N % 2 MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The T16 default flip (yt=4) left the dispatch gate at the donor's N%2, documented for YTILE=2. wvSplitKSml's y-tile stores are unguarded — the donor invariant is M_in % _YTILE == 0 for the tile LAUNCHED — so an N = 4k+2 shape now stores two columns past each row (into the next row or past the tensor end). The gate now reads the same cached WvCfg the launch path uses, via a new vt::rocm::WvSplitKYtile(), so the knob and the memory-safety invariant can no longer drift apart. Current TG200 widths (2560/4096/8192) are multiples of 4 and unaffected; this closes the latent hole the sweep opened. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:OMEN-ALPHA [OMP] --- src/vt/rocm/rocm_matmul_hipblaslt.hip | 14 ++++++++++---- src/vt/rocm/rocm_skinny_gemm.hip | 10 ++++++++++ 2 files changed, 20 insertions(+), 4 deletions(-) diff --git a/src/vt/rocm/rocm_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index a31e44fc2b..af816a4c55 100644 --- a/src/vt/rocm/rocm_matmul_hipblaslt.hip +++ b/src/vt/rocm/rocm_matmul_hipblaslt.hip @@ -619,6 +619,10 @@ void MatmulKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) { void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, int N, int K, int device); +// Effective YTILE after the T16 sweep knobs (VT_WVSPLIT_YTILE, default 4) — +// the gate must test the SAME tile the kernel will run. +int WvSplitKYtile(); + // Row-major trick: gemm(OP_T, OP_N, N, M, K, B, K, A, a_rs, C, N) // BLAS: C = op(A)*op(B) with opA=T => A is KxN in col form = row B[N,K] // opB=N, B is KxM col = row A[M,K] with ld=a_rs @@ -673,14 +677,16 @@ void MatmulBTKernelRocm(Queue& q, Tensor& out, const Tensor& a, const Tensor& b) // * N > 8 (upstream `m > 8`, utils.py:181 — the feature-dim lower bound; // at N==1 the kernel's y=1 store writes C[1] past the end of a 1-wide // output); - // * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0`; YTILE=2 — - // the last wave's y=1 store lands at C[N] on odd N); + // * N % YTILE == 0 (skinny_gemms.cu:1217 `M_in % _YTILE == 0` — the last + // wave's stores land past C[N] otherwise; the tile is the T16-effective + // one, NOT the donor's fixed 2: with the YT=4 default an N ≡ 2 mod 4 + // shape would write two columns past the row under the old N%2 gate); // * wave32 arch only (the gfx9 ROW_BCAST reduction arm is not ported); // * K % 8 == 0 and the LDS-fit (K*M activation fits the staging buffer); // everything else stays on the BLAS path. if (bf16 && out.dtype == DType::kBF16 && M >= 1 && M <= 4 && (K % 8) == 0 && - N > 8 && (N % 2) == 0 && a.stride[0] == K && K * M <= 32768 && - SkinnyGemmEnabled() && + N > 8 && (N % vt::rocm::WvSplitKYtile()) == 0 && a.stride[0] == K && + K * M <= 32768 && SkinnyGemmEnabled() && vt::rocm::SkinnyGemmArchOk(q.device.index, vt::rocm::DeviceArchName)) { WvSplitKBT(s, out.data, a.data, b.data, static_cast(M), static_cast(N), static_cast(K), q.device.index); diff --git a/src/vt/rocm/rocm_skinny_gemm.hip b/src/vt/rocm/rocm_skinny_gemm.hip index d644aa640e..0ddb3c83d6 100644 --- a/src/vt/rocm/rocm_skinny_gemm.hip +++ b/src/vt/rocm/rocm_skinny_gemm.hip @@ -252,6 +252,16 @@ void WvSplitKBTLaunch(hipStream_t s, void* out, const void* a, const void* b, } // namespace +// Effective YTILE after the T16 sweep knobs, cached identically to the launch +// path's static cfg. The dispatch gate (rocm_matmul_hipblaslt.hip) tests +// N % YTILE == 0 against THIS value: the kernel's unguarded y-tile stores are +// memory-safe only for N a multiple of the tile actually launched, and the +// T16 default flip to 4 made the old N%2 gate too weak. +int WvSplitKYtile() { + static const WvCfg cfg = WvCfgFromEnv(); + return cfg.yt; +} + // out[M,N] = a[M,K] @ b[N,K]^T, bf16 in/out. Only called for the decode-skinny // gate below (M in 1..4, K%8==0); the caller keeps every other shape on the // BLAS path. From 98435184a5fd0949d0364f0265f406e178942cdf Mon Sep 17 00:00:00 2001 From: ghazni Date: Sun, 6 Sep 2026 09:03:19 +0000 Subject: [PATCH 4/4] fix(rocm): restore wvSplitK launch error check and M refusal Restores the hipGetLastError check after the wvSplitK kernel launch and the throw for unsupported M values, both deleted by the YTILE=4 adoption without mention in the PR body. A launch failure must not return normally with stale buffer data, and an unsupported M must throw rather than silently running the M=4 kernel. Addressing localai-org-maint-bot review on #2787. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:GLM-5-2 [OMP] --- src/vt/rocm/rocm_skinny_gemm.hip | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/src/vt/rocm/rocm_skinny_gemm.hip b/src/vt/rocm/rocm_skinny_gemm.hip index 0ddb3c83d6..c77ee1250c 100644 --- a/src/vt/rocm/rocm_skinny_gemm.hip +++ b/src/vt/rocm/rocm_skinny_gemm.hip @@ -214,14 +214,18 @@ void WvSplitKBTDispatch(hipStream_t s, void* out, const void* a, const void* b, default: wvSplitKSml<3, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; } break; - default: + case 4: switch (yt) { case 1: wvSplitKSml<4, 1, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; case 2: wvSplitKSml<4, 2, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; default: wvSplitKSml<4, 4, OutT><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); break; } break; + default: + throw std::runtime_error("vt rocm: wvSplitK unsupported M=" + std::to_string(M)); } + if (hipGetLastError() != hipSuccess) + throw std::runtime_error("vt rocm: wvSplitK launch failed"); } template