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 7af460e65..b5f7320a8 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 @@ -229,84 +229,3 @@ 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 c3995c167..c14913aa0 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 2bd04cdf5..ec37c6976 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 bea2c21e7..e735fc9e7 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_matmul_hipblaslt.hip b/src/vt/rocm/rocm_matmul_hipblaslt.hip index a31e44fc2..af816a4c5 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 b4d7dd755..c77ee1250 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,31 +157,69 @@ 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 = 4; // T16 sweep: YT=4 wins 5/5 paired (+1.8% median), bit-identical + 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); + 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; case 4: - wvSplitKSml<4><<>>(K, K, K, N, B, A, C, wvPrGrp, cu); + 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)); @@ -172,4 +228,58 @@ void WvSplitKBT(hipStream_t s, void* out, const void* a, const void* b, int M, i 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 + +// 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. +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