perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in) by justinchuby · Pull Request #996 · justinchuby/onnx-genai

justinchuby · 2026-08-15T06:43:22Z

Summary

Opt-in fp16 int4 M=1 wide GEMV that matches ORT's MatMulFloat4BitsKernelM1 arithmetic — full-fp16 (half2) K accumulate, fp32 only in the final cross-lane shuffle — stacked on the merged #992 multicol NC=4 activation-reuse layout. This is the honest base-vs-base, fp16-vs-fp16 lever: ORT's int4 M=1 kernel is fp16, our merged path is fp32-accumulate, which structurally caps us below ORT. Ships permanently opt-in via ONNX_GENAI_GEMV_FP16=1; the default (fp32) path is untouched and byte-identical to main.

Merge disposition (Chew 🟡 land-as-opt-in): the fp16 numerics are sound (matches ORT's own fp16 scheme, error well under bound) but the kernel is not byte-identity-safe vs the fp32 path — glm has intermittent knife-edge argmax flips (see Accuracy gate) — so fp16 stays opt-in and is not a default-on candidate. This deliberately does not flip the default.

Result (glm-4-9b-int4, H200, `--steady --tokens 160 --decode-skip 40 --runs 3`)

variant	tok/s
multicol NC=4 #992 (fp32)	199.4
fp16 full-accumulate (this PR)	212.8 (+6.7%)

ORT base ≈ 250–252 → native-vs-ORT base gap 1.30× → 1.18×.

ncu (glm gate_up, `--graph-profiling node`, launch-skip 240)

	fp32 multicol	fp16 (this PR)	ORT
kernel	34.3 µs	26.4 µs	24.9 µs
DRAM	1.76 TB/s	2.29 TB/s	2.42 TB/s
regs/occ	64 / 41%	64 / 41%	32 / —

Kernel is now at ~ORT parity (26.4 vs 24.9 µs; 2.29 vs 2.42 TB/s).

Honest findings from the ORT source read

ORT's LOP3 dequant + prmt activation reorder idioms were already present in this kernel (int4x8_to_half2x4_sub, decode_activation8_h2). The only arithmetic delta was per-chunk-fp32 → full-fp16 accumulate.
The "full-fp16 is faster" hypothesis was DISPROVEN at the kernel level: full-fp16 vs per-chunk-fp32 is a speed tie (26.4 vs 27.0 µs; 212.8 vs 212.4 tok/s) and 2× less precise (f64 max_rel 1.16e-2 vs 6.3e-3). The limiter is the multicol WIDE_NC=4 register state + wave/tail quantization (achieved occupancy 41% << register-theoretical), not the accumulator. Shipped full-fp16 to match ORT's arithmetic exactly (cleanest "same error class as ORT" gate story).
WIDE_NC_FP16=2 (fewer cols → fewer regs) rejected: occupancy moved only 41→43% (wave-bound, not register-bound) and the kernel got slower.
Reaching ORT's 24.9 µs would need ORT's 1-col/warp 32-reg high-occupancy layout, not the multicol layout requested. The remaining e2e gap to 250 is other ops (lm_head cuBLASLt is owned separately).

Accuracy gate (opt-in, non-byte-identical by design)

f64 oracle (fp16_mixed_gemv_matches_f64_oracle_glm_decode): fp16 max_rel = 1.16e-2, 4.3× under the 5e-2 envelope (same envelope the reviewed fp32 int4 path meets). Full numerics suite 8/8. Chew confirmed: genuinely mirrors ORT's fp16 scheme, error class ≤ ORT by construction.
Token behavior (corrected — earlier "byte-identical" claim was wrong): because the fp16 K-accumulate is not bit-identical to the fp32 path, glm greedy argmax can differ vs fp32 on knife-edge (near-tie) tokens — Chew observed intermittent flips (English↔Chinese at token idx 3, a dup at idx 12; ~2 in the first 4 runs, then 0 across 70+ subsequent runs). qwen is fully byte-identical (block-32 never routes fp16). This is expected for a non-byte-identical fp16 path and is exactly why it ships opt-in gated on accuracy (≤ ORT's fp16 error class), not on bit-identity. Whether the residual glm flips are pure argmax knife-edge (deterministic given inputs) vs a low-rate kernel nondeterminism is being checked as a non-blocking follow-up.
qwen no-regression: block-32 never routes fp16 (gate lives in the block_size != 32 arm); tokens identical flag on/off.

Portability / capture

Opt-in; default-off path untouched & byte-identical to main. __hfma2 is SM53+ (no arch guard). Static grid, capture-safe.

Reviewer asks

Chew (accuracy): ✅ addressed — "≤ same-envelope-as-fp32-path" stands in for "≤ ORT" (ORT can't link into the ep-cuda test crate); full-fp16 matches ORT's arithmetic ⇒ same error class by construction.
Gaff (capture/perf): confirm capture-safety and the +6.7% (212.8) reproducibility. Note: not a default-on candidate — ships opt-in.

Co-authored-by: Copilot 223556219+Copilot@users.noreply.github.com

…opt-in) Adds an opt-in fp16 sibling of the column register-blocked wide GEMV (matmul_nbits_gemv_f16_general_bs_wide_multicol_fp16), gated by ONNX_GENAI_GEMV_FP16=1. The per-chunk MAC runs in fp16 __hfma2 (two fused MACs/instruction) instead of fp32 FFMA, matching ORT's fp16 arithmetic — the equal-conditions fp16-vs-fp16 path. PRECISION CONTRACT: only the 32 products of ONE quant chunk (single per-block scale) are summed in fp16 into a __half2 accumulator, then reduced to fp32, scaled, and accumulated across chunks in fp32. The K=4096 reduction stays fp32 (NOT a full-fp16 K accumulate). Not byte-identical to fp32 by construction -> opt-in, gated on accuracy (error <= ORT vs f64 oracle), not bit-identity. ncu (glm gate_up): 34.3->27.0us, DRAM 1.76->2.24 TB/s (ORT 2.42), L1/TEX 28->36%, SM 60->54% (now occupancy-bound at 41%). glm base decode 198->210 tok/s (+6.2%); greedy tokens identical to fp32 over 160 tokens on generic + repetitive prompts. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Adds fp16_mixed_gemv_matches_f64_oracle_glm_decode: runs the fp16 kernel (ONNX_GENAI_GEMV_FP16=1) on glm's real block-128 M=1 projection shapes (qkv/o/gate_up/down x sym+asym zp x fp16+fp32 scales) against the same f64 dequant->GEMM oracle, and holds it to the SAME justified Envelope the reviewed fp32 int4 path satisfies (no weakening). Also runs the fp32 path on each problem and reports the fp16-vs-fp32 oracle-error ratio, and asserts the fp16 output diverges from fp32 (proving the fp16 entry -- not the fp32 fallback -- was selected). Pins ONNX_GENAI_GENERAL_SPLITK=0 for deterministic entry selection. Measured: fp16 max_rel=6.3e-3 (8x under the 5e-2 bound); symmetric ~1.05-1.20x the fp32 path, asymmetric-zp 9-13x on rel but abs error stays <<bound (small output magnitude). fp16 keeps the K reduction in fp32 so it is at least as precise as any full-fp16 kernel. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

codecov · 2026-08-15T06:51:33Z

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 83.56%. Comparing base (07a0443) to head (dff4a1f).
⚠️ Report is 8 commits behind head on main.

Additional details and impacted files

@@           Coverage Diff           @@
##             main     #996   +/-   ##
=======================================
  Coverage   83.56%   83.56%           
=======================================
  Files          12       12           
  Lines        5270     5270           
  Branches     5270     5270           
=======================================
  Hits         4404     4404           
  Misses        682      682           
  Partials      184      184

Flag	Coverage Δ
cli-ort-linux	`83.52% <ø> (ø)`
cli-ort-windows	`83.11% <ø> (ø)`

Flags with carried forward coverage won't be shown. Click here to find out more.

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

github-actions · 2026-08-15T07:09:56Z

🔴 Benchmark Regression Detected

Comparison of criterion micro-benchmarks: PR head vs merge-base, measured on the same runner in the same job (base first → PR second).

ℹ️ Absolute times are informational only — they vary with runner load. The % change column is the reliable signal because both sides ran under identical conditions.

Status	Scenario	Base	PR	Change
🔴	`block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024`	57.15 µs	145.04 µs	+153.8%
🔴	`block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024`	48.32 µs	110.80 µs	+129.3%
🔴	`block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024`	506.09 µs	1.13 ms	+123.4%
🔴	`matmul/medium_generic_f32_threads=8/32x512x512`	1.04 ms	1.98 ms	+90.0%
🔴	`matmul/large_generic_bf16_threads=8/32x1024x1024`	1.34 ms	2.50 ms	+86.9%
🔴	`matmul/medium_generic_bf16_threads=8/32x512x512`	373.01 µs	687.24 µs	+84.2%
🔴	`matmul/large_generic_f32_threads=8/32x1024x1024`	3.67 ms	6.48 ms	+76.4%
🔴	`matmul/medium_generic_f16_threads=8/32x512x512`	29.78 µs	49.03 µs	+64.6%
🔴	`add/small_f32_threads=1-internal/1024`	207.9 ns	302.2 ns	+45.3%
🔴	`matmul/medium_generic_bf16_threads=1/32x512x512`	540.64 µs	776.21 µs	+43.6%
🔴	`qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline`	3.96 ms	5.47 ms	+38.3%
🔴	`matmul/large_generic_bf16_threads=1/32x1024x1024`	2.01 ms	2.70 ms	+34.4%
🔴	`matmul/medium_generic_f16_threads=1/32x512x512`	30.75 µs	40.61 µs	+32.1%
⚠️	`matmul/large_generic_f16_threads=8/32x1024x1024`	84.67 µs	109.96 µs	+29.9%
⚠️	`matmul/small_generic_f16_threads=1/1x256x256`	30.49 µs	39.53 µs	+29.6%
⚠️	`add/medium_bf16_threads=1-internal/262144`	2.61 ms	3.33 ms	+27.7%
⚠️	`add/small_f16_threads=1-internal/1024`	12.70 µs	15.78 µs	+24.2%
⚠️	`add/medium_f16_threads=1-internal/262144`	2.54 ms	3.13 ms	+23.3%
⚠️	`matmul/large_generic_f32_threads=1/32x1024x1024`	9.33 ms	11.19 ms	+20.0%
⚠️	`add/medium_f32_threads=1-internal/262144`	2.50 ms	2.99 ms	+19.4%
⚠️	`matmul/large_generic_f16_threads=1/32x1024x1024`	81.95 µs	97.60 µs	+19.1%
⚠️	`add/large_f32_threads=1-internal/4194304`	40.47 ms	47.37 ms	+17.1%
⚠️	`matmul/small_generic_bf16_threads=1/1x256x256`	32.91 µs	38.36 µs	+16.5%
✅	`add/small_bf16_threads=1-internal/1024`	12.88 µs	14.61 µs	+13.5%
✅	`add/large_bf16_threads=1-internal/4194304`	42.02 ms	46.82 ms	+11.4%
✅	`gather/medium_bf16_threads=1-internal/32768`	3.34 µs	3.67 µs	+9.9%
✅	`reduce_mean/small_f32_threads=1-internal/4096`	15.12 µs	16.42 µs	+8.6%
✅	`reduce_mean/medium_f32_threads=1-internal/65536`	249.56 µs	270.33 µs	+8.3%
✅	`sampling_latency/min_p_per_token`	234.41 µs	251.01 µs	+7.1%
✅	`tokenization/decode_tokens_per_second`	7.19 ms	7.65 ms	+6.5%
✅	`sampling_latency/top_k_per_token`	61.83 µs	65.22 µs	+5.5%
✅	`qwen3_sampling_processors/top_k_partial_selection`	163.77 µs	172.50 µs	+5.3%
✅	`tokenization/encode_tokens_per_second`	423.74 µs	442.15 µs	+4.3%
✅	`qwen3_sampling_processors/top_p_fast_after_top_k`	620.34 µs	644.80 µs	+3.9%
✅	`matmul/medium_generic_f32_threads=1/32x512x512`	2.63 ms	2.72 ms	+3.5%
✅	`add/large_f16_threads=1-internal/4194304`	43.30 ms	44.75 ms	+3.4%
✅	`gather/medium_f16_threads=1-internal/32768`	3.13 µs	3.19 µs	+2.0%
✅	`reduce_mean/large_f32_threads=1-internal/262144`	1.13 ms	1.13 ms	-0.1%
✅	`block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1`	99.50 µs	98.41 µs	-1.1%
✅	`gather/large_f32_threads=1-internal/131072`	40.66 µs	39.87 µs	-1.9%
✅	`matmul/small_generic_f32_threads=1/1x256x256`	48.52 µs	47.26 µs	-2.6%
✅	`sampling_latency/top_p_per_token`	461.67 µs	449.41 µs	-2.7%
✅	`logit_processing/seven_processor_chain_per_step`	375.11 µs	364.60 µs	-2.8%
✅	`kv_cache/alloc_dealloc_pages`	45.38 µs	43.85 µs	-3.4%
✅	`qwen3_sampling_processors/top_k_full_sort_baseline`	2.46 ms	2.37 ms	-3.6%
✅	`gather/small_f32_threads=1-internal/4096`	783.9 ns	752.9 ns	-4.0%
✅	`qwen3_sampling_processors/top_k_top_p_fast`	764.78 µs	725.58 µs	-5.1%
✅	`qwen3_sampling_processors/top_k_top_p_full_sort_baseline`	6.75 ms	6.11 ms	-9.5%
✅	`gather/medium_f32_threads=1-internal/32768`	5.41 µs	4.87 µs	-10.0%
✅	`gather/small_f16_threads=1-internal/4096`	617.6 ns	552.9 ns	-10.5%
✅	`grammar_masking/llguidance_compute_mask/32`	92.89 µs	80.99 µs	-12.8%
🟢	`gather/large_bf16_threads=1-internal/131072`	18.20 µs	15.44 µs	-15.2%
🟢	`gather/small_bf16_threads=1-internal/4096`	648.2 ns	512.4 ns	-21.0%
🟢	`sampling_latency/greedy_per_token`	4.69 µs	3.57 µs	-23.8%
🟢	`gather/large_f16_threads=1-internal/131072`	20.89 µs	15.58 µs	-25.4%
🟢	`matmul/small_generic_f16_threads=8/1x256x256`	57.99 µs	41.64 µs	-28.2%
🟢	`block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1`	672.40 µs	474.99 µs	-29.4%
🟢	`matmul/small_generic_bf16_threads=8/1x256x256`	72.43 µs	41.50 µs	-42.7%
🟢	`matmul/small_generic_f32_threads=8/1x256x256`	114.49 µs	57.45 µs	-49.8%

Visual flags: ⚠️ ≥ 15% slower, 🔴 ≥ 30% slower — calibrated against measured runner noise (~27% worst-case on multi-threaded matmul)

Host info

CPU: Apple M1 (Virtual)
Cores: 3
OS: Darwin 25.5.0 arm64
Rust: rustc 1.97.1 (8bab26f4f 2026-07-14)
Load avg: { 2.44 2.87 4.06 }

What this cannot catch

Regressions in code paths not covered by these benchmarks (e.g., end-to-end decode with a real model)
Sub-threshold regressions that compound over multiple PRs
Performance changes that only manifest under GPU execution
Latency changes in the ORT integration path (these benchmarks exercise the native Rust kernels)

Replace the per-chunk fp32 flush in the multicol fp16 int4 GEMV with a full-fp16 half2 accumulate across the lane's K chunks, folding the per-block scale in with __hfma2 and deferring fp32 to only the final cross-lane warp_sum. This mirrors ORT's MatMulFloat4BitsKernelM1 arithmetic exactly. Safe because the per-lane fp16 reduction is a wide, shallow tree: with 32 lanes striding K by 32, each lane folds only ~K/1024 chunks (~4 for K=4096), so fp16 accumulation depth is tens, not thousands — negligible mantissa loss (unlike a naive deep single-accumulator fp16 sum, which flips tokens). f64 oracle: fp16 max_rel=1.16e-2, 4.3x under the 5e-2 envelope (same error class as the reviewed fp32 int4 path). Steady glm base decode 212.8 vs fp32 199.4 tok/s. Speed is a tie with the per-chunk-fp32 version (kernel 26.4 vs 27.0us): the limiter is multicol register state + wave/tail quantization (achieved occupancy 41%), not the accumulator, so ORT's "full-fp16 is faster" holds only marginally at the kernel level. Kept behind ONNX_GENAI_GEMV_FP16=1 for the validation phase. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Reflect the final shipped kernel (full-fp16 half2 accumulate matching ORT's MatMulFloat4BitsKernelM1) and the honest findings from the ORT source read: LOP3+prmt idioms were already present; full-fp16 vs per-chunk-fp32 is a speed tie (occupancy/wave-bound, not accumulator- bound); WIDE_NC_FP16=2 rejected; f64 max_rel 1.16e-2; glm 212.8 tok/s. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…6 K-accum The const doc comment for GEMV_F16_GENERAL_BS_WIDE_MULTICOL_FP16_ENTRY and the accuracy-gate test doc comment both still described the fp16 GEMV as summing each chunk in fp16 then "reduced to fp32 and accumulated across chunks in fp32 (the K reduction stays fp32)". That contradicts the shipped kernel: the per-lane K reduction runs ENTIRELY in fp16 __half2 accumulators (per-block scale folded via __hfma2), matching ORT's MatMulFloat4BitsKernelM1, with fp32 used ONLY in the final cross-lane warp-shuffle reduction. It is safe (no token-flipping mantissa loss) because each lane strides K by 32 and folds only a handful of chunks — a wide, shallow fp16 tree of depth ~tens, not K. Comments only; no code or arithmetic change. The device-function block comment (gemv_int4_fp16_lane_dot_multicol) was already accurate and is unchanged. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… flip) Correct the default-on decision drop: withdraw the "one-liner hook to default-on fp16 GEMV later / fp16 fixes fp32 nondeterminism" plan. Chew's #996 verdict is land-as-opt-in — glm has intermittent knife-edge argmax flips vs fp32, so fp16 is not byte-identity-safe for default. Out-of-box GEMV stays fp32 (byte- identical); the honest out-of-box default-on number is ~207.7 short / ~171.7 KV2048 (fp16 OFF). 232 requires user-set ONNX_GENAI_GEMV_FP16=1 and is not the default. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

Promote the #991 cuBLASLt lm_head decode MatMul to default-on and harden its CUDA-graph capture guard from a single cached shape to a shape-keyed plan cache, so the out-of-box glm base decode banks the cuBLASLt logits projection without an env flag. Capture guard (Gaff's #991 blocker): `launch_dense_capturable` now keeps a shape-keyed `Vec<DenseGemmPlan>` (LRU, MRU-promote on hit, bounded by DENSE_PLAN_CACHE_CAP=8) instead of a single `Option<DenseGemmPlan>`. Every distinct (dtype, M, K, N) the node runs keeps its own preselected cuBLASLt algorithm + persistent workspace, so alternating shapes (prefill M>1, decode M==1, a speculative M=K verify width) all replay with no per-call heuristic query inside the captured region. The reject-cold-miss-during-capture contract is unchanged (now enforced per shape); eviction only ever runs off the non-capturing creation path, and the hot decode plan is kept at the MRU front, so a plan baked into a live captured graph is never freed. Strictly safer than the single-slot path (evicts less; same recapture-on-change discipline). Default-on: `lmhead_cublaslt_enabled()` inverts to default-TRUE-unless-disabled; set ONNX_GENAI_LMHEAD_CUBLASLT=0/false/off for the hand-GEMV escape hatch. #999 GQA flash-decoding is already default-on (its env vars are rollback-only knobs), so no flip is needed there. GEMV fp16 (#996) stays untouched (not in this base) and is a one-liner default flip once it merges. Measured (glm-4-9b-int4, H200, graph-on, fallbacks=0): - default-on short decode 207.72 tok/s (vs escape-hatch-off hand-GEMV 202.24) - default-on KV2048 171.69 tok/s, fallbacks=0 across 8 bucket-growth recaptures - lm_head cuBLASLt(default) == hand-GEMV(off) top-40 logprobs BYTE-IDENTICAL - qwen no regression (152.16 vs 153.60, within noise), fallbacks=0 - f64 oracle 7/7, GQA fp16 2/2, fmt + clippy clean Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

… flip) Correct the default-on decision drop: withdraw the "one-liner hook to default-on fp16 GEMV later / fp16 fixes fp32 nondeterminism" plan. Chew's #996 verdict is land-as-opt-in — glm has intermittent knife-edge argmax flips vs fp32, so fp16 is not byte-identity-safe for default. Out-of-box GEMV stays fp32 (byte- identical); the honest out-of-box default-on number is ~207.7 short / ~171.7 KV2048 (fp16 OFF). 232 requires user-set ONNX_GENAI_GEMV_FP16=1 and is not the default. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…1004) ## Default-on: lm_head cuBLASLt + multi-shape capture guard Promotes #991's cuBLASLt `lm_head` decode MatMul to **default-on** (out-of-box glm base), after clearing Gaff's #991 blocker: the single-shape cuBLASLt plan cache is hardened into a **shape-keyed multi-shape plan cache** so plan creation never happens inside a captured CUDA-graph region across prefill/decode/verify shape switches. Rebased onto main `ad141370` (post-#996); disjoint from #996's opt-in fp16 (which stays OFF by default). ### Capture guard — path (a): shape-keyed multi-shape plan cache #991 already had a warm-once / reject-cold-miss-during-capture contract, but for a SINGLE cached shape (a shape switch evicted + re-ran the heuristic search). This PR generalizes it to a `Vec<DenseGemmPlan>` keyed by `(dtype,M,K,N)`, MRU-ordered, bounded by `DENSE_PLAN_CACHE_CAP = 8`: - **hit** → MRU-promote, launch preselected algo (no heuristic query); - **cold miss while capturing** → `Err` (reject-before-warm, now per shape); - **cold miss not capturing** → create + insert-front + LRU-truncate. - Eviction runs **only** off the non-capturing path; the hot decode plan (M==1) stays MRU-front and is never the LRU victim. Strictly safer than the single slot (evicts far less; same recapture-on-change discipline). **Shape-transition proof** (temporary `DENSE_PLAN_SEARCH` counter, glm KV2048 run = prompt → decode → **8 KV-bucket-growth recaptures**, 8623 replays): - **exactly 2 plan searches for the whole run** — `M=3` prefill (once) + `M=1` decode (once), each on the non-capturing path; - **0 searches** during any of the 8623 replays or 8 recaptures; - **fallbacks=0** (a cold-miss-inside-capture would surface as a fallback; none did). This is the "no per-transition plan search, not just steady-state" guarantee. Counter reverted (not committed). ### Default-on flip (byte-identical wins only) - `lmhead_cublaslt_enabled()` → default-TRUE-unless-disabled; escape hatch `ONNX_GENAI_LMHEAD_CUBLASLT=0|false|off` → hand fp16 GEMV. - **#999 GQA flash-decoding already default-on** (audited — its env vars are rollback-only knobs); no flip needed. - **GEMV fp16 (#996) stays PERMANENTLY opt-in — NOT flipped.** Chew's #996 verdict is land-as-opt-in: the fp16 K-accumulate is not byte-identity-safe vs fp32 (glm knife-edge argmax flips), so it is not a default-on candidate. Out-of- box GEMV stays fp32 (byte-identical); fp16 via `ONNX_GENAI_GEMV_FP16=1`. ### Measured (glm-4-9b-int4, H200, graph-on, median, no env flags) | config | short tok/s | KV2048 tok/s | fallbacks | |---|---|---|---| | **default-on (no flags)** | **~206–208** | **~166–172** | **0** | | escape-hatch `=0` (hand GEMV) | 202.24 | — | 0 | (short 207.72 GPU3 / 206.44 GPU4; KV2048 171.69 / 168.06 / 165.59 across GPU3/6/4 — cross-GPU + shared-host variance; all fallbacks=0.) - Byte-identity: cuBLASLt lm_head (default) vs hand-GEMV (off) top-40 token-0 logprobs **byte-identical** (both select token 315). - qwen no regression (152.16 vs 153.60, noise), fallbacks=0. - f64 oracle 7/7; GQA fp16 2/2; fmt + clippy clean. **Out-of-box number is ~206–208 short / ~166–172 KV2048 and that is final for this PR.** The ~232 combined figure required fp16 GEMV opt-in ON; since fp16 ships permanently opt-in, 232 is only reachable via `ONNX_GENAI_GEMV_FP16=1`. ⚠️ Please gate with fresh Chew (byte-identity / f64 oracle) + Gaff (capture-safe, fallbacks=0, win real). --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

justinchuby and others added 3 commits August 15, 2026 06:28

docs(squad): decision drop for fp16 mixed-precision int4 GEMV

cbf6366

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

justinchuby and others added 2 commits August 15, 2026 07:17

justinchuby changed the title ~~perf(cuda): fp16 mixed-precision int4 wide GEMV — glm 199→212 tok/s, accuracy-gated (opt-in)~~ perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in) Aug 15, 2026

justinchuby marked this pull request as ready for review August 15, 2026 09:26

justinchuby merged commit ad14137 into main Aug 15, 2026
5 of 13 checks passed

justinchuby deleted the squad/int4-gemv-fp16-mixed branch August 15, 2026 09:26

Navigation Menu

Repository navigation

perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in) - #996

perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in)#996
justinchuby merged 6 commits into
mainfrom
squad/int4-gemv-fp16-mixed

justinchuby commented Aug 15, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Aug 15, 2026 •

edited

Loading

Uh oh!

github-actions Bot commented Aug 15, 2026 •

edited

Loading

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Conversation

justinchuby commented Aug 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Result (glm-4-9b-int4, H200, --steady --tokens 160 --decode-skip 40 --runs 3)

ncu (glm gate_up, --graph-profiling node, launch-skip 240)

Honest findings from the ORT source read

Accuracy gate (opt-in, non-byte-identical by design)

Portability / capture

Reviewer asks

Uh oh!

codecov Bot commented Aug 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

github-actions Bot commented Aug 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

🔴 Benchmark Regression Detected

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

justinchuby commented Aug 15, 2026 •

edited

Loading

Result (glm-4-9b-int4, H200, `--steady --tokens 160 --decode-skip 40 --runs 3`)

ncu (glm gate_up, `--graph-profiling node`, launch-skip 240)

codecov Bot commented Aug 15, 2026 •

edited

Loading

github-actions Bot commented Aug 15, 2026 •

edited

Loading