From 7cd028625196ff3406770343c21323ad088ccaf5 Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Sat, 15 Aug 2026 07:49:11 +0000 Subject: [PATCH 1/2] =?UTF-8?q?perf(cuda):=20cp.async=20double-buffered=20?= =?UTF-8?q?multicol=20int4=20GEMV=20=E2=80=94=20measured=20NO-GO,=20defaul?= =?UTF-8?q?t=20off?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Re-test of the sota/TRT-LLM roadmap's #1 lever (cp.async weight prefetch + L2::evict_first) on the current multicol NC=4 M=1 decode GEMV, prompted by fresh stall data showing Long-Scoreboard is now the top stall (2.35) on this kernel — unlike the prior narrow-kernel no-go (#978 follow-on) which was register/warp-bound. New entry matmul_nbits_gemv_f16_general_bs_wide_multicol_cpasync double-buffers each lane's WIDE_NC 128-bit weight words through shared memory with cp.async.cg.shared.global.L2::evict_first, prefetching K-tile k+1 while tile k computes. Keeps the exact per-lane depth striding + per-column fp32 accumulation order, so output is BYTE-IDENTICAL (glm greedy tokens 0 diffs over 160 tokens, generic + repetitive). SM80+ with a byte-identical direct-load #else fallback (Rule 11). Dispatch gate use_gemv_cpasync(cc): SM80+, int4/block!=32 multicol only, DEFAULT OFF, opt-in via ONNX_GENAI_GEMV_CPASYNC=1. qwen (block-32) and the default decode path are unchanged. Result: NO-GO. glm decode 199.9 -> 191.6 tok/s (-4.2%). ncu on the gate_up matrix: cp.async DID hide global-load latency (Long-Scoreboard 2.35 -> 1.88, -20%) exactly as designed, but it is a NET LOSS — the cvta-to-shared + pipeline bookkeeping raised registers 64 -> 71, dropping achieved occupancy 41.6% -> 33.3%, and the added smem round-trip pushed kernel time 33.9 -> 40.2 us (+18%) and DRAM 1.78 -> 1.50 TB/s. At M=1 this GEMV hides load latency primarily via warp-level parallelism (occupancy), so the register cost of a software prefetch pipeline outweighs the latency it hides. This is the third convergent prefetch negative (register depth-4, prior cp.async, this multicol cp.async): the M=1 decode GEMV is occupancy-bound, not prefetch-bound. Kept behind the env flag as a reproducible negative result; no regression ships by default. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../src/kernels/matmul_nbits.rs | 297 +++++++++++++++++- 1 file changed, 294 insertions(+), 3 deletions(-) diff --git a/crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs b/crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs index 6439208a61..e4306fc7de 100644 --- a/crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs +++ b/crates/onnx-runtime-ep-cuda/src/kernels/matmul_nbits.rs @@ -113,6 +113,15 @@ const GEMV_F16_GENERAL_BS_SPLITK_WIDE_ENTRY: &str = "matmul_nbits_gemv_f16_gener /// unchanged). Non-split-K only (used on the wide, occupancy-filled gate_up). const GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY: &str = "matmul_nbits_gemv_f16_general_bs_wide_multicol"; +/// cp.async double-buffered counterpart of +/// [`GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY`]. Same launch geometry, ABI and +/// byte-identical output, but each lane's next K-tile of WIDE_NC 128-bit weight +/// words is prefetched through shared memory with `cp.async.cg` + L2 evict-first +/// while the current tile computes, to hide the Long-Scoreboard global-load +/// stall. SM80+ only; opt-in via `ONNX_GENAI_GEMV_CPASYNC=1` (default off — kept +/// as a reproducible measurement until proven a win). +const GEMV_F16_GENERAL_BS_WIDE_MULTICOL_CPASYNC_ENTRY: &str = + "matmul_nbits_gemv_f16_general_bs_wide_multicol_cpasync"; /// Output columns each warp emits in the register-blocked wide GEMV. Must match /// `#define WIDE_NC` in the CUDA source. A 256-thread CTA (8 warps) therefore /// covers `8 * GEMV_F16_WIDE_MULTICOL_NC` columns. @@ -915,6 +924,44 @@ __device__ __forceinline__ float warp_sum(float value) return value; } +// ---- cp.async software-pipeline primitives (SM80+) ------------------------- +// `cp.async` streams a 128-bit global weight load straight into shared memory +// without occupying a destination register while the ~300-cycle HBM transfer is +// in flight, so a lane can prefetch its NEXT K-tile weights while the current +// tile's LOP3 dequant + FMA run. This targets the Long-Scoreboard global-load +// stall that is the top stall reason of the multicol M=1 GEMV. The `.cg` variant +// bypasses L1, and the `L2::evict_first` cache policy tells L2 to drop these +// (single-use) weight bytes immediately so the reusable activation vector stays +// resident (the Marlin `cp_async4_stream` idiom). Emitted only for compute +// capability >= 8.0; every caller pairs it with a byte-identical `#else` +// direct-load fallback so the same source still runs on < SM80 (Rule 11). +#if (__CUDA_ARCH__ >= 800) +#define ONNX_GENAI_CP_ASYNC 1 +// Async-copy one aligned 16-byte (uint4) weight word global->shared, hinting L2 +// to evict it first (weights are single-use; the activation is what we keep hot). +__device__ __forceinline__ void cp_async_cg16_evict(unsigned smem_addr, const void* gmem_ptr) +{ + asm volatile( + "{\n" + " .reg .b64 p;\n" + " createpolicy.fractional.L2::evict_first.b64 p, 1.0;\n" + " cp.async.cg.shared.global.L2::cache_hint [%0], [%1], 16, p;\n" + "}\n" ::"r"(smem_addr), + "l"(gmem_ptr)); +} +// Close the current async-copy group so it can be waited on independently. +__device__ __forceinline__ void cp_async_commit() +{ + asm volatile("cp.async.commit_group;\n" ::); +} +// Block until at most `N` (immediate) async-copy groups remain in flight. +template +__device__ __forceinline__ void cp_async_wait() +{ + asm volatile("cp.async.wait_group %0;\n" ::"n"(N)); +} +#endif // __CUDA_ARCH__ >= 800 + // Fp16 GEMV bias epilogue. The fp32 accumulator is always rounded to fp16 for // the base output. When a bias is present: // * `bias_post_round == 0` (native MatMulNBits bias): add in fp32 and round @@ -4143,7 +4190,228 @@ extern "C" __global__ void matmul_nbits_gemv_f16_general_bs_wide_multicol( } } -// Model-agnostic fp16 int4/int8 prefill GEMM supporting any power-of-two +// cp.async software-pipelined counterpart of the multicol wide GEMV (int4 only). +// Keeps the EXACT per-lane depth striding, per-column fp32 accumulation order and +// dequant arithmetic of `gemv_int4_wide_lane_dot_multicol` — so the output is +// byte-identical — but double-buffers each lane's WIDE_NC 128-bit weight words +// through shared memory with `cp.async.cg` + `L2::evict_first`, prefetching the +// next K-tile while the current tile computes. On < SM80 the `#else` branch is +// the byte-identical direct-load multicol loop (Rule 11 portability). +extern "C" __global__ void matmul_nbits_gemv_f16_general_bs_wide_multicol_cpasync( + const __half* __restrict__ activation, + const unsigned char* __restrict__ packed, + const void* __restrict__ scales, + const unsigned char* __restrict__ zero_points, + const __half* __restrict__ bias, + __half* __restrict__ output, + const int k, + const int n, + const int block_size, + const int k_blocks, + const int blob_size, + const int zp_row_bytes, + const int scales_fp16, + const int bias_post_round, + const int bits) +{ + (void)bits; // int4-only entry; dispatch guarantees bits == 4. + const int tid = (int)threadIdx.x; + const int lane = tid & 31; + const int warp = tid >> 5; + const int warps_per_block = (int)blockDim.x >> 5; + const long col_base = ((long)blockIdx.x * warps_per_block + warp) * (long)WIDE_NC; + + float values[WIDE_NC]; + bool valid[WIDE_NC]; + long col_kb[WIDE_NC]; +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + values[c] = 0.0f; + const long column = col_base + c; + valid[c] = (column < n); + col_kb[c] = column * (long)k_blocks; + } + + const int depth0 = lane * 32; + const int warp_stride = 32 * 32; // 1024 + +#if defined(ONNX_GENAI_CP_ASYNC) + // Per-warp double buffer: [stage][WIDE_NC columns][32 lanes] of uint4. + // 8 warps * 2 * 4 * 32 * 16B = 32 KiB / CTA. Lane is the fastest index so + // each thread owns a distinct 16-byte slot (no bank-conflict cross-talk) and + // reads back only its own slot, so cp.async.wait alone orders the copy. + __shared__ uint4 wbuf[8][2][WIDE_NC][32]; + + // Prefetch a full K-tile (all WIDE_NC weight words for one `depth`) into a + // shared-memory slot with cp.async, or fall through when the tile is out of + // range for this lane. + auto prefetch_tile = [&](int stage, int depth) { + if (depth + 32 > k) { + return; + } + const int block = depth / block_size; + const int within = depth - block * block_size; +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + const long base = (col_kb[c] + block) * (long)blob_size; + const unsigned smem = + (unsigned)__cvta_generic_to_shared(&wbuf[warp][stage][c][lane]); + cp_async_cg16_evict(smem, packed + base + (within >> 1)); + } + cp_async_commit(); + }; + + // Prologue: launch the first tile (t = 0). + prefetch_tile(0, depth0); + + int depth = depth0; + int t = 0; + while (depth + 32 <= k) { + const int stage = t & 1; + const int next_depth = depth + warp_stride; + const bool has_next = (next_depth + 32 <= k); + if (has_next) { + prefetch_tile((t + 1) & 1, next_depth); + cp_async_wait<1>(); // keep the just-issued next tile in flight + } else { + cp_async_wait<0>(); // drain: this is the last in-flight tile + } + __syncwarp(); + + const int block = depth / block_size; + float scale[WIDE_NC]; + unsigned int sub2[WIDE_NC]; + uint4 w[WIDE_NC]; +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + w[c] = wbuf[warp][stage][c][lane]; + if (scales_fp16) { + scale[c] = __half2float(reinterpret_cast(scales)[col_kb[c] + block]); + } else { + scale[c] = reinterpret_cast(scales)[col_kb[c] + block]; + } + sub2[c] = int4_zero_point_sub2( + int4_block_zero_point(zero_points, col_base + c, block, zp_row_bytes)); + } + + float a8[8]; +#pragma unroll + for (int s = 0; s < 4; ++s) { + decode_activation8(activation + depth + s * 8, a8); +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + const unsigned int word = + (s == 0) ? w[c].x : (s == 1) ? w[c].y : (s == 2) ? w[c].z : w[c].w; + values[c] += scale[c] * dot_int4x8_f16_sub_act(word, a8, sub2[c]); + } + } + depth = next_depth; + ++t; + } +#else + // < SM80: byte-identical direct-load multicol main loop. + int depth = depth0; + while (depth + 32 <= k) { + const int block = depth / block_size; + const int within = depth - block * block_size; + uint4 w[WIDE_NC]; + float scale[WIDE_NC]; + unsigned int sub2[WIDE_NC]; +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + const long base = (col_kb[c] + block) * (long)blob_size; + w[c] = *reinterpret_cast(packed + base + (within >> 1)); + if (scales_fp16) { + scale[c] = __half2float(reinterpret_cast(scales)[col_kb[c] + block]); + } else { + scale[c] = reinterpret_cast(scales)[col_kb[c] + block]; + } + sub2[c] = int4_zero_point_sub2( + int4_block_zero_point(zero_points, col_base + c, block, zp_row_bytes)); + } + float a8[8]; +#pragma unroll + for (int s = 0; s < 4; ++s) { + decode_activation8(activation + depth + s * 8, a8); +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + const unsigned int word = + (s == 0) ? w[c].x : (s == 1) ? w[c].y : (s == 2) ? w[c].z : w[c].w; + values[c] += scale[c] * dot_int4x8_f16_sub_act(word, a8, sub2[c]); + } + } + depth += warp_stride; + } +#endif // ONNX_GENAI_CP_ASYNC + + // Partial trailing chunk (`depth < k < depth + 32`): direct-load tail, + // identical to the multicol device fn so the K-tail stays byte-identical. + if (depth < k) { + const int block = depth / block_size; +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + if (!valid[c]) { + continue; + } + float scale; + if (scales_fp16) { + scale = __half2float(reinterpret_cast(scales)[col_kb[c] + block]); + } else { + scale = reinterpret_cast(scales)[col_kb[c] + block]; + } + const int zero_point = + int4_block_zero_point(zero_points, col_base + c, block, zp_row_bytes); + const unsigned int sub2 = int4_zero_point_sub2(zero_point); + const long blob_base = (col_kb[c] + block) * (long)blob_size; + for (int off = 0; off < 32 && depth + off < k; off += 8) { + const int d = depth + off; + const int within = d - block * block_size; + const int valid_n = min(8, k - d); + const unsigned int packed_word = + *reinterpret_cast(packed + blob_base + (within >> 1)); + if (valid_n == 8) { + values[c] += scale * dot_int4x8_f16_sub(packed_word, activation + d, sub2); + } else { + float partial = 0.0f; +#pragma unroll + for (int i = 0; i < 8; ++i) { + if (i < valid_n) { + const int q = (int)((packed_word >> (i * 4)) & 15u) - zero_point; + partial += (float)q * __half2float(activation[d + i]); + } + } + values[c] += partial * scale; + } + } + } + } + +#pragma unroll + for (int c = 0; c < WIDE_NC; ++c) { + const float reduced = warp_sum(values[c]); + const long column = col_base + c; + if (lane == 0 && column < n) { + output[column] = fold_bias_f16(reduced, bias, column, bias_post_round); + } + } +} + + // block_size. Identical 16x16 tiling and fp32 accumulation as the tuned // block-32 GEMM, but the reduction walks K in fixed 32-wide tiles and derives // the block index from the real block_size (block = depth / block_size), so the @@ -4439,6 +4707,22 @@ fn use_gemv_wide_multicol() -> bool { ) } +/// Whether the multicol wide GEMV should take the cp.async double-buffered entry +/// ([`GEMV_F16_GENERAL_BS_WIDE_MULTICOL_CPASYNC_ENTRY`]) instead of the plain +/// direct-load multicol entry. `cp.async` is SM80+ only, so on older devices this +/// always returns false and the byte-identical direct-load kernel is used +/// (Rule 11). Opt-in and DEFAULT OFF via `ONNX_GENAI_GEMV_CPASYNC=1` — the +/// prefetch is a reproducible measurement, not yet a proven default-on win. +fn use_gemv_cpasync(compute_capability: (u32, u32)) -> bool { + if compute_capability.0 < 8 { + return false; + } + matches!( + std::env::var("ONNX_GENAI_GEMV_CPASYNC").ok().as_deref(), + Some("1") | Some("true") | Some("on") + ) +} + fn use_f16_symmetric_splitk( k: usize, n: usize, @@ -7210,7 +7494,11 @@ impl MatMulNBitsKernel { } } else if use_gemv_wideload(self.bits, self.block_size, self.k) { if use_gemv_wide_multicol() { - GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY + if use_gemv_cpasync(capabilities.compute_capability()) { + GEMV_F16_GENERAL_BS_WIDE_MULTICOL_CPASYNC_ENTRY + } else { + GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY + } } else { GEMV_F16_GENERAL_BS_WIDE_ENTRY } @@ -7301,7 +7589,9 @@ impl MatMulNBitsKernel { (threads / 32) as usize / GENERAL_BS_SPLITK } else if use_scales_f16_splitk { (threads / 32) as usize / GEMV_F16_SCALES_F16_ZP_SPLITK - } else if entry == GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY { + } else if entry == GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY + || entry == GEMV_F16_GENERAL_BS_WIDE_MULTICOL_CPASYNC_ENTRY + { // Each warp emits WIDE_NC columns, so a `threads/32`-warp CTA // covers `warps * WIDE_NC` output columns; the grid shrinks by // WIDE_NC accordingly. @@ -7337,6 +7627,7 @@ impl MatMulNBitsKernel { || entry == GEMV_F16_GENERAL_BS_WIDE_ENTRY || entry == GEMV_F16_GENERAL_BS_SPLITK_WIDE_ENTRY || entry == GEMV_F16_GENERAL_BS_WIDE_MULTICOL_ENTRY + || entry == GEMV_F16_GENERAL_BS_WIDE_MULTICOL_CPASYNC_ENTRY { builder.arg(&bits); } From 2002ad90559bb02e8d108c57fca7d971fda1449a Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Sat, 15 Aug 2026 07:50:12 +0000 Subject: [PATCH 2/2] =?UTF-8?q?docs(squad):=20decision=20drop=20=E2=80=94?= =?UTF-8?q?=20cp.async=20GEMV=20prefetch=20measured=20NO-GO?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fresh-data re-test of TRT-LLM roadmap PORT-3 on the current multicol kernel: byte-identical but -4.2% (occupancy 41.6->33.3% from +7 registers outweighs the 20% Long-Scoreboard reduction). Third convergent prefetch negative; M=1 GEMV is occupancy-bound. Recommends dropping PORT-3/PORT-5 and pointing at PORT-1 (occupancy via cheaper dequant) or landing #996 fp16 + lm_head as the honest path to beat ORT base. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../decisions/inbox/deckard-cpasync-nogo.md | 96 +++++++++++++++++++ 1 file changed, 96 insertions(+) create mode 100644 .squad/decisions/inbox/deckard-cpasync-nogo.md diff --git a/.squad/decisions/inbox/deckard-cpasync-nogo.md b/.squad/decisions/inbox/deckard-cpasync-nogo.md new file mode 100644 index 0000000000..247b0e48a6 --- /dev/null +++ b/.squad/decisions/inbox/deckard-cpasync-nogo.md @@ -0,0 +1,96 @@ +# Decision drop — cp.async weight prefetch on the M=1 int4 GEMV: measured NO-GO (fresh data) + +**Author:** Deckard (Systems Dev, CUDA/decode-perf) +**Branch:** `squad/int4-gemv-cpasync-v2` (off main `3f826936`) +**Date:** 2026-08-15 +**Status:** DRAFT PR — reproducible NEGATIVE result, default OFF, byte-identical +**Refs:** TRT-LLM port plan PORT-3; prior no-go `2131046e` (narrow kernel); #986/#992/#996 + +## TL;DR +The sota/TRT-LLM roadmap flagged **cp.async double-buffered weight prefetch + +`L2::evict_first`** as the #1 lever (15–30%, byte-identical). I built it on the +current multicol NC=4 GEMV and **measured it: NO-GO, −4.2% glm decode**. It is +byte-identical and it *does* hide load latency as designed — but it is +occupancy-negative, which dominates at M=1. **Recommend dropping PORT-3 (and +PORT-5 PDL, same class).** Fresh data + mechanism below. + +## Why I re-tested (the prior no-go was on a different kernel) +cp.async was already a measured no-go once (`2131046e`, follow-on to #978) — but +on the **narrow** pre-multicol kernel, which was register/warp-bound at ~80% +occupancy and 8–12% DRAM. The current kernel is very different, so I re-profiled +first. Fresh ncu on the **current default fp32 multicol** gate_up GEMV (glm, +`--graph-profiling node`, launch-skip 240): + +| stall reason | warps stalled / issue-active | +|---|---| +| **Long Scoreboard (global-load latency)** | **2.35** ← cp.async's target, now #1 | +| Wait (math/FMA latency) | 1.80 | +| Short Scoreboard (smem/L1) | 0.20 | +| No instruction | 0.15 | + +Long-Scoreboard IS now the top stall (precondition for cp.async **met**, unlike +the prior no-go) — so a genuine re-test was warranted. + +## What I built +`matmul_nbits_gemv_f16_general_bs_wide_multicol_cpasync`: double-buffers each +lane's WIDE_NC=4 128-bit weight words through shared memory with +`cp.async.cg.shared.global.L2::evict_first` (the Marlin `cp_async4_stream` +idiom, 16-byte), prefetching K-tile k+1 while tile k computes. Exact per-lane +depth striding + per-column fp32 accumulation order preserved ⇒ byte-identical. +SM80+ with a byte-identical direct-load `#else` fallback (Rule 11). Gate +`use_gemv_cpasync(cc)`: SM80+, int4/block≠32 multicol only, **DEFAULT OFF**, +opt-in `ONNX_GENAI_GEMV_CPASYNC=1`. + +## Result — NO-GO +**glm decode: 199.9 → 191.6 tok/s (−4.2%)**, greedy tokens BYTE-IDENTICAL +(0 diffs / 160 tokens, generic + repetitive). qwen unaffected (block-32 never +routes here). f64 numerics 7/7 (default path unchanged). fmt + clippy clean. + +### ncu — cp.async worked, but is occupancy-negative +| metric | direct-load (default) | **cp.async** | delta | +|---|---|---|---| +| kernel time | 33.9 µs | **40.2 µs** | **+18%** | +| DRAM | 1.78 TB/s | 1.50 TB/s | −16% | +| **Long Scoreboard** | 2.35 | **1.88** | **−20% (hid load latency ✓)** | +| Wait (math latency) | 1.80 | 2.07 | +15% | +| registers/thread | 64 | **71** | +7 | +| **achieved occupancy** | 41.6% | **33.3%** | **−8 pp** | + +**Mechanism:** cp.async did exactly what it promises — it cut the global-load +stall 20%. But the `cvta`-to-shared address + pipeline bookkeeping raised +registers 64→71, dropping occupancy 41.6%→33.3%. At M=1 this GEMV hides load +latency **primarily via warp-level parallelism (occupancy)**, not per-warp +prefetch depth; losing 8 pp of occupancy removes more latency-hiding than the +prefetch adds. Net: kernel +18%, DRAM −16%, e2e −4.2%. The short (~4-trip) K-loop +and the kernel's existing 4-way weight ILP (all WIDE_NC loads issued up front) +leave little for a software pipeline to add. + +This is the **third convergent prefetch negative**: (1) register depth-4 prefetch +(occupancy loss), (2) prior cp.async on the narrow kernel, (3) this multicol +cp.async. **The M=1 decode GEMV is occupancy-bound, not prefetch-bound.** + +## Recommendation on the roadmap +- **PORT-3 (cp.async prefetch): DROP.** Disproven with fresh, mechanistically + explained data on the current kernel. +- **PORT-5 (PDL / programmatic launch): deprioritize.** Same latency-hiding class; + the kernel isn't launch-bubble bound (issue-active 69.5%). +- **The real lever is OCCUPANCY**, which is register-bound at 64 (WIDE_NC=4). + Anything that cuts the kernel's register/instruction footprint *without* losing + the multicol activation reuse would raise occupancy → more warp-parallel latency + hiding. That points at **PORT-1** (offline bias+interleave → cheaper runtime + dequant, fewer instrs/regs) as the better byte-identical next lever — but PORT-1 + needs an **offline weight-layout repack + a versioned layout flag** (loader-side, + model-artifact-compat risk), a much larger surface than a pure kernel change. +- **Strategic honesty:** per #996's measurements, this GEMV is already at ~95% of + ORT's DRAM (2.29 vs 2.42 TB/s) and near ORT's kernel time (fp16 26.4 vs 24.9 µs). + The remaining e2e gap to ORT ~250 base is dominated by **other ops (lm_head)**, + owned separately by deckard-4. The banked honest wins to beat ORT base are the + **fp16 GEMV (#996, +6.7%)** + **lm_head cuBLASLt**, not further GEMV prefetch. + +## Ask +- Coordinator: accept the NO-GO; decide whether to **merge this as a documented + negative** (default-off, like `2131046e`) or take the decision drop and close + the PR. Either is fine — no regression ships by default. +- Next: your call between (a) me scoping **PORT-1** (offline repack + layout flag, + the occupancy lever) vs (b) me landing #996 fp16 default-on after Chew's gate and + handing the remaining gap to the lm_head work.