From edc3918064616de58e4fb801aa35f82f9b3e6f69 Mon Sep 17 00:00:00 2001 From: Squad Scribe Date: Tue, 11 Aug 2026 05:20:04 +0000 Subject: [PATCH 1/2] chore(squad): persist session state (GLM/DeepSeek ship, megakernel Phase 0 NO-SHIP, CI fix) Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .squad/agents/cohaagen/history-archive.md | 97 +++++ .squad/agents/cohaagen/history.md | 103 +----- .squad/agents/harry/history-archive.md | 92 +++++ .squad/agents/harry/history.md | 96 +---- .squad/agents/quaid/history.md | 5 + .squad/agents/roper/history.md | 10 + .squad/decisions-archive/2026-08.md | 344 ++++++++++++++++++ .squad/decisions.md | 10 + .../inbox/cohaagen-qmoe-fc2-fusion.md | 34 ++ .../inbox/cohaagen-qmoe-profile-nextlever.md | 32 ++ .squad/decisions/inbox/harry-728-c1-review.md | 21 ++ .squad/decisions/inbox/harry-728-rereview.md | 19 + .squad/decisions/inbox/harry-728-review3.md | 18 + .squad/decisions/inbox/harry-728-review4.md | 24 ++ .squad/decisions/inbox/harry-728-review5.md | 24 ++ .squad/decisions/inbox/harry-728-review6.md | 23 ++ .squad/decisions/inbox/harry-728-review7.md | 28 ++ .squad/decisions/inbox/harry-728-review8.md | 50 +++ .../inbox/harry-glm-deepseek-review.md | 82 +++++ .../inbox/quaid-megakernel-phase0.md | 161 ++++++++ .../decisions/inbox/quaid-qmoe-int4-dp4a.md | 114 ++++++ 21 files changed, 1206 insertions(+), 181 deletions(-) create mode 100644 .squad/agents/harry/history-archive.md create mode 100644 .squad/agents/roper/history.md create mode 100644 .squad/decisions/inbox/cohaagen-qmoe-fc2-fusion.md create mode 100644 .squad/decisions/inbox/cohaagen-qmoe-profile-nextlever.md create mode 100644 .squad/decisions/inbox/harry-728-c1-review.md create mode 100644 .squad/decisions/inbox/harry-728-rereview.md create mode 100644 .squad/decisions/inbox/harry-728-review3.md create mode 100644 .squad/decisions/inbox/harry-728-review4.md create mode 100644 .squad/decisions/inbox/harry-728-review5.md create mode 100644 .squad/decisions/inbox/harry-728-review6.md create mode 100644 .squad/decisions/inbox/harry-728-review7.md create mode 100644 .squad/decisions/inbox/harry-728-review8.md create mode 100644 .squad/decisions/inbox/harry-glm-deepseek-review.md create mode 100644 .squad/decisions/inbox/quaid-megakernel-phase0.md create mode 100644 .squad/decisions/inbox/quaid-qmoe-int4-dp4a.md diff --git a/.squad/agents/cohaagen/history-archive.md b/.squad/agents/cohaagen/history-archive.md index ee43c8dcc1..d96f1d3b09 100644 --- a/.squad/agents/cohaagen/history-archive.md +++ b/.squad/agents/cohaagen/history-archive.md @@ -34,3 +34,100 @@ _Archived by Scribe round 9, 2026-07-31T10:24:07Z._ - PR #529 merged: qwen3.5-0.8b hybrid places 100% on CUDA — split package embedding.onnx 24 nodes + text.onnx 1265 = 1289 nodes, 0 declines (after #480/#484/#525). Regression lock `qwen35_0_8b_placement_lock`. - E2e decode still BLOCKED on the loader: `Engine::from_dir` rejects the 3-onnx split; `from_pipeline_dir` refuses during vision `smart_resize` admission. Parity harness `qwen35_0_8b_hybrid_native_cuda_e2e` graceful-skips until fixed. - In flight (cohaagen-4): loader-unblock — admit the text-only split hybrid for decode, flip the e2e parity harness active. + +## 2026-08-11T03:25:00Z: Scribe compaction from live history + +## 2026-07-31T03:03:15Z — PR #535 merged: text-only decode pipeline synthesis (hybrid loader unblock) + +- PR #535 merged: unblocks the split VLM package whose image preprocessing is unrepresentable (`smart_resize`). New `GenAiConfigError::UnrepresentablePreprocessing` → `to_strict_text_only_pipeline_metadata` synthesizes an embedding→decoder AR pipeline with NO vision component. Modality-driven, NOT a model-name case. Flips the `qwen35_0_8b_hybrid_native_cuda_e2e` parity harness active. +- Weight-offload chain (#63): #444 first increment merged — `page_lazy_weight` dispatch-layer seam + `CudaWeightResidency` LRU, gated behind `ONNX_GENAI_WEIGHT_OFFLOAD=1`. NEXT: #87 async page-in (PLAN-ONLY, awaiting Justin green-light). +- Guardrail proven: o_proj 2-way split-K (K_SPLIT=2) REGRESSES the 7B o_proj GEMV (−0.59%) — **do NOT re-try**. A K_SPLIT>2 new kernel is the future candidate. + +## 2026-07-31T08:48:28Z — PR #544 async page-in MERGED; #552 observability MERGED; GQA capture validated + +- PR #544 merged (Harry test fix; Melina RC→APPROVE): async `cuMemcpyHtoDAsync`+`compute_wait_fence` page-in; eviction re-serializes (WAR); ~1% win at 2 MiB only; byte-identical all budgets. `ONNX_GENAI_WEIGHT_OFFLOAD_ASYNC_PAGEIN` default ON. +- PR #552 merged (Lori APPROVE): `NativeDecodeSession::load_with_resolved_io` — genai_config decoders now emit `cuda_graph captures/fallbacks/fallback_report` + `--trace` capture-reject reasons. `profile_native` simple path now fully observable. +- GQA capture default-on validated (verification-only, no code change): Qwen2.5-0.5B int4 GQA — ON 2.14× eager, byte-identical, zero declines. Standing question answered. Gemma-4-E2B unmeasurable: stale HF export + GAP-3. Capture default-on for multi-component deferred until mary's native pipeline + re-export. + +## 2026-07-31T10:24:07Z — 27B decode roofline profile (profile-only, no code changed) + +- Profiled Qwen3.6-27B int4 native-CUDA resident decode on H200. Steady-state: 168 ms/tok (~35× off roofline). +- Root cause: `Scan` (48 LinearAttention blocks/step) = **56.5%** of decode step — structurally un-capturable control-flow. MatMulNBits already at roofline (4.4 ms, 2%) — **do NOT touch for this**. +- Capture engages (fallbacks=0); Transpose per-call sync 9.2% is next lever in Cohaagen's lane (separate from Scan fix). +- Fix is Justin's / pipeline-control-flow area. This was a profile+identify task only; no worktree remains. +- The Foundry baseline reports native CUDA ahead of ORT; 7B tracing localized o_proj at 19.5% of kernel time. +- Reverted the two-way o_proj split-K gate after a repeatable 0.59% 7B regression; do not retry that lever without a new higher-split kernel experiment. + +## 2026-08-02T10:05:00+0000 — Fused LinearAttention + Inc-1b default wave + +- Authored #594 (EP-driven keep-as-op and dual-domain fused LinearAttention) and #592 (removed `ONNX_GENAI_DECODE_INLINE_SCAN`; automatic graph-property gate). Both merged to main. +- Reported 27B fused LinearAttention decode around 41.3 ms/tok with 48 fused ops / 0 Scans; composition oracle on final main was in flight at Scribe time. + +## 2026-08-02T11:40:00+0000 — Final fused LinearAttention validation + +- Ran the final-main #592 + #594 composition oracle: 27B byte-exact greedy PASS (CUDA == CPU == expected ids) and structural 48 fused `LinearAttention` / 0 Scans PASS. +- Confirmed steady decode on merged main at 40.8 ms/tok (24.5 tok/s); fused-LinearAttention lane complete. + +## 2026-08-02T15:45:00+0000 — Foundry native-vs-ORT sweep and 7B scope + +- Ran dense Foundry native-CUDA vs ORT-CUDA steady sweep: native beat ORT on every dense model (0.5b 1.75×, 1.5b 1.66×, Phi-4-mini 1.35×, 7b 1.14×); no native loss found. +- Scoped qwen2.5-7b GQA and found no bounded win: Nsight showed MatMulNBits/HBM-bound decode (~88% GPU kernel time); tested GQA tweak regressed and was reverted. +- Adjudicated qwen2.5-1.5b divergence as native-correct and authored merged PR #597 regression lock. + +## 2026-08-02T16:45:00+0000 — DeepSeek/GLM post-#434 metadata validation + +- Ran the DeepSeek/GLM sweep and confirmed initial broad failures were stale pre-#434 artifact metadata, not product bugs. +- Regenerated sidecars with current Mobius #434 metadata: DeepSeek-V2 tiny native+ORT admit, DeepSeek-V2-Lite real int4 native/ORT tokens match with native ~250× faster, and GLM-4-9B native matches golden greedy lock. +- GLM-5.2 tiny/q4/qmoe remain blocked on unmerged Mobius #404; add a real DeepSeek-V2-Lite non-tiny-vocab lock before treating it as full correctness coverage. + +## 2026-08-02T19:00:00+0000 — Thread-3 hetero legalization design + #602 + +- Authored the Thread-3 pipeline map and relocation design; key finding: the multi-EP kept-function correctness hole is latent because `hetero.rs` is not wired into the default single-EP session path. +- Implemented bounded Phase 0+1 in #602: post-assignment legalization fixpoint in `hetero::plan`, fail-closed ambiguous `(domain, op_type)`, and CPU-only fake-provider tests. Locked out after Harry's round-1 review; Deckard took the revision. + +## 2026-08-03T02:40:00+0000 — PR #606 Thread-3 Phase 3 scaffold + +- Authored and merged #606 after finding cross-EP tensor movement is only in standalone `hetero::execute`, not the stateful session `Executor`; did not fake per-op hetero execution. +- Shipped opt-in `ONNX_GENAI_HETERO` fail-closed placement guard, `SessionError::HeterogeneousExecutionUnsupported`, C API mapping, and 8 new CPU-fake hetero tests; single-EP flag-off path remains byte-identical. +- Coordinator resolved the stale decode-inline option flip and pivoted from unused integrated Phase-3 execution to the Qwen3.6-35B-A3B `PackedMultiHeadAttention` admission bug. + +## 2026-08-03T03:10:00+0000 — 35B-A3B PackedMHA root cause + +- Root-caused the Qwen3.6/Qwen3.5 35B-A3B `vision_encoder` admission failure as a mobius export bug, not an onnx-genai loader bug. +- Confirmed ORT `PackedMultiHeadAttention` requires the optional `bias` positional slot before `token_offset`/`cumulative_sequence_length`; mobius call sites emitted 6 inputs but the fallback function declared only 5 formals, so onnx-genai correctly rejected `6 > 5`. + +## 2026-08-03T06:40:00+0000 — 35B-A3B validation and fairness vet + +- Validated Qwen3.6-35B-A3B under strict apples-to-apples methodology: ORT-CUDA crashes, while native is blocked by fp16 MoE-gate `TopK`, rank-3 mRoPE positions, and unimplemented native pipeline decode, so 35B is a capability gap on both engines today. +- Fairness-vetted native-vs-ORT claims for issue #610: report only same-artifact steady-state oracle-correct comparisons as throughput multipliers; report crashes, graph rejects, and CPU/different-kernel fallbacks as capability gaps. +- Started fp16 TopK CUDA enablement on `squad/cuda-fp16-topk` to unblock dense_fallback MoE routers. + +## 2026-08-03T07:40:00+0000 — PR #612 fp16 TopK conformance and GAP-3 kickoff + +- Authored merged PR #612 after finding fp16/bf16 CUDA TopK was already on main via #445; kept the work test-only and added router-shape fp16 GPU==CPU, non-final-axis, and EP-claim coverage. +- Validated 15/15 GPU tests, CUDA clippy `-D warnings`, and fmt; flagged latent CPU EP non-final-axis TopK ordering as a separate follow-up. +- Began GAP-3 native pipeline decode scoping for Qwen3.6-35B-A3B. + +## 2026-08-03T09:00:00Z — 35B-A3B origin/main revalidation correction + +- Cohaagen-24 found GAP-3 native pipeline decode was already landed on `origin/main`; PR #613 merged as docs truth-up/design note instead of duplicating implementation. +- Cohaagen-25 revalidated 35B-A3B on fresh `origin/main` @ `0a5ac3c5`: full native pipeline decode is output-correct vs native-CPU oracle, and fp16 TopK/Softmax execute on CUDA EP. +- Current native GPU blocker is now a cuDNN f16/bf16 `ReduceSum` compute-type bug plus decoder device-wiring; ORT-CUDA still hard-crashes. Cohaagen-26 is in flight to fix and measure. + +## 2026-08-03T10:00:00Z — PR #616 35B-A3B native GPU unblock + +- Authored merged PR #616 fixing cuDNN f16/bf16 reduce compute-type behavior and native pipeline decoder device wiring. +- Measured Qwen3.6-35B-A3B native GPU decode on H200 at 2726 ms/tok / 0.37 tok/s, correct vs native-CPU oracle and GPU-confirmed; ORT-CUDA still crashes, so this is a native capability number. +- Follow-up perf work is host/sync overhead and low GPU utilization, not correctness/unblock. + +## 2026-08-03T12:30:00Z — 35B-A3B Lever A reduce capture + +- Profiled 35B native CUDA decode and found capture shredded by 10240 fp16 `ReduceSum` seams from dense_fallback MoE, causing host/sync-bound 2725 ms/tok decode. +- Authored merged PR #618 making cuDNN float ReduceSum/Mean capture-eligible via descriptor/workspace caching, sync gating, warmed axes, and shape-stable capture-safe marking. +- Measured 35B decode improvement to 405 ms/tok / 2.47 tok/s (~6.7×), byte-exact vs CPU oracle; Lever B now targets ReduceSumSquare/Split/LinearAttention seams. + +## 2026-08-04T00:40:00Z — PR #625 native loader and 35B QMoE follow-through + +- Delivered the #625 native loader fix with GraphIo/GraphIoMetadata and rebased onto origin/main; locked out after Harry found the initializer-input leak. +- Confirmed ORT 1.28 still rejects fp16-activation/fp32-scale QMoE, making config-B a capability gap rather than a loader issue. +- Left GPU watcher PID 1060559 for config-A measurement while external vLLM occupies the GPU. diff --git a/.squad/agents/cohaagen/history.md b/.squad/agents/cohaagen/history.md index 017732dbbb..15bcf8d5ce 100644 --- a/.squad/agents/cohaagen/history.md +++ b/.squad/agents/cohaagen/history.md @@ -5,100 +5,8 @@ CUDA performance and kernel specialist. Owns CUDA EP op-coverage (#67), weight-o _Entries before 2026-07-31T03:03:15Z archived to `history-archive.md` (Scribe round 9). Archived: PR #380 re-review; 7B perf (o_proj split-K revert); #63 inc2 (#444 weight pager); #480 (CausalConvWithState+GBQ); #484+#525 (LinearAttention+RoPE-contrib/NonZero); #529 (qwen3.5 100% CUDA placement)._ -## 2026-07-31T03:03:15Z — PR #535 merged: text-only decode pipeline synthesis (hybrid loader unblock) - -- PR #535 merged: unblocks the split VLM package whose image preprocessing is unrepresentable (`smart_resize`). New `GenAiConfigError::UnrepresentablePreprocessing` → `to_strict_text_only_pipeline_metadata` synthesizes an embedding→decoder AR pipeline with NO vision component. Modality-driven, NOT a model-name case. Flips the `qwen35_0_8b_hybrid_native_cuda_e2e` parity harness active. -- Weight-offload chain (#63): #444 first increment merged — `page_lazy_weight` dispatch-layer seam + `CudaWeightResidency` LRU, gated behind `ONNX_GENAI_WEIGHT_OFFLOAD=1`. NEXT: #87 async page-in (PLAN-ONLY, awaiting Justin green-light). -- Guardrail proven: o_proj 2-way split-K (K_SPLIT=2) REGRESSES the 7B o_proj GEMV (−0.59%) — **do NOT re-try**. A K_SPLIT>2 new kernel is the future candidate. - -## 2026-07-31T08:48:28Z — PR #544 async page-in MERGED; #552 observability MERGED; GQA capture validated - -- PR #544 merged (Harry test fix; Melina RC→APPROVE): async `cuMemcpyHtoDAsync`+`compute_wait_fence` page-in; eviction re-serializes (WAR); ~1% win at 2 MiB only; byte-identical all budgets. `ONNX_GENAI_WEIGHT_OFFLOAD_ASYNC_PAGEIN` default ON. -- PR #552 merged (Lori APPROVE): `NativeDecodeSession::load_with_resolved_io` — genai_config decoders now emit `cuda_graph captures/fallbacks/fallback_report` + `--trace` capture-reject reasons. `profile_native` simple path now fully observable. -- GQA capture default-on validated (verification-only, no code change): Qwen2.5-0.5B int4 GQA — ON 2.14× eager, byte-identical, zero declines. Standing question answered. Gemma-4-E2B unmeasurable: stale HF export + GAP-3. Capture default-on for multi-component deferred until mary's native pipeline + re-export. - -## 2026-07-31T10:24:07Z — 27B decode roofline profile (profile-only, no code changed) - -- Profiled Qwen3.6-27B int4 native-CUDA resident decode on H200. Steady-state: 168 ms/tok (~35× off roofline). -- Root cause: `Scan` (48 LinearAttention blocks/step) = **56.5%** of decode step — structurally un-capturable control-flow. MatMulNBits already at roofline (4.4 ms, 2%) — **do NOT touch for this**. -- Capture engages (fallbacks=0); Transpose per-call sync 9.2% is next lever in Cohaagen's lane (separate from Scan fix). -- Fix is Justin's / pipeline-control-flow area. This was a profile+identify task only; no worktree remains. -- The Foundry baseline reports native CUDA ahead of ORT; 7B tracing localized o_proj at 19.5% of kernel time. -- Reverted the two-way o_proj split-K gate after a repeatable 0.59% 7B regression; do not retry that lever without a new higher-split kernel experiment. - -## 2026-08-02T10:05:00+0000 — Fused LinearAttention + Inc-1b default wave - -- Authored #594 (EP-driven keep-as-op and dual-domain fused LinearAttention) and #592 (removed `ONNX_GENAI_DECODE_INLINE_SCAN`; automatic graph-property gate). Both merged to main. -- Reported 27B fused LinearAttention decode around 41.3 ms/tok with 48 fused ops / 0 Scans; composition oracle on final main was in flight at Scribe time. - -## 2026-08-02T11:40:00+0000 — Final fused LinearAttention validation - -- Ran the final-main #592 + #594 composition oracle: 27B byte-exact greedy PASS (CUDA == CPU == expected ids) and structural 48 fused `LinearAttention` / 0 Scans PASS. -- Confirmed steady decode on merged main at 40.8 ms/tok (24.5 tok/s); fused-LinearAttention lane complete. - -## 2026-08-02T15:45:00+0000 — Foundry native-vs-ORT sweep and 7B scope - -- Ran dense Foundry native-CUDA vs ORT-CUDA steady sweep: native beat ORT on every dense model (0.5b 1.75×, 1.5b 1.66×, Phi-4-mini 1.35×, 7b 1.14×); no native loss found. -- Scoped qwen2.5-7b GQA and found no bounded win: Nsight showed MatMulNBits/HBM-bound decode (~88% GPU kernel time); tested GQA tweak regressed and was reverted. -- Adjudicated qwen2.5-1.5b divergence as native-correct and authored merged PR #597 regression lock. - -## 2026-08-02T16:45:00+0000 — DeepSeek/GLM post-#434 metadata validation - -- Ran the DeepSeek/GLM sweep and confirmed initial broad failures were stale pre-#434 artifact metadata, not product bugs. -- Regenerated sidecars with current Mobius #434 metadata: DeepSeek-V2 tiny native+ORT admit, DeepSeek-V2-Lite real int4 native/ORT tokens match with native ~250× faster, and GLM-4-9B native matches golden greedy lock. -- GLM-5.2 tiny/q4/qmoe remain blocked on unmerged Mobius #404; add a real DeepSeek-V2-Lite non-tiny-vocab lock before treating it as full correctness coverage. - -## 2026-08-02T19:00:00+0000 — Thread-3 hetero legalization design + #602 - -- Authored the Thread-3 pipeline map and relocation design; key finding: the multi-EP kept-function correctness hole is latent because `hetero.rs` is not wired into the default single-EP session path. -- Implemented bounded Phase 0+1 in #602: post-assignment legalization fixpoint in `hetero::plan`, fail-closed ambiguous `(domain, op_type)`, and CPU-only fake-provider tests. Locked out after Harry's round-1 review; Deckard took the revision. - -## 2026-08-03T02:40:00+0000 — PR #606 Thread-3 Phase 3 scaffold - -- Authored and merged #606 after finding cross-EP tensor movement is only in standalone `hetero::execute`, not the stateful session `Executor`; did not fake per-op hetero execution. -- Shipped opt-in `ONNX_GENAI_HETERO` fail-closed placement guard, `SessionError::HeterogeneousExecutionUnsupported`, C API mapping, and 8 new CPU-fake hetero tests; single-EP flag-off path remains byte-identical. -- Coordinator resolved the stale decode-inline option flip and pivoted from unused integrated Phase-3 execution to the Qwen3.6-35B-A3B `PackedMultiHeadAttention` admission bug. - -## 2026-08-03T03:10:00+0000 — 35B-A3B PackedMHA root cause - -- Root-caused the Qwen3.6/Qwen3.5 35B-A3B `vision_encoder` admission failure as a mobius export bug, not an onnx-genai loader bug. -- Confirmed ORT `PackedMultiHeadAttention` requires the optional `bias` positional slot before `token_offset`/`cumulative_sequence_length`; mobius call sites emitted 6 inputs but the fallback function declared only 5 formals, so onnx-genai correctly rejected `6 > 5`. - -## 2026-08-03T06:40:00+0000 — 35B-A3B validation and fairness vet - -- Validated Qwen3.6-35B-A3B under strict apples-to-apples methodology: ORT-CUDA crashes, while native is blocked by fp16 MoE-gate `TopK`, rank-3 mRoPE positions, and unimplemented native pipeline decode, so 35B is a capability gap on both engines today. -- Fairness-vetted native-vs-ORT claims for issue #610: report only same-artifact steady-state oracle-correct comparisons as throughput multipliers; report crashes, graph rejects, and CPU/different-kernel fallbacks as capability gaps. -- Started fp16 TopK CUDA enablement on `squad/cuda-fp16-topk` to unblock dense_fallback MoE routers. - -## 2026-08-03T07:40:00+0000 — PR #612 fp16 TopK conformance and GAP-3 kickoff - -- Authored merged PR #612 after finding fp16/bf16 CUDA TopK was already on main via #445; kept the work test-only and added router-shape fp16 GPU==CPU, non-final-axis, and EP-claim coverage. -- Validated 15/15 GPU tests, CUDA clippy `-D warnings`, and fmt; flagged latent CPU EP non-final-axis TopK ordering as a separate follow-up. -- Began GAP-3 native pipeline decode scoping for Qwen3.6-35B-A3B. - -## 2026-08-03T09:00:00Z — 35B-A3B origin/main revalidation correction - -- Cohaagen-24 found GAP-3 native pipeline decode was already landed on `origin/main`; PR #613 merged as docs truth-up/design note instead of duplicating implementation. -- Cohaagen-25 revalidated 35B-A3B on fresh `origin/main` @ `0a5ac3c5`: full native pipeline decode is output-correct vs native-CPU oracle, and fp16 TopK/Softmax execute on CUDA EP. -- Current native GPU blocker is now a cuDNN f16/bf16 `ReduceSum` compute-type bug plus decoder device-wiring; ORT-CUDA still hard-crashes. Cohaagen-26 is in flight to fix and measure. - -## 2026-08-03T10:00:00Z — PR #616 35B-A3B native GPU unblock - -- Authored merged PR #616 fixing cuDNN f16/bf16 reduce compute-type behavior and native pipeline decoder device wiring. -- Measured Qwen3.6-35B-A3B native GPU decode on H200 at 2726 ms/tok / 0.37 tok/s, correct vs native-CPU oracle and GPU-confirmed; ORT-CUDA still crashes, so this is a native capability number. -- Follow-up perf work is host/sync overhead and low GPU utilization, not correctness/unblock. - -## 2026-08-03T12:30:00Z — 35B-A3B Lever A reduce capture - -- Profiled 35B native CUDA decode and found capture shredded by 10240 fp16 `ReduceSum` seams from dense_fallback MoE, causing host/sync-bound 2725 ms/tok decode. -- Authored merged PR #618 making cuDNN float ReduceSum/Mean capture-eligible via descriptor/workspace caching, sync gating, warmed axes, and shape-stable capture-safe marking. -- Measured 35B decode improvement to 405 ms/tok / 2.47 tok/s (~6.7×), byte-exact vs CPU oracle; Lever B now targets ReduceSumSquare/Split/LinearAttention seams. - -## 2026-08-04T00:40:00Z — PR #625 native loader and 35B QMoE follow-through - -- Delivered the #625 native loader fix with GraphIo/GraphIoMetadata and rebased onto origin/main; locked out after Harry found the initializer-input leak. -- Confirmed ORT 1.28 still rejects fp16-activation/fp32-scale QMoE, making config-B a capability gap rather than a loader issue. -- Left GPU watcher PID 1060559 for config-A measurement while external vLLM occupies the GPU. +## Historical context +Older detailed dated entries through 2026-08-04T00:40:00Z — PR #625 native loader and 35B QMoE follow-through were moved to `history-archive.md` during Scribe compaction on 2026-08-11T03:25:00Z. Keep this live file focused on current routing-relevant context; full chronology is preserved in the archive. ## 2026-08-06T00:00:00Z — 35B-A3B native sparse QMoE shipped @@ -112,6 +20,7 @@ _Entries before 2026-07-31T03:03:15Z archived to `history-archive.md` (Scribe ro - Fixed #695 by disabling native host/device KV-mirror prefix reuse whenever `has_recurrent_state()` is true, forcing full recompute for hybrid Mamba/attention decoders. - Kept single-shot byte identity and added always-on gate coverage plus an env-gated GPU continuation regression where reused argmax matches the fresh oracle token `33803`. - PR #700 merged and closed #695; ORT paged-reuse residual tracked separately as #701. + ## 2026-08-06T12:30:27Z — PR #684 QMoE router parallelization - Cohaagen-37 profiled 35B-A3B QMoE decode and proved `qmoe_route` was the roofline limiter: 65.3% GPU time, rows=1 row-parallel top-k, GPU effectively idle. @@ -123,3 +32,9 @@ _Entries before 2026-07-31T03:03:15Z archived to `history-archive.md` (Scribe ro - Shipped PR #708/C3 making GatedDeltaNet Split capture-safe (resolved output-shape sizes, no host-read/sync): 13.415 → 12.132 ms/tok, 184 → 154 segments, token@119 `33803`; rejected unsafe C2 sync elision and moved on to pinned-vs-growing symbol classification after strict-C1 proved a no-op. - 2026-08-07: C1 build-time growing-symbol classifier mechanically collapsed capture segments (154→34, +4.3% tok/s) but was shelved behind #722 due to a 35B QMoE fp16 captured-vs-eager near-tie; revive after fp32-teacher-forcing oracle re-anchor. + +## 2026-08-11T03:25:00Z — GLM/DeepSeek native load gaps fixed + +- Fixed GLM-4-9B native CUDA load by honoring runtime CUDA KV capacity instead of reserving metadata `max_sequence_length=131072`; PR #770 auto-merge pending CI. +- Fixed DeepSeek-V2-Lite QMoE by accepting Cast-backed scale inputs as fp32 scale sources rather than requiring direct initializers; PR #771 auto-merge pending CI. +- Validation: GLM native greedy lock passes at 93.6 tok/s, DeepSeek lock passes at 52.5 tok/s, CPU/CUDA first-token top-40 sets match; no model-name gates, and root cause was not partial rotary or MLA. diff --git a/.squad/agents/harry/history-archive.md b/.squad/agents/harry/history-archive.md new file mode 100644 index 0000000000..74f816eb30 --- /dev/null +++ b/.squad/agents/harry/history-archive.md @@ -0,0 +1,92 @@ +# harry — History Archive + +## 2026-08-11T03:25:00Z: Scribe compaction from live history + +## 2026-07-29T05:55:00+0000 — PR #388 approved + +- Verified schema constraints are consumed by llguidance, and checked constraint precedence and validation behavior. +- Targeted HTTP tests and clippy passed in a scratch worktree. +- Confirmed two unrelated full-suite context-limit failures already occur on origin/main. + +## 2026-07-30T09:16:00Z — cfg-gated LoopStatePair hotfix + +- Landed PR #441, repairing the cfg-gated `LoopStatePair` import. + +## 2026-07-30T15:20:00Z — PR #477 merged (shape-inference container types + Sequence) + +- PR #477 merged (Lori APPROVED): shape-inference IR container-type model + Sequence foundation (#449). Additive `ValueType` layer, byte-identical tensor path, 4 Sequence ops, 300 tests. Unblocks the previously-deferred Sequence/Optional/Map/ZipMap propagation. + +## 2026-07-30T21:15:00Z — PR #486 merged (#449 inc2 sequence ops) + +- PR #486 merged: #449 inc2 — SequenceInsert/SequenceErase/SplitToSequence/ConcatFromSequence + seq↔tensor conversion; op catalog 213→217, shape-inference 258→262. +- In flight: #449 inc3 (#527 in review) and inc4. + +## 2026-07-31T00:25:00Z — PR #531 merged (#449 inc4); #534 held + +- PR #531 merged: #449 inc4 — SequenceMap + Scan container support + cross-subgraph capture. CLOSED issue #449. Container-type shape inference COMPLETE: additive `ValueType{Tensor|Sequence|Optional|Map}`, byte-identical tensor path guaranteed (gated on non-empty container map). Catalog 217 ops/262 entries. Deferred non-load-bearing: Optional/Map handlers, IR-persistence of `ValueType`. +- PR #534 (server contracts #481/#482 — `build_dirty` Option present-as-null; truncated predicate uses actual scan size): Melina APPROVED but HELD — targets Justin's active branch `feat/genai-demo-dashboard` (PR #476); code exists only there, not main. +- In flight (harry-5): generalize ORT `clone_value`/`clone_owned` to all POD dtypes (unblocks Bool / gemma-3n audio mask). + +## 2026-07-31T03:03:15Z — PR #540 merged: generalize ORT value-clone to all POD dtypes + +- PR #540 merged (requested by Justin): cloning an ORT cached `Value` now covers ALL POD dtypes via one dtype-agnostic raw-byte fallback — do not re-add per-dtype bail arms. `decode/values.rs::clone_value` and `onnx-genai-ort::value.rs::clone_owned` terminal arms use `Value::from_raw_bytes(value.as_raw_bytes()?.to_vec(), shape, dtype)` (typed f32/f16/bf16/i64 fast paths kept). 11 tests. +- Rule: use `as_raw_bytes()` (host-guarded — precise `InvalidArgument` on a device tensor), NEVER `to_raw_bytes()` (reads the data pointer blind). Unblocks the gemma-3n Bool audio mask. + +## 2026-07-31T08:48:28Z — #544 test deterministic fix; independent APPROVE #554 + +- Revised #544 flaky `async_pagein_fence_orders_weight_page_in_consumer` test (bf345904): root cause was the NEGATIVE poison-arm — pure wall-clock race. Fix: event-orders negative transfer after consumer via `record_compute_fence`+`copy_wait_fence`. Test green 5/5 parallel, fails 3/3 without the fix. Production code unchanged. +- Independent APPROVE for PR #554 (Mary's session-reuse recurrent-state fix). MERGED. + +## 2026-08-02T10:05:00+0000 — Reviews for #592 and #594 + +- Approved #592 after verifying the removed flag stayed gone and the no-Scan/no-sibling graph-property gate was mutation-proven. +- Rejected #594 only for rustfmt, confirmed logic/design were otherwise good, and assigned Deckard as revision author under reviewer lockout. + +## 2026-08-02T11:40:00+0000 — #595 review + +- Independently reviewed #595 and approved it after confirming the dangling reset caller, mutation-verifying reset coverage, checking hot-path behavior, and confirming bench builds. + +## 2026-08-02T15:45:00+0000 — PR #597 review + +- Approved #597 after independently rebuilding the fp32 acc-1 oracle, confirming native token 1909 matches oracle while ORT device-greedy token 821 is a near-tie artifact. +- Verified the 27-token stream lock is non-vacuous, resolved the ORT-logits-vs-stream puzzle, and merged the regression-lock PR. + +## 2026-08-02T19:00:00+0000 — PR #602 review + +- Rejected round 1 because the IR inliner failed open for attribute-parameterized functions by not binding call-site/ref_attr_name attrs; Add/Relu tests hid the byte-identity defect. +- Approved round 2 after Deckard added fail-closed attribute-parameter guards, metadata preservation, and mutation-proven `ParamLeakyRelu` coverage; #602 auto-merge armed. + +## 2026-08-02T19:50:00+0000 — PR #604 review + +- Approved #604 after confirming the phase-profile flake fix was test-only and scoped assertions to a unique phase name instead of the process-global stats map. +- Mutation-proved reset coverage by making `reg.clear()` a no-op and seeing the test fail; reproduced 30/30 clean full-parallel lib runs plus fmt/clippy clean. + +## 2026-08-03T02:40:00+0000 — PR #606 review + +- Approved #606 after verifying flag-off byte identity, fail-closed mixed-plan behavior, no stateful call path to standalone `hetero::execute`, and honest deferral of integrated execution to #603. +- Mutation-proved the guard by making the `Heterogeneous` arm return `Ok` and seeing `guard_enabled_mixed_fails_closed` fail; fmt, clippy, and C API checks were clean. + +## 2026-08-03T03:10:00+0000 — mobius PR #449 review + +- Approved #449 after verifying positional wiring, inert bias behavior, and all three call sites passing `None` in slot 4. +- Mutation-proved the new test by moving the bias formal to the end and seeing it fail; confirmed the diff is mobius-only with no onnx-genai changes. + +## 2026-08-03T07:40:00+0000 — PR #612 review + +- Approved #612 after verifying fp16 TopK writes back original raw fp16 values and only upcasts for total-order compare. +- Confirmed CUDA k-major non-final-axis order is ONNX-spec-correct and mutation-proved tie-break coverage; CPU non-final-axis ordering is a latent separate bug. + +## 2026-08-03T10:00:00Z — PR #616 review + +- Approved #616 after mutation-verifying both the cuDNN comp-type bug and native device fallback, and checking FFI safety, f32 byte-identity, bf16 NVRTC fallback, and shared dispatch generality. +- Gates clean: fmt/clippy for both crates, GPU parity, engine device tests, and EP reduce lib tests. + +## 2026-08-03T12:30:00Z — PR #618 review + +- Approved Lever A after mutation-verifying cache-key shape coverage and sync gating, and checking miss-during-capture, warm-before-capture, shape-change aborts, f32/f16 byte behavior, and bf16 NVRTC exclusion. +- Verified fmt/clippy clean and 6/6 capture + 3/3 parity GPU tests; noted non-blocking need for no-cuDNN test guards. + +## 2026-08-04T00:40:00Z — PR #625 loader review + +- Rejected #625 rev1 for a major initializer-input leak in the native loader metadata path and assigned revision away from the locked-out author. +- Approved rev2 after Quaid added initializer exclusion mirroring `graph_builder.rs` plus metadata==Session KV-geometry parity coverage. diff --git a/.squad/agents/harry/history.md b/.squad/agents/harry/history.md index a79fcf21c0..2eb5f3b8b0 100644 --- a/.squad/agents/harry/history.md +++ b/.squad/agents/harry/history.md @@ -1,92 +1,7 @@ # Harry — History -## 2026-07-29T05:55:00+0000 — PR #388 approved - -- Verified schema constraints are consumed by llguidance, and checked constraint precedence and validation behavior. -- Targeted HTTP tests and clippy passed in a scratch worktree. -- Confirmed two unrelated full-suite context-limit failures already occur on origin/main. - -## 2026-07-30T09:16:00Z — cfg-gated LoopStatePair hotfix - -- Landed PR #441, repairing the cfg-gated `LoopStatePair` import. - -## 2026-07-30T15:20:00Z — PR #477 merged (shape-inference container types + Sequence) - -- PR #477 merged (Lori APPROVED): shape-inference IR container-type model + Sequence foundation (#449). Additive `ValueType` layer, byte-identical tensor path, 4 Sequence ops, 300 tests. Unblocks the previously-deferred Sequence/Optional/Map/ZipMap propagation. - -## 2026-07-30T21:15:00Z — PR #486 merged (#449 inc2 sequence ops) - -- PR #486 merged: #449 inc2 — SequenceInsert/SequenceErase/SplitToSequence/ConcatFromSequence + seq↔tensor conversion; op catalog 213→217, shape-inference 258→262. -- In flight: #449 inc3 (#527 in review) and inc4. - -## 2026-07-31T00:25:00Z — PR #531 merged (#449 inc4); #534 held - -- PR #531 merged: #449 inc4 — SequenceMap + Scan container support + cross-subgraph capture. CLOSED issue #449. Container-type shape inference COMPLETE: additive `ValueType{Tensor|Sequence|Optional|Map}`, byte-identical tensor path guaranteed (gated on non-empty container map). Catalog 217 ops/262 entries. Deferred non-load-bearing: Optional/Map handlers, IR-persistence of `ValueType`. -- PR #534 (server contracts #481/#482 — `build_dirty` Option present-as-null; truncated predicate uses actual scan size): Melina APPROVED but HELD — targets Justin's active branch `feat/genai-demo-dashboard` (PR #476); code exists only there, not main. -- In flight (harry-5): generalize ORT `clone_value`/`clone_owned` to all POD dtypes (unblocks Bool / gemma-3n audio mask). - -## 2026-07-31T03:03:15Z — PR #540 merged: generalize ORT value-clone to all POD dtypes - -- PR #540 merged (requested by Justin): cloning an ORT cached `Value` now covers ALL POD dtypes via one dtype-agnostic raw-byte fallback — do not re-add per-dtype bail arms. `decode/values.rs::clone_value` and `onnx-genai-ort::value.rs::clone_owned` terminal arms use `Value::from_raw_bytes(value.as_raw_bytes()?.to_vec(), shape, dtype)` (typed f32/f16/bf16/i64 fast paths kept). 11 tests. -- Rule: use `as_raw_bytes()` (host-guarded — precise `InvalidArgument` on a device tensor), NEVER `to_raw_bytes()` (reads the data pointer blind). Unblocks the gemma-3n Bool audio mask. - -## 2026-07-31T08:48:28Z — #544 test deterministic fix; independent APPROVE #554 - -- Revised #544 flaky `async_pagein_fence_orders_weight_page_in_consumer` test (bf345904): root cause was the NEGATIVE poison-arm — pure wall-clock race. Fix: event-orders negative transfer after consumer via `record_compute_fence`+`copy_wait_fence`. Test green 5/5 parallel, fails 3/3 without the fix. Production code unchanged. -- Independent APPROVE for PR #554 (Mary's session-reuse recurrent-state fix). MERGED. -## 2026-08-02T10:05:00+0000 — Reviews for #592 and #594 - -- Approved #592 after verifying the removed flag stayed gone and the no-Scan/no-sibling graph-property gate was mutation-proven. -- Rejected #594 only for rustfmt, confirmed logic/design were otherwise good, and assigned Deckard as revision author under reviewer lockout. - -## 2026-08-02T11:40:00+0000 — #595 review - -- Independently reviewed #595 and approved it after confirming the dangling reset caller, mutation-verifying reset coverage, checking hot-path behavior, and confirming bench builds. - -## 2026-08-02T15:45:00+0000 — PR #597 review - -- Approved #597 after independently rebuilding the fp32 acc-1 oracle, confirming native token 1909 matches oracle while ORT device-greedy token 821 is a near-tie artifact. -- Verified the 27-token stream lock is non-vacuous, resolved the ORT-logits-vs-stream puzzle, and merged the regression-lock PR. - -## 2026-08-02T19:00:00+0000 — PR #602 review - -- Rejected round 1 because the IR inliner failed open for attribute-parameterized functions by not binding call-site/ref_attr_name attrs; Add/Relu tests hid the byte-identity defect. -- Approved round 2 after Deckard added fail-closed attribute-parameter guards, metadata preservation, and mutation-proven `ParamLeakyRelu` coverage; #602 auto-merge armed. - -## 2026-08-02T19:50:00+0000 — PR #604 review - -- Approved #604 after confirming the phase-profile flake fix was test-only and scoped assertions to a unique phase name instead of the process-global stats map. -- Mutation-proved reset coverage by making `reg.clear()` a no-op and seeing the test fail; reproduced 30/30 clean full-parallel lib runs plus fmt/clippy clean. - -## 2026-08-03T02:40:00+0000 — PR #606 review - -- Approved #606 after verifying flag-off byte identity, fail-closed mixed-plan behavior, no stateful call path to standalone `hetero::execute`, and honest deferral of integrated execution to #603. -- Mutation-proved the guard by making the `Heterogeneous` arm return `Ok` and seeing `guard_enabled_mixed_fails_closed` fail; fmt, clippy, and C API checks were clean. - -## 2026-08-03T03:10:00+0000 — mobius PR #449 review - -- Approved #449 after verifying positional wiring, inert bias behavior, and all three call sites passing `None` in slot 4. -- Mutation-proved the new test by moving the bias formal to the end and seeing it fail; confirmed the diff is mobius-only with no onnx-genai changes. - -## 2026-08-03T07:40:00+0000 — PR #612 review - -- Approved #612 after verifying fp16 TopK writes back original raw fp16 values and only upcasts for total-order compare. -- Confirmed CUDA k-major non-final-axis order is ONNX-spec-correct and mutation-proved tie-break coverage; CPU non-final-axis ordering is a latent separate bug. - -## 2026-08-03T10:00:00Z — PR #616 review - -- Approved #616 after mutation-verifying both the cuDNN comp-type bug and native device fallback, and checking FFI safety, f32 byte-identity, bf16 NVRTC fallback, and shared dispatch generality. -- Gates clean: fmt/clippy for both crates, GPU parity, engine device tests, and EP reduce lib tests. - -## 2026-08-03T12:30:00Z — PR #618 review - -- Approved Lever A after mutation-verifying cache-key shape coverage and sync gating, and checking miss-during-capture, warm-before-capture, shape-change aborts, f32/f16 byte behavior, and bf16 NVRTC exclusion. -- Verified fmt/clippy clean and 6/6 capture + 3/3 parity GPU tests; noted non-blocking need for no-cuDNN test guards. - -## 2026-08-04T00:40:00Z — PR #625 loader review - -- Rejected #625 rev1 for a major initializer-input leak in the native loader metadata path and assigned revision away from the locked-out author. -- Approved rev2 after Quaid added initializer exclusion mirroring `graph_builder.rs` plus metadata==Session KV-geometry parity coverage. +## Historical context +Older detailed dated entries through 2026-08-04T00:40:00Z — PR #625 loader review were moved to `history-archive.md` during Scribe compaction on 2026-08-11T03:25:00Z. Keep this live file focused on current routing-relevant context; full chronology is preserved in the archive. ## 2026-08-06T00:00:00Z — PR #676 review @@ -98,6 +13,7 @@ - Approved #700 after verifying the recurrent-state gate disables host/device KV-mirror reuse for hybrid decoders without regressing single-shot behavior. - Confirmed the env-gated GPU continuation regression compares reused continuation argmax against the fresh oracle (`33803`). - Flagged a minor residual on the ORT paged-reuse path (`kv_bridge.rs:407`); coordinator filed #701. + ## 2026-08-06T12:30:27Z — Reviews #684/#692 and cache-reuse bug - Approved PR #684 after verifying the parallel `qmoe_route` top-k reduction is byte-exact with the old serial total-order scan and leaves softmax aggregation in original order. @@ -107,3 +23,9 @@ ## 2026-08-06T19:40:00Z — PR #708 review - Approved #708 after verifying capture-safe Split sizes are statically inferred and byte-equivalent to the old host-read path; left nits for dead fallback docs and redundant re-locking. + +## 2026-08-11T03:25:00Z — Reviewing GLM/DeepSeek native fixes + +- Dispatched to review PR #770 and PR #771. +- Focus areas: KV precedence soundness for GLM native CUDA capacity, and Cast-backed QMoE scale placement dtype/region correctness for DeepSeek. +- Current status is running; no review decision recorded yet. diff --git a/.squad/agents/quaid/history.md b/.squad/agents/quaid/history.md index 2c3ed90d8f..8f5070244e 100644 --- a/.squad/agents/quaid/history.md +++ b/.squad/agents/quaid/history.md @@ -10,3 +10,8 @@ - Fixed the #676 oracle test in merged PR #692 by using a fresh engine for teacher-forced logits after autoregressive decode. - Root cause: reused-engine teacher forcing restored attention KV but not Mamba conv/recurrent state, causing wrong argmax 279; fresh-engine teacher forcing correctly returns oracle token 33803. - Flagged the underlying hybrid-Mamba prefix-cache-reuse engine bug, now filed as issue #695. + +## 2026-08-11T03:25:00Z — Megakernel Phase 0 dispatched + +- Dispatched for persistent single-op QMoE decode kernel work: counter-synchronized FC1→FC2 pipeline, remove four scratch DRAM round-trips, preserve fp32 accumulation order. +- Current status is running; continuation gates are oracle margin 0.09375 and >=3% wall-clock improvement. diff --git a/.squad/agents/roper/history.md b/.squad/agents/roper/history.md new file mode 100644 index 0000000000..32b12e8db6 --- /dev/null +++ b/.squad/agents/roper/history.md @@ -0,0 +1,10 @@ +# roper — History + +## Role +Research and feasibility specialist. Produces read-only technical scoping, competitive analysis, and go/no-go recommendations. + +## 2026-08-11T03:25:00Z — Megakernel feasibility scoping delivered + +- Delivered a read-only scoping study concluding a whole-step/persistent megakernel is the only remaining lever that attacks batch-1 decode GPU-side bubbles after CUDA graph capture. +- Found vLLM full CUDA graph and llama.cpp are capture/per-op systems, not true megakernels; Mirage MPK is blueprint-only for this Rust/ONNX/int4-QMoE stack. +- Recommended Phase 0 only first: persistent single-op QMoE decode, preserving fp32 accumulation order, gated on oracle margin 0.09375 and >=3% model wall-clock improvement. diff --git a/.squad/decisions-archive/2026-08.md b/.squad/decisions-archive/2026-08.md index 2a33971d0e..10517b026e 100644 --- a/.squad/decisions-archive/2026-08.md +++ b/.squad/decisions-archive/2026-08.md @@ -4852,3 +4852,347 @@ restores attention KV only). This is fine for the test (fresh engines), but the prefix/decode-cache reuse producing wrong next-token logits for hybrid models in a continuation scenario may warrant a separate engine-correctness look. +## 2026-08-11T03:25:00Z: Roper megakernel feasibility inbox drop + +# Roper: Whole-Step Megakernel Decode — Feasibility Scoping Study (NOT a build) + +Date: 2026-08-11 +Author: Roper (feasibility scoping; research + read-only codebase analysis) +Requested by: Justin (@justinchuby) +Scope: Decision-support for a go/no-go on a multi-day megakernel investment. + NO kernels written, NO code PRs. Explainer: GitHub issue #769. +GPU target: H200 (132 SMs, ~4.8 TB/s HBM3e, 228 KB smem/SM, 65536 regs/SM). +Model: 35B-A3B QMoE (hybrid linear-attn + GQA), int4 experts, decode batch=1. +Baseline: ~11.1 ms/tok (~90 tok/s), CUDA-graph capture already ON. + +--- + +## 0. Problem framing (what is actually true) + +Three surgical QMoE-decode experiments have proven the decode GEMV is +**occupancy / MLP / barrier bound at M=1**, NOT compute/scratch/launch-bound: +- FC2/down+combine fusion: +0.08% (noise) — `cohaagen-qmoe-fc2-fusion.md` +- ILP-2 K-unroll: regressed (active warps 13.2→9.7) — `cohaagen-qmoe-latency-floor.md` +- 128-bit vectorized int4 weight read: **+4.9% regression** (fewer warps → less MLP) + — `quaid-qmoe-int4-dp4a.md` + +Consistent signal: any transform that trades away active warps for per-thread +width/ILP loses. NCU on `qmoe_linear_f32`: DRAM ~4% of peak, barrier ~37%, +no-eligible ~22%. Measured QMoE kernels are ~2.8 ms/tok vs an int4 **bandwidth +floor of ~0.10 ms/tok** (packed weights) / ~0.13–0.24 ms/tok (incl. scales/zp) +— i.e. **~10–20× above roofline** (`cohaagen-qmoe-latency-floor.md`). + +CPU launch overhead is already removed (CUDA graph capture on). The residual gap +is therefore GPU-side and structural: (a) inter-kernel DRAM round-trips through +scratch, and (b) per-kernel internal barrier/occupancy stalls + straggler-block +drain between kernels. The ONLY lever that attacks both simultaneously is a +persistent/tiled or whole-step **megakernel**. This study scopes that. + +--- + +## 1. Competitive verdict + +**Do the incumbents have a true megakernel? No.** + +- **vLLM** `full_cuda_graph=True` (CUDAGraphMode FULL) is **graph capture** of the + whole forward (attention included, only for FA3/Triton/FlashMLA backends), NOT + a fused megakernel. Data still round-trips through global memory between kernels; + it only removes launch/CPU-submit overhead — which we *already* have via + `ONNX_GENAI_CUDA_GRAPH=1`. + Source: vLLM CUDA Graphs design doc + https://github.com/vllm-project/vllm/blob/main/docs/design/cuda_graphs.md + (Note: some secondary summaries loosely call FULL mode a "megakernel" — that is + imprecise; it is capture/replay of many kernels, not a single fused kernel.) +- **llama.cpp** uses CUDA Graphs + a few small (bias+activation) fusions, but is + fundamentally **one-kernel-per-op with intermediates in global memory** — not a + megakernel. + Source: NVIDIA "Optimizing llama.cpp AI Inference with CUDA Graphs" + https://developer.nvidia.com/blog/optimizing-llama-cpp-ai-inference-with-cuda-graphs/ + and llama.cpp issue #6763 https://github.com/ggml-org/llama.cpp/issues/6763 + +**Where a true megakernel wins — and where it does NOT:** +- WINS: **batch-1 / low-latency decode** (memory- and latency-bound). Hazy Research + reports 78% of H100 BW vs ~50% for vLLM/SGLang, sub-1ms Llama-1B forward, 1.5× + vs SGLang / 2.5× vs vLLM (3.5× vs vLLM on B200). + Source: https://hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles +- Mirage MPK reports **1.2–6.7× decode-latency** wins vs kernel-per-op frameworks, + again concentrated at small batch. + Source: arXiv:2512.22219 https://arxiv.org/abs/2512.22219 ; + https://github.com/mirage-project/mirage (mpk branch) +- Does NOT help **large-batch throughput** (compute-bound; kernels already saturate + SMs and tensor cores — a megakernel's straggler/launch savings vanish). Our target + is exactly batch-1 decode, so the regime is favorable. + +**Differentiation verdict:** A working int4-QMoE whole-step megakernel would be a +genuine, defensible differentiator in the batch-1 latency regime that neither vLLM +nor llama.cpp occupies. It is worthless for throughput serving. Our stated goal +(~11 ms/tok single-sequence decode) is squarely in the winning regime. + +--- + +## 2. Build vs. reuse: can we adopt Mirage MPK? + +**Verdict: cannot adopt directly; hand-roll, using MPK/Hazy as a design blueprint.** + +MPK input assumptions (from README + source, mpk branch): +- **Framework:** PyTorch / HuggingFace model zoo; graph is defined in **Python** + via `mi.PersistentKernel(...)` + chained fused ops (`rmsnorm_linear_layer`, + `silu_mul_linear`, `attention`, `moe_linear`). Demo target is Qwen3-8B. +- **Runtime:** a Python-driven persistent kernel; `num_workers`/schedulers must + equal physical SM counts; multi-GPU via **nvshmem** (`world_size`, `mpi_rank`). +- **Dtype/quant:** tasks are **bf16** MMA. `ampere/moe_linear.cuh` uses + `SM80_16x8x16_F32BF16BF16F32` tensor-core MMA — **no int4 path**. The only + quantization in-tree is `blackwell/fp8_group_gemm_sm100.cuh` (fp8, sm100/Blackwell). + There is **no int4 / block-affine dequant MoE** kernel. + +Why it doesn't fit us: +1. Our model is **not** a PyTorch/HF graph — it's an ONNX graph run by a pure-Rust + ORT-style per-op executor with a Rust CUDA EP. MPK's Python runtime + task-graph + builder is a foreign execution model. +2. Our experts are **int4 block-affine quantized** (block_size 16, optional zero + points), decoded scalar→fp32. MPK has no int4 task; we'd have to author int4 + QMoE task kernels inside *their* CUDA template framework anyway. +3. MPK is oriented toward multi-GPU (nvshmem) and requires SM-count-matched worker + config; it would need to coexist with our CUDA-graph capture + KV-cache/paging, + which it knows nothing about. +4. The 35B fp32 oracle lock (token 33803, margin 0.09375) demands byte-exact fp32 + accumulation order; adopting MPK's bf16 tensor-core MoE would change numerics + and almost certainly break the lock. + +**Reusable from MPK/Hazy (blueprint, not code):** +- **SM-level task/instruction graph** compiled once, replayed for many steps + (Hazy: schedule computed Python-side, reused for hundreds of forwards). +- **Counter-based cross-instruction synchronization** in global memory (increment + on complete, spin-wait to target) instead of a grid barrier — avoids the + cooperative-launch co-residency constraint. (Hazy Issue 3/3.) +- **Shared-memory paging** (Hazy: 13×16KiB pages on H100, explicit request/release) + to pipeline next-instruction weight loads while the previous stores — this is the + mechanism that turns ~4% DRAM utilization into ~78%. +- **Chunked producer/consumer** of MLP intermediates (Hazy: 4 chunks with per-chunk + counters so down-proj starts on chunk 0) — directly maps to our FC1→SwiGLU→FC2. +- MPK's instruction menu maps 1:1 to our chain: `norm_linear`, `rotary_embedding`, + `single_batch_gqa` + `merge_splitkv`, `silu_mul_linear`, `moe_linear`. + +Effort tradeoff: adopting MPK would still require writing int4 QMoE task kernels +*and* bridging Python↔Rust↔ONNX↔our KV cache — strictly more integration than +hand-rolling one persistent kernel in our existing NVRTC pipeline, with less +control over the oracle-critical fp32 accumulation. **Hand-roll.** + +--- + +## 3. Codebase reality: the decode op-chain and where DRAM round-trips are + +Per transformer layer, decode (rows=1) chain and backing kernels: + +| Step | Op | Kernel (file) | DRAM round-trip? | +|---|---|---|---| +| 1 | RMSNorm (input) | `skip_rmsnorm_f16_*` (`normalization.rs`) | writes normed activ | +| 2 | QKV proj | `matmul_nbits_gemv_*_splitk` (`matmul_nbits.rs`) | Q,K,V to DRAM | +| 3 | RoPE | `rotary_embedding.rs` | rewrites Q,K | +| 4 | Attention | `LinearAttention` + GQA split-K decode + merge (`gqa_decode.rs`, `linear_attention`, `group_query_attention.rs`) | KV cache + attn out | +| 5 | O-proj | `matmul_nbits_gemv_*` | attn-proj to DRAM | +| 6 | RMSNorm (post-attn) | `skip_rmsnorm_*` | + residual Add | +| 7 | Router | `matmul_nbits` + `qmoe_route` (`qmoe.rs:63`) | logits, selected experts/weights | +| 8 | QMoE FC1 | `qmoe_linear_*` (`qmoe.rs:1367`) | **`fc1_output` scratch** | +| 9 | QMoE FC3 (opt) | `qmoe_linear_*` (`qmoe.rs:1377`) | **`fc3_output` scratch** | +| 10 | SwiGLU/act | `qmoe_activate` (`qmoe.rs:1389`) | **`activated` scratch** | +| 11 | QMoE FC2/down | `qmoe_linear_*` (`qmoe.rs:1390`) | **`route_output` scratch** | +| 12 | Combine (router-weighted) | `qmoe_combine_*` (`qmoe.rs:1401`) | expert-weighted sum | +| 13 | Residual Add | `elementwise.rs` | hidden state | + +QMoE decode alone = **5 kernel launches** (FC1, [FC3], activate, FC2, combine) with +**4 global scratch buffers** (`fc1_output`, `fc3_output`, `activated`, +`route_output`) allocated from a warmed `ScratchPool` (`qmoe.rs:1270-1291`), each a +full write+read DRAM round-trip. The non-grouped rows==1 branch is `qmoe.rs:1367-1408`. +Per Cohaagen's Nsight grouping: QMoE route/linear/activate/combine = **32.8% / +2.84 ms/tok**; MatMulNBits GEMV = 23.5% / 2.05 ms/tok; elementwise/copy/act = 22.2% +/ 1.94 ms/tok; norms = 11.9% / 1.02 ms/tok; LinearAttention 6.2% / 0.54 ms/tok +(`cohaagen-qmoe-profile-nextlever.md`). + +These per-op DRAM round-trips + straggler-block drain between ~7 launches/layer × +~40 layers are precisely the "memory pipeline bubbles" a megakernel removes. + +--- + +## 4. Occupancy budget & the hardest integration point + +**Cooperative-launch co-residency risk.** A grid-wide barrier +(`cudaLaunchCooperativeKernel` + `grid.sync()`) requires *all* resident blocks to +co-reside on the SMs simultaneously or it **deadlocks**. On H200: 132 SMs, 65536 +regs/SM, 228 KB smem/SM. Our `qmoe_linear` runs 256 threads/block at ~13 active +warps/block. A cooperative kernel must cap total blocks at +`132 × (blocks that fit per SM)` and cannot exceed occupancy — this fights the very +resource (many resident warps for MLP) the decode GEMV depends on. + +**Recommended mitigation (matches Hazy/MPK):** do NOT use a grid barrier. Use the +**counter-in-global-memory** persistent-worker model: launch exactly +`num_workers ≈ 132` persistent blocks, each pulling instructions and spinning on +dependency counters. This sidesteps the co-residency deadlock entirely and is what +both Hazy and MPK actually do. + +**Coexistence with our CUDA-graph capture — the hardest integration point.** +- Today: 14 captured segments / 13 eager seams (`cohaagen-qmoe-profile-nextlever.md`), + seams from Squeeze copies + GQA growing-KV symbols. Capture eligibility gated by + `subgraph_graph_capturable` (`capture.rs:10`), and kernels are capture-safe + (no sync, no per-call malloc, fixed geometry — see `gqa_decode.rs:14-22`). +- A persistent megakernel is itself **one capturable node** IF it obeys the same + contract (fixed launch geometry, device-resident length selection, no host sync). + Good news: our kernels already follow this discipline, so a persistent kernel can + be **recorded into the existing graph** rather than replacing capture. +- KV-cache/paging: the megakernel must read/write the same KV layout the GQA + split-K decode uses (`gqa_decode.rs`, `sparse_kv_gather.rs`) with device-resident + valid-length selecting active splits — the persistent worker must consume the same + device-side length, not a captured constant. This is the integration crux for + Phase 2 and the main reason Phase 2 is high-risk. + +**There is NO cooperative/persistent infrastructure today.** Confirmed: zero +`cudaLaunchCooperativeKernel`/`grid.sync`/`cooperative_groups` usage in +`crates/onnx-runtime-ep-cuda/src` (the two `qmoe.rs` "cooperative" hits at lines 63 +and 1856 are comments describing block-wide argmax, not grid cooperation). This is +greenfield infrastructure. + +--- + +## 5. What MUST stay byte-exact (correctness gate) + +- **Primary lock:** 35B fp32 teacher-forced oracle — argmax token **33803**, + margin `logprob(33803) − logprob(5342) = 0.09375`. This is the ship gate. +- **Key de-risking finding (from `quaid-qmoe-int4-dp4a.md`):** reduction *restructuring* + held the margin **bit-identical** (0.09375) — the vectorized int4 read changed + block-reduction fan-in yet kept teacher-forced numerics exact. Therefore the + correctness risk of a megakernel is **integration, not numerics**, PROVIDED we + preserve per-thread fp32 accumulation order within each expert GEMV. +- **Tripwire caveat:** the *autoregressive* #722 tripwire (benign set {33803, 46283}) + is a documented fp16 near-tie coin-flip that flips under *any* tiny perturbation. + Any megakernel experiment should gate on the **teacher-forced** margin (robust), + and treat the autoregressive tripwire as informational, matching how Quaid's + bit-identical change still "failed" only the narrow autoregressive set. + +--- + +## 6. Phased plan (each phase independently sheddable, with a measurable gate) + +### Phase 0 — Persistent single-op QMoE decode kernel (smallest proof) +- **What:** Replace the 5-launch QMoE decode (FC1→[FC3]→act→FC2→combine) with ONE + persistent kernel keeping FC1/SwiGLU intermediates in shared memory/registers, + eliminating `fc1_output`/`fc3_output`/`activated`/`route_output` DRAM round-trips. + Use counter-sync producer/consumer for the FC1→FC2 chunk dependency (Hazy's 4-chunk + MLP trick). Keep per-thread fp32 accumulation order for the oracle. +- **Ceiling (bounded):** QMoE is 32.8% / 2.84 ms/tok of GPU kernel time. Removing + the 4 scratch round-trips + 4 launch/drain seams cannot exceed the QMoE slice. + Realistic: recover the ~4 inter-kernel bubbles + scratch traffic on a + latency-bound (not BW-bound) kernel → **~15–25% of the QMoE slice**, i.e. + ~0.4–0.7 ms/tok, **~4–6% model-level** (11.1 → ~10.5 ms/tok, ~90 → ~95 tok/s). + This is the "~MoE-only +23% ceiling" reference bound; treat +4–6% wall-clock as + the likely outcome and >3% as the ship gate. +- **Effort:** ~2–4 person-days (one persistent kernel, reuses NVRTC pipeline + + existing ScratchPool + capture contract; no cross-layer state). +- **Top integration risk:** holding fp32 accumulation order byte-exact while keeping + enough resident warps (the ILP/vec experiments show this kernel punishes any + occupancy loss). If the persistent structure reduces active warps, it can regress + like ILP-2 did — this is the real uncertainty Phase 0 measures. +- **Gate:** 35B teacher-forced margin holds at 0.09375 AND ≥3% steady-state + wall-clock win (profile_native --pipeline --steady). Fail → shed, stop. + +### Phase 1 — Single transformer BLOCK megakernel (off the capture path) +- **What:** Fuse ONE layer's norm + QKV + RoPE + attention + o-proj + norm + MoE + + residual into a persistent instruction stream (Hazy instruction menu), run + standalone (NOT integrated with capture) and measured against the current per-op + path for one layer. Proves the SM-level scheduler + shared-memory paging + counter + sync end-to-end without touching production capture/KV integration. +- **Expected gain (if extrapolated to all layers):** attacks the full per-layer + bubble budget (norms 11.9% + elementwise/copy 22.2% + QMoE 32.8% inter-op seams). + Hazy's 78% vs ~50% BW ≈ 1.5× suggests a **1.3–1.6× decode ceiling** if fully + realized (11.1 → ~7–8.5 ms/tok, ~90 → ~120–140 tok/s). Phase 1 only measures the + single-block fraction; it validates the mechanism, not the full number. +- **Effort:** ~8–15 person-days (instruction interpreter, smem pager, counter sync, + int4 QMoE + GQA split-K instructions, all fp32-order-preserving). +- **Top integration risk:** the on-SM interpreter + shared-memory paging is new + infrastructure; getting attention (split-K + online-softmax merge) into an + instruction while preserving the GQA decode numerics is the hardest single piece. +- **Gate:** single-block megakernel matches per-op output within oracle tolerance + AND single-block GPU time beats the summed per-op block time by a margin that + extrapolates to >10% model-level. Fail → shed; Phase 0 value stands alone. + +### Phase 2 — Whole-step megakernel integrated with capture + KV-cache +- **What:** All ~40 layers as one persistent kernel recorded as a single graph node, + consuming device-resident KV valid-length, coexisting with the existing capture + lifecycle (`graph.rs`, `capture.rs`) and KV layout (`gqa_decode.rs`, + `sparse_kv_gather.rs`). +- **Expected gain:** the full **1.3–1.6× decode** target (11.1 → ~7–8.5 ms/tok, + ~120–140 tok/s), i.e. the payoff that justifies the whole effort. This is the + number to sell a go/no-go on. +- **Effort:** ~15–30 person-days on top of Phase 1 (capture coexistence, KV/paging, + growing-sequence symbol handling that currently causes the 10 GQA eager seams, + multi-layer weight streaming/pipelining, full oracle + attention-only gates). +- **Top integration risk:** the persistent kernel vs. CUDA-graph-capture coexistence + and dynamic KV length — today there are 13 eager seams (10 from growing-KV GQA); + a whole-step kernel must handle variable context length device-side without + re-capture. Highest integration risk in the plan. +- **Gate:** full 35B oracle PASS + attention-only model gate + ≥1.25× steady-state + decode win. Fail → the Phase 0/1 wins may still ship independently. + +--- + +## 7. Recommendation + +- **Start Phase 0 only.** It is bounded (~2–4 person-days), reuses existing NVRTC + + ScratchPool + capture infrastructure, has a clean ship gate (oracle margin + + >3% wall-clock), and directly tests the single biggest open uncertainty: + *can a persistent, scratch-free QMoE decode structure keep enough resident warps + to win on a latency/occupancy-bound kernel, while holding fp32 accumulation order + byte-exact?* Every prior surgical attempt failed on the occupancy trade; Phase 0 + is the smallest experiment that decisively answers whether the persistent + structure escapes that trap. +- **Do NOT green-light Phases 1–2 yet.** They are the real prize (1.3–1.6× decode) + but are 8–30 person-day builds of greenfield SM-scheduler/interpreter + capture/KV + integration. Gate them on a **positive Phase 0** result: if Phase 0 cannot even + recover the QMoE scratch/bubble slice without an occupancy regression, the + whole-step kernel will hit the same wall at far higher cost. +- **Smallest uncertainty-reducing experiment:** Phase 0's persistent FC1→SwiGLU→FC2 + fusion with counter-synced chunking, measured on the teacher-forced oracle. If it + clears +3% with the margin at 0.09375, greenlight Phase 1. If it regresses like + ILP-2/vec-read, the megakernel thesis is materially weakened and the multi-day + investment should be declined. +- **Build, don't reuse:** Mirage MPK is a design reference, not an adoptable backend + (PyTorch/bf16/fp8, Python runtime, no int4, no ONNX). Hand-roll on our NVRTC path. + +--- + +## Sources (web) +- Hazy Research, "Look Ma, No Bubbles! Designing a Low-Latency Megakernel for + Llama-1B": https://hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles + (78% vs ~50% BW; sub-1ms H100; 1.5× SGLang / 2.5× vLLM; instruction interpreter; + shared-memory paging; counter synchronization; 4-chunk MLP producer/consumer; + CUDA-graph launch cost still ~1.3µs, PDL too coarse.) +- Hazy Research Megakernels code: https://github.com/HazyResearch/Megakernels +- Mirage MPK paper, "Mirage Persistent Kernel: A Compiler and Runtime for + Mega-Kernelizing Tensor Programs," arXiv:2512.22219 + https://arxiv.org/abs/2512.22219 (SM-level task graph; event/counter scheduling; + 1.2–6.7× decode latency; small-batch regime). +- Mirage repo (mpk branch): https://github.com/mirage-project/mirage/tree/mpk + (Python PersistentKernel API; bf16 tasks; moe_linear.cuh uses SM80 bf16 MMA, no + int4; only fp8 group gemm on Blackwell; nvshmem multi-GPU; SM-count-matched workers.) +- vLLM CUDA Graphs design (full_cuda_graph = capture, not megakernel): + https://github.com/vllm-project/vllm/blob/main/docs/design/cuda_graphs.md +- NVIDIA, "Optimizing llama.cpp AI Inference with CUDA Graphs" (graphs group + per-op kernels; not a megakernel): + https://developer.nvidia.com/blog/optimizing-llama-cpp-ai-inference-with-cuda-graphs/ +- llama.cpp CUDA Graphs issue #6763: https://github.com/ggml-org/llama.cpp/issues/6763 + +## Internal citations +- QMoE decode 5-launch chain + 4 scratch round-trips: + onnx-genai:crates/onnx-runtime-ep-cuda/src/kernels/qmoe.rs:1270-1408 +- QMoE route kernel (block-wide argmax comment, not grid cooperation): + onnx-genai:crates/onnx-runtime-ep-cuda/src/kernels/qmoe.rs:63-73,1856-1858 +- GQA decode capture-safety contract (no sync / no per-call malloc / fixed geometry): + onnx-genai:crates/onnx-runtime-ep-cuda/src/kernels/gqa_decode.rs:14-56 +- Capture eligibility gate: onnx-genai:crates/onnx-runtime-ep-cuda/src/capture.rs:10-39 +- Graph capture lifecycle: onnx-genai:crates/onnx-runtime-ep-cuda/src/graph.rs:1-60 +- No cooperative-launch infra anywhere in the CUDA EP (verified via grep). +- Profiling/attribution: .squad/decisions/inbox/cohaagen-qmoe-profile-nextlever.md +- Latency-floor/bandwidth roofline: (worktree) cohaagen-qmoe-latency-floor.md +- Reduction-restructure is bit-identical (numerics = low risk): + .squad/decisions/inbox/quaid-qmoe-int4-dp4a.md +- FC2/combine fusion ~0.08% (scratch removal alone insufficient): + .squad/decisions/inbox/cohaagen-qmoe-fc2-fusion.md diff --git a/.squad/decisions.md b/.squad/decisions.md index 2786ca365c..bd42153714 100644 --- a/.squad/decisions.md +++ b/.squad/decisions.md @@ -54,6 +54,16 @@ Archive by SIZE, not age. Age-only no-ops during high-volume campaigns because m - Dense prefetch stays eviction-neutral: executor-driven lazy-weight prefetch is scoped to dense MatMulNBits weights, and CUDA admits prefetch only when it fits without eviction or lease growth so MoE behavior and cache-victim selection are unchanged. - Native CUDA KV bindings may reserve the full-context VMM address range while exposing bucketed physical strides; growth commits the next bucket and repacks valid prefixes in place. Full-context strides committed one granule per head and worsened arena pressure, so floor claims must wait for #694 because ledger-refusal passes made the sweep non-monotonic. +### 2026-08-11: Whole-step megakernel is the next bounded batch-1 latency lever + +**By:** Roper + +**What:** Roper's read-only scoping concluded that a whole-step/persistent megakernel is the only remaining lever that attacks the current batch-1 decode regime's GPU-side bubbles after CUDA graph capture. vLLM full CUDA graph and llama.cpp are capture/per-op systems, not true megakernels; Mirage MPK is not directly adoptable for this Rust/ONNX/int4 QMoE stack. + +**Why:** Prior FC2 fusion, ILP, and vectorized-int4 attempts show QMoE decode is occupancy/MLP/barrier bound rather than launch or bandwidth bound. Recommendation is Phase 0 only first: a persistent single-op QMoE decode kernel that removes FC1/FC3/activation/FC2/combine scratch round-trips while preserving fp32 accumulation order. Gate continuation on oracle margin `0.09375` and at least a 3% model wall-clock win. + +**Detail:** Full feasibility memo archived in `.squad/decisions-archive/2026-08.md` under "Roper megakernel feasibility inbox drop". + ## Extension contract standing directive (#524) **By:** Justin Chu / contract audit diff --git a/.squad/decisions/inbox/cohaagen-qmoe-fc2-fusion.md b/.squad/decisions/inbox/cohaagen-qmoe-fc2-fusion.md new file mode 100644 index 0000000000..75d67324f7 --- /dev/null +++ b/.squad/decisions/inbox/cohaagen-qmoe-fc2-fusion.md @@ -0,0 +1,34 @@ +# Cohaagen: QMoE FC2/down+combine decode fusion experiment + +Date: 2026-08-11 +Branch: `squad/moe-fc2-combine-fusion` +GPU: H200, `CUDA_VISIBLE_DEVICES=1` + +## Design tried + +I implemented the safe remainder of the decode-path fusion for `rows == 1 && routes <= 16`: fuse FC2/down with weighted combine. The experimental `qmoe_down_combine_*` kernel kept the existing #765 FC1/SwiGLU fused activation global scratch, then for each output feature computed each selected expert's FC2 GEMV in slot order, applied the router weight, and wrote the final combined output. This removed the separate `qmoe_combine_*` launch and the `route_output` global scratch write/read. + +I did **not** attempt a single full FC1+SwiGLU+FC2 kernel because the current CTA shape computes one `(route, intermediate)` FC1 reduction per CTA. Keeping all `inter` activations private while preserving FC2 parallelism would require either cross-CTA shared state (not available), recomputing gate/up for each hidden output (prohibitively expensive), or collapsing a whole route/expert into one/few CTAs (massive under-parallelization for hidden=2048). That is a larger persistent/tiled design, not a surgical extension. + +## Correctness + +The experimental down+combine fusion was deterministic and passed validation: + +- New rows==1 fused-path tests: `qmoe_decode_fused_swiglu_down_combine_matches_cpu` and `qmoe_decode_fused_silu_fc3_down_combine_matches_cpu` passed. +- Full QMoE GPU suite with the experiment: 31/31 passed. +- 35B oracle lock passed: teacher-forced token 33803 held; measured `logprob(33803)-logprob(5342) = 0.09375`. + +## Profile result + +`profile_native --pipeline --steady --warmups 1 --runs 3 --tokens 128` with `ONNX_GENAI_PROFILE_OPS=1`: + +| Build | median ms/token | QMoE op | MatMulNBits op | LinearAttention op | RMSNorm op | +| --- | ---: | ---: | ---: | ---: | ---: | +| #765 tip baseline | 11.105 | 67.7-68.0 ms, ~12.1-13.2% | 69.3-79.6 ms, ~13.5-14.1% | 29.3-30.7 ms, ~5.4-5.7% | 24.2-26.8 ms, ~4.7% | +| down+combine experiment | 11.096 | 71.0-71.3 ms, ~13.6-13.8% | 70.1-72.2 ms, ~13.5-13.9% | 29.1-29.3 ms, ~5.6% | 24.9-25.0 ms, ~4.8% | + +Wall-clock delta was only ~0.08% (11.105 -> 11.096 ms/token), far below the 1% ship gate and within shared-host noise. Per-op QMoE time did not improve; it slightly increased in the steady profiled passes, likely because the fused down+combine kernel serializes the route loop per output CTA and adds synchronization inside a heavier kernel, offsetting the saved combine launch/scratch traffic. + +## Decision + +Do **not** ship this FC2/down+combine fusion. Although it is oracle-stable and test-clean, it does not clear the performance gate and adds a new reduction kernel on a numerically sensitive path. The remaining meaningful lever is not route-output scratch removal alone; it needs a bigger tiled/persistent QMoE decode design that can keep FC1 activations local while still exposing enough FC2 output-feature parallelism. diff --git a/.squad/decisions/inbox/cohaagen-qmoe-profile-nextlever.md b/.squad/decisions/inbox/cohaagen-qmoe-profile-nextlever.md new file mode 100644 index 0000000000..5f0b3e87e9 --- /dev/null +++ b/.squad/decisions/inbox/cohaagen-qmoe-profile-nextlever.md @@ -0,0 +1,32 @@ +# Cohaagen finding: QMoE decode next lever after capture fixes + +Date: 2026-08-10 +Worktree: `/home/justinchu/onnx-genai-wt-moeprof` +Commit: `a3856b6a06be8092477a9fa8398f4feb74562ba9` +GPU: H200, `CUDA_VISIBLE_DEVICES=1` + +## Measurement + +Command required `ONNX_GENAI_CUDA_KV_MAX_LEN=262144` because the pipeline decoder directory does not carry the parent `model.max_sequence_length`; without it, native CUDA resolves an unbounded KV cap and fails the mask reservation overflow guard. + +Clean steady-state median: **11.563 ms/token** over 3 runs (128 tokens, warmups=1, decode_skip=8). +Capture segmentation with `ONNX_GENAI_LOG_CAPTURE_SEGMENTS=1`: **14 captured segments, 13 eager seams**. Residual seams are 3 `Squeeze` copy-path seams and 10 `GroupQueryAttention` seams due growing KV/total-sequence symbols. + +## Attribution + +Per-op measured blocks rank: `MatMulNBits` 13.6%, `QMoE` 12.5%, `GroupQueryAttention` 10.0%, `Add` 8.3%, `Cast` 7.4%, `Softplus` 6.7%, `ReduceSumSquare` 6.0%, `LinearAttention` 5.6%, `RMSNormalization` 4.6%, `SkipSimplifiedLayerNormalization` 4.5%. + +Nsight Systems kernel grouping over 256 decode forwards: +- QMoE route/linear/activate/combine: **32.8%**, 2.84 ms/token of GPU kernel time. +- MatMulNBits dense GEMV: **23.5%**, 2.05 ms/token. +- Elementwise/copy/activation: **22.2%**, 1.94 ms/token. +- Norm/RMS components: **11.9%**, 1.02 ms/token. +- LinearAttention: **6.2%**, 0.54 ms/token. + +Top individual kernels: `matmul_nbits_gemv_f16_scales_f16_zp_splitk` 20.9%; `qmoe_linear_f32` 14.3%; `qmoe_linear_f16` 11.0%; `skip_rmsnorm_f16_warp_half4` 6.7%; `linear_attention_f16` 6.2%; `qmoe_route` 5.9%. + +## Conclusion + +The leading norm-fusion hypothesis is not the top lever on current main: norms are real but only about **1.0 ms/token / 11.9%** of GPU kernel time. The dominant model-level lever is QMoE decode, especially the two expert linear kernels plus routing. NCU on `qmoe_linear_f32` shows low DRAM throughput (~3.5% peak) and high barrier/no-eligible stalls (~46%, CTA barrier ~48% of issue cycles), so this is latency/synchronization-bound rather than bandwidth-bound. NCU on the top MatMulNBits GEMV likewise shows low bandwidth (DRAM ~17% peak) and L1TEX scoreboard stalls, consistent with M=1 GEMV latency behavior. + +Recommended next work: design and validate a QMoE decode-specialized expert linear path that reduces block-level barrier/imbalance or groups the k=8 active expert work without reordering fp32 accumulations. This is numerically high-risk and should not be shipped without the 35B oracle lock and an attention-only model gate. diff --git a/.squad/decisions/inbox/harry-728-c1-review.md b/.squad/decisions/inbox/harry-728-c1-review.md new file mode 100644 index 0000000000..a6c07d4f2d --- /dev/null +++ b/.squad/decisions/inbox/harry-728-c1-review.md @@ -0,0 +1,21 @@ +### 2026-08-07: PR #728 C1 capture-classifier review +**By:** Harry +**Verdict:** REJECT + +## Blocking findings + +1. **The growing-symbol gate is not closed under shape inference, so a growing KV axis can be marked capture-safe.** `node_capture_seq_independent` checks only exact `SymbolId` membership in output shapes (`crates/onnx-runtime-session/src/executor/kernel_cache.rs:147-157`). Elementwise broadcast inference can replace two distinct symbolic inputs with the lower-ID representative (`crates/onnx-runtime-shape-inference/src/context.rs:468-482`). Thus an op broadcasting `[seq_kv, D]` with `[batch, D]` (runtime batch `1`) may emit the `batch` symbol even though its output extent follows `seq_kv`; the new gate returns true and admits a stale launch geometry. The negative test only covers direct propagation of the same KV symbol (`executor/tests.rs:845-855`), so it misses this alias case. At minimum, this pointwise-family classifier must reject when any input or output references a growing symbol, with a regression test for symbol unification/aliasing. + +2. **The growing-source collector is incomplete for a supported stateful CUDA op.** It recognizes only GQA, default `Attention`, and `IndexShare` (`kernel_cache.rs:41-59`). `CompressedSparseAttention` derives growing cache-record dimensions from `total_sequence_length`, minting fresh symbols for outputs 1/3/5 (`crates/onnx-runtime-shape-inference/src/handlers/custom_ops.rs:115-129,166-190`), but none enter the set. A pointwise op on such a cache-shaped value is therefore classified independent, especially when no recognized attention node exists and the set is empty (`kernel_cache.rs:94-127,143-156`). This contradicts the PR’s all-model safety claim. + +3. **The oracle re-anchor no longer guards the changed captured autoregressive path.** The “primary lock” teacher-forces one fixed multi-token context on a fresh engine (`crates/onnx-genai-engine/tests/qwen36_35b_a3b_qmoe_divergence.rs:247-282`); that is prefill and does not exercise the C1 decode replay being shipped. The only autoregressive run accepts and logs *any* token at index 119 (`:284-312`). Therefore an actual captured-decode corruption to 5342, 279, or an unrelated token still passes as long as fixed-context prefill remains correct. The dense and fp32 checks are also optional on artifact presence (`:320-371`). Preserve the #722 allowance, but assert a bounded, independently adjudicated set/invariant for the live captured stream rather than making it wholly non-fatal. + +## Confirmed sound parts + +- The EP trait addition is default no-op (`onnx-runtime-ep-api/src/kernel.rs:449-475`) and is set immediately after kernel creation, before insertion/warm execution (`kernel_cache.rs:260-283`), so CPU and unrelated kernels remain unaffected. +- The native engine routes only `token_ids.len() == 1` through the capture state machine and sends multi-token input to eager execution (`onnx-genai-engine/src/native_decode/cuda.rs:430-448,481-553`), so the requested seq_len>1 guard is present. +- `git diff --check` passed. Targeted session classifier tests passed (2/2). + +**Revision owner:** Deckard should revise; Cohaagen is locked out for this rejected revision. + +**REJECT: the classifier can lose or omit growing-symbol dependencies, and the re-anchored test does not fail on captured autoregressive corruption.** diff --git a/.squad/decisions/inbox/harry-728-rereview.md b/.squad/decisions/inbox/harry-728-rereview.md new file mode 100644 index 0000000000..c898b846a4 --- /dev/null +++ b/.squad/decisions/inbox/harry-728-rereview.md @@ -0,0 +1,19 @@ +### 2026-08-07: PR #728 revision re-review +**By:** Harry +**Verdict:** REJECT + +## Blocking findings + +1. **Finding 1 remains open for downstream consumers; input+output checking is not transitive.** Shape inference deliberately replaces two distinct broadcast symbols with the lower-ID representative (`crates/onnx-runtime-shape-inference/src/context.rs:468-482`), and the existing regression proves the named/low-ID symbol wins (`crates/onnx-runtime-shape-inference/tests/op_rules.rs:1721-1738`). Binary inference writes that representative to the output (`handlers/elementwise.rs:20-28`), and graph inference persists it in the IR (`infer.rs:435-452`). A downstream unary/pointwise op then copies that already-aliased shape (`handlers/elementwise.rs:11-16`), so both its input and output contain only the pinned-looking representative. `node_capture_seq_independent` performs exact `SymbolId` membership on those edges (`kernel_cache.rs:234-252`) and therefore wrongly returns true. Runtime execution recomputes the aliasing op's real broadcast extent from concrete inputs (`executor/dynamic_shapes.rs:3-32`, `executor/dispatch.rs:610-650`), while replay launches captured segments without redispatching their nodes (`executor/dispatch.rs:169-182,212-250`), making the stale downstream geometry scenario reachable. The new test covers only the first aliasing op, not its consumer (`executor/tests.rs:889-909`). Require union-find/equivalence-class canonicalization (or equivalent transitive propagation) before growing-set membership, plus a downstream-consumer regression. + +2. **Finding 2 is improved but not fully closed for CSA output 5.** The collector adds CSA outputs 1/3 only (`kernel_cache.rs:88-96`) and extracts only the penultimate axis (`:101-114`). Dynamic ratio-4 CSA independently mints a fresh, growing `selections` symbol on output 5's last axis (`crates/onnx-runtime-shape-inference/src/handlers/custom_ops.rs:181-195`); the shape test confirms this symbolic output (`tests/op_rules.rs:3491-3505`). The generic `past…`/`present…` rank-4 boundary scan does not cover this internal selected-indices value. The added CSA test exercises output 1 only (`executor/tests.rs:943-985`). Collect output 5's growing axis and test a pointwise consumer. + +## Closed / confirmed + +- **Finding 3 is closed:** `C1_CAPTURE_TOKEN` is 46283 (`qwen36_35b_a3b_qmoe_divergence.rs:114-121`) and the fatal assertion permits only 33803 or 46283 (`:326-332`). +- The generic rank-4 symbolic-penultimate KV scan excludes static-penultimate recurrent/GDN state (`kernel_cache.rs:171-190`); the targeted recurrent-state test passes, so GDN pointwise eligibility and the reported 34-segment collapse remain consistent. +- Targeted classifier tests, the symbolic representative test, and `git diff --check` passed. + +**Revision owner:** Leon should produce the next revision. Cohaagen and Deckard are locked out for this artifact revision cycle. + +**REJECT: growing-symbol aliasing still escapes into downstream capturable consumers, and CSA output 5's growing selection axis is still omitted.** diff --git a/.squad/decisions/inbox/harry-728-review3.md b/.squad/decisions/inbox/harry-728-review3.md new file mode 100644 index 0000000000..f5776f75b4 --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review3.md @@ -0,0 +1,18 @@ +### 2026-08-07: PR #728 round-3 re-review +**By:** Harry +**Verdict:** REJECT + +## Blocking finding + +1. **Finding 1 is fixed only for the enumerated elementwise family; the claimed unification closure is not complete.** `op_broadcasts_elementwise` admits only the default-domain elementwise/comparison/variadic/`Where` list (`crates/onnx-runtime-session/src/executor/kernel_cache.rs:320-348`), but other shape handlers invoke the same representative-selecting broadcast logic. In particular, `MatMul` broadcasts its batch dimensions and writes that result into the output (`crates/onnx-runtime-shape-inference/src/handlers/linalg.rs:157-180`), using `broadcast_dim`'s lower-ID symbolic representative (`crates/onnx-runtime-shape-inference/src/context.rs:468-482`). Therefore `[seq_kv, M, K] @ [batch, K, N] -> [batch, M, N]` can erase the raw growing symbol from the output, and a downstream pointwise consumer carrying only `batch` is still classified capture-safe. I reproduced this on `817eee53`: a focused MatMul-alias consumer probe failed because the closed set remained `{SymbolId(1)}` and omitted `batch`. `Einsum` also calls `broadcast` for ellipses (`handlers/einsum.rs:121-125`); `Expand` and `Concat` directly call `broadcast_dim` (`handlers/movement/transform.rs:354-407`, `handlers/movement/concat_slice.rs:62-80`). The closure must cover every shape-inference path that can substitute a growing symbol into another representative, not only elementwise ops. + +## Confirmed closed / sound + +- Within the enumerated elementwise family, the union-find is transitive across arbitrary alias chains (`kernel_cache.rs:360-398`). Symbolic-vs-`1` correctly preserves the symbolic ID, while symbolic-vs-static-non-`1` produces the static extent rather than another symbolic representative (`context.rs:437-442,464-467`), so those cases need no union. +- CSA output 5 is `[query[0], index_heads, query_seq, selections]`; `selections` is its last axis (`handlers/custom_ops.rs:181-195`). `last_axis_outputs = &[5]` and collection at `kernel_cache.rs:99-111,210-218` close Finding 2. +- Both new tests pass on HEAD. The downstream-alias test asserts the consumer, not merely the aliasing op (`executor/tests.rs:958-990`), and I verified it fails on pre-closure commit `f8876baa`. The CSA test consumes output 5 and asserts eager (`executor/tests.rs:1077-1165`). +- Moving the alias case into its dedicated regression does not remove the prior pinned/direct-growing coverage. + +**Revision owner:** Batty should revise; Cohaagen, Deckard, and Leon remain locked out. + +**REJECT: symbol-unification closure still omits non-elementwise broadcast handlers, with a reproduced MatMul downstream-consumer escape.** diff --git a/.squad/decisions/inbox/harry-728-review4.md b/.squad/decisions/inbox/harry-728-review4.md new file mode 100644 index 0000000000..e4bf7510f6 --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review4.md @@ -0,0 +1,24 @@ +### 2026-08-07: PR #728 round-4 re-review +**By:** Harry +**Verdict:** REJECT + +## Blocking finding + +1. **`broadcast_dim` is not the single chokepoint where a growing dependency can disappear from IR symbols.** A valid `Reshape([-1])` over `[seq_kv, 8]` forms the derived expression `seq_kv * 8` (`crates/onnx-runtime-shape-inference/src/handlers/movement/transform.rs:71,120-138`). When shapes are written back, `SymbolInterner::lower` interns that non-bare expression as a fresh `SymbolId` without recording its source symbols (`crates/onnx-runtime-shape-inference/src/context.rs:246-269`). The new record only appends inside the two-bare-symbol broadcast arm (`context.rs:506-518`), so `Graph::symbol_unifications` contains no edge from `seq_kv` to the fresh output symbol. A downstream unary op copies only that fresh shape (`handlers/elementwise.rs:11-16`), while the executor closes only over `symbol_unifications` and performs exact membership (`crates/onnx-runtime-session/src/executor/kernel_cache.rs:334-355,395-412`), admitting the downstream op to capture even though its launch extent grows every decode step. + + I reproduced this on `571ea0d9` with real inference: `past_key` supplied growing `seq_kv`; `Reshape([seq_kv,8], [-1])` produced `[SymbolId(0x80000000)]`; `growing == {seq_kv}` and `symbol_unifications == []`; a downstream `Sigmoid` was classified capture-safe. The focused assertion failed exactly on that classification. `Flatten` has the same unrecorded derived-product path (`transform.rs:154-173`). This is the same stale-geometry/silent-corruption class, not an invalid symbolic-shape scenario. + + The authoritative record must include symbol lineage for lowered derived expressions (or the classifier must conservatively propagate growing dependencies through them), with a downstream-consumer regression for `Reshape([-1])` and preferably `Flatten`. + +## Confirmed sound in this revision + +- Recording placement is otherwise correct: `1`, equal dims, and symbolic-vs-static return before the distinct-bare-symbol arm; inferred representatives are unchanged. The full shape-inference suite passed. +- `record_unification` does not deduplicate, but each inference run overwrites the graph record and duplicates are bounded by broadcast calls, so this is non-blocking. +- The executor union-find is transitive, input-order-independent, and treats pairs bidirectionally. The Add and MatMul downstream-alias regressions pass. +- `Expand` and symbolic non-concat `Concat` funnel through `broadcast_dim`; removing the old elementwise mirror did not lose broadcast-alias coverage. +- Reconciliation keeps the inferred symbolic dim rather than replacing it with a declared symbolic name; Squeeze/Unsqueeze preserve surviving symbol IDs. +- Unioning a growing symbol with a global `batch` symbol can over-poison many nodes and threaten the 34-segment result, but only conservatively. The GPU growing-set/segment re-measure remains required. + +**Revision owner:** Sebastian should own the next revision; Cohaagen, Deckard, Leon, and Batty remain locked out. + +**REJECT: derived symbolic dimensions can erase growing-symbol lineage outside `broadcast_dim`, and a reproduced `Reshape([-1])` downstream consumer is still classified capture-safe.** diff --git a/.squad/decisions/inbox/harry-728-review5.md b/.squad/decisions/inbox/harry-728-review5.md new file mode 100644 index 0000000000..056015fc04 --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review5.md @@ -0,0 +1,24 @@ +### 2026-08-07: PR #728 round-5 re-review +**By:** Harry +**Verdict:** REJECT + +## Blocking finding + +1. **The fail-safe classifier is not a hard capture veto, so both new round-4 regressions remain capturable through the legacy runtime-shape heuristic.** The executor passes `node_capture_seq_independent(...)` only as the boolean `set_capture_seq_independent` hint (`crates/onnx-runtime-session/src/executor/kernel_cache.rs:602-625`). CUDA unary and binary pointwise kernels then compute eligibility as `capture_shape_eligible(self.capture_seq_independent, shape)` (`crates/onnx-runtime-ep-cuda/src/kernels/elementwise.rs:539-551,812-826`). That helper is an **OR**: `seq_independent || numel == 1 || is_fixed_decode_shape(shape)`, and `is_fixed_decode_shape` accepts any shape whose dimensions except the last multiply to one (`elementwise.rs:1038-1059`). + + Consequently, classifier `false` does not mean eager. The exact new `Reshape([-1])` regression produces rank 1, for which the empty leading-dimension product is 1; the `Flatten` regression produces `[batch, seq_kv*8]`, which is accepted when decode batch is 1. Their downstream `Sigmoid` kernels therefore warm a supported capture signature even though `node_capture_seq_independent` returns false. Replay then retains the warmed `n`/grid while `seq_kv*8` grows—the same stale-geometry corruption the provenance fix is intended to close. The tests at `crates/onnx-runtime-session/src/executor/tests.rs:1110-1251` assert only the metadata predicate, not the kernel's actual `capture_support`, so they are false-positive safety tests. + + Make fail-safe disqualification a hard veto distinct from “metadata did not add eligibility” (for example, a tri-state policy or explicit `capture_forbidden` flag checked before the legacy heuristic), and add an EP/capture-planner regression proving disqualified rank-1 and `[1,N]` growing shapes remain outside capture. + +## Confirmed + +- The provenance graph itself is coherent: broadcast edges are undirected, derivations flow source→derived, and the combined worklist closure is transitive/order-independent (`kernel_cache.rs:412-442`). +- `inference_symbol_floor` is seeded above all graph symbols and current-pass mints are at/above it (`shape-inference/src/infer.rs:69-75,1527-1542`); repeated inference preserved the Reshape classification in a focused probe. +- Derived-only-from-pinned roots are not reverse-poisoned. Over-eager classification changes performance, not numerical semantics. +- Both new regressions independently failed on `571ea0d9` with only `{SymbolId(1)}` and pass on `0fd87df3`. +- `lower` returns the prior dimensions; added recording is bookkeeping-only. Full shape-inference tests passed (16 + 41 + 275 + doctests), targeted fail-safe/derived tests passed, and `git diff --check` passed. +- The switch-gated denylist residual is non-blocking for this verdict, but consider removing or more strongly gating it. + +**Revision owner:** Roy should revise; Cohaagen, Deckard, Leon, Batty, and Sebastian remain locked out. + +**REJECT: the shipped fail-safe set is computed correctly but is only an additive hint; the CUDA kernels can still capture disqualified growing rank-1 / `[1,N]` pointwise shapes via the legacy heuristic.** diff --git a/.squad/decisions/inbox/harry-728-review6.md b/.squad/decisions/inbox/harry-728-review6.md new file mode 100644 index 0000000000..d5d8bd0bb0 --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review6.md @@ -0,0 +1,23 @@ +### 2026-08-07: PR #728 round-6 final re-review +**By:** Harry +**Verdict:** REJECT + +## Blocking finding + +1. **The classifier is still bypassed by three shape-dependent CUDA pointwise kernels.** The revised `capture_shape_eligible` is now a hard veto for every caller, and the signature-gated unary/binary/predicate/bitwise/PRelu/SiluMul paths correctly return `Unsupported` when `capture_seq_independent == false`. However, `UnaryMathKernel` (`Abs`, `Neg`, `Exp`, `Sign`, etc.) and `NotKernel` still return `CaptureSupport::Supported` unconditionally and do not store or override `set_capture_seq_independent` (`crates/onnx-runtime-ep-cuda/src/kernels/pointwise.rs:304-387,418-490`). `BitwiseNotKernel` has the same bypass (`crates/onnx-runtime-ep-cuda/src/kernels/bitwise.rs:470-564`). A classifier-disqualified growing rank-1 value consumed by any of these kernels can therefore still enter capture with a launch grid/count baked from the warmed extent. + + This directly fails the requested “ANY remaining path” check. Wire these kernels through the same authoritative flag/signature gate (or enforce the veto centrally before kernel-specific `capture_support`). + +2. **The claimed capture-level regression does not exercise `UnaryKernel::run` or `capture_support()`.** `disqualified_growing_reshape_consumer_yields_no_capture_signature` manually calls `capture_shape_eligible` and constructs a `Vec` (`elementwise.rs:1231-1248`); it is a helper test, not an EP/capture-admission test, and cannot catch the bypasses above. Its fail-pre claim is genuine—the old helper admitted `[320]`—but an actual kernel-path regression remains necessary. + +## Confirmed + +- `capture_shape_eligible(false, shape)` is now always false; all six existing callers are centrally fixed. +- Their struct defaults are fail-safe, and `KernelCache::get_or_create` sets the classifier flag before insertion/return (`kernel_cache.rs:602-631`). +- `is_fixed_decode_shape` has no remaining code references. +- `cargo test -p onnx-runtime-ep-cuda --lib --quiet`: 280 passed, 2 ignored. +- `git diff --check 0fd87df3..f6b6203c`: clean. + +**Revision owner:** Sapper. Roy, Cohaagen, Deckard, Leon, Batty, and Sebastian remain locked out. + +**REJECT: classifier-disqualified nodes can still reach capture through unconditional `Supported` paths in `UnaryMathKernel`, `NotKernel`, and `BitwiseNotKernel`.** diff --git a/.squad/decisions/inbox/harry-728-review7.md b/.squad/decisions/inbox/harry-728-review7.md new file mode 100644 index 0000000000..282c8f037e --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review7.md @@ -0,0 +1,28 @@ +### 2026-08-07: PR #728 round-7 review +**By:** Harry +**Verdict: REJECT** + +## Blocking finding + +1. **The “first check” veto corrupts the public seam classification for host-driven nodes.** `Executor::node_capture_reason` returns `ClassifierDisqualified` before invoking the EP structural policy (`crates/onnx-runtime-session/src/executor/capture.rs:556-615`). Therefore an `If`/`Loop`/`Scan` or Sequence node whose input/output references any disqualifying symbol is reported as `ClassifierDisqualified`, whose `path_kind()` is `EagerDeviceSeam` (`capture.rs:100-110`), even though `ExecutionProvider::plan_capture_region` defines that node as `HostControlFlowOrSequence` (`crates/onnx-runtime-ep-api/src/provider.rs:455-477`) and it actually executes through the host path. + + Concrete regression: a LongRoPE-style `If` with a fail-safe-disqualified symbolic output now appears in the public capture-segmentation report as an eager device seam rather than a host seam. This does not numerically corrupt tensors, but it corrupts the structured capture-path contract and any diagnostics/measurement that distinguish device seams from host round trips. Preserve host-control-flow/sequence precedence, then apply the classifier veto before any kernel admission check. + +**Fresh revision owner:** Gaff. Roy, Cohaagen, Deckard, Leon, Batty, Sebastian, and Sapper remain locked out for this artifact. + +## Confirmed sound + +- `plan_capture_segments` is the sole schedule builder and maps every plan node through `node_capture_reason`; the only stream-capture call is reached through that schedule. +- The veto uses the build-time `capture_growing_symbols` produced by `compute_capture_disqualifying_symbols`, identical to kernel-cache/dispatch classification. +- `node_capture_seq_independent` checks every input and every output. +- Warm-seeded capture planning still routes through the veto; clearing the set in the isolated Range seeding test matches that fixture's denylist-mode state. +- The negative mock faithfully models unconditional-`Supported` kernels and would be admitted on `f6b6203c`; the positive companion remains admitted. +- `ClassifierDisqualified` correctly prevents unconditional kernel support from re-admitting a disqualified device kernel. + +## Validation + +- `cargo test -p onnx-runtime-session --lib` — **133 passed, 0 failed**. +- `cargo test -p onnx-runtime-ep-cuda --lib --features cuda` — **285 passed, 0 failed, 2 ignored**. +- `git diff --check f6b6203c..9555b354` — clean. + +**REJECT: the central veto closes the capture-admission hole, but its precedence newly misreports host control-flow/sequence seams as eager device seams.** diff --git a/.squad/decisions/inbox/harry-728-review8.md b/.squad/decisions/inbox/harry-728-review8.md new file mode 100644 index 0000000000..20ff430e2d --- /dev/null +++ b/.squad/decisions/inbox/harry-728-review8.md @@ -0,0 +1,50 @@ +### 2026-08-07: PR #728 round-8 review +**By:** Harry + +**Verdict: APPROVE** + +## Confirmed sound + +- `Executor::node_capture_reason` now has the required precedence: quarantine + (`capture.rs:566-577`), EP structural policy (`capture.rs:578-593`), resolved-shape + contract assertion (`capture.rs:595-598`), classifier hard veto + (`capture.rs:614-623`), kernel warmth (`capture.rs:640-647`), then kernel admission + (`capture.rs:648`). Thus a disqualified `If`/`Loop`/`Scan`/Sequence node reports + `HostControlFlowOrSequence` / `HostSeam`, while a resolved plain device node still + reports `ClassifierDisqualified`. +- The classifier veto remains before `kernel_capture_decline`: the veto returns at + `capture.rs:614-623`; kernel admission is not consulted until `capture.rs:648`. + An unconditional-`Supported` kernel therefore cannot re-admit a disqualified node. +- Moving the veto after the assertion does not create a panic path in-tree. + `ExecutionProvider::plan_capture_region` must decline unresolved shapes and its + default implementation returns `UnresolvedOutputShape`/`UnresolvedInputShape` + before admission (`provider.rs:455-477`). The only in-tree override delegates to + that policy. Any node reaching the assertion and veto has resolved inputs and + outputs. +- `disqualified_control_flow_node_reports_host_seam_not_device_seam` calls the real + `node_capture_reason` path. It makes a default-domain node an `If`, proves the + structural helper recognizes it, adds a growing symbolic output and proves the + classifier disqualifies it, installs a warmed unconditional-`Supported` kernel, + and asserts `HostControlFlowOrSequence` plus `HostSeam`. On `9555b354`, the + veto-first implementation returns `ClassifierDisqualified`, so the test fails; + on `f310ed9f`, structural precedence makes it pass. +- Sapper's round-7 negative and positive device-node tests are unchanged: the + disqualified `Identity` remains vetoed despite `Supported`, while the qualified + `Identity` remains admitted. +- No bypass admission path remains. Every whole-graph or segmented capture schedule + maps every plan node through `node_capture_reason` (`capture.rs:660-680`); + warm-seeded planning uses that same schedule; capture recording consumes only its + captured segments (`dispatch.rs:183-246`). Control-flow child executors call + `run_scoped`, which explicitly uses `RunMode::Eager` (`run.rs:26-35`), so nested + bodies do not bypass this audit. + +## Validation + +- `cargo test -p onnx-runtime-session --lib` — **148 passed, 0 failed**. +- `cargo test -p onnx-runtime-ep-cuda --lib --features cuda` — + **289 passed, 0 failed, 2 ignored**. +- `git diff --check 9555b354..f310ed9f` — clean. +- GPU/oracle integration tests were not run, as requested. + +**APPROVE: the structural seam precedence is restored without reopening the +classifier admission hole or introducing another capture path.** diff --git a/.squad/decisions/inbox/harry-glm-deepseek-review.md b/.squad/decisions/inbox/harry-glm-deepseek-review.md new file mode 100644 index 0000000000..21b363c1b0 --- /dev/null +++ b/.squad/decisions/inbox/harry-glm-deepseek-review.md @@ -0,0 +1,82 @@ +# Review: native CUDA load unblock for GLM-4-9B & DeepSeek-V2-Lite + +Date: 2026-08-11 +Reviewer: Harry (code reviewer) +PRs: #770 (honor native CUDA KV capacity), #771 (accept Cast-backed QMoE scales) +Context: both OPEN, auto-merge on; reviews are not a merge gate. Genuine correctness audit. + +## PR #770 — native KV reservation uses runtime CUDA hard cap first + +Verdict: **APPROVE** (correct for the target models) with one recommended hardening follow-up. + +### What was verified +1. **Precedence is sound for the enforced ceiling.** `hard_max_len == capacity.max_len` + (`cuda.rs:2330`). The decode path rejects any request beyond that ceiling at every + entry point (`total_len > state.capacity.max_len` at cuda.rs:469, 604, 669, 815, 843) + and KV growth refuses past it (`ensure_capacity`: `if required > self.capacity.max_len` + at cuda.rs:1056). So reserving at `hard_max_len` cannot admit a sequence the device + later fails to carry — the original "discover mid-generation" hazard is avoided. The + env/programmatic caps are themselves clamped by metadata in `resolve_cuda_kv_capacity` + (cuda.rs:2771+), so `hard_max_len <= metadata max_sequence_length` always. +2. **Fallback path is correct.** When `cuda_kv_debug_stats()` is None (CPU/non-CUDA), + `native_kv_reservation_max_context` falls back to metadata `max_sequence_length` + (source "model.max_sequence_length"). Covered by `native_kv_reservation_falls_back_to_metadata`. +3. **No regression for normal declared-context models.** With a declared + `max_sequence_length` and no env/programmatic cap, `resolve_cuda_kv_capacity` sets + `hard_max_len == metadata` value, so the new path reserves exactly what the old path + reserved. Identical behavior. With an env cap it reserves the smaller enforced value — + strictly safer. + +### Finding (LOW severity, narrow, recommended fix — non-blocking for GLM/DeepSeek) +`resolve_cuda_kv_capacity` returns `max_len = usize::MAX`, source +`"unbounded (model.max_sequence_length unavailable)"` when a CUDA model declares **no** +`max_sequence_length` and no env/programmatic cap is set (cuda.rs ~2807-2814). In that case +`cuda_kv_debug_stats()` is `Some`, so the new `native_kv_reservation_max_context` returns +`Some((usize::MAX, "unbounded..."))` and load.rs calls `kv_reservation(usize::MAX)`, which +saturates to `u64::MAX`. Behavior then splits by allocator: +- VMM CUDA path (`commits_on_demand() == true`): only warns, still loads. +- **cuMemAlloc CUDA path (`commits_on_demand() == false`, provider.rs:965 / vmm_allocator.rs:945):** + `governor.reserve_on(NativeKvCache, Device, u64::MAX)` fails → **load fails**. + +Previously this case hit the `None` arm (load.rs:529) which deliberately warns-and-loads +("the model runs fine"). So for an undeclared-context model on the non-VMM CUDA path this +is a genuine load regression. + +Recommended guard: treat an unbounded/`usize::MAX` `hard_max_len` as "no cap" and fall +through to metadata (i.e. `native_capacity.filter(|(len, _)| *len != usize::MAX)`), which +reproduces the old warn-and-load behavior. Neither GLM-4-9B nor DeepSeek-V2-Lite hits this +(both declare context and GLM uses `ONNX_GENAI_CUDA_KV_MAX_LEN=4096`), so it does not block +Justin's priority — but it should be fixed before relying on this path for arbitrary models. + +## PR #771 — QMoE placement accepts one-hop Cast(initializer) scales + +Verdict: **APPROVE**. All three scrutiny points hold. + +1. **No dtype/region corruption.** `plan_static_weight_placement` is explicitly advisory + (placement.rs:22-27 doc): it produces a plan/report (device/host byte counts, explanation) + and does **not** feed bytes into the QMoE kernel. The runtime kernel still consumes its + graph inputs — the fp32 Cast output — unchanged. Region classification reads the backing + fp16 initializer only for dimensions/regions; `blocks_per_row = scale_dims[2]` is a count + and shape is Cast-invariant (Cast preserves shape), so the layout math is correct. The + fp16 backing bytes are never handed to the kernel as fp32. Cohaagen's claim is verified. + Forward-looking note (not a bug today): if/when static placement becomes *enforcing*, the + classified region points at the fp16 backing tensor while the kernel reads the fp32 Cast + output — a divergence to track — but the scale tensors are tiny (per-block) so accounting + impact is negligible. +2. **one-hop / domain guard is right.** `initializer_or_initializer_cast` accepts only a + default-domain (`node.domain.is_empty()`) `Cast` whose single input is a direct + initializer. A multi-hop `Cast(Cast(init))` is correctly rejected (the inner value is not + in `graph.initializers`). A `com.microsoft` Cast is rejected — correct, since Cast is a + standard-domain op. Not checking the Cast `to` dtype is harmless here because placement + uses the backing initializer's real dims and the runtime consumes whatever the Cast emits. +3. **DRY.** No model-name gates; generic runtime/placement behavior. Confirmed. + +### Validation plausibility +GLM top-40 identical, max logprob delta ~0.00794; DeepSeek top-40 identical (low-rank +reorder only), max delta ~0.02848; both golden-greedy locks pass. Deltas are consistent with +fp accumulation-order differences on the CUDA path, not a numeric defect. Plausible. + +## Bottom line +- #771: APPROVE, no changes needed. +- #770: APPROVE for the priority models; recommend guarding the `usize::MAX` unbounded case + so undeclared-context models on the cuMemAlloc CUDA path don't regress from load to failure. diff --git a/.squad/decisions/inbox/quaid-megakernel-phase0.md b/.squad/decisions/inbox/quaid-megakernel-phase0.md new file mode 100644 index 0000000000..0d8c98ee18 --- /dev/null +++ b/.squad/decisions/inbox/quaid-megakernel-phase0.md @@ -0,0 +1,161 @@ +# Megakernel Phase 0 — Persistent Single-Op QMoE Decode Kernel + +**Author:** Quaid (CUDA kernel engineer) +**Date:** 2026-08-11 +**Branch:** `squad/megakernel-phase0-qmoe` (off origin/main, #766 merged) +**GPU:** CUDA_VISIBLE_DEVICES=4 (H100/Hopper, 132 SM) +**Verdict:** ⛔ **NO-SHIP** — oracle holds byte-exact, but decode **regresses +7.4%**. Occupancy is the trap, exactly as predicted. +**Phase 1 recommendation:** **Do NOT greenlight Phase 1 as specified.** See bottom. + +--- + +## 1. What I built + +Replaced the QMoE **rows==1 decode** chain's two heaviest launches with ONE persistent +fused kernel, keeping the FC1/SwiGLU intermediate (`activated`) in the kernel and +eliminating its global scratch round-trip, using a **counter-based (global-atomic) +producer/consumer** for the FC1→FC2 chunk dependency (the Hazy "4-chunk MLP" trick) — +**NOT** a `grid.sync()` cooperative barrier (per Roper §4: grid barrier risks co-residency +deadlock and fights occupancy). + +Important correction to the feasibility spec: the "5-launch / 4-scratch" description in +`roper-megakernel-feasibility.md` §3 is **stale**. Current origin/main (#766) already fuses +FC1+FC3+activate into `qmoe_gate_up_activate`. The **real** decode chain is 4 launches / +2 scratch: + +``` +route -> gate_up_activate (FC1+FC3+SwiGLU) -> FC2 (qmoe_linear) -> combine + |-- activated (16 KB) --| |-- route_output (64 KB) --| +``` + +Phase 0 fused `gate_up_activate` + `FC2` into one persistent kernel producing +`route_output`; `combine` stays separate (it sums across routes → cross-block reduction, +must stay separate to keep byte-exact accumulation order). + +Design: +- Fixed grid of co-resident blocks (occupancy-capped for deadlock-freedom), cached in an + `AtomicU32` so warmup == capture geometry. +- **Phase A producers** compute `activated` with accumulation *identical* to + `gate_up_activate` (same `qmoe_int4_chunk`/`block_sum`/`swiglu_value`, 256-thread stride), + then `atomicAdd` a per-route done-counter. +- **Phase B consumers** spin on the per-route counter (volatile read + `__threadfence`), + then compute FC2 *identical* to `qmoe_linear`. +- Tiny `qmoe_fused_reset` kernel zeroes counters each launch (capturable). +- Env toggle `ONNX_GENAI_QMOE_PERSISTENT` (default on) → clean A/B on one binary. +- DRY: driven off tensor shapes/attrs only, **no model-name gates**; multi-row + (prefill/grouped) path untouched, falls back to per-op for rows>1. + +Files: `crates/onnx-runtime-ep-cuda/src/kernels/qmoe.rs` (kernel + launch + toggle), +`crates/onnx-runtime-ep-cuda/tests/qmoe_gpu.rs` (3 new A/B + CPU tests). + +--- + +## 2. Gate 1 — Oracle byte-exact: ✅ HOLDS + +35B-A3B teacher-forced fp32 oracle, persistent **ON**: + +``` +QUAID-PHASE0 QMoE teacher-forced margin = 0.09375 (argmax=33803) +``` + +- **argmax = 33803**, **margin logprob(33803)−logprob(5342) = 0.09375 EXACTLY.** ✔ +- Behavior is **byte-identical** between persistent ON and OFF: same #722 note + (autoregressive token@119 = 46283 = C1_CAPTURE_TOKEN, the documented benign fp16 + coin-flip — informational, not a fail), same everything. +- Unit tests: **32 pass** (29 existing + 3 new). New tests assert on `to_bits()` → + the persistent path is **bit-identical** to the per-op path AND matches CPU. + +**Environmental note (NOT a Phase-0 failure):** the oracle test's *step-3 DENSE int4 +cross-check* aborts with a pre-existing **cuDNN** error +(`CUDNN_STATUS_SUBLIBRARY_VERSION_MISMATCH` at node 80 `mlp/gate/Softmax`, a Float16 +Softmax op on a *different* model). This reproduces **identically with persistent OFF** and +has nothing to do with the QMoE kernel (my kernel touches neither cuDNN nor Softmax). The +QMoE **primary lock (step 1)** and **#722 tripwire (step 2)** both pass byte-identically +with the fused kernel. The correctness gate — the teacher-forced QMoE margin — is GREEN. + +This matches my int4/DP4A prior: reduction restructures that preserve per-thread fp32 +accumulation order hold the oracle bit-exact. **Numerics are low-risk.** + +--- + +## 3. Gate 2 — Wall-clock ≥3% decode win: ❌ FAILS (regresses) + +Steady-state decode, `profile_native --pipeline --steady --warmups 1 --runs 3 --tokens 128`, +`ONNX_GENAI_CUDA_GRAPH=1 ONNX_GENAI_CUDA_KV_MAX_LEN=262144`, GPU 4, median of 3: + +| Config | decode ms/tok | throughput | Δ vs baseline | +|---|---|---|---| +| **Baseline** (persistent OFF, per-op) | **11.109** | 90.01 tok/s | — | +| **Phase 0** (persistent ON, fused) | **11.932** | 83.81 tok/s | **+7.4% SLOWER** | + +Runs were tight (±0.02 ms across 3), so this is a real **regression**, not noise. The +≥3%-win gate required ≤10.77 ms/tok; we got 11.93. Decisive miss. + +--- + +## 4. Why — occupancy collapse (the mechanism) + +Launch geometry, decode (routes=8, inter=512, hidden=2048): + +| Kernel | resident blocks | note | +|---|---|---| +| per-op `gate_up_activate` | **4096** | one block per (route, inter-feature), uncapped (routes≤16) | +| per-op `FC2` (`qmoe_linear`) | **2112** | capped at min(tasks, 132 SM × 16) | +| **Phase 0 persistent fused** | **660** | `occupancy_max_active_blocks=5/SM × 132 SM`, **register-limited** | + +``` +QUAID-PHASE0 persistent_grid: per_sm=5 sm_count=132 grid=660 block_dim=256 smem=1024 +``` + +Fusing FC1-decode + FC2 into one kernel makes it register-heavy → only **5 blocks/SM** +(1280 of 2048 max threads/SM ≈ 62% occupancy). The bandwidth-bound **FC2 GEMV** therefore +runs with **~3.2× fewer** concurrent blocks (2112 → 660) driving HBM. Decode QMoE is +**HBM-bandwidth-bound** (~503 MB/token weight reads); fewer concurrent blocks → lower +memory-level parallelism → lower achieved bandwidth → the op slows. + +The thing Phase 0 *saved* — the `activated` scratch round-trip (~16 KB/token, plus the two +launch/drain seams) — is **negligible** vs 503 MB/token of weight traffic. This is fully +consistent with Cohaagen's FC2+combine fusion measuring only **+0.08%** from removing a +seam: at decode the QMoE cost is weight bandwidth, not launch overhead or scratch DRAM. + +**One-line:** the fusion that removes the seams also destroys the occupancy the +bandwidth-bound GEMVs need to saturate HBM, and the seams were never the bottleneck. + +--- + +## 5. Verdict & Phase 1 recommendation + +**NO-SHIP.** Oracle byte-exact (0.09375, 33803) — the persistent counter-synced structure +is numerically safe and bit-identical to per-op. But it **regresses decode 7.4%** because +decode QMoE is occupancy/bandwidth-bound, not seam-bound. Reverting the branch code; keeping +this memo. + +**Phase 1 (the real 1.3–1.6× prize) — recommend DO NOT greenlight as specified.** Phase 0 +is the cheapest instance of the whole-step-megakernel thesis, and it inverts: naive fusion +trades occupancy for seam-removal on the heaviest decode op, and loses. Phases 1–2 (fusing +attention + all MoE across the whole decoder step into one persistent kernel) face the SAME +occupancy-vs-fusion tension at **larger** scale — more weights + more state in one kernel → +even higher register/smem pressure → even lower occupancy — while the payoff (seam removal) +stays small because decode is bandwidth-bound. A megakernel could only win here if it holds +**≥ current per-op occupancy**, which requires a real **tile scheduler** (MPK/Hazy-style: +many small tiles dynamically load-balanced across a persistent grid, NOT one-block-per-task +fusion). That is a much larger, still-speculative effort. My recommendation: shelve the +whole-step megakernel; the higher-EV decode lever remains **weight bandwidth** (int4 layout / +DP4A / vectorized weight loads), not launch/seam fusion. + +--- + +## Reproduce + +```bash +source .cudaenv.sh # CUDA_VISIBLE_DEVICES=4, LD_LIBRARY_PATH=/home/tlwu/cudnn9.19_cuda13/lib +# A/B wall-clock: +ONNX_GENAI_CUDA_GRAPH=1 ONNX_GENAI_CUDA_KV_MAX_LEN=262144 ONNX_GENAI_QMOE_PERSISTENT={0,1} \ + ./target/release/profile_native --model /home/justinchu/qwen36-35b-a3b-qmoe-artifacts \ + --pipeline --ep cuda --steady --warmups 1 --runs 3 --tokens 128 +# Oracle (primary lock = teacher-forced QMoE margin; dense step-3 hits unrelated cuDNN env error): +ONNX_GENAI_CUDA_GRAPH=1 ONNX_GENAI_CUDA_KV_MAX_LEN=4096 ONNX_GENAI_QMOE_PERSISTENT=1 \ + cargo test -q -p onnx-genai-engine --features "cuda native-backend" \ + --test qwen36_35b_a3b_qmoe_divergence \ + qwen36_35b_a3b_qmoe_native_cuda_matches_fp32_oracle -- --ignored --nocapture +``` diff --git a/.squad/decisions/inbox/quaid-qmoe-int4-dp4a.md b/.squad/decisions/inbox/quaid-qmoe-int4-dp4a.md new file mode 100644 index 0000000000..02075dad17 --- /dev/null +++ b/.squad/decisions/inbox/quaid-qmoe-int4-dp4a.md @@ -0,0 +1,114 @@ +# Quaid memo: QMoE int4 DP4A / vectorized weight-read decode experiment + +Date: 2026-08-11 +Branch: `squad/moe-int4-dp4a-experiment` (fresh off `origin/main` @ 4dd5119b, i.e. #765 tip) +GPU: H200, `CUDA_VISIBLE_DEVICES=4` +Authorized by: Justin (@justinchuby) as a risk/reward branch experiment. NO-SHIP is an acceptable outcome; the goal is DATA. + +## TL;DR — NO-SHIP (perf regression, oracle test fails on the #722 tripwire) + +The highest-ceiling remaining QMoE decode lever named in Cohaagen's latency-floor memo +(*"int4 DP4A / vectorized weight-read"*) was prototyped as a **128-bit (`uint4`) vectorized +int4 weight-stream** for the `rows==1 && routes<=16` decode GEMV (FC1 fused gate/up + +SwiGLU, and FC2/down). It **regressed** the model by **+4.90% ms/token** and grew the QMoE +op by **+25%**. The teacher-forced correctness lock was numerically **unchanged** (token +33803 holds, margin **0.09375** — bit-identical to baseline), but the oracle test still +**FAILS** because the *autoregressive* #722 tripwire drifted to token 13 (outside the narrow +benign set {33803, 46283}). Both ship gates (oracle PASS + >3% ms/tok win) are missed, so +the code is reverted and only this memo is kept. + +## What was built + +A vectorized int4 GEMV chunk (`qmoe_int4_vec32`): one aligned **128-bit `uint4` load streams +32 packed int4 weights** (four 8-value sub-chunks) per memory transaction, reusing each quant +block's scale/zero-point across the sub-chunks it covers. Accumulation is **byte-order-identical +within a thread** to four sequential scalar `qmoe_int4_chunk` calls, so the only numerical +change is the block-reduction fan-in (fewer, wider per-thread partials). + +Wired behind new NVRTC entries (`qmoe_linear_vec_*`, `qmoe_gate_up_activate_vec_*`, templated on +a `bool Vectorized`), guarded to 4-bit layouts with `in_features % 32 == 0`, and applied to the +decode path only: fused FC1 gate/up (in=hidden=2048) and FC2/down (in=inter=512). Launch width +drops to one thread per 32-weight superchunk, rounded to whole warps and clamped `[32, 256]` +(FC1 → 64 threads, FC2 → 32 threads) so the warp-shuffle block reduction stays complete. + +Sub-experiment coverage: +- **(2) memory-coalesce (vectorized wide weight read):** BUILT + MEASURED — this memo's result. +- **(1) DP4A compute-quant (int8 activation × int8 weight, `__dp4a`):** NOT built. Analytically + dominated on this kernel (see "Why DP4A was not pursued"). +- **int4 tensor-core / mma:** NOT built (same reasoning as DP4A; multiply is not the bottleneck). + +## Risk / reward table + +| Metric | Baseline (#765 tip) | After (vec int4) | Delta | +| --- | ---: | ---: | ---: | +| Decode ms/token (steady, 3 runs) | **11.157** | **11.704** | **+0.547 ms / +4.90% (REGRESSION)** | +| QMoE op, per-token window | 67.5–67.9 ms, ~13.15% | 84.5 ms, ~15.8% | **+25% QMoE op time** | +| MatMulNBits op (control) | ~70.0 ms, ~13.1% | ~70.0 ms, ~13.1% | unchanged | +| Oracle: teacher-forced argmax @119 | **33803** | **33803** | HOLDS | +| Oracle: primary margin logprob(33803)−logprob(5342) | **0.09375** | **0.09375** | **UNCHANGED (bit-identical)** | +| Oracle: autoregressive token @119 (#722 tripwire) | 33803 / 46283 (benign coin-flip) | **13** | **out of benign set → test FAILS** | +| QMoE GPU unit suite | 31/31 | 31/31 (incl. 2 new vec parity tests) | green | + +Profile: `profile_native --pipeline --steady --warmups 1 --runs 3 --tokens 128`, +`ONNX_GENAI_PROFILE_OPS=1 ONNX_GENAI_CUDA_KV_MAX_LEN=262144`, model +`/home/justinchu/qwen36-35b-a3b-qmoe-artifacts`. +Oracle: `qwen36_35b_a3b_qmoe_native_cuda_matches_fp32_oracle` (`ONNX_GENAI_CUDA_GRAPH=1 +ONNX_GENAI_CUDA_KV_MAX_LEN=4096`). The change-invariant dense-int4 (step 3) and full-fp32-CPU +(step 4) cross-checks were skipped in the after-run to save ~40 min of CPU oracle time — they +do not touch the CUDA QMoE kernel and passed on baseline; steps 1–2 (the QMoE CUDA path) ran. + +## Why it regressed (root cause) + +Cohaagen's NCU data already showed `qmoe_linear_*` is **latency / occupancy / barrier bound, +not bandwidth bound** (DRAM ~4% of peak, barrier ~37%, no-eligible ~22%). At `in_features=2048` +the scalar path already issues a coalesced 32-bit (`uint32` = 8 int4) load per thread with 256 +threads/block — the weight read is *already* coalesced. Going to a 128-bit `uint4` load forces +**fewer threads** (64 for FC1, 32 for FC2) to cover the same K, which **lowers memory-level +parallelism (fewer outstanding requests) and occupancy** — exactly the resource this +latency-bound kernel depends on to hide load latency. The wider transactions do not compensate. +Net: QMoE op +25%. This mirrors the earlier ILP-2 unroll regression (#764 follow-up): any +transform that trades away warps/MLP for per-thread work loses on this M=1 GEMV. + +## Why DP4A was not pursued + +DP4A (int8×int8→int32 `__dp4a`) accelerates the **multiply**, but the multiply is not the +bottleneck: the kernel is barrier/occupancy/latency bound with DRAM at ~4% and low useful math +intensity. DP4A would (a) not reduce the dominant int4 **weight-read** stream (~503 MB/token, +still read as int4 then unpacked), (b) add per-row int8 activation-quant overhead and dynamic +scale bookkeeping, and (c) introduce accumulation-order / dequant-scale numerical risk on a +teacher-forced-locked path. On a kernel that is not multiply-bound, it is analytically dominated +by the memory-coalesce attempt that already regressed. Building it would consume another ~1 h +oracle/profile cycle for a predetermined-negative result. + +## Numerical-risk finding (the useful positive datum) + +The **capture-independent primary lock is completely insensitive** to this class of change: the +vectorized reduction reordering left `logprob(33803)−logprob(5342) = 0.09375` **bit-identical** +to baseline, and the new `rows==1` GPU parity tests conform to the CPU reference within the +2e-5 abs / 1e-4 rel band (block_size=16 with a 32-weight superchunk straddling two quant blocks +is exercised). So the *teacher-forced* numerical risk of int4 GEMV reduction restructuring is +effectively **zero** here. The oracle test failure is confined to the **autoregressive** #722 +tripwire (token 13 ∉ {33803, 46283}) — a manifestation of the documented fp16 near-tie +coin-flip whose winner shifts under any tiny decode perturbation; the allowed-set is simply too +narrow to include this perturbation's outcome. It is not evidence of a corrupt kernel, but it +does mean the oracle **does not PASS** as written, so the ship gate is not met regardless of perf. + +## Recommendation + +**This lever is NOT worth a bigger investment in its current form.** Three independent surgical +attempts on the existing per-output-element block-reduction CTA shape have now all failed to +clear the >3% gate: FC2/down+combine fusion (~0.08%, Cohaagen), ILP-2 K-unroll (regressed, +Cohaagen), and this vectorized 128-bit weight read (**+4.9% regression, Quaid**). The consistent +signal is that the QMoE decode GEMV is **occupancy/MLP/barrier bound**, and any transform that +reduces active warps to buy per-thread width or ILP loses. + +The remaining real headroom (the kernel is ~10× above the ~0.1–0.24 ms/token int4 bandwidth +floor) requires a **fundamentally different decode kernel structure**, not a chunk-level tweak: +a persistent / tiled QMoE decode worker that keeps many warps resident and busy across the +k=8 experts and the FC1→SwiGLU→FC2 chain (amortizing launch/graph-node overhead and the +per-output block reduction), while preserving fp32 accumulation order enough to hold the 33803 +lock. That is a multi-day design with high integration risk against the ORT-style per-op +executor + graph-capture model — schedule it deliberately, not as a surgical follow-up. + +DP4A / int4 tensor-core specifically should be **shelved** until the kernel is first made +compute-bound by such a restructure; on today's memory/latency-bound shape it cannot help. From dc8a07815934e90c4d93b2d380f7f69dbbb7508e Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Tue, 11 Aug 2026 14:54:14 +0000 Subject: [PATCH 2/2] fix(cpu-ep): k-major output layout for non-final-axis TopK CPU-EP TopK wrote outputs in push order [outer][inner][k]; ONNX requires k-major [outer][k][inner] (input shape with shape[axis]=k). The layouts match only when inner==1 (final axis), so non-final-axis TopK with k>1 produced wrong values+indices ordering vs CUDA EP / ONNX Runtime. Write each winner to its strided destination (outer*k + rank)*inner + i instead of pushing sequentially. Adds non-final-axis and first-axis regression tests; final-axis path unchanged. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../coordinator-topk-nonfinal-axis-fix.md | 33 ++++++++ .../src/kernels/selection.rs | 75 ++++++++++++++++--- 2 files changed, 96 insertions(+), 12 deletions(-) create mode 100644 .squad/decisions/inbox/coordinator-topk-nonfinal-axis-fix.md diff --git a/.squad/decisions/inbox/coordinator-topk-nonfinal-axis-fix.md b/.squad/decisions/inbox/coordinator-topk-nonfinal-axis-fix.md new file mode 100644 index 0000000000..c8dd5b363d --- /dev/null +++ b/.squad/decisions/inbox/coordinator-topk-nonfinal-axis-fix.md @@ -0,0 +1,33 @@ +# CPU-EP TopK: k-major output layout for non-final axes + +Date: 2026-08-11 +Owner: Squad (coordinator; sub-agent spawn hit a canary false-positive and returned 0 useful turns, so fixed inline per fallback rule) + +## Bug +`TopKKernel::execute` (crates/onnx-runtime-ep-cpu/src/kernels/selection.rs) built the +values/indices outputs by sequential `push` in loop order `outer -> inner(i) -> rank`, +producing memory layout `[outer][inner][k]`. ONNX TopK output keeps the input shape with +`shape[axis]=k`, i.e. the correct layout is k-major `[outer][k][inner]`. The rank-`r` +winner for slot (outer, i) must land at flat index `(outer*k + r)*inner + i`. + +The two layouts are EQUAL only when `inner == numel(shape[axis+1..]) == 1`, i.e. the +reduced axis is the final axis. So the bug was latent for final-axis TopK (the common +router/argmax-style case) but produced wrong values+indices ordering for any +**non-final-axis TopK with k>1** — dtype-independent, pre-existing. CUDA EP / ORT emit +k-major. + +## Fix +Pre-size the output vectors to `numel(outputs[0].shape)` and write each selected element +to its strided destination `(outer*k + rank)*inner + i` instead of pushing sequentially. +Removed the now-unused `sorted` struct field (the kernel always emits sorted winners, +which satisfies both `sorted=1` and `sorted=0`); documented the intent in the factory. + +## Validation +- New tests: `topk_non_final_axis_uses_k_major_layout` (shape [2,3,2], axis=1, k=2) and + `topk_first_axis_is_k_major` (shape [3,2], axis=0, k=2) — both assert the k-major + values AND indices. Old push-order code would emit e.g. `[5,3,4,2]` vs correct + `[5,4,3,2]`, so the tests are discriminating. +- Final-axis path unchanged: existing `topk_bfloat16_values_match_widened_reference` and + `topk_and_nonzero` still pass verbatim (inner==1). +- `cargo test -p onnx-runtime-ep-cpu --lib selection`: 22 passed / 0 failed. +- `cargo fmt --all`; `cargo clippy -p onnx-runtime-ep-cpu --lib` clean. diff --git a/crates/onnx-runtime-ep-cpu/src/kernels/selection.rs b/crates/onnx-runtime-ep-cpu/src/kernels/selection.rs index 6e77428505..f22c14c879 100644 --- a/crates/onnx-runtime-ep-cpu/src/kernels/selection.rs +++ b/crates/onnx-runtime-ep-cpu/src/kernels/selection.rs @@ -328,7 +328,6 @@ impl Kernel for ArgKernel { pub struct TopKKernel { axis: i64, largest: bool, - sorted: bool, } /// ONNX `NonMaxSuppression` for the standard `[batch, boxes, 4]` / @@ -550,7 +549,9 @@ impl KernelFactory for TopKFactory { Ok(Box::new(TopKKernel { axis: node.attr("axis").and_then(|a| a.as_int()).unwrap_or(-1), largest: node.attr("largest").and_then(|a| a.as_int()).unwrap_or(1) != 0, - sorted: node.attr("sorted").and_then(|a| a.as_int()).unwrap_or(1) != 0, + // `sorted` is intentionally not stored: this kernel always emits the + // k winners in sorted order, which satisfies both `sorted=1` (must be + // sorted) and `sorted=0` (order unspecified). })) } } @@ -584,8 +585,15 @@ impl Kernel for TopKKernel { )); } let inner = numel(&inputs[0].shape[axis + 1..]); - let mut values = Vec::with_capacity(numel(outputs[0].shape)); - let mut indices = Vec::with_capacity(numel(outputs[1].shape)); + // TopK output shares the input shape with `shape[axis] = k`, so the + // reduced axis stays in place: the rank-`r` winner for a given + // (outer, inner) slot lands at `(outer * k + r) * inner + i`, NOT at the + // sequential push position. These differ whenever `inner > 1` (a + // non-final axis), so writing to the strided destination is required to + // match the ONNX (and CUDA-EP) k-major layout `[outer][k][inner]`. + let out_len = numel(outputs[0].shape); + let mut values = vec![0f32; out_len]; + let mut indices = vec![0i64; out_len]; for outer in 0..numel(&inputs[0].shape[..axis]) { for i in 0..inner { let mut candidates: Vec = (0..width).collect(); @@ -598,12 +606,10 @@ impl Kernel for TopKKernel { self.largest, ) }); - if !self.sorted { - candidates.truncate(k); - } - for d in candidates.into_iter().take(k) { - values.push(x[(outer * width + d) * inner + i]); - indices.push(d as i64); + for (rank, d) in candidates.into_iter().take(k).enumerate() { + let dst = (outer * k + rank) * inner + i; + values[dst] = x[(outer * width + d) * inner + i]; + indices[dst] = d as i64; } } } @@ -932,7 +938,6 @@ mod tests { TopKKernel { axis: -1, largest: true, - sorted: true, } .execute( &[input.view(), count.view()], @@ -976,7 +981,6 @@ mod tests { TopKKernel { axis: -1, largest: true, - sorted: true, } .execute(&[x.view(), k.view()], &mut [v.view_mut(), i.view_mut()]) .unwrap(); @@ -990,6 +994,53 @@ mod tests { assert_eq!(o.to_i64(), vec![0, 1, 1, 0]); } + #[test] + fn topk_non_final_axis_uses_k_major_layout() { + // Shape [2,3,2], axis=1 (width=3, inner=2), k=2. This is the case the + // old push-order layout got wrong: with inner>1 and k>1 the winners must + // land in k-major `[outer=2][k=2][inner=2]` order, not `[outer][inner][k]`. + // Column i=0 down axis 1: [1,5,3] -> top2 values [5,3] idx [1,2]. + // Column i=1 down axis 1: [2,0,4] -> top2 values [4,2] idx [2,0]. + // outer=1 mirrors with +10 so the two outers are distinguishable. + let x = Owned::f32( + &[2, 3, 2], + &[ + 1., 2., 5., 0., 3., 4., // outer 0: rows (1,2),(5,0),(3,4) + 11., 12., 15., 10., 13., 14., // outer 1 + ], + ); + let k = Owned::i64(&[], &[2]); + let mut v = Owned::zeros_f32(&[2, 2, 2]); + let mut i = Owned::zeros(DataType::Int64, &[2, 2, 2]); + TopKKernel { + axis: 1, + largest: true, + } + .execute(&[x.view(), k.view()], &mut [v.view_mut(), i.view_mut()]) + .unwrap(); + // k-major: [ [rank0_i0, rank0_i1], [rank1_i0, rank1_i1] ] per outer. + assert_eq!(v.to_f32(), vec![5., 4., 3., 2., 15., 14., 13., 12.]); + assert_eq!(i.to_i64(), vec![1, 2, 2, 0, 1, 2, 2, 0]); + } + + #[test] + fn topk_first_axis_is_k_major() { + // Shape [3,2], axis=0 (width=3, inner=2), k=2. + // Column 0: [2,4,1] -> top2 [4,2] idx [1,0]. Column 1: [5,0,3] -> [5,3] idx [0,2]. + let x = Owned::f32(&[3, 2], &[2., 5., 4., 0., 1., 3.]); + let k = Owned::i64(&[], &[2]); + let mut v = Owned::zeros_f32(&[2, 2]); + let mut i = Owned::zeros(DataType::Int64, &[2, 2]); + TopKKernel { + axis: 0, + largest: true, + } + .execute(&[x.view(), k.view()], &mut [v.view_mut(), i.view_mut()]) + .unwrap(); + assert_eq!(v.to_f32(), vec![4., 5., 2., 3.]); + assert_eq!(i.to_i64(), vec![1, 0, 0, 2]); + } + #[test] fn non_max_suppression_filters_overlaps_per_class() { let boxes = Owned::f32(