From 813b5fa52bd3ff1d0562263c87d59a6d42c8285e Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Thu, 20 Aug 2026 04:04:34 +0000 Subject: [PATCH 1/2] fix(ci): account for the deleted ONNX_GENAI_GEMV_FOLDSCALE knob verify_documented_env_vars is failing on main. #1527 deleted the fold-scale GEMV and gemv_foldscale_enabled() after the +2.7% failed to reproduce, but dense-decode-megakernel-feasibility.md still names the env var that selected it, so the guard sees a documented knob no crate reads -- which is exactly what it exists to catch. The prose is already honest: 9.2 records the re-measurement (a 2.9% cost against a 28.62/28.62 baseline) and the repro note says the entries were removed. Two of the three mentions carry that caveat; the Phase C lever paragraph did not, and now does. So this is an allowlist entry, not a rewrite. The script requires each KNOWN_UNIMPLEMENTED name to stay accompanied by the statement in the prose, and after this commit all three mentions carry it. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .github/scripts/verify_documented_env_vars.py | 4 ++++ docs/research/dense-decode-megakernel-feasibility.md | 4 +++- 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/.github/scripts/verify_documented_env_vars.py b/.github/scripts/verify_documented_env_vars.py index aa060b00c0..048f7db554 100644 --- a/.github/scripts/verify_documented_env_vars.py +++ b/.github/scripts/verify_documented_env_vars.py @@ -70,6 +70,10 @@ "dense-decode-megakernel-feasibility.md §9.2; the cp.async double-buffered " "weight-load path regressed (−13%, 4 B/lane granularity) and was reverted, " "not shipped", + "ONNX_GENAI_GEMV_FOLDSCALE": "retracted experiment knob documented in " + "dense-decode-megakernel-feasibility.md; the fold-scale GEMV it selected was " + "deleted by #1527 after the +2.7% failed to reproduce (re-measured as a 2.9% " + "cost), and the document says so where it is named", "NXRT_AUTO_INSTALL_CUDA": "referenced in prose about a proposed installer step", "NXRT_SQNBIT_PREFILL_MIN": "referenced in prose about a proposed tuning knob", "ONNX_GENAI_BASE_URL": "client-side variable read by external tooling, not by this workspace", diff --git a/docs/research/dense-decode-megakernel-feasibility.md b/docs/research/dense-decode-megakernel-feasibility.md index a03f332fef..268706aef4 100644 --- a/docs/research/dense-decode-megakernel-feasibility.md +++ b/docs/research/dense-decode-megakernel-feasibility.md @@ -739,7 +739,9 @@ full 128-bit sector — that is a from-scratch kernel, out of scope for this bou LOP3 (4×`lop3.b32` + debias) — that is *why* the ALU pipe is at 65%. The one remaining ALU lever is to **fold the per-block scale into the dequant's zp-subtract**: replace the plain `q=(code-zp)` then a separate `__hmul2(q, scale)` with a single `fma(code, scale, -zp*scale)`, dropping **4 `__hmul2` per 8 -weights** (~20% fewer fp16 ALU ops in the MAC). Env `ONNX_GENAI_GEMV_FOLDSCALE`. +weights** (~20% fewer fp16 ALU ops in the MAC). Env `ONNX_GENAI_GEMV_FOLDSCALE` — **retracted**: +the kernel this selected was deleted by #1527 after the win failed to reproduce, so the variable +no longer exists and setting it does nothing. See the re-measurement in §9.2. | variant | tok/s | Δ | dominant-kernel µs | DRAM% | | --- | --- | --- | --- | --- | From cdfeeb8f584701c3d4777479bac61aac8f1fef2d Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Thu, 20 Aug 2026 04:10:08 +0000 Subject: [PATCH 2/2] Point the fold-scale retraction at 9.3, where the re-measurement is The caveat sits inside 9.2, so 'see 9.2' left the reader in the section they were already in -- and the next numbers there are the retracted +2.7% table. The re-measurement that retracts it is in 9.3. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/research/dense-decode-megakernel-feasibility.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/research/dense-decode-megakernel-feasibility.md b/docs/research/dense-decode-megakernel-feasibility.md index 268706aef4..1431fb815a 100644 --- a/docs/research/dense-decode-megakernel-feasibility.md +++ b/docs/research/dense-decode-megakernel-feasibility.md @@ -741,7 +741,7 @@ LOP3 (4×`lop3.b32` + debias) — that is *why* the ALU pipe is at 65%. The one separate `__hmul2(q, scale)` with a single `fma(code, scale, -zp*scale)`, dropping **4 `__hmul2` per 8 weights** (~20% fewer fp16 ALU ops in the MAC). Env `ONNX_GENAI_GEMV_FOLDSCALE` — **retracted**: the kernel this selected was deleted by #1527 after the win failed to reproduce, so the variable -no longer exists and setting it does nothing. See the re-measurement in §9.2. +no longer exists and setting it does nothing. See the re-measurement in §9.3. | variant | tok/s | Δ | dominant-kernel µs | DRAM% | | --- | --- | --- | --- | --- |