[justinchuby]

Summary — measured NO-GO, default off, byte-identical (reproducible negative)

Re-test of the sota/TRT-LLM roadmap's #1 lever — PORT-3 cp.async double-buffered weight prefetch + L2::evict_first — on the current multicol NC=4 M=1 int4 decode GEMV. Result: NO-GO, −4.2% glm decode. Kept behind ONNX_GENAI_GEMV_CPASYNC=1 (default OFF), byte-identical, so nothing regresses by default. This documents the negative with fresh, mechanistically-explained data.

Why re-test (the prior no-go was a different kernel)

cp.async was already a measured no-go (2131046e) — but on the narrow pre-multicol kernel (register-bound, ~80% occupancy, 8–12% DRAM). The current multicol kernel is different, so I re-profiled first: on the current default kernel Long-Scoreboard is now the #1 stall (2.35) — exactly cp.async's target — so the precondition IS met here. That justified a genuine re-test.

What it does

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 (Marlin idiom), 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: SM80+, int4/block≠32 multicol only, default off.

Result

• 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 ✓)
registers/thread	64	71	+7
achieved occupancy	41.6%	33.3%	−8 pp

Mechanism: cp.async cut the global-load stall 20% as designed, but the cvta-to-shared + pipeline state raised registers 64→71, dropping occupancy 41.6%→33.3%. At M=1 this GEMV hides latency primarily via warp-level parallelism (occupancy), so losing 8 pp of occupancy removes more hiding than the prefetch adds. Third convergent prefetch negative (register depth-4, prior cp.async, this one): the M=1 decode GEMV is occupancy-bound, not prefetch-bound.

Roadmap recommendation

• Drop PORT-3 (cp.async) and deprioritize PORT-5 (PDL) — same latency-hiding class.
• The real lever is occupancy (register-bound at 64 from WIDE_NC=4) → points at PORT-1 (offline bias+interleave → cheaper dequant → fewer regs), but that needs an offline weight repack + versioned layout flag (model-artifact-compat surface).
• Honest strategic note: this GEMV is already ~95% of ORT DRAM / near ORT kernel time (see perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in) #996). The remaining e2e gap to ORT ~250 base is dominated by other ops (lm_head). The banked honest wins are fp16 GEMV (perf(cuda): fp16 int4 wide GEMV matching ORT's M=1 arithmetic — glm 199→213 tok/s, accuracy-gated (opt-in) #996, +6.7%) + lm_head cuBLASLt.

Decision

Merge as a documented negative (like 2131046e) or take the decision drop and close — either way no regression ships by default. Decision drop: .squad/decisions/inbox/deckard-cpasync-nogo.md.

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

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 83.56%. Comparing base (3f82693) to head (2002ad9).

Additional details and impacted files

@@           Coverage Diff           @@
##             main    #1000   +/-   ##
=======================================
  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]

🔴 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_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1	72.99 µs	97.78 µs	+34.0%
⚠️	tokenization/encode_tokens_per_second	354.99 µs	429.18 µs	+20.9%
⚠️	add/medium_bf16_threads=1-internal/262144	2.44 ms	2.92 ms	+19.8%
⚠️	grammar_masking/llguidance_compute_mask/32	68.38 µs	79.36 µs	+16.1%
✅	qwen3_sampling_processors/top_k_top_p_fast	606.58 µs	662.85 µs	+9.3%
✅	qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline	3.30 ms	3.58 ms	+8.6%
✅	qwen3_sampling_processors/top_p_fast_after_top_k	479.17 µs	520.39 µs	+8.6%
✅	qwen3_sampling_processors/top_k_partial_selection	130.03 µs	141.05 µs	+8.5%
✅	tokenization/decode_tokens_per_second	5.70 ms	6.17 ms	+8.2%
✅	sampling_latency/top_p_per_token	364.41 µs	390.63 µs	+7.2%
✅	kv_cache/alloc_dealloc_pages	36.68 µs	39.24 µs	+7.0%
✅	logit_processing/seven_processor_chain_per_step	302.94 µs	318.61 µs	+5.2%
✅	sampling_latency/min_p_per_token	193.50 µs	202.61 µs	+4.7%
✅	qwen3_sampling_processors/top_k_top_p_full_sort_baseline	5.40 ms	5.65 ms	+4.7%
✅	qwen3_sampling_processors/top_k_full_sort_baseline	2.00 ms	2.09 ms	+4.3%
✅	add/large_f32_threads=1-internal/4194304	37.28 ms	38.63 ms	+3.6%
✅	add/medium_f32_threads=1-internal/262144	2.33 ms	2.42 ms	+3.6%
✅	add/small_f16_threads=1-internal/1024	11.85 µs	12.25 µs	+3.4%
✅	add/small_f32_threads=1-internal/1024	180.6 ns	186.6 ns	+3.3%
✅	add/medium_f16_threads=1-internal/262144	2.39 ms	2.46 ms	+2.8%
✅	matmul/large_generic_f16_threads=1/32x1024x1024	79.71 µs	81.72 µs	+2.5%
✅	reduce_mean/large_f32_threads=1-internal/262144	912.49 µs	926.80 µs	+1.6%
✅	add/large_bf16_threads=1-internal/4194304	38.84 ms	39.41 ms	+1.5%
✅	add/large_f16_threads=1-internal/4194304	37.98 ms	38.40 ms	+1.1%
✅	reduce_mean/medium_f32_threads=1-internal/65536	224.76 µs	226.79 µs	+0.9%
✅	gather/small_f16_threads=1-internal/4096	453.3 ns	452.4 ns	-0.2%
✅	gather/medium_f32_threads=1-internal/32768	3.47 µs	3.45 µs	-0.4%
✅	gather/small_f32_threads=1-internal/4096	621.7 ns	616.6 ns	-0.8%
✅	gather/small_bf16_threads=1-internal/4096	445.7 ns	440.0 ns	-1.3%
✅	gather/medium_f16_threads=1-internal/32768	2.26 µs	2.23 µs	-1.3%
✅	reduce_mean/small_f32_threads=1-internal/4096	14.02 µs	13.79 µs	-1.6%
✅	add/small_bf16_threads=1-internal/1024	12.20 µs	11.96 µs	-1.9%
✅	matmul/medium_generic_f16_threads=1/32x512x512	29.15 µs	28.14 µs	-3.5%
✅	matmul/large_generic_bf16_threads=1/32x1024x1024	2.02 ms	1.94 ms	-3.7%
✅	gather/medium_bf16_threads=1-internal/32768	2.42 µs	2.24 µs	-7.5%
✅	matmul/large_generic_f16_threads=8/32x1024x1024	93.88 µs	84.37 µs	-10.1%
✅	matmul/large_generic_f32_threads=1/32x1024x1024	9.60 ms	8.62 ms	-10.2%
✅	matmul/large_generic_f32_threads=8/32x1024x1024	4.00 ms	3.57 ms	-10.8%
✅	sampling_latency/greedy_per_token	3.52 µs	3.11 µs	-11.7%
✅	sampling_latency/top_k_per_token	59.22 µs	51.11 µs	-13.7%
🟢	gather/large_f16_threads=1-internal/131072	12.55 µs	10.15 µs	-19.1%
🟢	matmul/medium_generic_f32_threads=1/32x512x512	2.68 ms	2.16 ms	-19.6%
🟢	block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1	474.26 µs	380.01 µs	-19.9%
🟢	matmul/small_generic_f32_threads=1/1x256x256	43.05 µs	34.01 µs	-21.0%
🟢	gather/large_f32_threads=1-internal/131072	38.19 µs	30.15 µs	-21.0%
🟢	matmul/small_generic_f16_threads=1/1x256x256	36.73 µs	28.48 µs	-22.4%
🟢	matmul/medium_generic_bf16_threads=1/32x512x512	667.98 µs	507.18 µs	-24.1%
🟢	block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024	687.03 µs	519.72 µs	-24.4%
🟢	matmul/small_generic_bf16_threads=1/1x256x256	39.72 µs	29.57 µs	-25.6%
🟢	gather/large_bf16_threads=1-internal/131072	14.08 µs	10.41 µs	-26.0%
🟢	block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024	66.54 µs	48.81 µs	-26.7%
🟢	matmul/small_generic_bf16_threads=8/1x256x256	41.98 µs	30.20 µs	-28.0%
🟢	block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024	65.90 µs	46.66 µs	-29.2%
🟢	matmul/small_generic_f16_threads=8/1x256x256	40.69 µs	28.68 µs	-29.5%
🟢	matmul/large_generic_bf16_threads=8/32x1024x1024	1.88 ms	1.32 ms	-30.1%
🟢	matmul/medium_generic_f16_threads=8/32x512x512	46.70 µs	27.64 µs	-40.8%
🟢	matmul/small_generic_f32_threads=8/1x256x256	55.11 µs	31.91 µs	-42.1%
🟢	matmul/medium_generic_f32_threads=8/32x512x512	1.63 ms	893.37 µs	-45.1%
🟢	matmul/medium_generic_bf16_threads=8/32x512x512	784.09 µs	375.37 µs	-52.1%

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.95 3.13 4.27 }

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.