[justinchuby]

PORT-1 (offline bias+interleave) — measured NO-GO on the multicol asymmetric-zp GEMV

Stacked on #996 (fp16 multicol GEMV). Base branch: squad/int4-gemv-fp16-mixed.

PORT-1 from the TRT-LLM port plan was ranked as an occupancy lever: bake the signed→unsigned +8 bias and odd/even nibble interleave into the stored weight layout so the runtime dequant drops ORT's prmt.b32 → fewer registers → higher occupancy → faster. On our kernel that mechanism does not exist, and I proved it with ncu before building the large offline-repack/versioned-layout/loader surface (which would have delivered ~0 gain).

Method

Added template<bool NATURAL> to the fp16 multicol device fn. NATURAL=true swaps the 4× prmt.b32 activation reorder (decode_activation8_h2) for a prmt-free decode that assumes offline-interleaved weights — isolating exactly the one op PORT-1 removes for the asymmetric case. <false> is byte-identical codegen to the shipping #996 kernel (f64 oracle 8/8, registers + kernel time unchanged). Probe entry gated behind ONNX_GENAI_GEMV_FP16_NATURAL=1 (default OFF); its output is intentionally wrong vs current non-interleaved weights and is never a correctness path — kept as reproducible evidence only.

Result (ncu, glm gate_up, GPU6, fresh idle, --graph-profiling node)

Metric	Baseline fp16 (#996)	prmt-free (PORT-1 dequant)	Delta
Registers / thread	64	64	0
Achieved occupancy	41.17 %	40.92 %	~0 (noise)
Kernel time	26.69 us	26.53 us	-0.6 % (noise)
DRAM throughput	2.267 TB/s	2.272 TB/s	~0

Why the premise fails here

1. Asymmetric zero-points (glm & qwen both use zp): our dequant is already LOP3-based plus an irreducible 4x per-block zp-subtract. FastInterleaved's op-saving only finalizes a symmetric +8 dequant; with asymmetric zp the +8 bias just folds into the stored zp (same op count). Weight-dequant instruction count does not drop.
2. The activation prmt is amortized across WIDE_NC=4 columns and is register-transient, so removing it frees 0 registers and is already latency-hidden.
3. The 64-reg footprint is structural to WIDE_NC=4 column state (w[4] uint4=16, col_kb[4] long=8, total/acc/scale2/sub2). Occupancy won't rise without dropping WIDE_NC — which loses the multicol L1 win.

Bottom line

Third convergent "occupancy is the wall and it doesn't move" finding (after depth-4 prefetch and cp.async #1000). The fp16 GEMV is already ~95% of ORT's DRAM and near per-kernel parity; the remaining base-vs-ORT gap is dominated by other ops (lm_head — deckard-4). Recommendation: don't build PORT-1's offline-repack surface for this kernel; land #996, let lm_head land, re-measure the gap.

Draft / documented-negative PR — not intended to merge as a feature. Kept for reproducibility + decision record (.squad/decisions/inbox/deckard-port1-nogo.md).

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 83.56%. Comparing base (6e37403) to head (3563044).

Additional details and impacted files

@@                     Coverage Diff                     @@
##           squad/int4-gemv-fp16-mixed    #1002   +/-   ##
===========================================================
  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]

✅ Benchmarks — No Regression

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
✅	add/small_f16_threads=1-internal/1024	13.83 µs	15.01 µs	+8.5%
✅	matmul/large_generic_bf16_threads=8/32x1024x1024	1.35 ms	1.45 ms	+7.4%
✅	grammar_masking/llguidance_compute_mask/32	74.78 µs	77.13 µs	+3.1%
✅	add/large_f32_threads=1-internal/4194304	40.28 ms	41.38 ms	+2.7%
✅	gather/medium_f32_threads=1-internal/32768	3.71 µs	3.80 µs	+2.5%
✅	add/medium_f16_threads=1-internal/262144	2.56 ms	2.62 ms	+2.4%
✅	logit_processing/seven_processor_chain_per_step	317.45 µs	322.62 µs	+1.6%
✅	matmul/large_generic_f16_threads=8/32x1024x1024	84.51 µs	85.73 µs	+1.4%
✅	block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1	77.28 µs	78.16 µs	+1.1%
✅	qwen3_sampling_processors/top_k_top_p_full_sort_baseline	5.63 ms	5.66 ms	+0.6%
✅	block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1	378.49 µs	379.77 µs	+0.3%
✅	block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024	516.58 µs	518.00 µs	+0.3%
✅	qwen3_sampling_processors/top_k_top_p_fast	658.91 µs	658.74 µs	-0.0%
✅	matmul/large_generic_bf16_threads=1/32x1024x1024	2.01 ms	2.00 ms	-0.1%
✅	add/medium_f32_threads=1-internal/262144	2.54 ms	2.53 ms	-0.5%
✅	add/medium_bf16_threads=1-internal/262144	2.66 ms	2.64 ms	-0.5%
✅	matmul/large_generic_f32_threads=8/32x1024x1024	3.96 ms	3.94 ms	-0.6%
✅	gather/small_f16_threads=1-internal/4096	484.1 ns	479.7 ns	-0.9%
✅	matmul/large_generic_f16_threads=1/32x1024x1024	80.15 µs	79.20 µs	-1.2%
✅	kv_cache/alloc_dealloc_pages	39.02 µs	38.55 µs	-1.2%
✅	matmul/large_generic_f32_threads=1/32x1024x1024	9.44 ms	9.30 ms	-1.6%
✅	add/large_bf16_threads=1-internal/4194304	42.67 ms	41.93 ms	-1.7%
✅	reduce_mean/medium_f32_threads=1-internal/65536	246.93 µs	242.31 µs	-1.9%
✅	qwen3_sampling_processors/top_p_fast_after_top_k	528.53 µs	517.60 µs	-2.1%
✅	gather/medium_bf16_threads=1-internal/32768	2.46 µs	2.41 µs	-2.3%
✅	reduce_mean/small_f32_threads=1-internal/4096	15.32 µs	14.96 µs	-2.4%
✅	qwen3_sampling_processors/top_k_partial_selection	144.89 µs	138.71 µs	-4.3%
✅	block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024	46.65 µs	44.52 µs	-4.5%
✅	qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline	3.74 ms	3.53 ms	-5.5%
✅	add/large_f16_threads=1-internal/4194304	43.31 ms	40.64 ms	-6.2%
✅	reduce_mean/large_f32_threads=1-internal/262144	1.04 ms	974.85 µs	-6.4%
✅	matmul/medium_generic_f32_threads=1/32x512x512	2.54 ms	2.34 ms	-7.6%
✅	add/small_bf16_threads=1-internal/1024	14.05 µs	12.96 µs	-7.8%
✅	qwen3_sampling_processors/top_k_full_sort_baseline	2.31 ms	2.12 ms	-8.1%
✅	sampling_latency/greedy_per_token	3.52 µs	3.22 µs	-8.6%
✅	matmul/small_generic_bf16_threads=8/1x256x256	35.08 µs	32.02 µs	-8.7%
✅	sampling_latency/min_p_per_token	229.64 µs	207.65 µs	-9.6%
✅	gather/small_bf16_threads=1-internal/4096	523.7 ns	472.1 ns	-9.8%
✅	block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024	49.87 µs	44.65 µs	-10.5%
✅	sampling_latency/top_p_per_token	431.53 µs	384.32 µs	-10.9%
✅	gather/small_f32_threads=1-internal/4096	794.2 ns	691.6 ns	-12.9%
✅	sampling_latency/top_k_per_token	60.19 µs	52.40 µs	-12.9%
✅	gather/medium_f16_threads=1-internal/32768	2.84 µs	2.46 µs	-13.6%
✅	matmul/small_generic_f16_threads=1/1x256x256	35.77 µs	30.84 µs	-13.8%
🟢	matmul/medium_generic_bf16_threads=1/32x512x512	632.78 µs	536.81 µs	-15.2%
🟢	matmul/small_generic_bf16_threads=1/1x256x256	39.51 µs	33.12 µs	-16.2%
🟢	tokenization/encode_tokens_per_second	473.08 µs	373.70 µs	-21.0%
🟢	matmul/small_generic_f16_threads=8/1x256x256	39.78 µs	30.60 µs	-23.1%
🟢	tokenization/decode_tokens_per_second	8.12 ms	6.16 ms	-24.1%
🟢	matmul/small_generic_f32_threads=1/1x256x256	49.51 µs	37.13 µs	-25.0%
🟢	matmul/medium_generic_f16_threads=8/32x512x512	40.26 µs	29.72 µs	-26.2%
🟢	matmul/medium_generic_f16_threads=1/32x512x512	44.11 µs	31.25 µs	-29.1%
🟢	add/small_f32_threads=1-internal/1024	275.5 ns	192.5 ns	-30.1%
🟢	matmul/medium_generic_f32_threads=8/32x512x512	1.37 ms	919.22 µs	-32.8%
🟢	matmul/small_generic_f32_threads=8/1x256x256	54.35 µs	36.27 µs	-33.3%
🟢	matmul/medium_generic_bf16_threads=8/32x512x512	563.58 µs	370.74 µs	-34.2%
🟢	gather/large_f16_threads=1-internal/131072	21.96 µs	12.99 µs	-40.9%
🟢	gather/large_f32_threads=1-internal/131072	48.05 µs	26.26 µs	-45.3%
🟢	gather/large_bf16_threads=1-internal/131072	30.58 µs	13.50 µs	-55.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: { 3.43 3.16 5.59 }

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)

[justinchuby]

Triage, not a rejection. The measured negative is valuable and I want it in the tree; the code half I cannot validate.

Why I am not merging as-is: this PR carries a CUDA kernel variant alongside its decision drop, and this machine no longer has a CUDA runtime -- a recursive search finds no cublasLt64_*.dll and no cudart64*.dll, and a full native run reports 0 MiB of GPU memory. I cannot reproduce "default off, byte-identical", and I am not going to merge kernel code on the strength of a claim I have no way to check. That standard has earned its keep today: a PR whose speed claim I did reproduce turned out to triple peak memory, and the accounting fix for it under-reported by 2.4x on its first attempt.

What I would like: split the decision drop out and I will land it immediately, the way #1011 just landed. A measured no-go deserves permanent shelf space -- the whole point is that the next person to propose this lever finds the numbers instead of re-running the experiment. The kernel variant can then wait for a reviewer with a working CUDA host, where "default off" and "byte-identical" can actually be demonstrated rather than asserted.

If you would rather not split it, say so and I will leave this open until I have GPU hardware again; that is a worse outcome only because the negative record stays unpublished in the meantime.