Skip to content

perf(cuda): QMoE fused gate/up occupancy lever (ONNX_GENAI_QMOE_OCC, default-OFF) - #1167

Merged
justinchuby merged 1 commit into
mainfrom
squad/qmoe-vec-occ
Aug 18, 2026
Merged

justinchuby merged 1 commit into
mainfrom
squad/qmoe-vec-occ

Conversation

@justinchuby

Copy link
Copy Markdown
Owner

QMoE gate/up SwiGLU expert-GEMV occupancy lever

Ports the dense #1139 gate/up occupancy win into the fused gate/up SwiGLU expert GEMV (the MoE analog). New flag ONNX_GENAI_QMOE_OCC, default-OFF.

ncu evidence (spill-verified)

regs theo occ achieved duration
default 54 50% 43.3% 42.3us
_occ (256,6) 40 (spill-free) 75% 63.8% 37.8us (−10.6%)

Register-granularity trap avoided: (256,8) → 32 reg/100% theoretical but spills 1.22 MB → regression. Shipped block config 6.

qmoe_linear (fc2/down) already at 32 reg/100% (no lever); grouped-GEMM M>1 path already dequants-once with TILE_M=8. Gate/up is the sole remaining occupancy target.

Parity — byte-identical

  • qmoe_occ_is_bit_identical: 0 differing bits (f16/bf16/f32 × rows{1,4,6,8})
  • Full qmoe_gpu suite: 32/32 pass (re-verified post-rebase onto current main)
  • E2E DeepSeek-V2-Lite tokens identical OFF vs ON, both equal the merged golden […207,16,24,1012,1712,5075…]

E2E

V2-Lite decode ~−1.0% median (56.6→57.2 tok/s, runs=10 warmups=2, 3 rounds). Small because gate/up ≈8% of V2-Lite decode (MLA/Attention ≈21% dominates a different lane); bigger-intermediate MoE would convert more.

Review

  • 🧪 Rachael: 🟡 approve-with-notes — byte-identity + ncu occupancy + spill-free config all independently verified. Note: flag selects _occ for any QMoE QuantLayout when enabled (not literally symmetric-int4-only); affine OCC smoke tests pass; default-OFF keeps it non-blocking.

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

…h_bounds (byte-identical)

The fused QMoE gate/up SwiGLU expert GEMV (qmoe_gate_up_activate) is the
occupancy-limited kernel in the MoE decode path. Re-measured on the corrected
DeepSeek-V2-Lite baseline (#1150, 64 experts, top-6, hidden=2048, inter=1408,
int4 block-16, f32 activations) on H200: the default entry runs at 54 reg/thread
-> 4 blocks/SM (50% theoretical, 43.3% achieved occupancy), DRAM only 9.1%,
Long-Scoreboard bound on the int4 weight loads.

Add opt-in _occ siblings with __launch_bounds__(256, 6): 54->40 reg/thread,
6 blocks/SM, spill-free (ncu local ld/st = 0 bytes). Occupancy 50->75%
theoretical / 43.3->63.8% achieved, kernel duration 42.3->37.8us (-10.6%).
(256, 8) reaches 32 reg/100% theoretical but spills 1.22 MB and regresses to
43.5us -- the register-granularity trap that also sank the down-projection
occupancy attempt -- so 6 is shipped.

Byte-identical by construction: __launch_bounds__ only constrains register
allocation, so the FMA accumulate order, both block_sum reductions, and the
SwiGLU are unchanged. New byte-identity A/B test asserts 0 differing bits
between the default and _occ paths across f16/bf16/f32 and rows in {1,4,6,8}.
E2E V2-Lite decode A/B (interleaved, runs=10, warmups=2): -1.0% median decode
(56.6->57.2 tok/s) over 3 clean rounds, generated token ids byte-identical to
the merged golden under both flag states. Default-OFF behind
ONNX_GENAI_QMOE_OCC=1; symmetric int4 only (asymmetric _zp path untouched).

Investigation finding: the other QMoE expert-GEMV kernels have no comparable
occupancy headroom -- qmoe_linear (fc2/down) is already 32 reg/100% theoretical
occupancy, and the M>1 batched path (qmoe_grouped_linear) already dequants each
weight once and reuses it across TILE_M=8 rows -- so gate/up is the sole
remaining occupancy lever in the family.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
@justinchuby
justinchuby merged commit 651901a into main Aug 18, 2026
9 of 14 checks passed
@justinchuby
justinchuby deleted the squad/qmoe-vec-occ branch August 18, 2026 02:58
@codecov

codecov Bot commented Aug 18, 2026 •

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 80.44%. Comparing base (8219fa2) to head (9a6768c).
⚠️ Report is 12 commits behind head on main.

Additional details and impacted files

Impacted file tree graph

@@            Coverage Diff            @@
##           main    #1167       +/-   ##
=========================================
+ Coverage      0   80.44%   +80.44%     
=========================================
  Files         0      359      +359     
  Lines         0   157323   +157323     
  Branches      0   157323   +157323     
=========================================
+ Hits          0   126564   +126564     
- Misses        0    26196    +26196     
- Partials      0     4563     +4563     
Flag Coverage Δ
mlas 85.36% <ø> (?)
offline 80.35% <ø> (?)

Flags with carried forward coverage won't be shown. Click here to find out more.
see 359 files with indirect coverage changes

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

@github-actions

Copy link
Copy Markdown

🔴 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
🔴 gather/large_f16_threads=1-internal/131072 9.81 µs 22.08 µs +125.1%
🔴 gather/large_f32_threads=1-internal/131072 25.44 µs 49.57 µs +94.8%
🔴 add/medium_f32_threads=1-internal/262144 22.78 µs 41.25 µs +81.1%
🔴 block_quantized_matmul_cached_dense/mxfp4_uncached_dequant_each_call/1x1024x1024 513.68 µs 918.68 µs +78.8%
🔴 matmul/small_generic_f16_threads=8/1x256x256 27.65 µs 46.74 µs +69.0%
🔴 matmul/large_generic_bf16_threads=8/32x1024x1024 1.26 ms 1.99 ms +57.7%
🔴 matmul/small_generic_bf16_threads=1/1x256x256 28.53 µs 44.74 µs +56.8%
🔴 matmul/medium_generic_f16_threads=8/32x512x512 27.98 µs 40.89 µs +46.1%
🔴 matmul/small_generic_bf16_threads=8/1x256x256 30.79 µs 44.04 µs +43.0%
🔴 add/small_bf16_threads=1-internal/1024 401.5 ns 571.1 ns +42.2%
🔴 matmul/small_generic_f32_threads=8/1x256x256 31.83 µs 44.92 µs +41.1%
🔴 reduce_mean/large_f32_threads=1-internal/262144 908.40 µs 1.28 ms +41.0%
🔴 block_quantized_matmul_cached_dense/mxfp4_cached_dense_repeated_call/1x1024x1024 40.15 µs 55.23 µs +37.5%
🔴 reduce_mean/medium_f32_threads=1-internal/65536 224.84 µs 295.32 µs +31.4%
🔴 gather/small_f32_threads=1-internal/4096 621.5 ns 810.2 ns +30.4%
🔴 add/small_f32_threads=1-internal/1024 181.1 ns 235.4 ns +30.0%
⚠️ gather/medium_f32_threads=1-internal/32768 3.40 µs 4.38 µs +28.9%
⚠️ add/large_f32_threads=1-internal/4194304 567.89 µs 729.95 µs +28.5%
⚠️ gather/medium_bf16_threads=1-internal/32768 2.31 µs 2.93 µs +26.9%
⚠️ matmul/small_generic_f16_threads=1/1x256x256 27.74 µs 35.08 µs +26.5%
⚠️ block_quantized_matmul_cached_dense/mxfp4_preexpanded_dense_oncelock_like_proxy/1x1024x1024 42.22 µs 53.26 µs +26.1%
⚠️ matmul/small_generic_f32_threads=1/1x256x256 34.63 µs 43.53 µs +25.7%
⚠️ matmul/large_generic_f32_threads=8/32x1024x1024 3.69 ms 4.64 ms +25.7%
⚠️ matmul/medium_generic_f32_threads=8/32x512x512 908.28 µs 1.13 ms +24.3%
⚠️ gather/large_bf16_threads=1-internal/131072 10.75 µs 13.23 µs +23.0%
⚠️ matmul/medium_generic_f16_threads=1/32x512x512 27.63 µs 33.88 µs +22.6%
⚠️ add/large_f16_threads=1-internal/4194304 1.54 ms 1.88 ms +21.9%
⚠️ add/medium_f16_threads=1-internal/262144 97.84 µs 119.28 µs +21.9%
⚠️ gather/small_f16_threads=1-internal/4096 437.7 ns 527.4 ns +20.5%
⚠️ add/small_f16_threads=1-internal/1024 414.6 ns 499.6 ns +20.5%
⚠️ add/large_bf16_threads=1-internal/4194304 1.50 ms 1.79 ms +18.8%
⚠️ matmul/medium_generic_bf16_threads=8/32x512x512 366.55 µs 433.55 µs +18.3%
⚠️ reduce_mean/small_f32_threads=1-internal/4096 13.83 µs 16.32 µs +18.0%
⚠️ matmul/large_generic_f16_threads=8/32x1024x1024 77.67 µs 91.15 µs +17.4%
⚠️ matmul/large_generic_f32_threads=1/32x1024x1024 8.74 ms 10.12 ms +15.7%
⚠️ matmul/medium_generic_bf16_threads=1/32x512x512 485.01 µs 560.21 µs +15.5%
✅ kv_cache/alloc_dealloc_pages 42.00 µs 48.29 µs +15.0%
✅ gather/small_bf16_threads=1-internal/4096 436.7 ns 500.4 ns +14.6%
✅ matmul/large_generic_f16_threads=1/32x1024x1024 73.50 µs 83.84 µs +14.1%
✅ add/medium_bf16_threads=1-internal/262144 94.22 µs 106.68 µs +13.2%
✅ matmul/medium_generic_f32_threads=1/32x512x512 2.17 ms 2.44 ms +12.7%
✅ gather/medium_f16_threads=1-internal/32768 2.27 µs 2.50 µs +10.3%
✅ block_quantized_moe_cached_dense/mxfp4_cached_dense_expert_repeated_call/rows=1,H=256,I=256,E=4,top_k=1 154.86 µs 170.40 µs +10.0%
✅ logit_processing/seven_processor_chain_per_step 321.22 µs 351.54 µs +9.4%
✅ block_quantized_moe_cached_dense/mxfp4_uncached_expert_dequant_each_call/rows=1,H=256,I=256,E=4,top_k=1 476.52 µs 521.24 µs +9.4%
✅ matmul/large_generic_bf16_threads=1/32x1024x1024 1.82 ms 1.98 ms +8.7%
✅ grammar_masking/llguidance_compute_mask/32 76.78 µs 82.12 µs +7.0%
✅ sampling_latency/top_p_per_token 400.06 µs 418.23 µs +4.5%
✅ sampling_latency/min_p_per_token 220.70 µs 225.53 µs +2.2%
✅ qwen3_sampling_processors/top_k_full_sort_baseline 2.17 ms 2.17 ms +0.3%
✅ qwen3_sampling_processors/top_k_partial_selection 151.75 µs 150.87 µs -0.6%
✅ sampling_latency/top_k_per_token 57.48 µs 54.69 µs -4.9%
✅ qwen3_sampling_processors/top_p_full_sort_after_top_k_baseline 4.00 ms 3.69 ms -7.9%
✅ qwen3_sampling_processors/top_k_top_p_fast 767.45 µs 703.66 µs -8.3%
✅ qwen3_sampling_processors/top_p_fast_after_top_k 594.89 µs 541.11 µs -9.0%
✅ tokenization/decode_tokens_per_second 7.64 ms 6.71 ms -12.1%
✅ tokenization/encode_tokens_per_second 436.71 µs 383.54 µs -12.2%
✅ sampling_latency/greedy_per_token 3.71 µs 3.21 µs -13.5%
🟢 qwen3_sampling_processors/top_k_top_p_full_sort_baseline 7.95 ms 6.32 ms -20.5%

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.50 3.04 3.32 }
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)

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant