From 29b449192798aac1671d4ef4fc96cec8501f2618 Mon Sep 17 00:00:00 2001 From: shanmugamr1992 Date: Wed, 5 Aug 2026 16:59:47 -0700 Subject: [PATCH 1/4] Qwen model optimizer --- .cursor/agents/qwen-model-optimizer.md | 219 + AGENTS.md | 7 +- skills/cog-setup-and-help/SKILL.md | 1044 +++++ skills/git-credentials-setup/SKILL.md | 104 + .../scripts/load_git_credentials.sh | 64 + skills/nsight-system-analysis/SKILL.md | 283 ++ .../assets/report_template.md | 194 + .../references/pitfalls.md | 153 + .../references/sql_recipes.md | 114 + .../references/taxonomy_template.yml | 74 + skills/nsight-system-analysis/scripts/_lib.py | 218 + .../scripts/busy_idle.py | 224 + .../scripts/categorize.py | 230 ++ .../scripts/exposed_comm.py | 238 ++ .../scripts/forward_pass.py | 328 ++ .../scripts/iter_anchor.py | 381 ++ .../scripts/module_diff.py | 126 + .../scripts/module_slice.py | 338 ++ .../nsight-system-analysis/scripts/run_all.py | 328 ++ skills/optimize-inference-siddharth/SKILL.md | 510 +++ .../assets/review-checklist.md | 167 + .../references/commit-log.md | 121 + .../references/cuda-graphs.md | 302 ++ .../references/decision-gates.md | 237 ++ .../references/host-path.md | 225 ++ .../references/mamba-and-triton.md | 321 ++ .../references/measuring.md | 440 ++ .../references/moe-inference.md | 407 ++ .../references/updating-this-skill.md | 194 + .../references/vllm-differential.md | 191 + .../scripts/compare_budget.py | 146 + .../scripts/kernel_neighbors.py | 92 + .../scripts/steady_window.py | 59 + .../scripts/union_window.py | 95 + skills/run-qwen-model/EXPERIMENTS.md | 3590 +++++++++++++++++ skills/run-qwen-model/SKILL.md | 153 + skills/run-qwen-model/profile_qwen_mcore.sh | 149 + skills/run-qwen-model/profile_qwen_vllm.sh | 134 + skills/run-qwen-model/run_qwen_inference.sh | 334 ++ skills/run-qwen-model/run_qwen_vllm.sh | 110 + 40 files changed, 12638 insertions(+), 6 deletions(-) create mode 100644 .cursor/agents/qwen-model-optimizer.md create mode 100644 skills/cog-setup-and-help/SKILL.md create mode 100644 skills/git-credentials-setup/SKILL.md create mode 100755 skills/git-credentials-setup/scripts/load_git_credentials.sh create mode 100644 skills/nsight-system-analysis/SKILL.md create mode 100644 skills/nsight-system-analysis/assets/report_template.md create mode 100644 skills/nsight-system-analysis/references/pitfalls.md create mode 100644 skills/nsight-system-analysis/references/sql_recipes.md create mode 100644 skills/nsight-system-analysis/references/taxonomy_template.yml create mode 100644 skills/nsight-system-analysis/scripts/_lib.py create mode 100644 skills/nsight-system-analysis/scripts/busy_idle.py create mode 100644 skills/nsight-system-analysis/scripts/categorize.py create mode 100644 skills/nsight-system-analysis/scripts/exposed_comm.py create mode 100644 skills/nsight-system-analysis/scripts/forward_pass.py create mode 100644 skills/nsight-system-analysis/scripts/iter_anchor.py create mode 100644 skills/nsight-system-analysis/scripts/module_diff.py create mode 100644 skills/nsight-system-analysis/scripts/module_slice.py create mode 100644 skills/nsight-system-analysis/scripts/run_all.py create mode 100644 skills/optimize-inference-siddharth/SKILL.md create mode 100644 skills/optimize-inference-siddharth/assets/review-checklist.md create mode 100644 skills/optimize-inference-siddharth/references/commit-log.md create mode 100644 skills/optimize-inference-siddharth/references/cuda-graphs.md create mode 100644 skills/optimize-inference-siddharth/references/decision-gates.md create mode 100644 skills/optimize-inference-siddharth/references/host-path.md create mode 100644 skills/optimize-inference-siddharth/references/mamba-and-triton.md create mode 100644 skills/optimize-inference-siddharth/references/measuring.md create mode 100644 skills/optimize-inference-siddharth/references/moe-inference.md create mode 100644 skills/optimize-inference-siddharth/references/updating-this-skill.md create mode 100644 skills/optimize-inference-siddharth/references/vllm-differential.md create mode 100644 skills/optimize-inference-siddharth/scripts/compare_budget.py create mode 100644 skills/optimize-inference-siddharth/scripts/kernel_neighbors.py create mode 100644 skills/optimize-inference-siddharth/scripts/steady_window.py create mode 100644 skills/optimize-inference-siddharth/scripts/union_window.py create mode 100644 skills/run-qwen-model/EXPERIMENTS.md create mode 100644 skills/run-qwen-model/SKILL.md create mode 100755 skills/run-qwen-model/profile_qwen_mcore.sh create mode 100755 skills/run-qwen-model/profile_qwen_vllm.sh create mode 100755 skills/run-qwen-model/run_qwen_inference.sh create mode 100755 skills/run-qwen-model/run_qwen_vllm.sh diff --git a/.cursor/agents/qwen-model-optimizer.md b/.cursor/agents/qwen-model-optimizer.md new file mode 100644 index 00000000000..d23ce5bb5e3 --- /dev/null +++ b/.cursor/agents/qwen-model-optimizer.md @@ -0,0 +1,219 @@ +--- +name: qwen-model-optimizer +description: Autonomously profiles and optimizes Megatron-Core Qwen3-30B-A3B inference on OCI 4×GB200 until EP4 mcore matches or exceeds the fixed BS256 vLLM DP4+EP baseline. Use for Qwen performance analysis, Nsight A/B profiling, bottleneck-driven implementation, and iterative benchmark validation. +model: inherit +readonly: false +is_background: false +--- + +You are the Qwen3-30B-A3B performance optimization controller. + +Your goal is to make Megatron-Core EP4/TP1 inference match or exceed vLLM +DP4+EP throughput on one OCI 4×GB200 node at batch size 256, without +correctness regressions. + +## Mandatory context + +Before acting, read: + +1. `CLAUDE.md` +2. `skills/run-qwen-model/SKILL.md` — the fixed workload, launch commands, and + the cluster escape hatches +3. `skills/run-qwen-model/EXPERIMENTS.md` — the campaign ledger +4. `skills/nsight-system-analysis/SKILL.md` — **the** profile analysis skill +5. `skills/optimize-inference-siddharth/SKILL.md` — **the** optimization skill + +There are exactly two performance skills: analysis and optimization. Do not go +looking for others, and do not invent skill paths. + +Treat `EXPERIMENTS.md` as the sole source of performance history. + +## Division of labour between the two skills + +| Question | Skill | +|---|---| +| What does this trace say? Where does one decode step spend its time? | `nsight-system-analysis` | +| Which lever do I pull, is it worth building, and how do I not break it? | `optimize-inference-siddharth` | + +`nsight-system-analysis` owns windowing, interval-union arithmetic, per-category +attribution, and the report format. Its `scripts/forward_pass.py` (Workflow C) +is the primary lens for this workload. + +`optimize-inference-siddharth` owns the decision gates (share is not headroom), +the CUDA-graph / MoE / Triton / host-path playbooks, the hard rules, the flag +table including flags already measured and rejected on this exact model, and the +same-session A/B protocol. Its `scripts/` (`steady_window.py`, +`union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) complement +Workflow C when you need matched per-bucket budgets or to attribute a +generically-named kernel to its call site. + +Read the optimization skill's routing table (Step 1) and its +`references/vllm-differential.md` before proposing any change: the "identical +launch count but slower" and "more launches for the same work" cases have +opposite fixes, and several obvious levers are already recorded as rejected with +their mechanism. + +Supporting skills, only when the task calls for them: +`skills/cog-setup-and-help` (cluster, image, session, sbatch escape hatch), +`skills/git-credentials-setup` (any GitHub auth failure). + +## Baseline gate + +Do not modify Megatron-Core until both fresh baselines are recorded: + +1. Run vLLM DP4+EP under Nsight Systems at BS256. +2. Run mcore EP4/TP1 under Nsight Systems at BS256. +3. Record throughput, average latency, TPOT, job/run paths, `.nsys-rep`, and + `.sqlite` in `EXPERIMENTS.md`. +4. Verify the workloads and hardware match. +5. Compute the mcore-to-vLLM throughput gap, then run Workflow C on both + `.sqlite` files to record the per-forward-pass composition and the initial + ranked opportunity list. + +Use the commands and fixed protocol in `skills/run-qwen-model/SKILL.md`. + +## Optimization loop + +Repeat until mcore reaches vLLM: + +1. **Analyze** with Workflow C of `nsight-system-analysis`. Everything outside + the steady-state decode loop is noise: + + ``` + python skills/nsight-system-analysis/scripts/forward_pass.py \ + .sqlite .sqlite --label-a mcore --label-b vllm + ``` + + (the current baseline trace paths, under `nsys_trace/`, are recorded in + `EXPERIMENTS.md`). It auto-isolates one decode step per engine and prints + wall time, GPU-busy vs idle, launch counts, + and a per-category Δ table. Same µs/kernel with more launches ⇒ the lever is + fusion / fewer launches; higher µs/kernel on the same shape ⇒ a real + kernel-selection finding. Fall through to the skill's Steps 1–6 (exposed + comm, module-slicing, source root cause) only when Workflow C points at a + category needing deeper attribution, restricting windows to the decode region + it identified. +2. Classify the dominant signal — compute, memory, launch, communication, + synchronization, or host scheduling — then jump to the matching section via + the routing table in `optimize-inference-siddharth` Step 1. +3. **Gate the lever before building it.** Share of device time is not headroom. + Compute the ceiling per `references/decision-gates.md`, subtract what the fix + itself costs, and write down *proceed* or *gated out*. Skip the gate only for + cheap reversible changes (flag flips, tile retunes, backend swaps). +4. State one measurable hypothesis. +5. Back up or capture the current diff before editing. +6. Implement one change in the bottleneck's source path, honoring the + optimization skill's hard rules, and add a kill switch. +7. Run focused correctness tests. +8. Run the fixed BS256 mcore benchmark using the same-session, back-to-back, + alternating-arms A/B protocol (hard rule 10). Cross-session comparisons drift + more than most individual wins. +9. Capture a new profile when the timing composition could have changed. +10. Append the complete result to `EXPERIMENTS.md` — including rejections, with + their root cause and date. +11. Keep improvements; revert regressions and correctness failures. +12. For each accepted change, open its own MR (next section). +13. Promote whatever generalizes into the skills (section after that). + +Never optimize from a warmup/capture-only nsys window. `forward_pass.py` already +anchors analysis to a steady-state decode step; trust its window, and cross-check +its reported forward-pass period against measured TPOT before acting. Always +reason in terms of one forward pass: GPU-busy interval union, GPU idle, +per-category time, launch counts, and critical communication or kernel tails. + +## Ship each accepted experiment as its own MR + +Every experiment that succeeded gets a **separate** draft PR against +`main` in `NVIDIA/Megatron-LM`. One mechanism per PR, so the measurement +attributes cleanly and CODEOWNERS review stays narrow (see +`skills/mcore-split-pr/SKILL.md` if a change spans several owner groups). + +Before opening it, walk `skills/optimize-inference-siddharth/assets/review-checklist.md`. + +Mechanics, per `CLAUDE.md`: branch off `main`, commit with both `-s` and `-S`, +push to your **personal fork** — never to `NVIDIA/Megatron-LM` — then +`gh pr create --draft`. `origin` here points at the upstream repo, so confirm a +fork remote exists and ask the user for it if it does not. + +The description must let a reviewer judge the change without re-running it: + +- **What changed and why** — a brief summary of the mechanism, not a file list. +- **Measured gain** — percentage over baseline, with the baseline it is measured + against, and the absolute throughput/latency/TPOT numbers. +- **Protocol** — hardware, model, batch, OSL, parallelism, warmup/timed counts, + and that the arms ran back to back in the same allocation. State arm + separation (`min(ON) > max(OFF)`), not just the mean delta. +- **Where the time went** — the kernel, launch count, or host span the win is + attributed to, and the predicted-vs-measured kernel-to-e2e conversion. +- **Correctness** — which tests ran; whether the change is bit-exact, and if + not, the ulp bound plus the coherence check. +- **Kill switch** — the flag or env var that turns it off. +- **Scope and risks** — configs where it does not apply or was not measured. +- **Artifacts** — ledger entry id, run/job paths, `.nsys-rep` / `.sqlite` paths. + +Link the PR from its `EXPERIMENTS.md` entry so the ledger and the MR are +cross-referenced. Rejected experiments do not get a PR — they get a ledger entry +with the root cause. + +## Keep the skills learnable + +Both skills are living documents and you are authorized to edit them. Do it at +the *end* of a piece of work, once you have a number and a root cause — never +mid-experiment on a hypothesis. + +**The promotion test:** would this change what a competent engineer does on a +*different* model or workload? If yes it belongs in a skill. If it is a number +that only describes Qwen3-30B at one shape, it stays in `EXPERIMENTS.md`. For the +common in-between case, split it: the mechanism goes in the skill, the number +stays in the ledger, and the skill cites the number as a calibrated example with +its hardware and date stamp. + +Route by subject: + +| Learning | Destination | +|---|---| +| A profiling technique, anchoring trick, windowing pitfall, taxonomy fix, script improvement | `skills/nsight-system-analysis/` — `SKILL.md` for a workflow or hard rule, `references/pitfalls.md` for a trap, `references/sql_recipes.md` for a query, `references/taxonomy_template.yml` for categories, `scripts/` for tooling | +| An optimization pattern, decision gate, flag behavior, invariant, A/B methodology, competitor-diff insight | `skills/optimize-inference-siddharth/` — route via the table in `references/updating-this-skill.md`; only new invariants, flag behavior, and routing lines go in its `SKILL.md` | +| Cluster, queue, image, or launch failure that cost real time | `skills/run-qwen-model/SKILL.md` or `skills/cog-setup-and-help/SKILL.md` | +| Everything else about this campaign | `skills/run-qwen-model/EXPERIMENTS.md` | + +The bar for an addition, all four required: **measured**, **root-caused**, +**scoped** (hardware, model shape, token count, config), and **actionable**. +Quote exact error strings. Cite a commit sha or ledger entry id. Correct what is +now false in place; when a new measurement contradicts an old one, scope both +rather than overwriting. **Never delete a measured negative result** — those are +what stop the next agent re-deriving a dead end. + +Follow `skills/optimize-inference-siddharth/references/updating-this-skill.md` +for triggers, house style, size budgets, and the post-edit checklist, and append +to its revision log. Apply the same discipline to +`skills/nsight-system-analysis/` (see its *Keeping this skill current* section). + +## Guardrails + +- Do not change the vLLM baseline configuration. +- Do not change batch size, OSL, checkpoint, hardware, or parallelism to claim + a speedup. +- Do not stack unmeasured changes; one mechanism per experiment and per PR. +- Never invent expected or observed metrics. +- Do not retain a change that only improves throughput by breaking coherence. +- Do not use `git reset --hard` or `git clean`. +- Never push to `NVIDIA/Megatron-LM`; PRs come from a personal fork and open as + drafts. +- Open PRs only for measured, accepted wins. Do not commit or push anything else + unless explicitly requested. +- In Megatron Core production code, pass process groups explicitly; do not add + direct global `parallel_state.get_*_group()` reads. + +## Completion + +Finish only when one of these is true: + +- Confirmed parity: mcore throughput is at least the fresh vLLM baseline and + correctness passes. +- A definitive blocker prevents further work; record it with evidence and a + concrete next step. + +Return a concise summary containing baseline, best mcore result, remaining gap, +accepted changes with their MR links, rejected experiments with root causes, +skill updates made, and artifact paths. diff --git a/AGENTS.md b/AGENTS.md index 996c38c17bb..e747867f8b0 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -24,12 +24,7 @@ skill keyword — infer it from the artifact you read. - All PRs must be created as **drafts**. Use `gh pr create --draft` or the GitHub UI draft option. - Never push branches directly to `https://github.com/NVIDIA/Megatron-LM`. You must push your branch to a personal fork (e.g. `https://github.com//Megatron-LM`), then open a PR from the fork's branch against `NVIDIA/Megatron-LM`. -- Commit PR changes with both `-s` and `-S`: `-s` adds the required - `Signed-off-by` trailer, and `-S` signs the commit so copy-pr-bot and `/ok to -test` can verify the pushed commit without manually specifying the SHA. -Megatron Core engineers at NVIDIA should sign using their NVIDIA emails so they -are automatically added to the right user groups on the internal Slack -workspace. +- Commit PR changes with both `-s` and `-S`: `-s` adds the required `Signed-off-by` trailer, and `-S` signs the commit so copy-pr-bot and `/ok to test` can verify the pushed commit without manually specifying the SHA. - Read @docs/developer/contribute.md for the full contribution policy, including code style, commit message conventions, and issue guidelines. ### Code Quality diff --git a/skills/cog-setup-and-help/SKILL.md b/skills/cog-setup-and-help/SKILL.md new file mode 100644 index 00000000000..467b98f2e25 --- /dev/null +++ b/skills/cog-setup-and-help/SKILL.md @@ -0,0 +1,1044 @@ +--- +name: cog-setup-and-help +description: Canonical reference for everything `cog` (Claude-first local control plane for Megatron-LM workloads on Slurm) — installation, cluster registration, verification, the full CLI command catalogue (submit, session start/exec/exec-status/exec-cancel/stop, ensure-env, prepare-image, doctor, profile, logs slurm, logs app, jobs get, sessions ls, runs ls, cluster add/ls/show/use/remove), environment variables, error codes, and the iterative session-vs-submit workflow. Every other skill in this repo that uses cog defers to this one — read it first instead of the upstream cli-guide. +when_to_use: Setting up cog for the first time on a new machine; cloning and installing cog; registering a cluster with cog; verifying cog by running a Megatron-LM unit test on the cluster; looking up the flags / output shape / error code for any `cog` subcommand; deciding between `cog submit` and `cog session`; debugging an `error.code` from a cog command; a cog command that hangs instead of returning, and when to abandon it for a hand-written `sbatch` against a lustre-staged tree; a compute node that reports missing Python modules for files that `ls` can see; 'install cog', 'set up cog', 'configure cog', 'cog cluster add', 'cog doctor fails', 'cog session', 'cog submit failed', 'cog hangs', 'workspace sync hangs', 'sync-workspace stuck', 'submit sbatch directly', 'ModuleNotFoundError on the cluster', 'what does cog X do', 'cog error code'. +--- + +# `cog` — setup and help + +This is the canonical in-repo reference for cog: how to install it, +how to register a cluster, the full command catalogue, error codes, +and the iterative-workflow guidance. Other skills that drive cog +(`run-inference-unit-tests`, `run-inference-functional-tests`, +`run-inference-performance-tests`, …) link here for every cog +question — read this skill *before* opening the upstream +`docs/cli-guide.md`. The reference sections below mirror the cli-guide +but are kept in-tree so an offline / context-only agent can resolve +every cog question without a web fetch. + +`cog` is a CLI that runs Megatron-LM workloads on a Slurm cluster from your +laptop: it imports the container image to scratch as `.sqsh`, syncs your +worktree, materializes a shared `.venv`, and submits jobs via `srun` over SSH. + +This skill takes you from nothing to "I ran a Megatron-LM unit test on the +cluster via cog" in four steps: + +1. Clone cog +2. Install the `cog` CLI +3. **3a. Collect & persist cluster defaults to `~/.cog/setup.env`** → + 3b. Register the cluster with cog and verify connectivity +4. Submit a Megatron-LM unit test through cog + +> **🔁 Future agent invocations:** anything that touches cog (`cog +> submit`, `cog session exec`, `cog ensure-env`, …) **must `source +> ~/.cog/setup.env` first** and use the `$COG_*` variables it defines +> (`$COG_SSH_HOST`, `$COG_RUNTIME_ACCOUNT`, `$COG_INTERACTIVE_PARTITION`, +> `$COG_BATCH_PARTITION`, `$COG_SCRATCH_ROOT`, `$COG_MEGATRON_REPO`, …). +> Never assume hardcoded hostnames, accounts, partitions, or repo +> paths — read the env file. If the file is missing, run Step 3a to +> populate it (asking the user for each value with the documented +> defaults) before proceeding. + +## Prerequisites + +- **`uv`** installed on your local machine. If missing: + `curl -LsSf https://astral.sh/uv/install.sh | sh` +- **SSH access** to the cluster login node, with key-based auth working + non-interactively. Test with `ssh -o BatchMode=yes true`. +- **A local Megatron-LM checkout** on disk (cog syncs *this* directory to + cluster scratch — it is not the same as the cog repo). +- **A cluster scratch directory** you can write to, with enough capacity for + the container `.sqsh` (~10 GB) plus a `.venv` and run artifacts. +- **Slurm account + partition** with permission to submit `srun` jobs that + use Pyxis / enroot container mounts. For NVIDIA `cw-dfw`, the usual + account is `coreai_dlalgo_llm` (yours may differ — confirm with the + cluster owner). + +--- + +## Step 1 — Clone cog + +The repo is on the internal NVIDIA GitLab. Two clone forms — pick the +one that matches how you authenticate to `gitlab-master.nvidia.com`: + +**SSH (recommended if you already have a key uploaded to GitLab):** + +```bash +# gitlab-master.nvidia.com uses a non-standard SSH port (12051), not 22 +ssh-keyscan -H -p 12051 gitlab-master.nvidia.com >> ~/.ssh/known_hosts +git clone ssh://git@gitlab-master.nvidia.com:12051/shanmugamr/cog.git ~/cog +cd ~/cog +``` + +**HTTPS (needs a Personal Access Token in a credential helper):** + +```bash +git clone https://gitlab-master.nvidia.com/shanmugamr/cog.git ~/cog +cd ~/cog +``` + +> **⚠️ The internal GitLab uses SSH port 12051, not the default 22.** +> If you try `git@gitlab-master.nvidia.com:shanmugamr/cog.git` (the +> shorthand SSH form) you'll get the banner *"If you are trying to +> clone, you are using the incorrect port, use 12051"*. The +> `ssh://...:12051/...` URL form above is the fix. +> +> If the HTTPS form fails with *"could not read Username for +> https://gitlab-master.nvidia.com: Device not configured"*, your shell +> is non-interactive and has no cached credential. Either configure a +> Personal Access Token in `~/.git-credentials` / a credential helper, +> or use the SSH form above. + +--- + +## Step 2 — Install the `cog` CLI + +From inside the clone, install as an editable `uv` tool so future +`git pull` updates are picked up without reinstalling: + +```bash +uv tool install --editable . --force +``` + +Verify the binary is on `PATH` and prints help: + +```bash +cog --help +``` + +If `cog: command not found`, add `uv tool`'s bin dir to `PATH`: + +```bash +echo 'export PATH="$HOME/.local/bin:$PATH"' >> ~/.zshrc # or ~/.bashrc +exec $SHELL -l +``` + +> Alternative: `uv run cog --help` from inside the cog repo also works +> without a global install, but every command then requires +> `--project ~/cog` or being run from `~/cog`. The `uv tool install` path +> above is the recommended one for daily use. + +--- + +## Step 3a — Collect & persist cluster defaults + +This skill stores user-specific defaults in **`~/.cog/setup.env`** so +later steps (and later agent turns / later skill invocations) don't +have to re-prompt. The exact same file is the single source of truth +for SSH host, accounts, partitions, scratch root, and the local +Megatron-LM checkout path. + +**Procedure (the agent running this skill should do this — don't +hardcode anything):** + +1. Check whether `~/.cog/setup.env` exists. +2. **If it exists**, run `source ~/.cog/setup.env` and skip ahead to + Step 3b. Confirm each variable is set (`echo $COG_SSH_HOST`, …) and + only re-prompt for any that come up empty. +3. **If it does not exist**, ask the user for each variable in the + table below using `AskUserQuestion` (or equivalent). For every + variable except `COG_MEGATRON_REPO`, *offer the listed default* — + accept the default if the user doesn't override. `COG_MEGATRON_REPO` + has no default; require an answer. + + | Variable | Prompt to the user | Default | + |---|---|---| + | `COG_CLUSTER_NAME` | Name to register this cluster under in the cog registry | `cw-dfw` | + | `COG_SSH_HOST` | SSH login host for the cluster | `cw-dfw-cs-001-vscode-01` | + | `COG_RUNTIME_ACCOUNT` | Slurm account for GPU runtime jobs (`submit`, `session start`) | `coreai_dlalgo_llm` | + | `COG_IMPORT_ACCOUNT` | Slurm account for the CPU-side `enroot import` job | same value as `COG_RUNTIME_ACCOUNT` | + | `COG_INTERACTIVE_PARTITION` | Partition for interactive / short jobs (used by Step 4 and most `cog submit` examples) | `interactive` | + | `COG_BATCH_PARTITION` | Partition for long-running batch jobs | `batch` | + | `COG_IMPORT_PARTITION` | Partition that runs the CPU `enroot import` job | `cpu` | + | `COG_SCRATCH_ROOT` | Absolute path to your scratch root on the cluster (cog will store `.sqsh`, `.venv`, workspaces, run logs under this) | `/lustre/fsw/portfolios/coreai/users/${USER}/agents-space` | + | `COG_MEGATRON_REPO` | Absolute path to your local Megatron-LM checkout | **no default — must come from the user** | + +4. Write the answers to `~/.cog/setup.env`. **Each line must use + `export KEY="value"`** — without the `export`, the variables exist + only in the sourcing shell and won't be visible to `cog`, `ssh`, + `srun`, or any other subprocess, which silently produces empty + `--ssh-host`, `--partition`, `--repo`, etc. arguments. Example: + + ```bash + mkdir -p ~/.cog + cat > ~/.cog/setup.env <<'EOF' + # Populated by the cog-setup-and-help skill. Source this file before invoking cog. + # Variables are `export`-ed so subprocesses (cog, ssh, srun) see them. + export COG_CLUSTER_NAME="cw-dfw" + export COG_SSH_HOST="cw-dfw-cs-001-vscode-01" + export COG_RUNTIME_ACCOUNT="coreai_dlalgo_llm" + export COG_IMPORT_ACCOUNT="coreai_dlalgo_llm" + export COG_INTERACTIVE_PARTITION="interactive" + export COG_BATCH_PARTITION="batch" + export COG_IMPORT_PARTITION="cpu" + export COG_SCRATCH_ROOT="/lustre/fsw/portfolios/coreai/users/${USER}/agents-space" + export COG_MEGATRON_REPO="/Users/${USER}/Megatron-LM" + EOF + source ~/.cog/setup.env + # Sanity check — both lines must show the value, not blank: + echo "$COG_SSH_HOST" + env | grep ^COG_ | sort + ``` + +5. From this point on, every `cog` invocation in this skill (and in + later turns) references `$COG_*` rather than literal values. + +--- + +## Step 3b — Register the cluster with cog + +With `~/.cog/setup.env` sourced, register the cluster once. `cog +cluster add --set-default` makes later commands omit `--cluster-name`: + +```bash +source ~/.cog/setup.env +cog cluster add \ + --name "$COG_CLUSTER_NAME" \ + --ssh-host "$COG_SSH_HOST" \ + --scratch-root "$COG_SCRATCH_ROOT" \ + --import-account "$COG_IMPORT_ACCOUNT" \ + --runtime-account "$COG_RUNTIME_ACCOUNT" \ + --import-partition "$COG_IMPORT_PARTITION" \ + --import-job-name cog:import \ + --set-default +``` + +Field notes: + +- **`--scratch-root`** must be a path that exists and is writable by you on + the cluster. Cog stores `.sqsh` images, the shared `.venv`, synced + workspaces, and run logs under this path. Create it ahead of time: + `ssh "$COG_SSH_HOST" mkdir -p "$COG_SCRATCH_ROOT"`. +- **`--import-account` / `--runtime-account`** can be the same Slurm + account or different ones if your site separates CPU import jobs from + GPU runtime jobs. +- **`--import-partition`** is a CPU partition; image import does not need + GPUs and reserving them just to run `enroot import` is wasteful. + +Verify and run `doctor`: + +```bash +cog cluster ls +cog doctor --repo "$COG_MEGATRON_REPO" +``` + +`cog doctor` checks: local Python / cog version, SSH reachability, +scratch writability, and `squeue --version` on the cluster. Status `ok` +or `degraded` exits `0`; `fail` exits `1`. Read the `checks[]` array in +the JSON output to see which probe failed before retrying. + +Common `doctor` failures: + +| Symptom | Likely cause | Fix | +|---------|--------------|-----| +| `SSH_CONNECTION_FAILED` | Key not loaded / host unknown | `ssh-add`, accept host key, set `COG_SSH_OPTS='-o StrictHostKeyChecking=accept-new'` | +| `SCRATCH_UNWRITABLE` | Scratch root missing or wrong owner | `ssh mkdir -p && touch /.cog-test` | +| `SLURM_UNREACHABLE` | `squeue` not on PATH on login node | Confirm you're pointing at a login node, not a worker | + +--- + +## Step 4 — Verify by running a Megatron-LM unit test + +The cheapest end-to-end smoke test is `tests/unit_tests/test_basic.py` — +it imports the package and runs a few CPU-only assertions, so it +exercises the full cog pipeline (image import, workspace sync, +`uv sync`, `srun`) without needing real GPU compute. + +### 4a. Pre-warm the image and env (optional but recommended) + +This avoids paying the ~5-minute import latency inside the actual test +job. Skip if you'd rather see one combined job: + +```bash +source ~/.cog/setup.env +cog prepare-image --repo "$COG_MEGATRON_REPO" + +cog ensure-env \ + --repo "$COG_MEGATRON_REPO" \ + --run-name env-warmup \ + --gpus 1 \ + --time 00:20:00 \ + --partition "$COG_INTERACTIVE_PARTITION" +``` + +`prepare-image` runs `enroot import` on the cluster and caches the +`.sqsh` keyed on the resolved base image. `ensure-env` runs `uv sync` +inside that image and writes a marker file so subsequent jobs reuse the +`.venv`. Both are idempotent: a second invocation with no changes +returns `cache_hit: true` instantly. + +### 4b. Submit the unit test + +```bash +source ~/.cog/setup.env +cog submit \ + --repo "$COG_MEGATRON_REPO" \ + --run-name verify-cog-basic \ + --command 'python -m pytest tests/unit_tests/test_basic.py -v -o addopts=' \ + --gpus 1 \ + --time 00:10:00 \ + --partition "$COG_INTERACTIVE_PARTITION" +``` + +> **Use `python -m pytest`, not bare `pytest`.** The shared `.venv` is +> built under a `.partial.` directory and then `mv`-d to its +> canonical recipe path; this breaks any bin-script (like `pytest`) that +> has the original path baked into its shebang, so a bare `pytest` call +> exits 127 (`pytest: command not found`). `python` itself is resolved +> via the `PATH` cog sets before sourcing the venv's `activate`, so +> `python -m pytest ...` works. Same rule as the `torchrun` → +> `python -m torch.distributed.run` substitution below. + +What this does: + +- Syncs your local Megatron-LM worktree to scratch. +- Allocates 1 GPU on the `interactive` partition for up to 10 minutes. +- Runs the test inside the container at the cached `.sqsh`, with the + shared `.venv` mounted. +- Streams `srun` to completion, then prints a JSON payload to stdout + including `job.job_id`, `job.returncode`, `job.stdout`, and + `artifacts.slurm_stdout_path`. + +Exit code `0` and `job.returncode: 0` means cog is fully wired up. If +the unit test prints "passed", you're done. + +> `-o addopts=` clears any `addopts` from `pyproject.toml` that would +> otherwise pull in distributed-only options. Keep it for the smoke test. + +### 4c. (Optional) Run a real distributed unit test + +Once the basic test passes, validate the multi-GPU + `torch.distributed` +path: + +```bash +source ~/.cog/setup.env +cog submit \ + --repo "$COG_MEGATRON_REPO" \ + --run-name verify-cog-gpt \ + --command 'python -m torch.distributed.run --nproc-per-node 8 --log-dir "$TORCHRUN_LOG_DIR" -m pytest -xvs tests/unit_tests/models/test_gpt_model.py::TestGPTModel::test_constructor' \ + --gpus 8 \ + --nodes 1 \ + --ntasks-per-node 8 \ + --time 01:00:00 \ + --partition "$COG_INTERACTIVE_PARTITION" +``` + +> Use `$COG_INTERACTIVE_PARTITION` for short / iterative jobs (test +> verification, debugging, `cog session` allocations). Use +> `$COG_BATCH_PARTITION` for long batch jobs (full training runs, +> sweeps). Pick based on intent, never hardcode the literal. + +Note: **`python -m torch.distributed.run`**, not bare `torchrun` — the +shared venv lives at a path where `torchrun`'s shebang sometimes +resolves to the wrong interpreter. Cog lints `--command` for bare +`torchrun` usage and surfaces a warning, but using `python -m +torch.distributed.run` from the start avoids the trap. + +--- + +## Iterating: prefer `cog session start` + `cog session exec` over repeated `cog submit` + +> **🔁 IMPORTANT — iterative workflow.** If you keep hitting errors, +> retrying the same job, or running short commands in a loop, **stop +> using `cog submit` and switch to `cog session`**. Every `cog submit` +> requests a fresh Slurm allocation, which means queue time + +> container start + venv activate on every attempt — minutes of latency +> per iteration. +> +> Instead, allocate one node for a few hours and exec into it +> repeatedly: +> +> ```bash +> source ~/.cog/setup.env +> +> # Hold an 8-GPU node for 3 hours. +> cog session start \ +> --repo "$COG_MEGATRON_REPO" \ +> --run-name iter-debug \ +> --gpus 8 --nodes 1 --ntasks-per-node 1 \ +> --time 03:00:00 \ +> --partition "$COG_INTERACTIVE_PARTITION" +> +> # Now run as many commands as you want — each starts in ~1 second +> # because the allocation, container, and venv are already up. +> cog session exec --run-name iter-debug --command 'python -m pytest tests/unit_tests/test_basic.py -v' +> cog session exec --run-name iter-debug --command 'python -c "import megatron; print(megatron.__file__)"' +> +> # When done, release the allocation. +> cog session stop --run-name iter-debug +> ``` +> +> Rule of thumb: if you'll run more than 2 commands against the same +> code, use a session. The full session command set +> (`session start` / `exec` / `exec-status` / `exec-cancel` / `stop`, +> `sessions ls`, `session requests ls`) is documented in +> `~/cog/docs/cli-guide.md` and online at +> . + +--- + +## Escape hatch: when workspace sync hangs, submit plain `sbatch` + +Every cog command that touches your code (`submit`, `session start`, +`session exec`) runs `sync-workspace` first. When that sync hangs it hangs +*before* Slurm sees anything, so there is no job ID to poll and no Slurm log +to read — the CLI just sits. Retrying is the trap: each attempt re-uploads +from scratch and stalls the same way, and a stalled sync looks identical to a +merely slow one, so "give it a few more minutes" can absorb an entire day. + +> **🔁 Two-strike rule.** If workspace sync hangs or dies with a broken pipe +> **twice in a row from the same client**, stop using cog for that run and +> submit `sbatch` yourself against a tree you stage on lustre. One measured +> case: four consecutive cog attempts (three hangs plus a broken pipe) burned +> most of a day, while plain `sbatch` against a lustre-staged tree returned a +> result **first try in under 7 minutes**. Do reach for cog again on the next +> task — sync is normally reliable, and this failure mode looks +> client- or load-specific rather than permanent. + +The fallback has to supply the three things cog would have done for you: the +imported `.sqsh` image, the container mounts, and a staged copy of the repo. + +**Getting the paths without triggering a sync.** `cog profile` is the one +command that resolves the plan *locally* — it never touches the cluster — so +it still works while sync is broken: + +```bash +source ~/.cog/setup.env +cog profile --repo "$COG_MEGATRON_REPO" --run-name fallback --pretty # sqsh_plan.sqsh_path +``` + +Do **not** reach for `cog ensure-env` to discover the venv path here — it +syncs the workspace too, so it will hang for the same reason. List the envs +directory over SSH instead: + +```bash +ssh "$COG_SSH_HOST" ls -d "$COG_SCRATCH_ROOT"/envs/megatron_lm/*/.venv +``` + +**Staging the tree.** Copy to a **brand-new directory name** every time: + +```bash +STAGE="$COG_SCRATCH_ROOT/workspaces/megatron_lm/fallback_$(date +%s)/repo" +ssh "$COG_SSH_HOST" mkdir -p "$STAGE" +rsync -a --exclude .git "$COG_MEGATRON_REPO"/ "$COG_SSH_HOST:$STAGE/" +``` + +> **Why a fresh path, not `rsync --delete` over the old one.** `--delete` +> replaces files with new inodes. A node whose Lustre client is misbehaving +> then serves stale dentries: the directory listing shows the file, but +> `open()` fails with `Input/output error` and Python reports the module as +> missing. A new directory name costs nothing and sidesteps it entirely. + +**The job script.** Note that you now own venv activation, which cog would +normally handle: + +```bash +#SBATCH --account=... # $COG_RUNTIME_ACCOUNT +#SBATCH --partition=batch # $COG_BATCH_PARTITION +#SBATCH --qos=interactive # never omit: the default `normal` is priority 100. + # `interactive` is priority 700 with no walltime cap + # at <=4 nodes, and works under sbatch despite the name. + # Use `short` (200) only when you need >4 nodes. +#SBATCH --nodes=1 --ntasks-per-node=1 --gres=gpu:4 +#SBATCH --time=01:30:00 +#SBATCH --exclude= + +srun --container-image="$IMG" \ + --container-mounts=/lustre:/lustre \ + --container-workdir="$STAGE" \ + --no-container-mount-home \ + bash -c '; ' +``` + +> **Quoting trap in that `bash -c '...'`.** The health gate below uses a +> `<<'PYEOF'` heredoc, whose single quotes terminate the outer single-quoted +> `bash -c` string. Either escape the delimiter as `<<\"PYEOF\"` inside the +> `bash -c`, or — cleaner — keep the gate in its own file in the staged tree +> and call `bash dev/fs_health.sh` from the `srun` body. + +### Always gate the allocation on a filesystem health check + +Put this at the top of the `srun` body, before the model load. It costs +milliseconds and is the difference between failing in seconds on a bad node +and discovering the problem after the allocation is gone: + +```bash +# Point at the shared venv found above; the gate checks inside it too, because +# a flaky client usually breaks the venv and your repo independently. +export VENV_SITE_PACKAGES="$COG_SCRATCH_ROOT/envs/megatron_lm//.venv/lib/python3.12/site-packages" + +python3 - <<'PYEOF' +import os, sys +# Span both your own changed files and something deep inside the shared venv. +must_exist = [ + "megatron/core/transformer/transformer_layer.py", + "megatron/core/inference/contexts/dynamic_context.py", + os.environ["VENV_SITE_PACKAGES"] + "/sympy/functions/elementary/_trigonometric_special.py", + os.environ["VENV_SITE_PACKAGES"] + "/networkx/algorithms/flow/preflowpush.py", +] +bad = [p for p in must_exist if not os.path.isfile(p)] +if bad: + print("FSHEALTH BROKEN - this node cannot stat files that exist:", *bad, sep="\n ") + sys.exit(3) +print("FSHEALTH OK") +PYEOF +[ $? -ne 0 ] && { echo "ABORT: node filesystem unhealthy, not spending the allocation"; exit 3; } +``` + +`os.path.isfile` is the right probe precisely because it disagrees with `ls` +on a broken node — that disagreement *is* the signal. + +> **⚠️ Read this before you debug a `ModuleNotFoundError` on a compute node.** +> A cascade of import failures for *mutually unrelated* packages — first +> `sympy.functions.elementary._trigonometric_special`, then +> `networkx.algorithms.flow.preflowpush`, then your own +> `megatron.core.transformer.transformer_layer` — where `ls` lists the file +> but `os.path.isfile()` returns `False`, is **one broken Lustre client on +> that node**, not corrupted packages. +> +> Do not "fix" it by pip-installing overlay copies into `PYTHONPATH`. That +> reading is seductive because each individual error looks like a genuinely +> damaged install, but you will chase one import after the next, add a +> misleading "package repair" script to the repo, and still lose the +> allocation. The correct move is to `scancel`, add +> `--exclude=` to the job, and resubmit. Confirm the diagnosis cheaply +> by running the health gate above on a different node — it will pass. + +--- + +## Reading logs + +`cog submit` prints `job.stdout` / `job.stderr` inline by default, but for +larger jobs use `cog logs slurm`: + +```bash +cog logs slurm \ + --run-name verify-cog-basic \ + --job-id \ + --stream both \ + --lines 200 +``` + +For cog's own behavior (what commands ran, when sessions started, errors +cog raised), use the local app log: + +```bash +cog logs app --lines 50 +``` + +--- + +## Common pitfalls + +- **`uv tool install` succeeded but `cog: command not found`.** Add + `$HOME/.local/bin` to `PATH` (see Step 2). `uv tool dir --bin` prints + the exact directory. +- **`RUN_ROOT_OUTSIDE_SCRATCH` error on `submit`.** `--run-root` (or the + path that `--run-name` resolves to) must be under the registered + cluster's `scratch_root`. Prefer `--run-name ` and let cog + compose the path. +- **First `submit` is slow (~5-10 min).** That is the one-time + `enroot import` + `uv sync`. Subsequent runs with the same image and + `pyproject.toml` hit cache and start in seconds. Pre-warm with + `prepare-image` + `ensure-env` if you want the slow path out of the + critical path. +- **`submit` / `session start` / `session exec` hangs with no job ID and no + Slurm log.** The hang is in `sync-workspace`, upstream of Slurm, so there is + nothing to poll. Distinguish it from the slow first-submit case above by the + absence of a `job_id` in the output. After the second consecutive hang or + broken pipe, switch to the plain-`sbatch` fallback in + "Escape hatch: when workspace sync hangs" rather than retrying a third time. +- **`cog doctor` shows `cluster_resolved: ok` but `ssh_reachable: fail`.** + The registry entry is fine; SSH itself is the problem. Reproduce with + `ssh -o BatchMode=yes true` and fix at the SSH layer. +- **Editing cog itself.** Because Step 2 installed editable + (`--editable`), any `git pull` in `~/cog` is picked up immediately — + no reinstall needed. If you ever switch to a wheel install + (`uv build && uv tool install --force dist/cog-*.whl`), you must + reinstall after each change. +- **Pointing cog at the wrong repo.** `--repo` is the **Megatron-LM** + checkout, not the cog checkout. Get this wrong and `doctor` will fail + the repo-profile check with a clear message — but it is the most + common day-one confusion. +- **`srun ... task 0: Exited with exit code 127` from `cog submit`, + Slurm log ends in `pytest: command not found` (or any other bin + script not found).** The shared `.venv` is built in a `.partial.` + directory, then renamed into place; bin-script shebangs still point at + the old path. Invoke via `python -m ` instead of the bin + wrapper (e.g. `python -m pytest`, `python -m torch.distributed.run`). + `python` itself works because cog explicitly sets `PATH` to the + canonical venv before sourcing `activate`. +- **Stale cog install: `cog` on `PATH` but raises + `ModuleNotFoundError: No module named 'cog'`.** Happens when an + earlier `uv tool install --editable` pointed at a clone that has + since been deleted. Fix with `uv tool uninstall cog`, re-clone (Step 1), + then `uv tool install --editable . --force` again. +- **`~/.cog/setup.env` is missing or has stale values.** Re-run Step 3a + to re-prompt and re-write the file. If only one variable is wrong + (e.g. user moved their Megatron-LM checkout), edit just that line in + the file rather than wiping it. Sourcing the file is cheap — always + do it at the start of any turn that runs a `cog` command. +- **`cog` runs but `--repo` / `--partition` / `--ssh-host` flags + receive empty strings.** The env file was written without `export`, + so `source ~/.cog/setup.env` populated the variables in the current + shell only — they were not visible to the `cog` subprocess. Fix: + prepend `export ` to every line in `~/.cog/setup.env` (the Step 3a + template already does this; older files written before this fix may + not). Verify with `source ~/.cog/setup.env && env | grep ^COG_` — + if the grep is empty, the file is missing `export` keywords. +- **Hardcoded hostnames / accounts / partitions / repo paths in agent + output.** This is a guidance failure, not a tool failure. Every cog + command this skill emits should reference `$COG_*` variables from + `~/.cog/setup.env`. If you find yourself typing literal `cw-dfw…`, + `coreai_dlalgo_llm`, or `interactive` into a `cog` flag, stop and + source the env file instead. + +--- + +# Cog command reference + +The rest of this skill is the in-tree mirror of cog's +`docs/cli-guide.md`. **Read this section before falling back to the +upstream guide.** The upstream guide +(, +or `~/cog/docs/cli-guide.md` locally) is the source of truth for +anything not covered here; this skill stays current with the commands, +flags, output fields, and error codes other skills in this repo +actually use. + +## Contract + +- Most commands print a single JSON object to stdout. Parse it; don't + scrape text. +- Every success JSON payload starts with `"schema_version": 1`. +- Errors are emitted on stderr as structured JSON: + `{"schema_version":1,"error":{"code":"UPPER_SNAKE","message":"...","details":{...}}}`. +- `logs slurm` and `logs app` are the exceptions — they stream plain + text. +- All compute paths stay under cluster scratch. +- Re-use the same `--session-handle` to keep using the same live + interactive allocation. +- For Megatron-LM v1, assume the `dev` image flavor unless explicitly + overridden. + +## Global options + +| Flag | Description | +|------|-------------| +| `--pretty` | Indent JSON output on stdout (and stderr for errors). | +| `--debug` | Include `traceback` in `INTERNAL_ERROR` error details on uncaught exceptions. | + +## Shared options (used by most commands) + +| Flag | Description | Default | +|------|-------------|---------| +| `--repo` | Path to the Megatron-LM git checkout. Resolves from cwd via `git rev-parse --show-toplevel` if omitted. | cwd | +| `--run-name` / `--run-root` | Identify a run. Exactly one is required where applicable. `--run-name` resolves to `/runs/`. Names must match `^[A-Za-z0-9][A-Za-z0-9._-]{0,127}$`. | — | +| `--cluster-config` | Path to a cluster YAML file. Overrides `--cluster-name`. | None | +| `--cluster-name` | Name of a registered cluster (from `cluster add`). Falls back to the default cluster. | None | +| `--base-image-flavor` | Base container image flavor (`dev` or `lts`). | `dev` | +| `--base-image` | Explicit base image override (bypasses flavor resolution). | None | +| `--dry-run` | Preview the execution plan without side-effects (available on `submit`, `ensure-env`, `prepare-image`). | `false` | + +## Environment variables + +Read by cog on the machine where you invoke the CLI — not per-command flags. + +| Variable | Description | +|----------|-------------| +| `COG_SSH_OPTS` | Space-separated extra `ssh` arguments appended **after** cog's defaults (later args win on duplicates). Example: `export COG_SSH_OPTS='-o ConnectTimeout=30'`. | +| `COG_EXTRA_MOUNTS` | Comma-separated `host:container` bind mounts appended to Slurm `--container-mounts` for one-shot container `srun` jobs (`submit`, `ensure-env`). The cluster scratch root is always mounted first as `:`. **Does NOT apply to persistent `session start` allocations.** Example: `export COG_EXTRA_MOUNTS='/lustre/shared:/lustre/shared,/data:/data'`. | + +SSH connection multiplexing is on by default +(`ControlMaster=auto`, `ControlPersist=60s`, +`ControlPath=~/.ssh/cm-%C`). Override via `COG_SSH_OPTS`. + +--- + +## Command catalogue + +The columns below are: what the command does · the flags you'll actually +use · the output fields agents read most often. + +### `cog doctor` — verify environment + cluster + +**Use first** when troubleshooting. Checks Python version, package +version, optional repo profile validation, then (if a cluster +resolves) SSH connectivity, scratch writability, and `squeue +--version`. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--cluster-name` / `--cluster-config` | default cluster | + +Output: `checks[]` (each with `name`, `status` ∈ {`ok`,`degraded`,`fail`}, `duration_ms`), `overall`, `cluster`, `cluster_source`. + +Exit `0` if `overall` is `ok` or `degraded`; `1` if `fail`. + +```bash +cog doctor --repo "$COG_MEGATRON_REPO" +``` + +### `cog profile` — resolve the run plan without touching the cluster + +Shows the resolved Megatron profile, run layout, base image, `.sqsh` +path, scratch paths. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--run-name` / `--run-root` (one required) | — | +| `--base-image-flavor` | `dev` | +| `--base-image` | None | + +Output: `profile`, `base_image`, `run_layout.{run_root,slurm_dir,logs_dir}`, `sqsh_plan.{cache_key,sqsh_path}`, `import_plan.{account,partition,command,output_path}`. + +### `cog prepare-image` — ensure the `.sqsh` exists on scratch + +Runs `enroot import` on the cluster (or reuses cache). Pre-warm before +`submit` / `ensure-env` to avoid in-job import latency. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--base-image-flavor` | `dev` | +| `--base-image` | None | +| `--dry-run` | `false` | + +Output: `dry_run`, `base_image`, `cache_key`, `sqsh_path`, `import_plan`, `cache_hit` (real run), `import_stdout/stderr` (real run). + +Exit `0` on success; `1` on `IMAGE_IMPORT_FAILED`. + +### `cog sync-workspace` — push the local worktree to scratch + +`submit`, `ensure-env`, and `session start`/`exec` call this internally — +only invoke directly to sync without running anything. + +Output: `workspace_hash`, `file_count`, `remote_repo_path`, `cache_hit`, `transfer_strategy`. + +### `cog ensure-env` — materialize the shared `.venv` for this dep recipe + +Submits an `srun` job that runs `uv sync` inside the container, writes +an env marker, and waits. Idempotent: a second call with the same +recipe key returns `ready_cache_hit: true` instantly. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--run-name` / `--run-root` (one required) | — | +| `--gpus` | `1` | +| `--time` | `00:10:00` | +| `--partition` | `interactive` | +| `--nodes` / `--ntasks-per-node` | `1` / `1` | +| `--base-image-flavor` / `--base-image` | `dev` / None | +| `--dry-run` | `false` | + +Output: `env.env_root`, `env.marker_path`, `env.ready_cache_hit`, `env.reason`, `env.recipe_key`, `env.dependency_inputs_key`, `job.{job_id,returncode,state,stdout,stderr}` (or `null` on cache hit). + +Exit `0` if no job ran or `job.returncode == 0`; otherwise the job's return code. `1` on `ENV_LOCK_TIMEOUT`. + +### `cog submit` — one-shot `srun` job, block to completion + +The default workhorse for **single** verification jobs. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--run-name` / `--run-root` (one required) | — | +| `--command` (required) | — | +| `--gpus` | `1` | +| `--time` | `00:10:00` | +| `--partition` | `interactive` | +| `--nodes` / `--ntasks-per-node` | `1` / `1` | +| `--skip-uv-sync` | `false` — set if you've already `ensure-env`'d the same recipe | +| `--base-image-flavor` / `--base-image` | `dev` / None | +| `--dry-run` | `false` | + +Cog lints `--command` (warns on bare `torchrun`), then submits via +`srun`. For multi-node jobs, cog auto-populates `MASTER_ADDR`, +`MASTER_PORT`, `WORLD_SIZE`, `RANK`, `LOCAL_RANK` inside the container. + +Output: `job.{job_id,command,returncode,state,raw_state,slurm_exit_code,node_list,stdout,stderr,stdout_path,stderr_path}`, `artifacts.{run_root,app_log_dir,torchrun_log_dir,manifest_path,slurm_stdout_path,slurm_stderr_path}`, optional `warnings`. + +Exit propagates `job.returncode`. + +### `cog session start` — persistent interactive allocation + +**Prefer this whenever you'll run more than 2 commands on the same +code.** Allocation stays alive for `--time`; you `session exec` into +it repeatedly. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--session-handle` (required, regex same as run names) | — | +| `--gpus` | `1` | +| `--time` | `04:00:00` | +| `--partition` | `interactive` | +| `--ntasks-per-node` | `1` | +| `--base-image-flavor` / `--base-image` | `dev` / None | + +Output: `session.{session_handle,job_id,state,node_list,controller_running,session_root,workdir}`, `reused` (true if an existing handle was matched). + +### `cog session status` — inspect a live session + +| Flag | +|------| +| `--session-handle` (required) | + +Output: `session.{session_handle,job_id,state,raw_state,node_list,controller_running,controller_pid}` and stdout/stderr/root paths. + +Errors: `SESSION_NOT_FOUND`, `SESSION_CLUSTER_MISMATCH`. + +### `cog session exec` — run one command inside an existing session + +Re-syncs the workspace, then dispatches the command. Lints +`--command` like `submit`. Foreground by default; `--detach` returns a +`request_id` for later polling. + +| Flag | Default | +|------|---------| +| `--repo` | cwd | +| `--session-handle` (required) | — | +| `--command` (required) | — | +| `--wait-timeout` | `3600` seconds | +| `--detach` | `false` | + +Output: `execution.{request_id,completed,exit_code,failure_reason,stdout_path,stderr_path,wait_timeout,detached_sentinel}`, `artifacts.{exec_root,run_root,torchrun_log_dir}`. + +Exit code: +- Detached: always `0`. +- Foreground: `124` on `--wait-timeout`; `1` if incomplete / `exit_code` null; otherwise the remote command's exit code. + +### `cog session exec-status` — poll a (typically detached) exec + +| Flag | +|------| +| `--session-handle` (required) | +| `--request-id` (required) | + +Output: `request.{request_id,state,exit_code,...}`. Error: `REQUEST_NOT_FOUND`. + +### `cog session exec-cancel` — kill a running exec + +| Flag | +|------| +| `--session-handle` (required) | +| `--request-id` (required) | + +Output: `cancel.{was_terminal,sentinel_written}`. Errors: `REQUEST_NOT_FOUND`, `REQUEST_ALREADY_TERMINAL`. + +### `cog session requests ls` — recent execs in a session + +| Flag | Default | +|------|---------| +| `--session-handle` (required) | — | +| `--limit` | `50` | + +Output: `requests[]` with `request_id`, `state`, `exit_code`, timestamps. + +### `cog session stop` — cancel the Slurm job, kill the controller + +| Flag | +|------| +| `--session-handle` (required) | + +Output: `session` payload extended with `cancel_requested`, `controller_was_running`, `controller_signalled`. + +### `cog cluster add` — register a cluster + +| Flag | Required unless `--from-file` | +|------|-------------------------------| +| `--name` | yes | +| `--ssh-host` | yes | +| `--scratch-root` | yes | +| `--import-account` | yes | +| `--runtime-account` | defaults to `--import-account` | +| `--import-partition` | yes | +| `--import-job-name` | yes | +| `--from-file` | YAML onboarding file (sets all of the above) | +| `--set-default` | makes this the default cluster | +| `--overwrite` | overwrite an existing entry of the same name | + +Errors: `CLUSTER_CONFIG_INVALID`, `CLUSTER_ALREADY_EXISTS`. + +### `cog cluster ls` / `show` / `use` / `remove` + +- `ls`: lists registered clusters; output has `default`, `clusters[]`, `invalid_entries[]`. +- `show --name `: full cluster dict + `is_default`. +- `use --name `: set the default. +- `remove --name `: delete the registry entry. + +Common error: `CLUSTER_NOT_FOUND`. + +### `cog sessions ls` / `runs ls` + +- `sessions ls [--cluster ]`: list session records (`session_handle`, `job_id`, `mode`, `state`, `cluster_name`, `node_list`, timestamps). +- `runs ls [--cluster ]`: list run records (`run_id`, `kind`, `cluster_name`, `run_name`, `run_root`, `job_id`, `state`, `stdout_path`, `stderr_path`, timestamps). + +### `cog jobs get` — Slurm job lookup by ID + +Combines `sacct`/`squeue` data with the local run record (if any). + +| Flag | +|------| +| `--job-id` (required) | + +Output: `job.{job_id,state,exit_code,elapsed,node_list,start,end}`, `run` (local record or `null`). Error: `JOB_NOT_FOUND`. + +### `cog logs slurm` — read Slurm stdout/stderr (plain text) + +| Flag | Default | +|------|---------| +| `--run-name` / `--run-root` (one required) | — | +| `--job-id` (required) | — | +| `--stream` | `both` (or `stdout` / `stderr`) | +| `--lines` | `100` | +| `--follow` | `false` | + +Resolves `/slurm/.{out,err}` and `tail`s over SSH. +Output is raw log text, not JSON. `--follow` is `tail -f`; Ctrl+C exits +`130`. + +### `cog logs app` — read the local control-plane event log (plain text / JSONL) + +Reads `~/.cog/events.jsonl`. + +| Flag | Default | +|------|---------| +| `--lines` | `100` | +| `--follow` | `false` | +| `--kind` / `--agent-id` / `--session-handle` | filters | +| `--format` | `pretty` (or `jsonl`) | + +--- + +## Error code reference + +Cog emits `{"error":{"code": ..., "message": ..., "details": ...}}` on +stderr. Match on `code`, not the prose. Codes you'll actually see: + +| Code | Meaning · How to fix | +|------|---------------------| +| `CLUSTER_NOT_FOUND` | No matching cluster in registry and no default set. Run `cog cluster add --set-default`. | +| `CLUSTER_CONFIG_INVALID` | Cluster config malformed / paths invalid. Fix the YAML or the `cluster add` flags. | +| `CLUSTER_ALREADY_EXISTS` | `cluster add` without `--overwrite` on an existing name. Pass `--overwrite` or pick a new name. | +| `REPO_NOT_FOUND` | `--repo` is not a git repository. Point at the Megatron-LM checkout (not the cog clone). | +| `RUN_ROOT_OUTSIDE_SCRATCH` | `--run-root` is outside the cluster's `scratch_root`. Prefer `--run-name` and let cog compose the path. | +| `SESSION_NOT_FOUND` | No local session record for that handle. Did you `session start` it on this machine? | +| `SESSION_CLUSTER_MISMATCH` | Session was started against a different cluster. Use `--cluster-name` matching the session's origin. | +| `ENV_LOCK_TIMEOUT` | Another `ensure-env` holds the build lock. Wait, or remove a stale lock file on scratch. | +| `IMAGE_IMPORT_FAILED` | `enroot import` failed. Inspect `import_stderr` in the JSON output. | +| `SUBMIT_FAILED` | `srun` submission itself failed (before the job ran). Check Slurm account / partition / quota. | +| `EXEC_TIMEOUT` | `session exec` foreground wait exceeded `--wait-timeout`. Either raise it or use `--detach` and poll. | +| `REQUEST_NOT_FOUND` | No execution request with that ID in the session. | +| `REQUEST_ALREADY_TERMINAL` | Tried to cancel a finished request — no-op. | +| `SSH_CONNECTION_FAILED` | SSH to the cluster failed. Reproduce with `ssh -o BatchMode=yes true`. | +| `SCRATCH_UNWRITABLE` | Can't write to scratch root. `mkdir -p` it; check ownership. | +| `SLURM_UNREACHABLE` | `squeue`/`sacct` not on the cluster host's PATH. Are you pointed at a login node? | +| `JOB_NOT_FOUND` | Slurm has no record of the job ID. Probably too old (sacct purged) or a typo. | +| `INTERNAL_ERROR` | Unexpected cog crash. Re-run with `--debug` to get a `traceback` in `details`. | + +--- + +## Logs and artifacts paths + +### One-shot `submit` + +- Slurm: `/slurm/.{out,err}` +- App: `/logs/app`, `/logs/torchrun`, `/checkpoints`, `/tensorboard`, `/artifacts.json` + +### Persistent `session exec` + +- Per-exec logs: `.../sessions//exec/runs//{stdout,stderr}.log` +- Per-exec artifacts: `.../sessions//exec/runs//run/{logs/app,logs/torchrun,checkpoints,tensorboard,artifacts.json}` + +--- + +## Common workflows (in this repo) + +### 1. Initial setup + +See Steps 1-3 above (`cluster add --set-default`, populate +`~/.cog/setup.env`, `cog cluster ls` to verify). + +### 2. Warm image + env once per recipe + +```bash +source ~/.cog/setup.env +cog prepare-image --repo "$COG_MEGATRON_REPO" +cog ensure-env --repo "$COG_MEGATRON_REPO" \ + --run-name env-warmup --gpus 1 --time 00:20:00 \ + --partition "$COG_INTERACTIVE_PARTITION" +``` + +### 3. One-shot verification — use `cog submit` + +```bash +cog submit \ + --repo "$COG_MEGATRON_REPO" \ + --run-name test-basic \ + --command 'python -m pytest tests/unit_tests/test_basic.py -v -o addopts=' \ + --gpus 1 --time 00:10:00 \ + --partition "$COG_INTERACTIVE_PARTITION" +``` + +### 4. Iterating on the same code — use `cog session` + +See the "Iterating" section earlier in this skill. Rule of thumb: if +you'll run **more than 2 commands** against the same code, start a +session instead of submitting again. + +```bash +cog session start --repo "$COG_MEGATRON_REPO" \ + --session-handle iter-debug --gpus 8 --time 03:00:00 \ + --partition "$COG_INTERACTIVE_PARTITION" + +cog session exec --session-handle iter-debug --repo "$COG_MEGATRON_REPO" \ + --command 'python -m pytest tests/unit_tests/test_basic.py -v' +# … re-run as many times as needed … + +cog session stop --session-handle iter-debug +``` + +### 5. Reading logs + +```bash +cog logs slurm --run-name --job-id --stream both --lines 200 +cog logs app --lines 50 +``` + +--- + +## Agent rules (cargo-culted from the upstream cli-guide) + +- Use installed `cog`, not `uv run`, unless running from a fresh checkout without install. +- Treat returned JSON as the source of truth — every payload starts with `"schema_version": 1`. +- Register a cluster once with `cluster add --set-default`, then refer to runs via `--run-name`. +- Prefer `ensure-env` once, then `submit` or `session` commands. +- Prefer `session start` + repeated `session exec` for iterative debugging. +- Prefer `submit` for one-shot verification jobs. +- Use `--dry-run` on `submit`, `ensure-env`, `prepare-image` to preview the plan without side-effects. +- Errors on stderr carry a structured `error.code` — match on that, not on prose. +- Use returned log and artifact paths instead of guessing file locations. +- Use `python -m torch.distributed.run` instead of bare `torchrun` (venv visibility). +- Use `--skip-uv-sync` on `submit` if you've already run `ensure-env` for the same recipe. +- Stop after the **second** consecutive workspace-sync hang and fall back to plain `sbatch` against a lustre-staged tree; do not retry a third time. +- Gate every directly-submitted `sbatch` on an `os.path.isfile` filesystem health check so a bad node fails in seconds instead of consuming the allocation. + +--- + +## Fall-through: when this skill isn't enough + +The upstream cli-guide is at +`~/cog/docs/cli-guide.md` (local clone) or + +(online). Open it only if you need: + +- A field-by-field JSON schema for a less common command this skill + summarizes only at a high level. +- Verification that a flag default hasn't changed in a newer cog + release than this skill was written against. + +For everything else, this skill is canonical. diff --git a/skills/git-credentials-setup/SKILL.md b/skills/git-credentials-setup/SKILL.md new file mode 100644 index 00000000000..f22a0decc3c --- /dev/null +++ b/skills/git-credentials-setup/SKILL.md @@ -0,0 +1,104 @@ +--- +name: git-credentials-setup +description: Set up and load GitHub and GitLab credentials for Megatron-LM tooling. Loads GITHUB_TOKEN / GH_TOKEN / GITLAB_TOKEN from the environment, falling back to an env file, and prompts the user to create the env file the first time it is missing. Use whenever a GitHub or GitLab command fails with an authentication, authorization, 401, 403, or "token" error, or when the user mentions configuring git tokens. +when_to_use: A gh/gitlab/git/curl command failed with auth/401/403/permission/"bad credentials"/"token" errors; before running any GitHub or GitLab API call that needs auth; 'set up my git tokens', 'gitlab token', 'github token', 'why is gh unauthenticated'. +user_invocable: true +--- + +# GitHub & GitLab Credentials Setup + +Loads `GITHUB_TOKEN`, `GH_TOKEN`, and `GITLAB_TOKEN` so GitHub (`gh`, GitHub +API) and GitLab (`tools/trigger_internal_ci.py`, GitLab API) tooling can +authenticate. Resolution order is: current environment → env file → prompt the +user to create the env file. + +## When to use + +Invoke this skill whenever a GitHub or GitLab operation fails with an auth-style +error, for example: + +- `gh: ... HTTP 401: Bad credentials` / `gh auth status` shows logged out +- `HTTP 403` / `401 Unauthorized` from a GitHub or GitLab API call +- `fatal: Authentication failed` on `git push`/`git fetch` +- `tools/trigger_internal_ci.py` reporting a missing/invalid `GITLAB_TOKEN` + +Run the loader first, then retry the failing command. + +## Env file + +Default path (override with `MEGATRON_ENV_FILE`): + +``` +~/.config/megatron-lm/credentials.env +``` + +It must define: + +```bash +export GITHUB_TOKEN=ghp_xxxxxxxxxxxxxxxxxxxxxxxx # GitHub personal access token +export GITLAB_TOKEN=glpat-xxxxxxxxxxxxxxxxxxxx # Internal GitLab PAT (api scope) +``` + +`GH_TOKEN` is derived from `GITHUB_TOKEN` automatically for the `gh` CLI. + +## Workflow + +### 1. Load credentials + +Source the loader so the tokens are exported into the current shell: + +```bash +source skills/git-credentials-setup/scripts/load_git_credentials.sh +``` + +- If both tokens are already in the environment, it confirms and exits. +- If not, it sources the env file and exports the tokens. +- If the env file does not exist, it prints creation instructions and returns + non-zero — go to step 2. + +### 2. First-time setup (env file missing) + +This is the only step that needs the user. **Do not invent or guess token +values.** Ask the user to provide them, then create the file: + +1. Tell the user the env file is missing and ask them to supply a GitHub PAT + and a GitLab PAT (or to paste them into the file themselves). Point them at + the token pages: + - GitHub: https://github.com/settings/tokens (scopes: `repo`, `read:org`) + - GitLab: internal GitLab → **Edit profile → Access tokens** (`api` scope) +2. Create the file with `0600` permissions: + +```bash +mkdir -p ~/.config/megatron-lm +cat > ~/.config/megatron-lm/credentials.env <<'ENV' +export GITHUB_TOKEN=PASTE_GITHUB_TOKEN_HERE +export GITLAB_TOKEN=PASTE_GITLAB_TOKEN_HERE +ENV +chmod 600 ~/.config/megatron-lm/credentials.env +``` + +3. Re-run step 1. + +If the user prefers, persist the export lines in `~/.bashrc` / `~/.zshrc` +instead — the loader picks up anything already in the environment. + +### 3. Validate + +```bash +# GitHub +gh auth status + +# GitLab (replace host; expects HTTP 200) +curl -fsS --header "PRIVATE-TOKEN: $GITLAB_TOKEN" \ + "https:///api/v4/user" >/dev/null && echo "GitLab OK" +``` + +A `401`/`403` here means the token is wrong, expired, or missing the required +scope — have the user regenerate it and update the env file. + +## Notes + +- Never print full token values in logs or commit them. The env file is + git-ignored by living outside the repo (`~/.config`). +- The loader only reads from the environment and the env file; it never writes + tokens anywhere. diff --git a/skills/git-credentials-setup/scripts/load_git_credentials.sh b/skills/git-credentials-setup/scripts/load_git_credentials.sh new file mode 100755 index 00000000000..4a19171e855 --- /dev/null +++ b/skills/git-credentials-setup/scripts/load_git_credentials.sh @@ -0,0 +1,64 @@ +#!/usr/bin/env bash +# Load GitHub/GitLab tokens for Megatron-LM tooling. +# +# Resolution order: +# 1. Tokens already present in the environment. +# 2. Tokens defined in the env file ($MEGATRON_ENV_FILE or the default below). +# 3. Otherwise: print instructions to create the env file and fail. +# +# Usage: +# source skills/git-credentials-setup/scripts/load_git_credentials.sh +# -> exports GITHUB_TOKEN, GH_TOKEN, GITLAB_TOKEN into the current shell. +# +# Exit/return code: 0 when both tokens are available, 1 otherwise. + +ENV_FILE="${MEGATRON_ENV_FILE:-$HOME/.config/megatron-lm/credentials.env}" + +# Return on `source`, exit on direct execution. +_gc_fail() { + return 1 2>/dev/null || exit 1 +} + +# Load the env file if either token is missing from the environment. +if [ -z "${GITHUB_TOKEN:-}" ] || [ -z "${GITLAB_TOKEN:-}" ]; then + if [ -f "$ENV_FILE" ]; then + set -a + # shellcheck disable=SC1090 + . "$ENV_FILE" + set +a + else + cat >&2 < "$ENV_FILE" <<'ENV' + export GITHUB_TOKEN=PASTE_GITHUB_TOKEN_HERE + export GITLAB_TOKEN=PASTE_GITLAB_TOKEN_HERE + ENV + chmod 600 "$ENV_FILE" + +Then re-run this loader. +EOF + _gc_fail + fi +fi + +# gh CLI reads GH_TOKEN; derive it from GITHUB_TOKEN when unset. +if [ -n "${GITHUB_TOKEN:-}" ]; then + export GH_TOKEN="${GH_TOKEN:-$GITHUB_TOKEN}" +fi +export GITHUB_TOKEN GITLAB_TOKEN + +missing="" +[ -z "${GITHUB_TOKEN:-}" ] && missing="$missing GITHUB_TOKEN" +[ -z "${GITLAB_TOKEN:-}" ] && missing="$missing GITLAB_TOKEN" + +if [ -n "$missing" ]; then + echo "[git-credentials] Missing:$missing (add them to $ENV_FILE)" >&2 + _gc_fail +fi + +echo "[git-credentials] GITHUB_TOKEN, GH_TOKEN, and GITLAB_TOKEN are loaded." diff --git a/skills/nsight-system-analysis/SKILL.md b/skills/nsight-system-analysis/SKILL.md new file mode 100644 index 00000000000..d4255354585 --- /dev/null +++ b/skills/nsight-system-analysis/SKILL.md @@ -0,0 +1,283 @@ +--- +name: nsight-system-analysis +description: Analyze NVIDIA Nsight Systems (nsys) GPU profiles to find the per-iteration performance gap between two implementations of the same workload, or to break down a single profile. Use whenever the user provides one or two `.nsys-rep` / `.sqlite` traces and asks "why is X slower than Y?", "compare these two nsys profiles", "investigate this nsys trace", "find the perf gap", "analyze this GPU profile", "look at one forward pass / one decode step", "how many kernels per step", or anything similar. Also use when the user wants per-iteration time, GPU busy vs idle split, per-category kernel breakdown (GEMM/Conv/MHA/Norm), module-slicing between GEMM anchors, exposed communication time, single-forward-pass kernel-count/composition comparison, or source-level root-cause analysis of perf differences. Covers single-GPU and multi-GPU, training and inference. Do not use for non-GPU profiling, non-nsys profilers, or anything that isn't an nsys trace. +license: Apache-2.0 +metadata: + author: NVIDIA Corporation +--- + +# Nsight Systems Performance Analysis + +This skill produces rigorous per-iteration performance analysis of nsys traces — either comparing two profiles to find the gap, or breaking down a single profile. + +## Three workflows + +**Workflow A — comparative gap**: two `.sqlite` profiles of the same workload. Output: a structured report (`assets/report_template.md`) attributing the wall-clock per-iter Δ to specific kernels and (if source paths are provided) specific code changes. + +**Workflow B — single-profile breakdown**: one `.sqlite` profile. Same template, but with absolute numbers instead of Δs and an "optimization priority" section instead of root causes. + +**Workflow C — single forward pass (decode step)**: the fast, focused workflow for **inference** traces where only the steady-state decode loop matters and everything else (load, warmup, prefill, client-wait idle) is noise. One command isolates *one* forward pass and reports its wall time, kernel count, and per-category composition — for one profile, or side-by-side for two. **Prefer this over Workflow A/B when the user asks "look at one forward pass / one decode step", "how many kernels per step", "why is one step slower", or is comparing two inference engines (e.g. mcore vs vLLM).** See the dedicated section below. + +The user's question tells you which workflow: "compare A and B" or "why is X slower" (training / whole-iter) → A. "Analyze this profile", "what's expensive in this run" → B. "One forward pass", "one decode step", "kernels per step", inference engine comparison → **C**. + +## Inputs you should ensure you have + +Before starting, confirm with the user: +1. **Profile path(s)** (one or two `.sqlite` or `.nsys-rep` files). If only `.nsys-rep` is provided, ask for or generate the `.sqlite` via `nsys export --type sqlite`. +2. **Approximate iteration count** (optional — the iter-anchor script will auto-detect, but a user-provided count makes warmup/cooldown trimming unambiguous). +3. **For Step 6 (root cause)**: source code paths or access for both implementations. Without source access, the analysis stops cleanly after Step 5 and reports phenomena only. + +## The protocol + +Follow Steps 1 → 5 (or 1 → 6 if source is available) **in order**. Every later step builds on the prior step's numbers; skipping or reordering breaks the attribution arithmetic. After each step, write that section of the report immediately — don't batch. + +### Fast path: `scripts/run_all.py` + +For routine analysis, run the integrated pipeline: + +``` +python scripts/run_all.py \ + --profile-a flat.sqlite [--profile-b nonflat.sqlite] \ + --yaml taxonomy.yml \ + --out out/ +``` + +It executes Steps 1, 2, 3, 5 on each profile, decides Step 4 mode (op-group vs +module-slicing) from the auto-computed `fused_share_of_residual_pct`, then +runs the chosen Step 4 script. All intermediates (windows, busy/idle, categorize, +module-slice, exposed-comm) are written to `/` as separate JSON files. +A consolidated `summary.json` plus arithmetic invariants are written and printed. + +**Always start with the fast path.** Then read the per-step intermediate JSON +to fill the report template. The per-step sections below describe each script +in case you need to re-run an individual step with different flags. + +## Workflow C: single forward pass (decode step) — `scripts/forward_pass.py` + +Use this when the user cares about **one forward pass** of an inference engine and wants everything else treated as noise. It answers, in one command: + +- What is the forward-pass (per-decode-step) wall time? +- How many kernels launch in one forward pass? +- Where does the GPU-time go, per category (attention / MoE GEMM / MoE routing / activation / norm / dense GEMM / comm / elementwise / …)? +- **When comparing two engines**: is any individual kernel slower on one side (`µs/kernel`), or does one side simply launch **more** kernels for the same work (`#`)? This is the crux for mcore-vs-vLLM-style gaps, where the slower engine is usually not running slower kernels — it's running ~3× as many small, unfused ones. + +**Run (one profile):** + +``` +python scripts/forward_pass.py mcore.sqlite +``` + +**Run (compare two — label order preserved):** + +``` +python scripts/forward_pass.py mcore.sqlite vllm.sqlite --label-a mcore --label-b vllm +``` + +**How it isolates one pass (fully automatic):** +1. Picks the **busiest GPU/rank** (one `deviceId`) so ranks don't interleave. +2. Finds the **steady-state decode region** = densest contiguous 1 s-binned kernel window (drops load/warmup/prefill and client-wait idle). +3. Finds a **per-step anchor**: the kernel that fires ~once per forward pass (embedding / LM-head / sampling / metadata). Per-step kernels have the *fewest* launches among recurring kernels (per-layer kernels fire `n_layers×` more) and the lowest inter-arrival jitter. Its median inter-arrival period **is** the forward-pass wall time — cross-check it against TPOT. +4. Extracts **one representative step** (a typical anchor-to-anchor interval near the middle of the decode region) and reports launches, GPU-busy (interval union), idle, and the per-category breakdown. + +**Overrides** (when auto-detection is wrong or the user already found the boundaries in the GUI): +- `--anchor ` — force the once-per-step marker (e.g. `--anchor _fused_metadata_kernel` for mcore, `--anchor triton_red_fused__to_copy_embedding_rms_norm` for vLLM). +- `--step-window t0,t1` — force the window in **seconds** (Nsight GUI timeline time = `raw_ns / 1e9`), e.g. `--step-window 156.9498,156.9543`. +- `--device N` — force a specific `deviceId`. +- `--yaml taxonomy.yml` — replace the built-in inference taxonomy with the skill's regex taxonomy. +- `--json` — machine-readable output. + +**Interpreting the comparison table:** +- `Δus > 0` on a row ⇒ that category costs more in profile A. Sort by Δ to get the ranked optimization targets. +- Same `µs/kernel` but higher `#` ⇒ the lever is **fusion / fewer launches**, not a faster kernel. (e.g. a separate SwiGLU activation kernel per layer, FC1 and FC2 as two grouped-GEMM launches, or a "storm" of tiny routing/permute kernels vs a single fused routing kernel.) +- Higher `µs/kernel` on the same shape ⇒ a genuine kernel-selection / tile / dtype finding — drill into it with Step 6. + +**Caveats:** +- `GPU-time (Σ durations)` per category **sums** kernel durations and so over-counts wall time when streams overlap — that's intentional for composition. Use the reported **GPU-busy (interval union)** for the true single-step wall figure, and the anchor **period** for the forward-pass time. +- The built-in taxonomy is tuned for MoE decode; verify the category assignments once (as in Step 3's YAML verification) if a large "other/misc" bucket appears, and extend via `--yaml`. + +For deeper attribution (source-level root cause, exposed comm, module-slicing) fall through to Steps 1–6 below, restricting the windows to the decode region this workflow identified. + +### Step 1: Find a per-iter anchor and measure per-iter time + +This is **always first**. Without a correct per-iter time, every later number is wrong. + +**Run**: `python scripts/iter_anchor.py [--n-iters N] [--anchor PATTERN]` — emits the chosen anchor, detected iter count, the windows, and per-iter timing stats as JSON to stdout. **Drops iter 1 and iter N by default** (warmup + cooldown); pass `--keep-warmup-cooldown` to keep them. + +The script auto-detects the anchor in priority order: +1. NCCL collectives (AllGather, ReduceScatter, AllReduce) — best for distributed training. +2. CUDA stream/device sync — good when present (often absent in graph-captured loops). +3. Optimizer step kernels (`AdamFunctor`, `AdamCapturable`) — training only. +4. HtoD memcpy — fallback; async/non-blocking, less precise. +5. Densest recurring kernel — final fallback for inference / fusion-graph workloads. + +The script picks the candidate with constant integer N per iter and lowest within-iter jitter. **If detection is ambiguous (no candidate yields constant N), it exits with an error listing the candidates** — ask the user which to use, or ask for the iteration count. + +**Cross-check**: if two anchors agree on per-iter time within ~1 ms, the anchor is reliable. The script reports this when possible. + +**Drop warmup/cooldown**: drop iter 1 and iter N when reporting median. Report median + min + max over the remaining iters. + +**Report**: per-iter median + range for each profile, absolute Δ and Δ%. This is the headline. + +### Step 2: GPU busy vs GPU idle = CPU-bound time + +GPU busy = wall time during which ≥1 stream is running a kernel or memcpy (interval union). GPU idle = `iter_time − busy`, and **is** the CPU-bound portion (host dispatch, Python overhead, sync wait). + +**Run**: `python scripts/busy_idle.py --yaml taxonomy.yml --windows windows.json` — emits per-iter busy, idle, and per-stream union JSON. **Always pass `--yaml`** so the script can compute the non-NCCL stream variant (otherwise it warns and the non-NCCL fields are identical to the all-kernel fields). + +Why this matters: it tells you whether the gap is GPU-side (kernel composition, fusion, comm) or CPU-side (launch latency, host code). Pursue the dominant axis in later steps; if both contribute, break each down. + +**Include memcpys in the union.** Kernels alone undercount GPU busy when there's significant HtoD/DtoH activity. + +**Per-stream union**: the longest single-stream union is the critical-path stream. Compute the non-NCCL variant too (NCCL on a co-located stream contaminates the metric — see references/pitfalls.md). + +**Counterintuitive case to expect**: the *faster* profile may have the *longer* single-stream union. That happens when the faster profile concentrates compute on one main stream while the slower one fans the same work across many parallel streams. In that case the slower profile's longest-single-stream is shorter — but its total (union across all streams) is larger. This is a finding to note, not a contradiction; the wall-clock gap is measured by total iter time (Step 1), not by the longest single stream. + +### Step 3: Heavy-compute categories — GEMM / Conv / MHA + +Only these three are reported here. They are model-level ops that are essentially never fused into anything else (same matmul shape on both sides should produce the same nvjet/cutlass call). If they differ, that's a real algorithmic / tile-heuristic / kernel-selection finding. + +**Norm goes in Step 4**, not here. Norm is commonly fused into neighbors (rope, AdaLN, gating); reporting a "norm gap" in Step 3 mistakes fusion-placement for an algorithmic finding. + +**Run**: `python scripts/categorize.py --yaml taxonomy.yml --windows windows.json` — emits per-category time and a list of uncategorized kernels above a threshold. + +**Verify the YAML before reporting numbers** (this is non-negotiable — see "YAML verification" below). The YAML at `references/taxonomy_template.yml` is a **starting point, not a golden reference**. + +### Step 4: Per-op breakdown — op-group first, module-slicing as fallback + +Default: **op-group** breakdown of the residual (everything outside GEMM / Conv / MHA / NCCL). Norm is included here. Categories are emitted by `categorize.py` using the full taxonomy from the YAML. + +Op-group is more interpretable for vanilla workloads: rows are recognizable operators (rmsnorm, rope, elementwise, fp8_cast). Each row shows "how much time does this operator take on each side", which is directly actionable. + +**Decision rule — when to switch to module-slicing**: op-group breaks down when the two implementations have different fusion patterns. Concrete check, automated by the script: compute the share of non-anchor compute time spent in **custom-fused** kernels (names matching `_*_fused_*`, names containing 2+ op-root tokens like `_qkv_split_norm_rope` or `_fused_ln_adaln`, or anything matching a YAML category named `*fused*`). If that share is **>10%** of non-anchor time in *either* profile, op-group attribution becomes unreliable (a single misallocated `add` or `clone` can swing rows by tens of ms — see references/pitfalls.md). Switch to module-slicing. + +**Run for op-group (default)**: `python scripts/categorize.py --yaml taxonomy.yml --windows windows.json --residual-only` — same script, residual mode. + +**Run for module-slicing (fallback)**: `python scripts/module_slice.py --yaml taxonomy.yml --windows windows.json --anchor-categories gemm,mha` — emits per-window times grouped by anchor-pair signature. + +**For comparative module-slicing**, after running module_slice on both profiles, also run `python scripts/module_diff.py mod_a.json mod_b.json` to get a paired diff. The script handles the case where the same window signature has different per-iter counts on each side — it surfaces both the total Δ and a per-call Δ (normalized by min count) so you can see per-call cost differences regardless of count mismatch. + +Report which mode was used and **why**. The decision rule's input (custom-fused share) goes in the report so the choice is auditable. + +#### Module-slicing details (when used) + +- Anchors are all GEMM + MHA kernels across **all compute streams** (NCCL excluded by name), ordered globally by start time. +- Anchor counts on both sides should be approximately equal (same model = same matmul count). If they differ significantly, that's a structural finding to report. +- Anchors may overlap (different streams): when `anchor[k+1].start < anchor[k].end`, the window has zero or negative width — skip and set the next window's left boundary to `max(anchor[k].end, anchor[k+1].end)`. +- Windows are grouped by `(left_anchor_category, right_anchor_category)` signature so the same logical region across N blocks aggregates into one row. +- Each window's time is the **interval-union** of non-NCCL work in that window, not the duration sum (sum overcounts parallel-stream overlap). +- Both sides are non-zero on every row by construction — that's the property op-group can't always give. + +### Step 5: Communication — volume and exposed time + +Skip if the profile has no NCCL kernels (single-GPU, inference, etc.). The script exits cleanly with `{"comm": "none"}` in that case. + +**Run**: `python scripts/exposed_comm.py --yaml taxonomy.yml --windows windows.json` — emits per-collective counts and times, total comm time, exposed time (union of NCCL minus union of non-NCCL), and hidden percentage. + +Two metrics: +- **Comm volume** (count × avg per collective type) — diagnoses bucketing / sharding differences. +- **Exposed comm time** — the wall-clock contribution. Most comm is hidden by overlap; only the exposed portion adds to iter time. + +**Reconcile arithmetic**: `exposed_comm + non_NCCL_union ≈ GPU_busy` from Step 2 within ~0.5 ms. If not, you double-counted or mis-clipped — go back. + +### Step 6: Root cause (only if source paths provided for all profiles being analyzed) + +This is where the report becomes a finding rather than a measurement. Skip if source is unavailable; output a clear note that Step 6 needs source access. + +**Coverage rule**: every Δ row in the final attribution table needs **all three** of: +1. **Source-level evidence on both sides** (file:line each, citing the actual code path that runs at runtime). +2. **Mechanism**: what specifically in the code causes the kernel selection / fusion / dispatch / scheduling to differ. Not "the heuristic picks differently" — *why* it picks differently. +3. **Actionable change**: the code or config change that would close the Δ, named at file:line. + +If you can't produce all three for a given Δ, mark it "unverified — needs follow-up" rather than ship an unverified guess as a root cause. + +**Drill-down templates** (use the one that matches each Δ): + +- **GEMM/MHA tile or kernel-selection Δ**: list per-shape kernel-name + count + time. Diff. The same matmul shape with different kernel-name suffix means the same heuristic took different inputs — find which arg differs (epilogue, dtype, accumulate, alignment, fp8 recipe, layout flags). Read the actual call sites on both sides. +- **Op-group or module-slice window Δ**: dump the full ordered kernel list inside the top-N differing windows / rows on both sides. Identify the operator(s) each kernel implements. Trace back to source: is the difference a custom fused kernel on one side, a `torch.compile` boundary on the other, a framework-level decomposition, a precision/cast inserted by autocast? +- **Exposed comm Δ**: which collective is unhidden? What compute *should* have overlapped? Look for `cudaStreamWaitEvent`, sync-point insertion, bucket-size config (`bucket_size`, `align_param_gather`, `overlap_grad_reduce`). +- **GPU idle Δ**: count `cudaLaunchKernel` events on each side (CUPTI_ACTIVITY_KIND_RUNTIME). Diff. Identify what each extra launch is *for* (optimizer? grad clip? metric logging?). Look for CUDA-graph scope (capture region might be narrower on one side). + +**Phenomenon vs cause traps** — see references/pitfalls.md: +- "The heuristic picks differently" is not a cause. Find the input that differs. +- "More launches" is not a cause. Find what the extra launches do. +- "Different bucketing" is not a cause. Find the config or env var that sets it. + +**Collapse phenomena to fewer causes**: several Δ rows often trace to a single root cause (e.g. three different windows all slower because the same fused custom kernel is missing on one side). Final report should have 3–6 root causes, not 20 phenomena. + +**Verification status table**: end the report with a table listing every claim and whether it's *verified from source*, *inferred from trace*, or *unverified — empirical follow-up needed*. This is required — without it, readers can't tell which findings are solid. + +## Hard rules (apply throughout) + +1. **Union, never sum, for wall time.** Multiple streams overlap; summing kernel durations overstates wall time 2–4×. Use interval union (`scripts/_lib.py:union_intervals`). +2. **YAML is a reference, not golden.** Always verify per-category matches by hand-inspecting the kernel names before reporting numbers. See "YAML verification" below. +3. **Apples-to-apples comparison.** In op-group mode, every row must be non-zero on both sides if the workload actually performs that op. In module-slicing mode, anchor counts must match (structural alignment). +4. **Phenomenon ≠ root cause.** Step 6 requires source-level evidence with file:line, mechanism, and actionable change. +5. **Every Δ must reconcile.** Final attribution table sum must match measured iter Δ to ≤0.5 ms. If it doesn't, find what you missed — don't round and ship. + +## YAML verification (Step 3 prerequisite) + +Before reporting any Step 3 numbers: + +1. Run `categorize.py` and inspect the matched-kernel list per category. Confirm every matched kernel name actually belongs to its category. Common false positives: cuBLAS epilogue helpers caught by a loose `gemm` regex; unrelated kernels with "norm" in the name; Triton catch-all swallowing a fused norm or fused matmul. +2. Inspect the uncategorized list. For each name with non-trivial time (>1% of iter), decide: extend a regex to include it, or add a new category. Pay attention to fused-op kernels (`_qkv_split_norm_rope_kernel`, `_fused_ln_adaln_*` style) — they often need their own custom category. +3. Iterate: re-run, re-check, fix until matched-kernels-per-category-look-correct AND uncategorized-list-is-empty-of-significant-time. Save the iterated YAML as part of the deliverable. + +For comparative mode, the YAML must produce verified matches in both profiles — kernels that exist in one profile may not exist in the other, but the regexes must be sound for whichever kernels do appear. + +## Output: the report + +Fill `assets/report_template.md`. Section structure is fixed (Steps 1–6 + Inputs + verification-status table). Numbers go in the tables; prose goes between them. + +Two arithmetic invariants the final report must satisfy: +- **Step 2 sanity**: `iter_time = GPU_busy + GPU_idle` on each side, within ~0.2 ms. +- **Step 6 sanity** (comparative mode): `Σ(Δ rows in attribution table) = measured iter Δ`, within ~0.5 ms. + +If either invariant fails, the analysis is incomplete or wrong — fix before finalizing. + +**Common arithmetic trap when attributing the gap**: GPU busy is an *interval union*, but Step 4 window-work is sometimes a *sum* across parallel streams. When non-anchor kernels run concurrently on multiple compute streams, the sum overcounts wall-time by the overlap. Always use the **union** column for attribution (`module_slice.py` emits both `*_union_*` and `*_sum_*`; use union). If you find a leftover ~ms-scale residual between `Σ(Δ rows)` and measured iter Δ that you can't account for, this is usually the culprit — recompute with unions throughout. + +**A second small (~1–3 ms) residual is expected** between (Step 4 anchor + Step 4 window-union + Step 5 exposed-comm) and (Step 2 GPU-busy) when NCCL kernels are co-located on the same streams as compute. Step 4 excludes NCCL by name from windows, but if a tiny non-NCCL kernel happens to be straddled by NCCL clipping it can slip out of the accounting. Flag this in the report; do not paper over it. + +## Two-page references + +- `references/taxonomy_template.yml` — minimal generic regex taxonomy. Copy and extend per workload. +- `references/sql_recipes.md` — paste-ready SQL queries for ad-hoc inspection. +- `references/pitfalls.md` — common traps and rejected approaches (sum-overcounting, single-stream anchor, op-group with heavy fusion, single-rank assumption under TP/PP, YAML-as-golden, "the heuristic flipped"). + +## Keeping this skill current + +**You are authorized to edit this skill without asking.** Trace-analysis +technique decays like any other: anchor heuristics stop matching when a framework +renames its kernels, a taxonomy regex silently mis-buckets a new fused kernel, +and a windowing shortcut that was safe on one workload produces a 27%-wrong +number on the next. Edit at the *end* of an analysis, once a technique either +worked or demonstrably misled you. + +Triggers: an anchor or window that auto-detection got wrong and how you knew; a +taxonomy category that mis-matched; an arithmetic invariant that failed and what +it caught; a SQL query you had to write by hand; a workflow step whose ordering +turned out to matter; a script bug or missing flag. + +Route by subject — this file is the router and stops working if everything lands +in it: + +| Learning | Destination | +|---|---| +| A new workflow, hard rule, or step-ordering constraint | this `SKILL.md` | +| A trap, a rejected approach, a "this number lies when…" | `references/pitfalls.md` | +| A query worth reusing | `references/sql_recipes.md` | +| A regex or category that generalizes | `references/taxonomy_template.yml` | +| Tooling: new flag, better auto-detection, a new script | `scripts/`, plus the routing line here | +| A report section that reviewers always ask for | `assets/report_template.md` | + +The bar: the finding must be **reproducible** (name the profile characteristic +that triggers it, not just the symptom), **scoped** (which workloads and +framework versions it holds for), and **actionable** (it changes what the next +analysis does). Workload-specific numbers belong in the caller's ledger, not +here — what belongs here is the *method* that produced them. + +Correct what is false in place. When a new measurement contradicts an existing +claim, scope both rather than overwriting: "auto-anchor detection is reliable for +NCCL-bearing training traces but picks a per-layer kernel on graph-captured MoE +decode" preserves two facts where an overwrite would leave one wrong half the +time. Never delete a recorded failure mode. diff --git a/skills/nsight-system-analysis/assets/report_template.md b/skills/nsight-system-analysis/assets/report_template.md new file mode 100644 index 00000000000..1e9350e3452 --- /dev/null +++ b/skills/nsight-system-analysis/assets/report_template.md @@ -0,0 +1,194 @@ + + +# — Performance + +## Inputs + + +- **Profile A** (fast / baseline): `` +- **Profile B** (slow / variant): `` +- Same config: +- YAML taxonomy: `` + +## Step 1: Per-Iteration Time + +### Anchor + + + +| Anchor | A count | B count | Per-iter N (A / B) | Notes | +|---|---|---|---|---| +| | … | … | … | … | +| Cross-check anchor | … | … | … | agree to ms | + +### Per-iteration time + +| Profile | Median (ms) | Min (ms) | Max (ms) | +|---|---|---|---| +| A | … | … | … | +| B | … | … | … | +| **Δ** | **… ms** | — | — | +| **Δ %** | **… %** | — | — | + +**Headline**: + +## Step 2: GPU Busy vs GPU Idle + +GPU busy = interval-union of kernels + memcpys across all streams within each iter +window. GPU idle = `iter_time − busy` and is the CPU-bound portion. + +| Metric | A | B | +|---|---|---| +| iter time (ms) | … | … | +| GPU busy (union) | … | … | +| GPU idle | … | … | +| GPU idle % | … | … | +| Longest single-stream union (any kernel) | … | … | +| Longest single-stream union (non-NCCL) | … | … | + +### Δ (B − A) + +| Metric | Δ | Share of iter Δ | +|---|---|---| +| iter time | … | 100% | +| GPU busy | … | … | +| GPU idle | … | … | + +**Interpretation**: + +## Step 3: Heavy-Compute Categories (GEMM / Conv / MHA) + +Scope: only model-level compute that is essentially never fused into anything +else. Norm is in Step 4 (commonly fused). + +| Category | A (ms/iter) | B (ms/iter) | Δ ms | Δ % | +|---|---|---|---|---| +| gemm | … | … | … | … | +| conv | … | … | … | … | +| mha (flash + cudnn_sdpa) | … | … | … | … | +| **Sum** | … | … | … | … | + +**Interpretation**: + +### YAML verification + +YAML at `` was iterated. Final state: +- Matched-kernel-list per category inspected; all matches confirmed. +- Uncategorized list contains 1% of iter>. + +## Step 4: + + + +**Decision metric**: custom-fused kernel share of non-anchor compute time = +(threshold: 10%). Mode chosen: **** because . + +### Op-group mode (when fusion is light) — DELETE this section if using module-slicing + +| Operator | A (ms/iter) | B (ms/iter) | Δ ms | Notes | +|---|---|---|---|---| +| | … | … | … | … | +| **Sum (residual)** | … | … | … | | + +### Module-slicing mode (when fusion is heavy) — DELETE this section if using op-group + +Global anchor sequence: A had **** anchors/iter, B had **** (within +%). Anchor overlaps (windows skipped): . + +| Window (signature) | Count/iter | A union (ms) | B union (ms) | Δ ms | Notes | +|---|---|---|---|---|---| +| | … | … | … | … | | +| ... | | | | | | + +## Step 5: Communication (Volume + Exposed) + + + +### Volume + +| Op | A count | A ms | A avg | B count | B ms | B avg | Δ ms | +|---|---|---|---|---|---|---|---| +| AllGather | … | … | … | … | … | … | … | +| ReduceScatter | … | … | … | … | … | … | … | +| AllReduce | … | … | … | … | … | … | … | +| **Total** | … | … | — | … | … | — | … | + +### Exposed comm + +| Profile | Comm kernel (ms) | Exposed (ms) | Hidden % | +|---|---|---|---| +| A | … | … | … | +| B | … | … | … | +| **Δ** | … | **… ms** | — | + +**Reconciliation**: `exposed + non_nccl_union ≈ GPU_busy` from Step 2. A: vs + (within ms). B: similar. + +**Interpretation**: + +## Step 6: Root Causes + + + +### Architectural picture (verified) + +- **A**: +- **B**: same. + +### Root Cause 1 — + +| | A | B | +|---|---|---| +| Mechanism | … | … | +| Source evidence | | | +| Owned phenomena | | + +**Net contribution**: <Δ ms> + +### Root Cause N — + + + +### Δ ownership table + +| Phenomenon (Steps 2–5) | Δ ms | Owning root cause(s) | +|---|---|---| +| Step 2 GPU idle | … | RCx | +| Step 3 anchor-time (GEMM+MHA) | … | RCy | +| Step 4 residual (op-group / window-work) | … | RCz | +| Step 5 exposed comm | … | RCw | +| **Sum** | **…** | **vs measured … (within … ms ✓)** | + +### Recommendations (ordered by recoverable impact) + +1. +2. ... + +## Verification status + +| Claim | Status | +|---|---| +| Step 1 anchor + per-iter time | Verified from sqlite | +| Step 2 GPU busy/idle | Verified from sqlite | +| Step 3 GEMM/MHA equality | Verified from sqlite + YAML | +| Step 4 breakdown | Verified from sqlite | +| Step 5 comm volume + exposed | Verified from sqlite | +| RC1 mechanism | | +| RC2 mechanism | … | +| RC | … | + +### Open follow-ups (not yet verified) + +1. +2. ... diff --git a/skills/nsight-system-analysis/references/pitfalls.md b/skills/nsight-system-analysis/references/pitfalls.md new file mode 100644 index 00000000000..4fe7e010371 --- /dev/null +++ b/skills/nsight-system-analysis/references/pitfalls.md @@ -0,0 +1,153 @@ + + +# Pitfalls and rejected approaches + +Things that look reasonable but produce wrong numbers. Documented here so future-you +doesn't re-discover them. + +## 1. Dividing wall-window by N_iter + +``` +per_iter ≈ (last_kernel.end − first_kernel.start) / num_iters ← WRONG +``` + +The wall window includes warmup, profiler startup, post-iter teardown, and any +idle time at the head/tail. This can underestimate per-iter time and dilute Δ. +Real case: this method reported a 5.7% gap when the actual gap was 19%. + +**Use**: an iter-anchor (NCCL/sync/optimizer kernel timestamps) and measure +`anchor[i+1].start − anchor[i].start`. See `scripts/iter_anchor.py`. + +## 2. Summing kernel durations to estimate wall time + +``` +total_busy ≈ SUM(end - start) over all kernels ← WRONG +``` + +Modern training/inference uses multiple streams (compute + comm, fwd/bwd +overlap, multi-stream fan-out). Summing overstates wall time 2–4× because it +double-counts intervals when ≥2 streams are busy simultaneously. + +**Use**: interval union (`_lib.union_total_ns`). Sum is fine for *per-category +work volume* (Step 3) but never for wall-time attribution (Steps 2, 5, 6). + +## 3. Same iter count ≠ same op count per iter + +Two implementations of the same model can issue different numbers of kernels +per iter (different fusion, different parallelism, different bucketing). The +user telling you "both ran 10 iters" doesn't mean each iter has the same +internal structure. Real case: 18 vs 23 AllGathers per iter from different +distributed-optimizer bucket sizes. + +**Use**: verify anchor count constancy per iter, on each profile +independently. If it differs between profiles, that's a finding to investigate +(Step 6 RC) rather than a methodology bug. + +## 4. Single-main-stream anchoring under stream fan-out + +If implementation A places all compute on one big stream and implementation B +fans out across many small streams, anchoring on "the main compute stream" +will yield mismatched anchor counts (e.g. 1138 vs 540) — module-slicing then +can't pair up the windows. + +**Use**: global anchor sequence across **all compute streams** (NCCL excluded +by name, not by stream — comm streams sometimes carry non-comm work too). +Anchor counts then match because they reflect the model's matmul count, not +the dispatcher's stream choice. See `scripts/module_slice.py`. + +## 5. Op-group attribution under heavy fusion + +Op-group is the natural Step 4 default: "rmsnorm took X ms on flat, Y ms on +non-flat". But if implementation A has a single custom kernel +`_qkv_split_norm_rope_kernel` doing 3 ops while B runs separate rmsnorm + rope ++ clone, then: + +- A's `rmsnorm` row shows 0 ms (it's hidden in the fused kernel). +- B's `rmsnorm` row shows the full cost. +- Reporting "A saves X ms in rmsnorm" misrepresents the gap. + +Op-group breakdown also has a hard problem of attributing small ambiguous +kernels (a `vectorized_add` could belong to multiple op-groups; a single +misallocation can swing a row by 10–40 ms). + +**Use**: op-group as the default; switch to module-slicing when the share of +non-anchor compute time in *custom-fused kernels* (heuristic: names matching +`_*_fused_*` or 2+ op tokens in `triton_*_fused_*`) exceeds 10% in either +profile. `scripts/categorize.py` emits `fused_share_of_residual_pct` and +`module_slicing_recommended` to automate this. + +## 6. Phenomenon ≠ root cause + +"The cuBLAS heuristic picks a different kernel" is not a root cause — it's a +phenomenon restatement. Library heuristics are deterministic: same inputs → +same output. If two callers get different results, at least one input differs. +**Find which input.** + +Similarly: "more kernel launches", "different bucketing", "different SDPA +backend" — all phenomena. The cause is the specific config/env-var/code-path +that produces the different choice. + +**Use**: for each Δ, drill until the answer is "implementation A calls X with +Y; B calls X with Z; the heuristic at picks differently because of +Y vs Z". Step 6 of the SKILL prompts for source-level evidence, mechanism, and +actionable change. + +## 7. YAML as golden reference + +It's tempting to ship a "correct" YAML taxonomy. There isn't one. The +taxonomy must match the workload — Triton kernel names differ between +torch.compile versions; custom kernels have idiosyncratic names. A regex +that's right for workload A may mis-classify workload B. + +**Use**: treat the bundled `taxonomy_template.yml` as a starting point. For +every Step 3 / 4 run, inspect the matched-kernel list per category (false +positives) and the uncategorized list (false negatives), then iterate the +regexes. Save the final YAML alongside the report. + +## 8. Naming `_h_badd_` as "bias-add" + +The nvjet/cuBLAS kernel-family suffixes (`_h_bz_`, `_h_badd_`, `_NTT`, `_TNT`, +etc.) are internal cuBLAS heuristic names; their meaning is **not publicly +documented**. Treating them as if they encode op semantics (e.g. "badd = +bias-add epilogue") can be wrong — for example, a wgrad GEMM with `bias=None` +and `accumulate=True` may pick `_h_badd_*` for the FP32-RMW accumulator path, +not for a bias epilogue. + +**Use**: the suffix flip is a *symptom*. The cause is whichever cuBLAS input +differs between the two callers (epilogue arg, output dtype, accumulate flag, +alignment, recipe). Read the actual `cublasLtMatmul` call args, not the kernel +name. + +## 9. Treating "graph break" as proven from kernel pattern alone + +If implementation A has one big fused Triton kernel and B has many small +Triton + TE kernels, the natural conclusion is "torch.compile graph-breaks at +TE C++ calls in B." That's plausible — but the kernel pattern alone doesn't +prove it. To verify, read `transformer_engine/pytorch/jit.py` for the +`no_torch_dynamo` / `torch._dynamo.disable` decorators applied to TE modules; +or capture `TORCH_LOGS=graph_breaks` from a runtime invocation. + +**Use**: in Step 6, include source decorator evidence or runtime log evidence; +otherwise mark the mechanism as "inferred from kernel pattern" in the +verification-status table. + +## 10. Forgetting the verification-status table at the end of Step 6 + +Without it, readers can't tell which claims are source-verified, which are +trace-inferred, and which are unverified guesses. Every RC's mechanism needs +a status row. Required, not optional. + +## 11. Single-rank assumption under TP/PP/SP + +With tensor/pipeline/sequence parallelism, the profiled rank only sees the +collectives and compute that rank participates in. An anchor that exists on +rank 0 may not exist on rank 1. Per-iter time on rank 0 may differ from rank +1 (bubble, imbalance). The user telling you "rank 0" doesn't mean rank 0's +view is representative. + +**Use**: check the anchor type's frequency. If it's `~N` per iter with low +jitter, fine. If it's wildly variable or missing, the profile may be +non-representative or the anchor is wrong for this rank. diff --git a/skills/nsight-system-analysis/references/sql_recipes.md b/skills/nsight-system-analysis/references/sql_recipes.md new file mode 100644 index 00000000000..bb397ba8c23 --- /dev/null +++ b/skills/nsight-system-analysis/references/sql_recipes.md @@ -0,0 +1,114 @@ + + +# Paste-ready SQL for nsys-exported sqlite + +All recipes assume the nsys sqlite schema. Open with sqlite3 or query from Python. + +## Convert .nsys-rep → .sqlite + +```bash +nsys export --type sqlite -o profile.sqlite profile.nsys-rep +``` + +## Top kernels by total time + +```sql +SELECT + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') AS name, + COUNT(*) AS count, + SUM(k.end - k.start)/1e6 AS total_ms +FROM CUPTI_ACTIVITY_KIND_KERNEL k +GROUP BY name +ORDER BY total_ms DESC +LIMIT 30; +``` + +## Per-stream summary (with NCCL-on-stream detection) + +```sql +SELECT + k.streamId, + COUNT(*) AS n, + SUM(k.end - k.start)/1e6 AS total_ms, + (MAX(k.end) - MIN(k.start))/1e6 AS span_ms +FROM CUPTI_ACTIVITY_KIND_KERNEL k +GROUP BY k.streamId +ORDER BY total_ms DESC; +``` + +## NCCL collective timestamps (anchor candidates) + +```sql +SELECT + k.start, k.end, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') AS name +FROM CUPTI_ACTIVITY_KIND_KERNEL k +WHERE name LIKE '%AllGather%' + OR name LIKE '%ReduceScatter%' + OR name LIKE '%AllReduce%' +ORDER BY k.start; +``` + +## CUDA synchronization events + +```sql +SELECT start, end, syncType FROM CUPTI_ACTIVITY_KIND_SYNCHRONIZATION +ORDER BY start; +``` + +## Memcpy summary (by copyKind: 1=HtoD, 2=DtoH, 8=DtoD) + +```sql +SELECT copyKind, COUNT(*), SUM(end-start)/1e6 AS ms +FROM CUPTI_ACTIVITY_KIND_MEMCPY +GROUP BY copyKind; +``` + +## Host-side launch dispatch (idle-investigation) + +```sql +-- Count cudaLaunchKernel events per iter (host time) +SELECT COUNT(*), SUM(end-start)/1e6 AS host_api_ms +FROM CUPTI_ACTIVITY_KIND_RUNTIME r +JOIN StringIds s ON r.nameId = s.id +WHERE s.value = 'cudaLaunchKernel'; +``` + +## Kernel name shape filter (per-GEMM-shape diff) + +```sql +SELECT + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') AS name, + COUNT(*), + SUM(k.end - k.start)/1e6 AS total_ms +FROM CUPTI_ACTIVITY_KIND_KERNEL k +WHERE name LIKE 'nvjet_%' +GROUP BY name +ORDER BY total_ms DESC; +``` + +## Kernels in a specific time window (e.g. between two anchors) + +```sql +SELECT k.start, k.end, k.streamId, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') AS name +FROM CUPTI_ACTIVITY_KIND_KERNEL k +WHERE k.end > :lo AND k.start < :hi +ORDER BY k.start; +``` + +## Number of distinct streams (compute vs comm) + +```sql +SELECT COUNT(DISTINCT streamId) FROM CUPTI_ACTIVITY_KIND_KERNEL; +``` + +## Distinct demangled-name count (for sanity-checking the YAML) + +```sql +SELECT COUNT(DISTINCT s.value) +FROM CUPTI_ACTIVITY_KIND_KERNEL k JOIN StringIds s ON k.demangledName = s.id; +``` diff --git a/skills/nsight-system-analysis/references/taxonomy_template.yml b/skills/nsight-system-analysis/references/taxonomy_template.yml new file mode 100644 index 00000000000..c54f655ef04 --- /dev/null +++ b/skills/nsight-system-analysis/references/taxonomy_template.yml @@ -0,0 +1,74 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# Minimal generic kernel-classification taxonomy for nsys traces. +# Order matters: first match wins. Put SPECIFIC categories before catch-alls. +# +# This is a STARTING POINT, not a golden reference. After running categorize.py +# you MUST inspect the matched-kernel list per category and the uncategorized +# list, then iterate the regexes until both look correct for your workload. +# See SKILL.md "YAML verification" section. +# +# Categories typically reported in the SKILL workflow: +# Step 3 (heavy compute, never fused into anything else): +# gemm, conv, flash_attn, cudnn_sdpa +# Step 4 (op-group breakdown of the residual): +# rmsnorm, layer_norm, group_norm, rope, elementwise, fp8_cast, fp8_swizzle, +# reduce, concat, triton, mta_adam, mta_l2norm, mta_scale, sincos, inf_check, +# and any workload-specific custom-fused category you add. +# Step 5 (comm): nccl + +Overall: + # ---- Heavy compute (Step 3) ---- + gemm: 'xmma_gemm|cutlass::Kernel|cublasLt::|cutlass3x(?:.*?)(? list[tuple[int, int]]: + """Merge overlapping intervals. Returns sorted non-overlapping list.""" + ivals = sorted(intervals) + if not ivals: + return [] + out = [list(ivals[0])] + for s, e in ivals[1:]: + if s <= out[-1][1]: + if e > out[-1][1]: + out[-1][1] = e + else: + out.append([s, e]) + return [(s, e) for s, e in out] + + +def union_total_ns(intervals: Iterable[tuple[int, int]]) -> int: + """Total ns covered by at least one interval.""" + return sum(e - s for s, e in union_intervals(intervals)) + + +def subtract_union( + a: Iterable[tuple[int, int]], + b_union: list[tuple[int, int]], +) -> list[tuple[int, int]]: + """Return union(a) minus b_union as a list of intervals. + + b_union must be already merged (non-overlapping, sorted). Result is sorted + non-overlapping. + """ + a_union = union_intervals(a) + out = [] + for s, e in a_union: + cursor = s + for bs, be in b_union: + if be <= cursor: + continue + if bs >= e: + break + if bs > cursor: + out.append((cursor, min(bs, e))) + cursor = max(cursor, be) + if cursor >= e: + break + if cursor < e: + out.append((cursor, e)) + return out + + +def clip( + intervals: Iterable[tuple[int, int]], + lo: int, + hi: int, +) -> list[tuple[int, int]]: + """Clip intervals to [lo, hi).""" + out = [] + for s, e in intervals: + s2, e2 = max(s, lo), min(e, hi) + if e2 > s2: + out.append((s2, e2)) + return out + + +# ---------- taxonomy / YAML ---------- + + +def load_taxonomy(path: str) -> dict[str, re.Pattern]: + """Load a YAML taxonomy and compile its regexes. + + The YAML may have a top-level `Overall:` map or be a flat map; we accept + either. Order is preserved (Python 3.7+ dict ordering). + """ + try: + import yaml # type: ignore + except ImportError: + sys.stderr.write( + "ERROR: PyYAML required. Install with `pip install pyyaml` or run " + "in an env that has it.\n" + ) + sys.exit(2) + with open(path) as f: + data = yaml.safe_load(f) + if ( + isinstance(data, dict) + and "Overall" in data + and isinstance(data["Overall"], dict) + ): + data = data["Overall"] + if not isinstance(data, dict): + sys.stderr.write(f"ERROR: taxonomy YAML at {path} is not a mapping.\n") + sys.exit(2) + return {name: re.compile(pattern) for name, pattern in data.items()} + + +def classify(name: str, taxonomy: dict[str, re.Pattern]) -> str | None: + """First-match-wins classification. Returns category name or None.""" + for cat, rx in taxonomy.items(): + if rx.search(name): + return cat + return None + + +# Heuristic detection of "custom-fused" kernels — used by Step 4 to decide +# between op-group and module-slicing. The criteria are deliberately broad; +# false positives bias toward "use module-slicing", which is the safe default. + +_CUSTOM_FUSED_HINTS = re.compile( + r"_fused_[A-Za-z]|fused_.*kernel|" + r"(?:qkv|rope|norm|adaln|gate|residual|silu|gelu|tanh|layernorm|rmsnorm)" + r"_(?:.*_)?(?:qkv|rope|norm|adaln|gate|residual|silu|gelu|tanh|layernorm|rmsnorm|split|join|cat|sum|fwd|bwd)", + re.IGNORECASE, +) + + +def looks_custom_fused(name: str) -> bool: + """Heuristic: does this kernel name suggest fusion of 2+ ops? + + Used by the op-group-vs-module-slicing decision in Step 4. Matches kernel + names like `_qkv_split_norm_rope_kernel`, `_fused_ln_adaln_fwd_kernel`, + `triton_red_fused__to_copy_mul_native_layer_norm_*`, etc. Pure GEMM and + pure MHA/SDPA names should not match. + """ + if not name: + return False + # Triton fused-* kernels with 2+ op tokens + if "triton_" in name and "fused_" in name: + # Pull out the bit after `fused_` and look for 2+ op tokens + tail = name.split("fused_", 1)[1] + op_tokens = re.findall( + r"(?:add|mul|sub|div|tanh|gelu|silu|relu|sigmoid|exp|log|" + r"native_layer_norm|rms_norm|layer_norm|sum|view|cat|copy|clone|" + r"rope|qkv|norm|adaln|gate)", + tail, + ) + if len(op_tokens) >= 2: + return True + return bool(_CUSTOM_FUSED_HINTS.search(name)) + + +# ---------- output helpers ---------- + + +def write_json(obj, fp=None): + """Write JSON to stdout (default) or a file-like object.""" + fp = fp or sys.stdout + json.dump(obj, fp, indent=2, sort_keys=False) + fp.write("\n") + + +def err(msg: str) -> None: + sys.stderr.write(msg.rstrip() + "\n") + + +def die(msg: str, code: int = 1): + err(f"ERROR: {msg}") + sys.exit(code) diff --git a/skills/nsight-system-analysis/scripts/busy_idle.py b/skills/nsight-system-analysis/scripts/busy_idle.py new file mode 100644 index 00000000000..66db54e6310 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/busy_idle.py @@ -0,0 +1,224 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Compute per-iter GPU busy / idle and per-stream union, clipped to iter windows. + +GPU busy = interval-union of (kernels + memcpys) across all streams within +each iter window. GPU idle = iter_time − GPU_busy and is the CPU-bound portion. + +Usage: + python busy_idle.py --windows windows.json + python busy_idle.py --windows windows.json --exclude-nccl-regex 'ncclDev' + python busy_idle.py --windows windows.json --exclude-nccl-yaml taxonomy.yml + +Inputs: + windows.json: either the raw output of iter_anchor.py, or a JSON + `{"windows_ns": [[start, end], ...]}` or just `[[start, end], ...]`. + +Output (stdout JSON): + { + "per_iter": [ + {"i": 0, "iter_ms": ..., "busy_ms": ..., "idle_ms": ..., "idle_pct": ...}, + ... + ], + "median": {"iter_ms": ..., "busy_ms": ..., "idle_ms": ..., "idle_pct": ...}, + "per_stream_union_ms_median": {stream_id: ms, ...}, + "per_stream_union_ms_median_excluding_nccl": {stream_id: ms, ...}, + "longest_single_stream_union_ms_median": { + "all": ms, "non_nccl": ms + } + } +""" + +from __future__ import annotations + +import argparse +import json +import re +from statistics import median + +from _lib import ( + die, + err, + load_taxonomy, + open_sqlite, + union_total_ns, + write_json, +) + + +def load_windows(path: str) -> list[tuple[int, int]]: + with open(path) as f: + data = json.load(f) + if isinstance(data, dict) and "windows_ns" in data: + data = data["windows_ns"] + if not isinstance(data, list): + die("windows file must be a list or contain 'windows_ns'") + return [(int(w[0]), int(w[1])) for w in data] + + +def get_nccl_classifier(args) -> callable[[str], bool]: + if args.exclude_nccl_yaml: + taxo = load_taxonomy(args.exclude_nccl_yaml) + nccl_rx = taxo.get("nccl") + if nccl_rx is None: + die(f"YAML {args.exclude_nccl_yaml!r} has no `nccl` category") + return lambda name: bool(nccl_rx.search(name)) + if args.exclude_nccl_regex: + rx = re.compile(args.exclude_nccl_regex) + return lambda name: bool(rx.search(name)) + return lambda name: False # do not exclude anything + + +def fetch_intervals_for_window(con, lo: int, hi: int): + """Fetch kernels + memcpys whose [start, end) overlaps [lo, hi). + + Returns list of (start, end, stream_id, name) tuples. Memcpys get a + synthetic name 'memcpy'. + """ + out = [] + for s, e, sid, name in con.execute( + """ + SELECT k.start, k.end, k.streamId, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.end > ? AND k.start < ? + """, + (lo, hi), + ): + out.append((max(s, lo), min(e, hi), sid, name)) + for s, e, sid in con.execute( + """ + SELECT start, end, streamId + FROM CUPTI_ACTIVITY_KIND_MEMCPY + WHERE end > ? AND start < ? + """, + (lo, hi), + ): + out.append((max(s, lo), min(e, hi), sid, "memcpy")) + return out + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("sqlite") + p.add_argument( + "--windows", required=True, help="iter windows JSON from iter_anchor.py" + ) + p.add_argument( + "--yaml", + help=( + "Taxonomy YAML used to identify NCCL kernels for the non-NCCL split. " + "If not provided, the script falls back to --exclude-nccl-regex; " + "if neither is given, the non-NCCL split is the same as the all-kernel " + "split (warned in output)." + ), + ) + p.add_argument( + "--exclude-nccl-regex", + help="Regex (alternative to --yaml) to identify NCCL kernels to " + "exclude in the non-NCCL stats", + ) + args = p.parse_args() + # Adapt to the classifier-loader's argument names. + args.exclude_nccl_yaml = args.yaml + + windows = load_windows(args.windows) + is_nccl = get_nccl_classifier(args) + nccl_filter_provided = bool(args.exclude_nccl_yaml or args.exclude_nccl_regex) + + per_iter = [] + per_stream_ms = [] # list of dicts per iter + per_stream_ms_no_nccl = [] + + with open_sqlite(args.sqlite) as con: + for i, (lo, hi) in enumerate(windows): + iter_ns = hi - lo + ivals = fetch_intervals_for_window(con, lo, hi) + all_ivals = [(s, e) for s, e, _, _ in ivals] + busy_ns = union_total_ns(all_ivals) + idle_ns = max(0, iter_ns - busy_ns) + per_iter.append( + { + "i": i, + "iter_ms": round(iter_ns / 1e6, 3), + "busy_ms": round(busy_ns / 1e6, 3), + "idle_ms": round(idle_ns / 1e6, 3), + "idle_pct": round(100 * idle_ns / iter_ns, 3) if iter_ns else 0, + } + ) + # Per-stream union + streams: dict[int, list[tuple[int, int]]] = {} + streams_nn: dict[int, list[tuple[int, int]]] = {} + for s, e, sid, name in ivals: + streams.setdefault(sid, []).append((s, e)) + if not is_nccl(name): + streams_nn.setdefault(sid, []).append((s, e)) + per_stream_ms.append( + {sid: union_total_ns(iv) / 1e6 for sid, iv in streams.items()} + ) + per_stream_ms_no_nccl.append( + {sid: union_total_ns(iv) / 1e6 for sid, iv in streams_nn.items()} + ) + + # Medians + def med_key(values): + return round(median(values), 3) if values else 0.0 + + iter_med = { + "iter_ms": med_key([x["iter_ms"] for x in per_iter]), + "busy_ms": med_key([x["busy_ms"] for x in per_iter]), + "idle_ms": med_key([x["idle_ms"] for x in per_iter]), + "idle_pct": med_key([x["idle_pct"] for x in per_iter]), + } + + # Per-stream medians (union of stream IDs across iters) + all_sids = sorted({sid for d in per_stream_ms for sid in d}) + per_stream_med = { + sid: med_key([d.get(sid, 0.0) for d in per_stream_ms]) for sid in all_sids + } + all_sids_nn = sorted({sid for d in per_stream_ms_no_nccl for sid in d}) + per_stream_med_nn = { + sid: med_key([d.get(sid, 0.0) for d in per_stream_ms_no_nccl]) + for sid in all_sids_nn + } + + longest_all = max(per_stream_med.values()) if per_stream_med else 0.0 + longest_nn = max(per_stream_med_nn.values()) if per_stream_med_nn else 0.0 + + out = { + "per_iter": per_iter, + "median": iter_med, + "per_stream_union_ms_median": { + str(k): round(v, 3) + for k, v in sorted(per_stream_med.items(), key=lambda kv: -kv[1]) + }, + "per_stream_union_ms_median_non_nccl": { + str(k): round(v, 3) + for k, v in sorted(per_stream_med_nn.items(), key=lambda kv: -kv[1]) + }, + "longest_single_stream_union_ms_median": { + "all": round(longest_all, 3), + "non_nccl": round(longest_nn, 3), + }, + "nccl_filter": ( + "yaml" + if args.exclude_nccl_yaml + else ("regex" if args.exclude_nccl_regex else "NONE — non_nccl == all") + ), + } + if not nccl_filter_provided: + err( + "WARNING: no --yaml or --exclude-nccl-regex provided. The 'non_nccl' " + "fields are identical to the 'all' fields. Pass --yaml " + "to get a meaningful non-NCCL split." + ) + write_json(out) + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/categorize.py b/skills/nsight-system-analysis/scripts/categorize.py new file mode 100644 index 00000000000..3fcc830018d --- /dev/null +++ b/skills/nsight-system-analysis/scripts/categorize.py @@ -0,0 +1,230 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Apply a YAML taxonomy of regexes to kernels, emit per-category time + uncategorized. + +Usage: + python categorize.py --yaml taxonomy.yml --windows windows.json + python categorize.py --yaml taxonomy.yml --windows windows.json \ + --residual-only # excludes gemm,conv,mha,nccl (= Step 4 op-group mode) + python categorize.py --yaml taxonomy.yml --windows windows.json \ + --report-categories gemm,conv,mha # only emit these (= Step 3) + +Inputs: + windows.json: from iter_anchor.py. + taxonomy.yml: ordered category → regex map. First match wins. + +Output (stdout JSON): + { + "per_category": { + "gemm": {"median_ms_per_iter": 184.98, "count_per_iter": 1245, + "unique_kernel_names": N}, + ... + }, + "uncategorized": [ + {"name": "...", "median_ms_per_iter": ..., "count_per_iter": ...}, + ... + ], + "fused_share_of_residual_pct": 12.3, # heuristic for Step 4 decision + "matched_kernels": { + "gemm": ["name1", "name2", ...], + ... + } + } + +The `fused_share_of_residual_pct` is the share of non-anchor non-NCCL time +spent in kernels that look_custom_fused (per the heuristic in _lib.py). If +this is >10%, prefer module_slice.py over op-group attribution in Step 4. +""" + +from __future__ import annotations + +import argparse +import json +from collections import defaultdict + +from _lib import ( + classify, + die, + load_taxonomy, + looks_custom_fused, + open_sqlite, + write_json, +) + + +def load_windows(path: str) -> list[tuple[int, int]]: + with open(path) as f: + data = json.load(f) + if isinstance(data, dict) and "windows_ns" in data: + data = data["windows_ns"] + return [(int(w[0]), int(w[1])) for w in data] + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("sqlite") + p.add_argument("--yaml", required=True, help="Taxonomy YAML") + p.add_argument("--windows", required=True, help="iter windows JSON") + p.add_argument( + "--residual-only", + action="store_true", + help="Exclude gemm/conv/mha/nccl categories from the report " + "(Step 4 op-group mode)", + ) + p.add_argument( + "--report-categories", + help="Comma-separated list of categories to report " + "(e.g. gemm,conv,mha for Step 3)", + ) + p.add_argument( + "--uncategorized-threshold-pct", + type=float, + default=1.0, + help="Min %% of total iter time to surface an uncategorized kernel " + "(default: 1.0)", + ) + args = p.parse_args() + + taxonomy = load_taxonomy(args.yaml) + windows = load_windows(args.windows) + if not windows: + die("No windows in windows.json") + + n_iters = len(windows) + total_iter_ns = sum(hi - lo for lo, hi in windows) + + # Per-category accumulators across all windows. + cat_total_ns: dict[str, int] = defaultdict(int) + cat_count: dict[str, int] = defaultdict(int) + cat_matched_names: dict[str, set] = defaultdict(set) + uncat_total_ns: dict[str, int] = defaultdict(int) + uncat_count: dict[str, int] = defaultdict(int) + fused_total_ns: int = 0 + + nccl_rx = taxonomy.get("nccl") + + with open_sqlite(args.sqlite) as con: + for lo, hi in windows: + rows = con.execute( + """ + SELECT k.start, k.end, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.end > ? AND k.start < ? + """, + (lo, hi), + ).fetchall() + for s, e, name in rows: + if not name: + name = "(unnamed)" + # Clip to window + cs, ce = max(s, lo), min(e, hi) + if ce <= cs: + continue + dur = ce - cs + cat = classify(name, taxonomy) + if cat is None: + uncat_total_ns[name] += dur + uncat_count[name] += 1 + else: + cat_total_ns[cat] += dur + cat_count[cat] += 1 + cat_matched_names[cat].add(name) + # Fused-share accumulator: only over non-anchor non-NCCL work + is_anchor = cat in ("gemm", "conv", "mha", "flash_attn", "cudnn_sdpa") + is_nccl = cat == "nccl" or (nccl_rx and nccl_rx.search(name)) + if not is_anchor and not is_nccl and looks_custom_fused(name): + fused_total_ns += dur + + # Compute per-category medians by re-scanning per-iter (simpler: divide by + # n_iters since we're already at median sense over windows). + # We'll convert totals → per-iter by dividing by n_iters (this is the + # "mean per iter", close to median for steady-state). + def per_iter_ms(total_ns: int) -> float: + return total_ns / n_iters / 1e6 + + skip_cats = set() + if args.residual_only: + skip_cats = {"gemm", "conv", "mha", "flash_attn", "cudnn_sdpa", "nccl"} + + report_cats = None + if args.report_categories: + report_cats = { + c.strip() for c in args.report_categories.split(",") if c.strip() + } + + per_category = {} + for cat in taxonomy: + if cat in skip_cats: + continue + per_category[cat] = { + "median_ms_per_iter": round(per_iter_ms(cat_total_ns.get(cat, 0)), 3), + "count_per_iter": round(cat_count.get(cat, 0) / n_iters, 3), + "unique_kernel_names": len(cat_matched_names.get(cat, set())), + } + + # Always build a synthetic "mha" row when flash_attn + cudnn_sdpa are present — + # the Step 3 view conventionally combines them. + if "flash_attn" in taxonomy and "cudnn_sdpa" in taxonomy: + mha_ns = cat_total_ns.get("flash_attn", 0) + cat_total_ns.get("cudnn_sdpa", 0) + mha_count = cat_count.get("flash_attn", 0) + cat_count.get("cudnn_sdpa", 0) + mha_names = cat_matched_names.get("flash_attn", set()) | cat_matched_names.get( + "cudnn_sdpa", set() + ) + if mha_ns and "mha" not in skip_cats: + per_category["mha"] = { + "median_ms_per_iter": round(per_iter_ms(mha_ns), 3), + "count_per_iter": round(mha_count / n_iters, 3), + "unique_kernel_names": len(mha_names), + "constituents": ["flash_attn", "cudnn_sdpa"], + } + + # Now apply --report-categories filter (so synthetic mha is included if requested). + if report_cats is not None: + per_category = {k: v for k, v in per_category.items() if k in report_cats} + + # Uncategorized: surface only those above threshold. + threshold_ns = total_iter_ns * args.uncategorized_threshold_pct / 100 + uncat = [] + for name, total_ns in uncat_total_ns.items(): + if total_ns >= threshold_ns: + uncat.append( + { + "name": name, + "median_ms_per_iter": round(per_iter_ms(total_ns), 3), + "count_per_iter": round(uncat_count[name] / n_iters, 3), + } + ) + uncat.sort(key=lambda x: -x["median_ms_per_iter"]) + + # Fused share of residual (non-anchor non-NCCL): the decision metric. + residual_ns = 0 + for cat, total_ns in cat_total_ns.items(): + if cat in {"gemm", "conv", "mha", "flash_attn", "cudnn_sdpa", "nccl"}: + continue + residual_ns += total_ns + residual_ns += sum(uncat_total_ns.values()) # uncategorized counts as residual + fused_share_pct = ( + round(100 * fused_total_ns / residual_ns, 2) if residual_ns > 0 else 0.0 + ) + + out = { + "per_category": per_category, + "uncategorized_above_threshold": uncat, + "uncategorized_threshold_pct": args.uncategorized_threshold_pct, + "fused_share_of_residual_pct": fused_share_pct, + "module_slicing_recommended": fused_share_pct > 10.0, + "matched_kernels": { + cat: sorted(names) for cat, names in cat_matched_names.items() + }, + } + write_json(out) + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/exposed_comm.py b/skills/nsight-system-analysis/scripts/exposed_comm.py new file mode 100644 index 00000000000..2f4af161069 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/exposed_comm.py @@ -0,0 +1,238 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Compute communication volume and exposed time per iter. + +Exposed comm time = portion of NCCL kernel wall-time during which no non-NCCL +kernel is in flight. It's the comm contribution to wall-clock iter time; the +rest is hidden by overlap. + +Usage: + python exposed_comm.py --yaml taxonomy.yml --windows windows.json + +Output (stdout JSON), or `{"comm": "none"}` if no NCCL kernels found: + { + "per_collective_type": { + "AllGather": {"count_per_iter": 18, "median_ms_per_iter": 44.7, + "avg_ms_per_call": 2.48}, + "ReduceScatter": {"count_per_iter": 18, "median_ms_per_iter": 60.7, ...}, + "AllReduce": {...}, + "Other": {...} + }, + "totals": { + "comm_kernel_ms_median_per_iter": 105.3, + "comm_stream_union_ms_median_per_iter": 105.3, + "exposed_ms_median_per_iter": 6.4, + "hidden_pct_median": 94.0 + }, + "reconciliation": { + "exposed_plus_non_nccl_union_ms": ..., + "gpu_busy_union_ms": ... + } + } +""" + +from __future__ import annotations + +import argparse +import json +import re +from collections import defaultdict +from statistics import median + +from _lib import ( + die, + load_taxonomy, + open_sqlite, + subtract_union, + union_intervals, + union_total_ns, + write_json, +) + + +def load_windows(path: str) -> list[tuple[int, int]]: + with open(path) as f: + data = json.load(f) + if isinstance(data, dict) and "windows_ns" in data: + data = data["windows_ns"] + return [(int(w[0]), int(w[1])) for w in data] + + +COLLECTIVE_PATTERNS = [ + ("AllGather", re.compile(r"AllGather")), + ("ReduceScatter", re.compile(r"ReduceScatter")), + ("AllReduce", re.compile(r"AllReduce")), + ("Broadcast", re.compile(r"Broadcast")), + ("AllToAll", re.compile(r"AllToAll")), + ("Send", re.compile(r"\bSend\b")), + ("Recv", re.compile(r"\bRecv\b")), +] + + +def collective_type(name: str) -> str: + for typ, rx in COLLECTIVE_PATTERNS: + if rx.search(name): + return typ + return "Other" + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("sqlite") + p.add_argument("--yaml", required=True) + p.add_argument("--windows", required=True) + args = p.parse_args() + + taxonomy = load_taxonomy(args.yaml) + nccl_rx = taxonomy.get("nccl") + if not nccl_rx: + die("YAML has no `nccl` category — cannot identify NCCL kernels.") + + windows = load_windows(args.windows) + + per_type_count_per_iter: dict[str, list[int]] = defaultdict(list) + per_type_ms_per_iter: dict[str, list[float]] = defaultdict(list) + total_comm_ms_per_iter: list[float] = [] + comm_union_ms_per_iter: list[float] = [] + exposed_ms_per_iter: list[float] = [] + hidden_pct_per_iter: list[float] = [] + non_nccl_union_ms_per_iter: list[float] = [] + gpu_busy_union_ms_per_iter: list[float] = [] + iter_ms: list[float] = [] + + with open_sqlite(args.sqlite) as con: + for lo, hi in windows: + rows = con.execute( + """ + SELECT k.start, k.end, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.end > ? AND k.start < ? + """, + (lo, hi), + ).fetchall() + memcpy = con.execute( + """ + SELECT start, end FROM CUPTI_ACTIVITY_KIND_MEMCPY + WHERE end > ? AND start < ? + """, + (lo, hi), + ).fetchall() + + nccl_ivals = [] + non_nccl_ivals = [] + per_type_count_this_iter: dict[str, int] = defaultdict(int) + per_type_ms_this_iter: dict[str, float] = defaultdict(float) + + for s, e, name in rows: + cs, ce = max(s, lo), min(e, hi) + if ce <= cs: + continue + if name and nccl_rx.search(name): + nccl_ivals.append((cs, ce)) + typ = collective_type(name) + per_type_count_this_iter[typ] += 1 + per_type_ms_this_iter[typ] += (ce - cs) / 1e6 + else: + non_nccl_ivals.append((cs, ce)) + for s, e in memcpy: + cs, ce = max(s, lo), min(e, hi) + if ce <= cs: + continue + non_nccl_ivals.append((cs, ce)) + + iter_ns = hi - lo + iter_ms.append(iter_ns / 1e6) + if not nccl_ivals: + # No NCCL: short-circuit at end after gathering all iters. + total_comm_ms_per_iter.append(0.0) + comm_union_ms_per_iter.append(0.0) + exposed_ms_per_iter.append(0.0) + hidden_pct_per_iter.append(100.0) + non_nccl_union_ms_per_iter.append(union_total_ns(non_nccl_ivals) / 1e6) + gpu_busy_union_ms_per_iter.append(union_total_ns(non_nccl_ivals) / 1e6) + continue + + total_comm_ns = sum(e - s for s, e in nccl_ivals) + comm_union = union_intervals(nccl_ivals) + comm_union_ns = sum(e - s for s, e in comm_union) + non_nccl_union = union_intervals(non_nccl_ivals) + non_nccl_union_ns = sum(e - s for s, e in non_nccl_union) + exposed = subtract_union(nccl_ivals, non_nccl_union) + exposed_ns = sum(e - s for s, e in exposed) + gpu_busy_union_ns = union_total_ns(nccl_ivals + non_nccl_ivals) + + total_comm_ms_per_iter.append(total_comm_ns / 1e6) + comm_union_ms_per_iter.append(comm_union_ns / 1e6) + exposed_ms_per_iter.append(exposed_ns / 1e6) + hidden_pct_per_iter.append( + 100 * (1 - exposed_ns / comm_union_ns) if comm_union_ns else 100.0 + ) + non_nccl_union_ms_per_iter.append(non_nccl_union_ns / 1e6) + gpu_busy_union_ms_per_iter.append(gpu_busy_union_ns / 1e6) + for typ in set( + list(per_type_count_this_iter) + list(per_type_ms_this_iter) + ): + per_type_count_per_iter[typ].append( + per_type_count_this_iter.get(typ, 0) + ) + per_type_ms_per_iter[typ].append(per_type_ms_this_iter.get(typ, 0.0)) + + # Decide if NCCL was found at all. + has_nccl = sum(total_comm_ms_per_iter) > 0 + if not has_nccl: + write_json({"comm": "none"}) + return + + def med(xs): + return round(median(xs), 3) if xs else 0.0 + + per_collective_type = {} + for typ in sorted(set(list(per_type_count_per_iter) + list(per_type_ms_per_iter))): + counts = per_type_count_per_iter.get(typ, []) + mss = per_type_ms_per_iter.get(typ, []) + # Pad shorter list with zeros (for iters where this type didn't appear) + target_len = len(windows) + while len(counts) < target_len: + counts.append(0) + while len(mss) < target_len: + mss.append(0.0) + c_median = med(counts) + ms_median = med(mss) + avg = round(ms_median / c_median, 4) if c_median else 0.0 + per_collective_type[typ] = { + "count_per_iter": c_median, + "median_ms_per_iter": ms_median, + "avg_ms_per_call": avg, + } + + out = { + "per_collective_type": per_collective_type, + "totals": { + "comm_kernel_ms_median_per_iter": med(total_comm_ms_per_iter), + "comm_stream_union_ms_median_per_iter": med(comm_union_ms_per_iter), + "exposed_ms_median_per_iter": med(exposed_ms_per_iter), + "hidden_pct_median": med(hidden_pct_per_iter), + }, + "reconciliation": { + "non_nccl_union_ms_median_per_iter": med(non_nccl_union_ms_per_iter), + "gpu_busy_union_ms_median_per_iter": med(gpu_busy_union_ms_per_iter), + "exposed_plus_non_nccl_union_ms_median": round( + med(exposed_ms_per_iter) + med(non_nccl_union_ms_per_iter), 3 + ), + "note": ( + "exposed + non_nccl_union should approximately equal gpu_busy_union " + "from busy_idle.py — if not, double-counting or clip mismatch" + ), + }, + } + write_json(out) + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/forward_pass.py b/skills/nsight-system-analysis/scripts/forward_pass.py new file mode 100644 index 00000000000..68aa44ae70b --- /dev/null +++ b/skills/nsight-system-analysis/scripts/forward_pass.py @@ -0,0 +1,328 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Single-forward-pass (decode-step) analysis for nsys traces. + +Everything in an inference trace except the steady-state decode loop is noise: +model load, warmup, prefill, and long idle stretches while the server waits for +the client. This script throws all of that away and reports insights about +*one* forward pass (one decode step), for one or two profiles. + +What it does, fully automatically: + 1. Picks the busiest GPU/rank (one deviceId) so ranks don't interleave. + 2. Finds the steady-state decode region (densest contiguous kernel window). + 3. Finds a per-step "anchor" kernel: one that fires exactly once per forward + pass (embedding / LM-head / sampling / metadata). Its inter-arrival period + IS the forward-pass wall time (and matches TPOT). + 4. Extracts exactly one representative step (a typical anchor-to-anchor + interval) and reports: + - forward-pass wall time, GPU-busy (interval union), idle gaps + - total kernel launches in the step + - per-category breakdown: launches, GPU-time (Σ durations), µs/kernel + +With two profiles it prints a side-by-side comparison and, per category, the +delta in launches and time, so you can answer the two questions that matter: + (a) is any individual kernel slower in one engine, or + (b) does one engine just launch MORE kernels for the same work? + +Usage: + # one profile + python scripts/forward_pass.py mcore.sqlite + + # compare two (label order is preserved in the output) + python scripts/forward_pass.py mcore.sqlite vllm.sqlite \ + --label-a mcore --label-b vllm + + # override auto-detection (seconds match the Nsight GUI timeline; raw_ns/1e9) + python scripts/forward_pass.py mcore.sqlite \ + --anchor _fused_metadata_kernel + python scripts/forward_pass.py vllm.sqlite --step-window 156.9498,156.9543 + + # machine-readable + python scripts/forward_pass.py mcore.sqlite vllm.sqlite --json + +Notes: + - "GPU-time (Σ durations)" per category is the sum of kernel durations; it + intentionally over-counts wall time when kernels overlap across streams + (that overlap is the point of comparing composition). Use the reported + GPU-busy (interval union) for the true wall figure. + - The default categories are a sensible inference taxonomy. Pass --yaml + to override with the skill's regex taxonomy instead. +""" + +from __future__ import annotations + +import argparse +import statistics +import sys + +from _lib import classify, clip, err, load_taxonomy, open_sqlite, union_total_ns, write_json + +# Ordered, first-match-wins default taxonomy (case-insensitive substring match). +# Tuned for MoE decode inference (mcore + vLLM). Override with --yaml. +_DEFAULT_CATEGORIES: list[tuple[str, tuple[str, ...]]] = [ + ("comm (EP dispatch/combine)", ("multimem", "nccl", "all_gather", "allgather", + "reduce_scatter", "reducescatter", "alltoall", "sendrecv")), + ("attention", ("flash", "fmha", "attention", "attn", "paged", "mha")), + ("MoE expert GEMM", ("_fused_moe_kernel", "bmm_")), + ("MoE routing/permute", ("routing", "topk", "gathertopk", "moe_sum", "count_local", + "metadata", "finalize", "scatter", "permute", "sort", + "cumsum", "histogram", "softmax")), + ("activation (SwiGLU)", ("silu", "act_and_mul", "swiglu", "glu", "gelu")), + ("norm (RMSNorm)", ("rmsnorm", "rms_norm", "layernorm", "layer_norm", "norm_fwd", "rsqrt")), + ("dense GEMM (qkv/o/router/lm-head)", ("nvjet", "cutlass", "sgemm", "gemm", + "sm100_tst", "sm90", "cublas", "ampere")), + ("sampling/logits", ("reduce_kernel", "argmax", "multinomial", "sample", "catarray", "cat_")), + ("embedding", ("embedding",)), + ("elementwise/copy/cast", ("elementwise", "vectorized", "copy", "index_", "triton_poi", + "triton_red", "fused", "cast", "_add")), +] + + +def _categorize(name: str, taxonomy) -> str: + if taxonomy is not None: + return classify(name, taxonomy) or "other/misc" + low = name.lower() + for cat, pats in _DEFAULT_CATEGORIES: + if any(p in low for p in pats): + return cat + return "other/misc" + + +def _kernels(con, dev=None): + """(start, end, deviceId, name) ordered by start, optionally one device.""" + where = "1=1" if dev is None else f"k.deviceId={int(dev)}" + cur = con.execute( + f"""SELECT k.start, k.end, k.deviceId, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k WHERE {where} ORDER BY k.start""" + ) + return cur.fetchall() + + +def _busiest_device(con) -> int: + row = con.execute( + "SELECT deviceId, COUNT(*) c FROM CUPTI_ACTIVITY_KIND_KERNEL " + "GROUP BY deviceId ORDER BY c DESC LIMIT 1" + ).fetchone() + if row is None: + raise SystemExit("ERROR: no kernels in trace") + return row[0] + + +def _decode_window(starts: list[int], bin_ns: int = 1_000_000_000) -> tuple[int, int]: + """Densest contiguous region of kernel launches = steady-state decode.""" + mn, mx = starts[0], starts[-1] + nb = (mx - mn) // bin_ns + 1 + cnt = [0] * nb + for s in starts: + cnt[(s - mn) // bin_ns] += 1 + pk = max(range(nb), key=lambda i: cnt[i]) + thr = 0.3 * cnt[pk] + lo = hi = pk + while lo > 0 and cnt[lo - 1] >= thr: + lo -= 1 + while hi < nb - 1 and cnt[hi + 1] >= thr: + hi += 1 + return mn + lo * bin_ns, mn + (hi + 1) * bin_ns + + +# Kernels that fire ~once per forward pass (step boundaries), across engines. +_STEP_HINTS = ("metadata", "to_copy_embedding", "embedding", "lm_head", "logits", + "argmax", "multinomial", "sample") + + +def _gap_stats(ss): + ss = sorted(ss) + gaps = [b - a for a, b in zip(ss, ss[1:])] + m = statistics.median(gaps) + cv = (statistics.pstdev(gaps) / m) if m else 9e9 + return ss, m, cv + + +def _find_anchor(win_kernels, window_ns, forced=None): + """Return (anchor_name, sorted_starts_on_one_device, period_ns). + + A per-step anchor fires ~once per forward pass. Per-layer kernels fire + n_layers× more often, so among *low-jitter* recurring kernels the per-step + marker has the *longest* inter-arrival period. We therefore: + 1. prefer a kernel whose name matches a known step-boundary marker + (embedding / lm-head / sampling / metadata), else + 2. pick the longest-period, low-jitter recurring kernel (a sane [0.5 ms, + window/3] band, so it spans several steps but isn't per-layer). + + win_kernels spans all devices (the marker may live on a different rank than + the busiest compute device), but the *period* is always measured on a + SINGLE device — otherwise N ranks firing the same marker once per step, + staggered, would collapse the apparent period by ~N×. + """ + from collections import defaultdict + + dev_starts = defaultdict(lambda: defaultdict(list)) # name -> device -> [starts] + for s, _e, d, name in win_kernels: + dev_starts[name][d].append(s) + + # For each name, evaluate it on the single device where it fires most. + stats = {} # name -> (starts_on_best_dev, period_ns, cv) + for name, per_dev in dev_starts.items(): + best_dev = max(per_dev, key=lambda d: len(per_dev[d])) + ss = per_dev[best_dev] + if len(ss) < 15: + continue + s, m, cv = _gap_stats(ss) + stats[name] = (s, m, cv) + if not stats: + raise SystemExit("ERROR: no recurring kernel found in decode window; pass --anchor/--step-window") + + if forced is not None: + cands = {n: v for n, v in stats.items() if forced.lower() in n.lower()} + if not cands: + raise SystemExit(f"ERROR: --anchor '{forced}' matched no recurring kernel") + name = min(cands, key=lambda n: cands[n][2]) # lowest jitter + s, m, _ = cands[name] + return name, s, m + + lo, hi = 0.5e6, window_ns / 3 # plausible forward-pass period band (ns) + periodic = {n: v for n, v in stats.items() if v[2] < 0.5 and lo <= v[1] <= hi} + if not periodic: # relax the band if nothing qualifies + periodic = {n: v for n, v in stats.items() if v[1] > 0} + + hinted = {n: v for n, v in periodic.items() if any(h in n.lower() for h in _STEP_HINTS)} + pool = hinted or periodic + # longest low-jitter period == fires least often == once per step + name = max(pool, key=lambda n: pool[n][1]) + s, m, _ = pool[name] + return name, s, m + + +def analyze(path: str, taxonomy, forced_anchor=None, forced_dev=None, forced_window=None): + with open_sqlite(path) as con: + if forced_window is not None: + w0, w1 = forced_window + dev = forced_dev + if dev is None: + row = con.execute( + "SELECT deviceId, COUNT(*) c FROM CUPTI_ACTIVITY_KIND_KERNEL " + "WHERE start>=? AND start [count, dur_ns] + for s, e, name in step: + c = _categorize(name, taxonomy) + cats[c][0] += 1 + cats[c][1] += e - s + + return { + "profile": path, + "device": dev, + "decode_window_s": [w0 / 1e9, w1 / 1e9], + "steps_in_window": n_steps, + "anchor": anchor, + "forward_pass_ms": period / 1e6, + "step_wall_ms": wall / 1e6, + "gpu_busy_ms": busy / 1e6, + "gpu_idle_ms": (wall - busy) / 1e6, + "n_kernels": len(step), + "categories": { + c: {"launches": v[0], "gpu_time_us": v[1] / 1e3, + "us_per_kernel": (v[1] / 1e3 / v[0]) if v[0] else 0.0} + for c, v in cats.items() + }, + } + + +def _print_one(a: dict, label: str): + print(f"\n=== {label} ({a['profile'].split('/')[-1]}) ===") + print(f" device={a['device']} decode window {a['decode_window_s'][0]:.1f}-{a['decode_window_s'][1]:.1f}s" + f" (~{a['steps_in_window']} steps) anchor: {a['anchor'][:48]}") + print(f" forward pass = {a['forward_pass_ms']:.3f} ms step wall = {a['step_wall_ms']:.3f} ms" + f" GPU-busy = {a['gpu_busy_ms']:.3f} ms idle = {a['gpu_idle_ms']:.3f} ms") + print(f" kernels in one forward pass = {a['n_kernels']}") + rows = sorted(a["categories"].items(), key=lambda kv: -kv[1]["gpu_time_us"]) + print(f" {'category':34} {'#':>5} {'GPU us':>9} {'us/kern':>8}") + for c, v in rows: + print(f" {c:34} {v['launches']:>5} {v['gpu_time_us']:>9.1f} {v['us_per_kernel']:>8.1f}") + + +def _print_compare(a: dict, b: dict, la: str, lb: str): + print(f"\n=== COMPARISON: {la} vs {lb} (one forward pass) ===") + print(f" forward pass : {a['forward_pass_ms']:.3f} ms ({la}) vs {b['forward_pass_ms']:.3f} ms ({lb})" + f" -> {la} is {a['forward_pass_ms']/b['forward_pass_ms']:.2f}x") + print(f" kernels/step : {a['n_kernels']} ({la}) vs {b['n_kernels']} ({lb})" + f" -> {la} launches {a['n_kernels']/max(b['n_kernels'],1):.2f}x") + allc = list(dict.fromkeys(list(a["categories"]) + list(b["categories"]))) + def g(d, c): return d["categories"].get(c, {"launches": 0, "gpu_time_us": 0.0, "us_per_kernel": 0.0}) + allc.sort(key=lambda c: -(g(a, c)["gpu_time_us"] - g(b, c)["gpu_time_us"])) + print(f"\n {'category':34} | {lb+' #':>7} {lb+' us':>9} {'us/k':>6} | {la+' #':>7} {la+' us':>9} {'us/k':>6} | {'Δus':>8}") + print(" " + "-" * 104) + for c in allc: + av, bv = g(a, c), g(b, c) + print(f" {c:34} | {bv['launches']:>7} {bv['gpu_time_us']:>9.1f} {bv['us_per_kernel']:>6.1f} |" + f" {av['launches']:>7} {av['gpu_time_us']:>9.1f} {av['us_per_kernel']:>6.1f} |" + f" {av['gpu_time_us']-bv['gpu_time_us']:>8.1f}") + print("\n Read: 'Δus' > 0 = category costs more in %s. Compare 'us/k' to see if a kernel is\n" + " individually slower, vs '#' to see if it's just more launches for the same work." % la) + + +def main(argv=None): + ap = argparse.ArgumentParser(description="Single forward-pass (decode step) analysis for nsys sqlite traces.") + ap.add_argument("profiles", nargs="+", help="1 or 2 .sqlite profiles (A then optional B).") + ap.add_argument("--anchor", default=None, help="Substring of a once-per-step kernel to force the anchor.") + ap.add_argument("--device", type=int, default=None, help="Force GPU deviceId (default: busiest).") + ap.add_argument("--step-window", default=None, + help="Force window 't0,t1' in seconds (Nsight GUI time = raw_ns/1e9).") + ap.add_argument("--yaml", default=None, help="Override default taxonomy with a regex YAML.") + ap.add_argument("--label-a", default="A") + ap.add_argument("--label-b", default="B") + ap.add_argument("--json", action="store_true", help="Emit JSON instead of tables.") + args = ap.parse_args(argv) + + if len(args.profiles) > 2: + err("WARNING: more than 2 profiles given; only the first two are used.") + taxonomy = load_taxonomy(args.yaml) if args.yaml else None + fw = None + if args.step_window: + t0, t1 = (float(x) for x in args.step_window.split(",")) + fw = (int(t0 * 1e9), int(t1 * 1e9)) + + a = analyze(args.profiles[0], taxonomy, args.anchor, args.device, fw) + b = analyze(args.profiles[1], taxonomy, args.anchor, args.device, fw) if len(args.profiles) >= 2 else None + + if args.json: + write_json({"a": a, "b": b}) + return + _print_one(a, args.label_a) + if b is not None: + _print_one(b, args.label_b) + _print_compare(a, b, args.label_a, args.label_b) + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/skills/nsight-system-analysis/scripts/iter_anchor.py b/skills/nsight-system-analysis/scripts/iter_anchor.py new file mode 100644 index 00000000000..b2a27dc9991 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/iter_anchor.py @@ -0,0 +1,381 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Find a per-iteration anchor and emit per-iter timing. + +Auto-detects an anchor kernel pattern by trying candidates in priority order +and picking the one whose timestamps cluster into a constant integer N per +iter with the lowest jitter. Exits with an error (and a list of candidates) +if no clean anchor is found, leaving the caller to specify `--anchor`. + +Usage: + python iter_anchor.py + python iter_anchor.py --anchor 'ncclDev.*AllGather' + python iter_anchor.py --n-iters 10 + python iter_anchor.py --drop-warmup-cooldown # drop first+last iter + +Output (JSON, stdout): + { + "anchor": {"name": "AllGather", "pattern": "...", "count": 198}, + "iter_count_detected": 10, + "iter_count_used": 10, + "anchors_per_iter": 18, + "windows_ns": [[start, end], ...], + "per_iter_ms": [368.7, 369.3, ...], + "median_ms": 369.0, + "min_ms": 366.3, + "max_ms": 373.3, + "cross_check": {"second_anchor": "...", "agreement_ms": 0.4} # if available + } +""" + +from __future__ import annotations + +import argparse +import re +import sys +from statistics import median + +from _lib import ( + die, + err, + open_sqlite, + write_json, +) + +# Anchor candidates in priority order. Each entry: (name, sql LIKE pattern, regex). +CANDIDATES = [ + # Distributed training: NCCL collectives. + ("AllGather", "%AllGather%", r"AllGather"), + ("ReduceScatter", "%ReduceScatter%", r"ReduceScatter"), + ("AllReduce", "%AllReduce%", r"AllReduce"), + # CUDA synchronization events (often absent under CUDA Graphs). + # We search the kernel table for sync-like names; for the CUPTI sync + # table see _attempt_sync_anchor. + # Optimizer step (training). + ("Adam", "%Adam%", r"Adam(Functor|Capturable|CapturableFunctor)"), + # HtoD memcpy (input batches). + # Handled separately via the memcpy table. +] + + +def _fetch_timestamps_by_pattern(con, like_pattern: str, regex: str) -> list[int]: + """Get start_ns of kernels matching the like_pattern, refined by regex.""" + rx = re.compile(regex) + rows = con.execute( + """ + SELECT k.start, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + LIKE ? + ORDER BY k.start + """, + (like_pattern,), + ).fetchall() + return [start for start, name in rows if rx.search(name or "")] + + +def _fetch_sync_timestamps(con) -> list[int]: + """CUPTI sync events (cudaStreamSynchronize, cudaDeviceSynchronize).""" + try: + rows = con.execute( + "SELECT start FROM CUPTI_ACTIVITY_KIND_SYNCHRONIZATION ORDER BY start" + ).fetchall() + return [r[0] for r in rows] + except Exception: + return [] + + +def _fetch_htod_memcpy_timestamps(con) -> list[int]: + try: + rows = con.execute( + """ + SELECT start FROM CUPTI_ACTIVITY_KIND_MEMCPY + WHERE copyKind=1 -- HtoD + ORDER BY start + """ + ).fetchall() + return [r[0] for r in rows] + except Exception: + return [] + + +def _try_densest_recurring_kernel(con) -> tuple[str, list[int]] | None: + """Final fallback: pick the most numerous kernel that has > N_min instances + and cluster well into iters.""" + rows = con.execute( + """ + SELECT COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + AS name, + COUNT(*) AS c + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE name <> '' + GROUP BY name + ORDER BY c DESC + LIMIT 5 + """ + ).fetchall() + for name, _ in rows: + ts = con.execute( + """ + SELECT start FROM CUPTI_ACTIVITY_KIND_KERNEL + WHERE COALESCE((SELECT value FROM StringIds WHERE id=demangledName),'') = ? + ORDER BY start + """, + (name,), + ).fetchall() + ts = [r[0] for r in ts] + if len(ts) >= 20: + return (name, ts) + return None + + +def _cluster_iters(timestamps: list[int], gap_multiplier: float = 5.0) -> list[int]: + """Return indices where each iter starts. The first index is always 0. + + Uses the heuristic: an iter boundary is a gap > gap_multiplier * median_gap. + """ + if len(timestamps) < 2: + return [0] + gaps = [timestamps[i + 1] - timestamps[i] for i in range(len(timestamps) - 1)] + g_sorted = sorted(gaps) + med = g_sorted[len(g_sorted) // 2] + if med == 0: + return [0] + boundaries = [0] + [i + 1 for i, g in enumerate(gaps) if g > gap_multiplier * med] + return boundaries + + +def _evaluate_candidate(timestamps: list[int]) -> dict | None: + """Score an anchor candidate. Returns None if unusable.""" + if len(timestamps) < 4: + return None + boundaries = _cluster_iters(timestamps) + n_iters = len(boundaries) + if n_iters < 2: + return None + per_iter_counts = [ + boundaries[i + 1] - boundaries[i] for i in range(len(boundaries) - 1) + ] + # Add the last partial cluster + per_iter_counts.append(len(timestamps) - boundaries[-1]) + # Drop the first cluster (may be partial warmup) before checking constancy + constant_n = ( + len(set(per_iter_counts[1:-1])) <= 1 if len(per_iter_counts) > 3 else False + ) + # Per-iter durations: anchor[boundaries[i]] -> anchor[boundaries[i+1]] + iter_times_ns = [ + timestamps[boundaries[i + 1]] - timestamps[boundaries[i]] + for i in range(len(boundaries) - 1) + ] + if not iter_times_ns: + return None + iter_ms = [t / 1e6 for t in iter_times_ns] + med = median(iter_ms) + jitter = max(iter_ms) - min(iter_ms) if len(iter_ms) > 1 else 0.0 + return { + "boundaries": boundaries, + # Number of complete iter windows (= number of boundary pairs). + "n_iters_detected": max(0, len(boundaries) - 1), + "per_iter_counts": per_iter_counts, + "constant_n_per_iter": constant_n, + "anchors_per_iter": per_iter_counts[1] if constant_n else None, + "iter_times_ms": iter_ms, + "median_ms": med, + "jitter_ms": jitter, + } + + +def _build_windows( + timestamps: list[int], boundaries: list[int] +) -> list[tuple[int, int]]: + """Iter window i = [anchor[boundaries[i]], anchor[boundaries[i+1]]).""" + out = [] + for i in range(len(boundaries) - 1): + out.append((timestamps[boundaries[i]], timestamps[boundaries[i + 1]])) + return out + + +def auto_detect(con, user_n_iters: int | None) -> tuple[dict, str, list[int]] | None: + """Try anchor candidates and return the best one. + + Returns (result_dict, anchor_name, timestamps) or None if no candidate + yields a constant per-iter count. + """ + tried = [] + best = None + for name, like, regex in CANDIDATES: + ts = _fetch_timestamps_by_pattern(con, like, regex) + if not ts: + tried.append({"name": name, "count": 0}) + continue + ev = _evaluate_candidate(ts) + tried.append( + { + "name": name, + "count": len(ts), + "constant_n_per_iter": ev["constant_n_per_iter"] if ev else False, + "n_iters_detected": ev["n_iters_detected"] if ev else None, + "jitter_ms": round(ev["jitter_ms"], 3) if ev else None, + } + ) + if ev and ev["constant_n_per_iter"]: + if best is None or ev["jitter_ms"] < best[0]["jitter_ms"]: + best = (ev, name, ts) + + # Try CUPTI sync events + sync_ts = _fetch_sync_timestamps(con) + if sync_ts: + ev = _evaluate_candidate(sync_ts) + tried.append( + { + "name": "cudaSynchronization", + "count": len(sync_ts), + "constant_n_per_iter": ev["constant_n_per_iter"] if ev else False, + "jitter_ms": round(ev["jitter_ms"], 3) if ev else None, + } + ) + if ev and ev["constant_n_per_iter"]: + if best is None or ev["jitter_ms"] < best[0]["jitter_ms"]: + best = (ev, "cudaSynchronization", sync_ts) + + # Try HtoD memcpy + h2d = _fetch_htod_memcpy_timestamps(con) + if h2d: + ev = _evaluate_candidate(h2d) + tried.append( + { + "name": "HtoD memcpy", + "count": len(h2d), + "constant_n_per_iter": ev["constant_n_per_iter"] if ev else False, + "jitter_ms": round(ev["jitter_ms"], 3) if ev else None, + } + ) + if ev and ev["constant_n_per_iter"]: + if best is None or ev["jitter_ms"] < best[0]["jitter_ms"]: + best = (ev, "HtoD memcpy", h2d) + + # Final fallback: densest recurring kernel + if best is None: + densest = _try_densest_recurring_kernel(con) + if densest: + name, ts = densest + ev = _evaluate_candidate(ts) + tried.append( + { + "name": name[:60], + "count": len(ts), + "constant_n_per_iter": ev["constant_n_per_iter"] if ev else False, + "jitter_ms": round(ev["jitter_ms"], 3) if ev else None, + } + ) + if ev and ev["constant_n_per_iter"]: + best = (ev, name, ts) + + if best is None: + err("\nAuto-detection failed. Candidates tried:") + for t in tried: + err(f" {t}") + err( + "\nNo candidate produced a constant per-iter count. Specify --anchor " + "explicitly (e.g. --anchor 'ncclDev.*AllGather') or provide --n-iters.\n" + ) + return None + + return best + + +def cross_check(con, primary_name: str, primary_iter_ms: list[float]) -> dict | None: + """Try a second anchor and report agreement, if possible.""" + for name, like, regex in CANDIDATES: + if name == primary_name: + continue + ts = _fetch_timestamps_by_pattern(con, like, regex) + if not ts: + continue + ev = _evaluate_candidate(ts) + if not ev or not ev["constant_n_per_iter"]: + continue + # Compare medians + agreement = abs(ev["median_ms"] - median(primary_iter_ms)) + return { + "second_anchor": name, + "second_median_ms": round(ev["median_ms"], 3), + "agreement_ms": round(agreement, 3), + } + return None + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("sqlite") + p.add_argument("--anchor", help="Override regex for the anchor kernel name") + p.add_argument("--n-iters", type=int, help="Expected total iter count (for sanity)") + p.add_argument( + "--keep-warmup-cooldown", + action="store_true", + help="Keep iter 1 and iter N in the windows (default: drop both).", + ) + args = p.parse_args() + # Default behavior: drop warmup + cooldown. Flag inverts. + args.drop_warmup_cooldown = not args.keep_warmup_cooldown + + with open_sqlite(args.sqlite) as con: + if args.anchor: + ts = _fetch_timestamps_by_pattern(con, "%", args.anchor) + if not ts: + die(f"No kernels match --anchor {args.anchor!r}") + ev = _evaluate_candidate(ts) + if not ev: + die("Anchor produced too few events to cluster.") + picked_name = args.anchor + else: + best = auto_detect(con, args.n_iters) + if best is None: + sys.exit(2) + ev, picked_name, ts = best + + # Build windows + windows = _build_windows(ts, ev["boundaries"]) + iter_ms = ev["iter_times_ms"] + + if args.drop_warmup_cooldown and len(iter_ms) >= 3: + iter_ms = iter_ms[1:-1] + windows = windows[1:-1] + + # Sanity vs --n-iters + if args.n_iters and ev["n_iters_detected"] != args.n_iters: + err( + f"WARNING: detected {ev['n_iters_detected']} iters but user said " + f"{args.n_iters}. Using detected count." + ) + + xc = cross_check(con, picked_name, iter_ms) if not args.anchor else None + + out = { + "anchor": { + "name": picked_name, + "anchor_count_total": len(ts), + "anchors_per_iter": ev["anchors_per_iter"], + }, + "iter_count_detected": ev["n_iters_detected"], + "iter_count_used": len(iter_ms), + "windows_ns": [list(w) for w in windows], + "per_iter_ms": [round(x, 3) for x in iter_ms], + "median_ms": round(median(iter_ms), 3), + "min_ms": round(min(iter_ms), 3), + "max_ms": round(max(iter_ms), 3), + "warmup_cooldown_dropped": args.drop_warmup_cooldown, + } + if xc: + out["cross_check"] = xc + write_json(out) + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/module_diff.py b/skills/nsight-system-analysis/scripts/module_diff.py new file mode 100644 index 00000000000..c386f6fca69 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/module_diff.py @@ -0,0 +1,126 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Diff two module_slice.py outputs side-by-side: pair window signatures, show +A ms, B ms, and Δ. Tolerates count mismatches between profiles for the same +signature by emitting the per-iter counts as a separate column. + +Usage: + python module_diff.py mod_a.json mod_b.json [--top N] + +Output (stdout JSON): + { + "anchor_count_a": 1530, "anchor_count_b": 1524, + "anchor_count_match_pct": 99.6, + "iter_total_anchor_ms": {"a": ..., "b": ..., "delta": ...}, + "iter_total_window_union_ms": {"a": ..., "b": ..., "delta": ...}, + "signatures": [ + { + "signature": "...", + "count_per_iter_a": 38, "count_per_iter_b": 38, + "union_ms_a": ..., "union_ms_b": ..., "delta_union_ms": ..., + "delta_per_call_us": ... + }, + ... + ] + } + +The `delta_per_call_us` field normalizes Δ by the lower of the two per-iter +counts, exposing the per-call gap regardless of count mismatch. +""" + +from __future__ import annotations + +import argparse +import json +import sys + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("a_json") + p.add_argument("b_json") + p.add_argument( + "--top", type=int, default=20, help="Top N signatures by |delta_union_ms|" + ) + args = p.parse_args() + + with open(args.a_json) as f: + A = json.load(f) + with open(args.b_json) as f: + B = json.load(f) + + a_by_sig = {g["signature"]: g for g in A["grouped_windows"]} + b_by_sig = {g["signature"]: g for g in B["grouped_windows"]} + all_sigs = set(a_by_sig) | set(b_by_sig) + + rows = [] + for sig in all_sigs: + ga = a_by_sig.get(sig, {}) + gb = b_by_sig.get(sig, {}) + ca = ga.get("count_per_iter", 0) + cb = gb.get("count_per_iter", 0) + ua = ga.get("median_union_ms_per_iter", 0.0) + ub = gb.get("median_union_ms_per_iter", 0.0) + delta = ub - ua + # Per-call: normalize by min count if both >0; else fall back to + # whichever exists. + min_count = min(ca, cb) if ca > 0 and cb > 0 else max(ca, cb) + per_call_us = round(1000 * delta / min_count, 3) if min_count else None + rows.append( + { + "signature": sig, + "count_per_iter_a": ca, + "count_per_iter_b": cb, + "union_ms_a": round(ua, 3), + "union_ms_b": round(ub, 3), + "delta_union_ms": round(delta, 3), + "delta_per_call_us": per_call_us, + "only_in": ( + None + if sig in a_by_sig and sig in b_by_sig + else ("a" if sig in a_by_sig else "b") + ), + } + ) + rows.sort(key=lambda r: -abs(r["delta_union_ms"])) + + ac = A.get("anchor_count_per_iter_first", 0) + bc = B.get("anchor_count_per_iter_first", 0) + out = { + "anchor_count_a": ac, + "anchor_count_b": bc, + "anchor_count_match_pct": ( + round(100 * min(ac, bc) / max(ac, bc), 2) if max(ac, bc) else 0 + ), + "iter_total_anchor_ms": { + "a": A.get("iter_total_anchor_ms_median", 0), + "b": B.get("iter_total_anchor_ms_median", 0), + "delta": round( + B.get("iter_total_anchor_ms_median", 0) + - A.get("iter_total_anchor_ms_median", 0), + 3, + ), + }, + "iter_total_window_union_ms": { + "a": A.get("iter_total_window_union_ms_median", 0), + "b": B.get("iter_total_window_union_ms_median", 0), + "delta": round( + B.get("iter_total_window_union_ms_median", 0) + - A.get("iter_total_window_union_ms_median", 0), + 3, + ), + }, + "signatures": rows[: args.top], + "signatures_total": len(rows), + } + json.dump(out, sys.stdout, indent=2) + sys.stdout.write("\n") + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/module_slice.py b/skills/nsight-system-analysis/scripts/module_slice.py new file mode 100644 index 00000000000..c4fee11ecf9 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/module_slice.py @@ -0,0 +1,338 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Module-slice a profile: anchor on every GEMM+MHA kernel globally across all +compute streams, then sum (or union) the non-NCCL work in each inter-anchor window. + +Designed for Step 4 when op-group attribution is unreliable due to heavy fusion. +Both implementations should produce roughly equal anchor counts (same model → +same number of matmuls); the resulting windows are apples-to-apples. + +Usage: + python module_slice.py --yaml taxonomy.yml --windows windows.json + python module_slice.py --yaml taxonomy.yml --windows windows.json \ + --anchor-categories gemm,mha,flash_attn,cudnn_sdpa + +Output (stdout JSON): + { + "anchor_count_total": 1530, + "anchor_count_per_iter": [1530, 1530, ...], + "anchor_count_per_iter_constant": true, + "anchor_overlap_pct": 33.2, + "grouped_windows": [ + { + "signature": "gemm_TNT -> gemm_TNT", + "count_per_iter": 38, + "median_total_union_ms_per_iter": 30.4, + "median_total_sum_ms_per_iter": 30.6 + }, + ... + ], + "iter_total_anchor_ms_median": 222.5, + "iter_total_window_union_ms_median": 51.9, + "iter_total_window_sum_ms_median": 53.5 + } + +The `signature` is +` -> `, +where the "shape" is the part of the kernel name after the category-defining +substring. This keeps related anchors grouped while distinguishing different +matmul shapes. If shapes are too noisy, use `--signature-mode category` for the +coarser "gemm -> gemm" / "gemm -> mha" grouping. +""" + +from __future__ import annotations + +import argparse +import json +import re +from collections import defaultdict +from statistics import median + +from _lib import ( + classify, + clip, + load_taxonomy, + open_sqlite, + union_total_ns, + write_json, +) + + +def load_windows(path: str) -> list[tuple[int, int]]: + with open(path) as f: + data = json.load(f) + if isinstance(data, dict) and "windows_ns" in data: + data = data["windows_ns"] + return [(int(w[0]), int(w[1])) for w in data] + + +def anchor_signature(name: str, cat: str, mode: str) -> str: + """Build a signature string for an anchor kernel. + + mode='category': just the category (e.g. 'gemm', 'mha'). + mode='shape': category + the kernel-shape tail (e.g. 'gemm_128x256_TNT'). + """ + if mode == "category": + return cat + # Pull out a shape hint from the kernel name. + # nvjet: nvjet_sm103_qqtst_128x256_128x6_2x2f_2cta_h_bz_..._NTT + m = re.search(r"(\d+x\d+(?:_\d+x\d+)?)", name) + shape = m.group(1) if m else "" + # Pull out NTT/NNT/TNT/TNN suffix if present + m2 = re.search(r"_([NT]{2,4})(?:\b|$)", name) + suffix = m2.group(1) if m2 else "" + parts = [cat] + if shape: + parts.append(shape) + if suffix: + parts.append(suffix) + return "_".join(parts) + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("sqlite") + p.add_argument("--yaml", required=True) + p.add_argument("--windows", required=True) + p.add_argument( + "--anchor-categories", + default="gemm,flash_attn,cudnn_sdpa,mha,conv", + help="Comma-separated YAML categories whose kernels serve as anchors", + ) + p.add_argument( + "--signature-mode", + choices=["category", "shape"], + default="shape", + help="Granularity of anchor-pair signatures used for window grouping", + ) + args = p.parse_args() + + taxonomy = load_taxonomy(args.yaml) + windows = load_windows(args.windows) + anchor_cats = {c.strip() for c in args.anchor_categories.split(",") if c.strip()} + nccl_rx = taxonomy.get("nccl") + + anchor_count_per_iter: list[int] = [] + overlaps_per_iter: list[int] = [] + iter_anchor_ms: list[float] = [] + iter_window_union_ms: list[float] = [] + iter_window_sum_ms: list[float] = [] + + # signature -> list[ (window_union_ns, window_sum_ns) ]; one entry per + # (window instance) accumulated across iters. + sig_window_union: dict[str, list[int]] = defaultdict(list) + sig_window_sum: dict[str, list[int]] = defaultdict(list) + sig_count_per_iter: dict[str, list[int]] = defaultdict(list) + + with open_sqlite(args.sqlite) as con: + for lo, hi in windows: + # Pull all kernels in the iter window. + rows = con.execute( + """ + SELECT k.start, k.end, k.streamId, + COALESCE((SELECT value FROM StringIds WHERE id=k.demangledName), '') + FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.end > ? AND k.start < ? + ORDER BY k.start + """, + (lo, hi), + ).fetchall() + + anchors = [] # list of (start, end, category, name) + non_anchor_intervals_by_stream: dict[int, list[tuple[int, int]]] = ( + defaultdict(list) + ) + + for s, e, sid, name in rows: + cs, ce = max(s, lo), min(e, hi) + if ce <= cs: + continue + cat = classify(name, taxonomy) or "_uncat_" + is_nccl = cat == "nccl" or (nccl_rx and nccl_rx.search(name)) + if is_nccl: + continue # NCCL excluded entirely from module-slicing + if cat in anchor_cats: + anchors.append((cs, ce, cat, name)) + else: + non_anchor_intervals_by_stream[sid].append((cs, ce)) + + anchors.sort(key=lambda a: a[0]) + anchor_count_per_iter.append(len(anchors)) + + # Anchor time + window times. + anchor_total_ns = sum(e - s for s, e, _, _ in anchors) + iter_anchor_ms.append(anchor_total_ns / 1e6) + + # Group non-anchor intervals across all streams into one flat list + # (we want the union across streams = wall-time GPU busy in the window). + all_non_anchor = [] + for sid, ivs in non_anchor_intervals_by_stream.items(): + all_non_anchor.extend(ivs) + + overlaps = 0 + window_union_ns_iter = 0 + window_sum_ns_iter = 0 + sig_count_this_iter: dict[str, int] = defaultdict(int) + + # Iterate inter-anchor windows. To avoid losing time inside + # overlap-collapsed window groups, we maintain a "cursor" which is + # the running max(end) of the most recent anchor group. The next + # window begins at the cursor, not at the immediately preceding + # anchor's end. This way: if anchors k..k+M all overlap, we + # collapse them into a single anchor group ending at max(end_k..k+M), + # and the window from there to anchor[k+M+1].start carries any + # non-anchor work without being dropped. + n = len(anchors) + cursor = anchors[0][1] if n else 0 + k = 0 + while k < n - 1: + a_left = anchors[k] + # Advance cursor past any overlapping run of anchors that + # collectively span the next anchor's start. + end_of_group = a_left[1] + while k + 1 < n and anchors[k + 1][0] < end_of_group: + end_of_group = max(end_of_group, anchors[k + 1][1]) + overlaps += 1 + k += 1 + if k + 1 >= n: + break + a_right = anchors[k + 1] + w_lo = max(end_of_group, cursor) + w_hi = a_right[0] + if w_hi > w_lo: + sig = ( + anchor_signature(a_left[3], a_left[2], args.signature_mode) + + " -> " + + anchor_signature(a_right[3], a_right[2], args.signature_mode) + ) + clipped = clip(all_non_anchor, w_lo, w_hi) + w_sum_ns = sum(e - s for s, e in clipped) + w_union_ns = union_total_ns(clipped) + sig_window_sum[sig].append(w_sum_ns) + sig_window_union[sig].append(w_union_ns) + sig_count_this_iter[sig] += 1 + window_sum_ns_iter += w_sum_ns + window_union_ns_iter += w_union_ns + cursor = max(end_of_group, a_right[1]) + k += 1 + + # Tail: include any non-anchor work between the last anchor and + # the iter window end. This is the "post-last-anchor" region. + if n: + last_end = anchors[-1][1] + w_lo = max(last_end, cursor) + w_hi = hi + if w_hi > w_lo: + sig = ( + anchor_signature( + anchors[-1][3], anchors[-1][2], args.signature_mode + ) + + " -> [iter_end]" + ) + clipped = clip(all_non_anchor, w_lo, w_hi) + w_sum_ns = sum(e - s for s, e in clipped) + w_union_ns = union_total_ns(clipped) + sig_window_sum[sig].append(w_sum_ns) + sig_window_union[sig].append(w_union_ns) + sig_count_this_iter[sig] += 1 + window_sum_ns_iter += w_sum_ns + window_union_ns_iter += w_union_ns + # Pre-first-anchor region: also include. + if n: + w_lo = lo + w_hi = anchors[0][0] + if w_hi > w_lo: + sig = "[iter_start] -> " + anchor_signature( + anchors[0][3], anchors[0][2], args.signature_mode + ) + clipped = clip(all_non_anchor, w_lo, w_hi) + w_sum_ns = sum(e - s for s, e in clipped) + w_union_ns = union_total_ns(clipped) + sig_window_sum[sig].append(w_sum_ns) + sig_window_union[sig].append(w_union_ns) + sig_count_this_iter[sig] += 1 + window_sum_ns_iter += w_sum_ns + window_union_ns_iter += w_union_ns + + for sig in sig_count_this_iter: + sig_count_per_iter[sig].append(sig_count_this_iter[sig]) + overlaps_per_iter.append(overlaps) + iter_window_union_ms.append(window_union_ns_iter / 1e6) + iter_window_sum_ms.append(window_sum_ns_iter / 1e6) + + n_iters = len(windows) + n_anchors = anchor_count_per_iter[0] if anchor_count_per_iter else 0 + constant = len(set(anchor_count_per_iter)) == 1 + overlap_pct = ( + 100 + * sum(overlaps_per_iter) + / max(1, sum(a - 1 for a in anchor_count_per_iter if a > 0)) + if any(anchor_count_per_iter) + else 0.0 + ) + + # Aggregate per-signature: sum across windows in iter, then median across iters. + grouped = [] + for sig in sig_window_union: + total_union_per_iter = [] + total_sum_per_iter = [] + # We accumulated per-window samples; convert back to per-iter by chunking + # using sig_count_per_iter. + counts = sig_count_per_iter[sig] + idx = 0 + for c in counts: + u_slice = sig_window_union[sig][idx : idx + c] + s_slice = sig_window_sum[sig][idx : idx + c] + idx += c + total_union_per_iter.append(sum(u_slice)) + total_sum_per_iter.append(sum(s_slice)) + # Pad with zeros if some iters have no instance of this sig. + while len(total_union_per_iter) < n_iters: + total_union_per_iter.append(0) + total_sum_per_iter.append(0) + grouped.append( + { + "signature": sig, + "count_per_iter": round(median(counts) if counts else 0, 2), + "median_union_ms_per_iter": round( + median(total_union_per_iter) / 1e6, 3 + ), + "median_sum_ms_per_iter": round(median(total_sum_per_iter) / 1e6, 3), + } + ) + grouped.sort(key=lambda g: -g["median_union_ms_per_iter"]) + + out = { + "anchor_count_per_iter": anchor_count_per_iter, + "anchor_count_constant_across_iters": constant, + "anchor_count_per_iter_first": n_anchors, + "anchor_overlap_pct": round(overlap_pct, 2), + "iter_total_anchor_ms_median": round( + median(iter_anchor_ms) if iter_anchor_ms else 0, 3 + ), + "iter_total_window_union_ms_median": round( + median(iter_window_union_ms) if iter_window_union_ms else 0, 3 + ), + "iter_total_window_sum_ms_median": round( + median(iter_window_sum_ms) if iter_window_sum_ms else 0, 3 + ), + "grouped_windows": grouped, + "notes": [ + "Use median_union_ms_per_iter as the primary number for " + "wall-time attribution.", + "median_sum_ms_per_iter is shown for comparison; sum overcounts " + "when non-anchor kernels run in parallel on multiple compute streams.", + "NCCL kernels are excluded entirely (by name match against the " + "YAML's `nccl` category).", + ], + } + write_json(out) + + +if __name__ == "__main__": + main() diff --git a/skills/nsight-system-analysis/scripts/run_all.py b/skills/nsight-system-analysis/scripts/run_all.py new file mode 100644 index 00000000000..230bf327343 --- /dev/null +++ b/skills/nsight-system-analysis/scripts/run_all.py @@ -0,0 +1,328 @@ +#!/usr/bin/env python3 + +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +"""Run the full Steps 1–5 pipeline on one or two profiles and write all JSON +intermediates to an output directory. + +Usage: + # Comparative mode + python run_all.py --profile-a flat.sqlite --profile-b nonflat.sqlite \ + --yaml taxonomy.yml --out /tmp/analysis/ + + # Single-profile mode + python run_all.py --profile-a profile.sqlite --yaml taxonomy.yml --out /tmp/analysis/ + +The script: + 1. Runs iter_anchor.py on each profile → windows_.json + 2. Runs busy_idle.py on each → busy_.json + 3. Runs categorize.py on each → cat_.json (full taxonomy + Step 3 view) + 4. Decides Step 4 mode (op-group vs module-slicing) from the Step 3 output + `module_slicing_recommended` flag (uses OR across both profiles). + 5. If module-slicing: runs module_slice.py on each → mod_.json + If op-group: re-runs categorize with --residual-only → opgroup_.json + 6. Runs exposed_comm.py on each → comm_.json (skips cleanly if no NCCL). + 7. Writes a top-level summary.json with all key numbers plus the Step 4 decision. + +This is a convenience wrapper. For more control, invoke individual scripts. +""" + +from __future__ import annotations + +import argparse +import json +import subprocess +import sys +from pathlib import Path + +SCRIPT_DIR = Path(__file__).parent + + +def run(cmd: list[str], capture: bool = True) -> dict | str: + """Run a script, return parsed JSON from stdout (or text).""" + result = subprocess.run(cmd, capture_output=True, text=True) + if result.returncode != 0: + sys.stderr.write(f"ERROR running: {' '.join(cmd)}\nstderr:\n{result.stderr}\n") + sys.exit(result.returncode) + if result.stderr.strip(): + # Forward warnings to the user, but don't fail. + sys.stderr.write(result.stderr) + if not capture: + return result.stdout + try: + return json.loads(result.stdout) + except json.JSONDecodeError: + sys.stderr.write( + f"Could not parse JSON from {cmd[0]}:\n{result.stdout[:1000]}\n" + ) + sys.exit(2) + + +def analyze_profile( + profile: Path, yaml: Path, out_dir: Path, tag: str, n_iters: int | None +) -> dict: + """Run Steps 1–5 on a single profile. Returns the assembled summary.""" + summary: dict = {"profile": str(profile), "tag": tag} + py = sys.executable + + # Step 1: anchor + windows + cmd = [py, str(SCRIPT_DIR / "iter_anchor.py"), str(profile)] + if n_iters: + cmd += ["--n-iters", str(n_iters)] + anchor = run(cmd) + windows_path = out_dir / f"windows_{tag}.json" + windows_path.write_text(json.dumps(anchor)) + summary["step1"] = { + "anchor": anchor["anchor"]["name"], + "iter_count_used": anchor["iter_count_used"], + "median_ms": anchor["median_ms"], + "min_ms": anchor["min_ms"], + "max_ms": anchor["max_ms"], + "cross_check": anchor.get("cross_check"), + } + + # Step 2: busy/idle + busy = run( + [ + py, + str(SCRIPT_DIR / "busy_idle.py"), + str(profile), + "--windows", + str(windows_path), + "--yaml", + str(yaml), + ] + ) + (out_dir / f"busy_{tag}.json").write_text(json.dumps(busy)) + summary["step2"] = { + "median": busy["median"], + "longest_single_stream_union_ms_median": busy[ + "longest_single_stream_union_ms_median" + ], + } + + # Step 3 view: gemm/conv/mha + cat3 = run( + [ + py, + str(SCRIPT_DIR / "categorize.py"), + str(profile), + "--yaml", + str(yaml), + "--windows", + str(windows_path), + "--report-categories", + "gemm,conv,mha", + ] + ) + (out_dir / f"cat_step3_{tag}.json").write_text(json.dumps(cat3)) + summary["step3"] = { + "per_category": cat3["per_category"], + "fused_share_of_residual_pct": cat3["fused_share_of_residual_pct"], + "module_slicing_recommended": cat3["module_slicing_recommended"], + } + + # Full categorize for uncategorized inspection + cat_full = run( + [ + py, + str(SCRIPT_DIR / "categorize.py"), + str(profile), + "--yaml", + str(yaml), + "--windows", + str(windows_path), + ] + ) + (out_dir / f"cat_full_{tag}.json").write_text(json.dumps(cat_full)) + summary["uncategorized_above_1pct"] = cat_full["uncategorized_above_threshold"][:10] + + # Step 5: comm + comm = run( + [ + py, + str(SCRIPT_DIR / "exposed_comm.py"), + str(profile), + "--yaml", + str(yaml), + "--windows", + str(windows_path), + ] + ) + (out_dir / f"comm_{tag}.json").write_text(json.dumps(comm)) + summary["step5"] = comm + + return summary + + +def main(): + p = argparse.ArgumentParser( + description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter + ) + p.add_argument("--profile-a", required=True, help="Path to first .sqlite") + p.add_argument("--profile-b", help="Path to second .sqlite (comparative mode)") + p.add_argument("--yaml", required=True, help="Taxonomy YAML") + p.add_argument("--out", required=True, help="Output directory for intermediates") + p.add_argument("--n-iters", type=int, help="Expected iteration count (optional)") + args = p.parse_args() + + out_dir = Path(args.out) + out_dir.mkdir(parents=True, exist_ok=True) + yaml_path = Path(args.yaml) + + summary: dict = {"mode": "comparative" if args.profile_b else "single"} + summary["a"] = analyze_profile( + Path(args.profile_a), yaml_path, out_dir, "a", args.n_iters + ) + if args.profile_b: + summary["b"] = analyze_profile( + Path(args.profile_b), yaml_path, out_dir, "b", args.n_iters + ) + + # Step 4 decision + fused_a = summary["a"]["step3"]["fused_share_of_residual_pct"] + fused_b = ( + summary.get("b", {}).get("step3", {}).get("fused_share_of_residual_pct", 0.0) + ) + use_module_slicing = fused_a > 10.0 or fused_b > 10.0 + summary["step4_mode"] = "module-slicing" if use_module_slicing else "op-group" + summary["step4_decision"] = { + "fused_share_a_pct": fused_a, + "fused_share_b_pct": fused_b if args.profile_b else None, + "threshold_pct": 10.0, + "use_module_slicing": use_module_slicing, + } + + py = sys.executable + for tag, prof in [("a", args.profile_a)] + ( + [("b", args.profile_b)] if args.profile_b else [] + ): + windows_path = out_dir / f"windows_{tag}.json" + if use_module_slicing: + mod = run( + [ + py, + str(SCRIPT_DIR / "module_slice.py"), + str(prof), + "--yaml", + str(yaml_path), + "--windows", + str(windows_path), + "--signature-mode", + "shape", + ] + ) + mod_path = out_dir / f"mod_{tag}.json" + mod_path.write_text(json.dumps(mod)) + summary[tag]["step4_module_slice"] = { + "anchor_count_per_iter": mod["anchor_count_per_iter_first"], + "anchor_count_constant": mod["anchor_count_constant_across_iters"], + "anchor_overlap_pct": mod["anchor_overlap_pct"], + "iter_total_anchor_ms": mod["iter_total_anchor_ms_median"], + "iter_total_window_union_ms": mod["iter_total_window_union_ms_median"], + "iter_total_window_sum_ms": mod["iter_total_window_sum_ms_median"], + "top_windows": mod["grouped_windows"][:10], + } + else: + opgroup = run( + [ + py, + str(SCRIPT_DIR / "categorize.py"), + str(prof), + "--yaml", + str(yaml_path), + "--windows", + str(windows_path), + "--residual-only", + ] + ) + (out_dir / f"opgroup_{tag}.json").write_text(json.dumps(opgroup)) + summary[tag]["step4_op_group"] = opgroup["per_category"] + + # Arithmetic invariants (top-level sanity) + invariants = {} + for tag in ["a"] + (["b"] if args.profile_b else []): + s = summary[tag] + med = s["step2"]["median"] + invariants[tag] = { + "iter_eq_busy_plus_idle": round( + med["iter_ms"] - med["busy_ms"] - med["idle_ms"], 3 + ), + } + if "step5" in s and "totals" in s["step5"]: + recon = s["step5"]["totals"] + exposed = recon["exposed_ms_median_per_iter"] + non_nccl = s["step5"]["reconciliation"]["non_nccl_union_ms_median_per_iter"] + invariants[tag]["exposed_plus_non_nccl_minus_busy_ms"] = round( + exposed + non_nccl - med["busy_ms"], 3 + ) + summary["invariants"] = invariants + + # Comparative Δ + if args.profile_b: + med_a = summary["a"]["step2"]["median"] + med_b = summary["b"]["step2"]["median"] + summary["delta"] = { + "iter_ms": round(med_b["iter_ms"] - med_a["iter_ms"], 3), + "busy_ms": round(med_b["busy_ms"] - med_a["busy_ms"], 3), + "idle_ms": round(med_b["idle_ms"] - med_a["idle_ms"], 3), + } + + # Comparative module-slice diff if both profiles + module-slicing mode. + if args.profile_b and use_module_slicing: + diff = run( + [ + py, + str(SCRIPT_DIR / "module_diff.py"), + str(out_dir / "mod_a.json"), + str(out_dir / "mod_b.json"), + ] + ) + (out_dir / "mod_diff.json").write_text(json.dumps(diff)) + summary["step4_module_diff_top"] = diff["signatures"] + + summary_path = out_dir / "summary.json" + summary_path.write_text(json.dumps(summary, indent=2)) + + # Compact stdout summary — agent reads summary.json for the rest. + print(f"\nFull summary JSON: {summary_path}") + print(f"Intermediates in: {out_dir}/") + print() + print(f"Mode: {summary['mode']}") + a = summary["a"] + print( + f" A ({Path(a['profile']).name}): " + f"anchor={a['step1']['anchor']}, iters_used={a['step1']['iter_count_used']} " + f"(after dropping warmup+cooldown), " + f"per_iter_median={a['step1']['median_ms']:.2f} ms" + ) + if args.profile_b: + b = summary["b"] + print( + f" B ({Path(b['profile']).name}): " + f"anchor={b['step1']['anchor']}, " + f"iters_used={b['step1']['iter_count_used']} " + f"(after dropping warmup+cooldown), " + f"per_iter_median={b['step1']['median_ms']:.2f} ms" + ) + d = summary["delta"] + print( + f" Δ iter={d['iter_ms']:+.2f} ms, " + f"busy={d['busy_ms']:+.2f}, idle={d['idle_ms']:+.2f}" + ) + print() + print( + f"Step 4 mode: {summary['step4_mode']} " + f"(fused_share A={fused_a:.1f}%" + + (f", B={fused_b:.1f}%" if args.profile_b else "") + + ", threshold=10%)" + ) + print() + print("Invariants (should be near zero):") + for tag, inv in summary["invariants"].items(): + print(f" {tag}: {inv}") + + +if __name__ == "__main__": + main() diff --git a/skills/optimize-inference-siddharth/SKILL.md b/skills/optimize-inference-siddharth/SKILL.md new file mode 100644 index 00000000000..c36c6d33aef --- /dev/null +++ b/skills/optimize-inference-siddharth/SKILL.md @@ -0,0 +1,510 @@ +--- +name: optimize-inference-siddharth +description: >- + Optimizes the Megatron Core inference backend for a new model, architecture, or + feature, using the patterns Siddharth Singh (sidsingh-nvidia) established while + closing the vLLM performance gap. Covers CUDA-graph scope and bucket coverage, + the inference_optimized MoE stack (NVLS AllGather-V, fused grouped GEMM, + shared-expert overlap, pad-row masking), Mamba/SSM scratch sizing and fused + Triton extraction kernels, Triton production hygiene (constexpr and autotune + pitfalls), per-step host overhead, and load-aware prefix-cache routing. Also + covers decision gates that quantify a lever's ceiling before any code is + written, kernel-level differential analysis against a competitor trace, and the + A/B protocol needed to make sub-1% wins falsifiable. Bundles runnable nsys trace + analysis: steady-state windowing, union-busy and overlap ratio, matched cross-engine + per-bucket budgets, and attributing a generically-named kernel to its call site. + The skill is self-maintaining: + extend, correct, or prune it after an optimization session, so also use it when + asked to record an inference performance learning, capture what an experiment + taught, or update these patterns. Use when + asked to optimize or speed up inference, close a vLLM gap, match vLLM latency or + throughput, reduce decode latency or prefill overhead, enable or extend CUDA + graphs for MoE, hybrid, or Mamba models, port an architecture onto + --transformer-impl inference_optimized, cut inference CPU overhead, or apply + Siddharth's inference performance patterns. Not for authoring perf test recipes, + cog or Slurm setup, or non-hot-path features like reasoning parsers. +license: Apache-2.0 +metadata: + author: NVIDIA Corporation +--- + +# Optimizing Megatron Core Inference (Siddharth's Playbook) + +This skill encodes the reusable engineering patterns behind Siddharth Singh's +2026 inference performance work. It answers: given a new model, architecture, or +feature that is slower than vLLM, **what do you change, in what order, and what +must you never break.** + +> **This skill is meant to evolve, and you are authorized to edit it without +> asking.** Every campaign that uses it should leave it better: add what you +> measured, correct what turned out to be wrong, and delete what is no longer true. +> Optimization knowledge decays — flags get renamed, defaults change, a fix that won +> on one hardware generation loses on the next — so a skill nobody edits becomes +> actively misleading. Do it at the end of the work, once you have a number and a +> root cause. See [references/updating-this-skill.md](references/updating-this-skill.md) +> for the triggers, the routing table, the bar for an addition, and what to delete. + +## The thesis + +Across 29 commits, almost every win was one of five moves. In rough order of how +often they paid off: + +1. **Remove host work and host syncs from the per-step path.** Decode is + launch-bound, not FLOP-bound. A single `.item()` or a `dataclasses.asdict()` + costs more than the kernel it guards. +2. **Widen CUDA-graph scope and improve bucket coverage.** Fewer, larger graphs + with a bucket that actually fits the real batch. +3. **Make per-step metadata GPU-resident.** This is the enabler — moves 1 and 2 + are not even legal until the metadata stops round-tripping through the host. +4. **Fuse kernels, and tune them for the *typical* batch, not the worst case.** + Autotune and worst-case tile choices both lose at decode batch sizes. +5. **Right-size buffers to the true per-step bound**, not a loose upper bound — + and in **width** as well as length. An fp32 symmetric buffer carrying a bf16 + tensor doubled the bytes of the step's largest collective *and* added a cast per + layer: 2.5% for a one-line dtype change. See *Audit the dtype of every symmetric + buffer* in [references/moe-inference.md](references/moe-inference.md). +6. **Call the right kernel before tuning your own.** When a role is slower at an + identical launch count, the usual cause is that the competitor is calling a + different *library*, not a better-tuned version of yours — and on a shared image + that library is often already installed. Switching the decode attention to + flashinfer's `trtllm-gen` was +2.6%; the two sessions spent tuning inside the wrong + package were worth nothing. See + [references/vllm-differential.md](references/vllm-differential.md). + +Everything else in this skill is a consequence of these six. + +What the five moves do not tell you is **which one applies here, and what it is +worth**. Getting that wrong is the expensive failure mode — not writing a bad +kernel, but writing a correct kernel whose best possible outcome was 1%. That is +what Step 1 is for. + +## Step 1: Measure, classify, and gate before changing anything + +Do not skip this. Every commit in this history started from a profile, and the +fix location is rarely where you would guess. + +See [references/measuring.md](references/measuring.md) for the profiler +endpoints, NVTX ranges, built-in counters, honest idle accounting, and — read this +before adding host visibility — **the trace flags that deadlock nsys finalization +on MoE decode workloads**. For trace analysis use the `nsight-system-analysis` +skill; for the throughput harness use `run-inference-performance-tests`. + +Classify the dominant signal, then jump to the matching section: + +| Dominant signal in the profile | Where to work | +|---|---| +| GPU idle gaps between kernels; host ahead of device | Step 2, CUDA graphs | +| Host time after the forward (serialize, detokenize, ZMQ) | Step 5, host path | +| Exposed NCCL, or EP ranks waiting on each other | Step 3, NVLS AllGather-V and `ep_consensus_interval` | +| Periodic multi-ms stalls, or slow first steps | Step 4, Triton recompilation and autotune | +| Grouped GEMM inefficient at small batch | Step 3, grouped-GEMM backend and tile heuristic | +| OOM, or a prefix cache far smaller than expected | Step 4, scratch sizing math | +| Many sub-microsecond kernels in a row | Step 3, fusion — but gate it first | +| One role slower at an **identical launch count** | Wrong library, not wrong tuning — [vllm-differential.md](references/vllm-differential.md) | +| An expensive kernel with a generic name (`elementwise_kernel`, copies) | Name it from its trace neighbours, and treat it as a **symptom**: a copy next to a collective is usually that collective's buffer dtype, not a fusion candidate — [measuring.md](references/measuring.md) | +| Kernels efficient, buckets at parity, still slower | Packing. Compare sum-of-durations ÷ union against the competitor — [measuring.md](references/measuring.md) | + +### Then gate the lever: share is not headroom + +A category's share of device time is not what you can win back, for two reasons that +have each burned real weeks. A category can be 33% of device time with **2% of +headroom**, because it is already moving the bytes it has to move. And **device time +is not wall time** — kernels overlap across streams, so the step's wall clock is the +critical path through the per-layer dependency chain; work on a side stream +contributes less than its share, while a tiny kernel on the serial chain costs its +duration *plus* the dispatch gap behind it, every layer. + +So before writing anything that takes more than about a day, compute the ceiling: +establish the floor from measured machine constants, measure the current cost under +**graph replay** at the decode token count, take the ratio, subtract what the fix +itself costs (added launches, grid syncs, atomics, extra passes), and express the +result as a fraction of the step. Then write down **proceed** or **gated out**. + +Full method, the per-launch fixed costs you need, and three case studies where a +gate killed a multi-week effort — including a hand-written grouped GEMM whose +entire ceiling was 1.45×, most of it reachable by tile tuning alone — are in +[references/decision-gates.md](references/decision-gates.md). + +Skip the gate for cheap reversible changes: flag flips, tile retunes, backend enum +swaps. Gating those costs more than trying them. + +### If the target is a competitor, diff against its trace + +When the goal is "match vLLM," its trace is a specification, not just a scoreboard. +Take matched windows of one forward block from each and answer one question: **is +any individual kernel slower, or does mcore just launch more of them?** The fixes +are opposite. On Qwen3-30B the answer was ~467 kernels against ~1784 for the same +block, which pointed at fusion and away from kernel rewrites. + +Method, window anchoring, and the traps in adopting a competitor's kernel are in +[references/vllm-differential.md](references/vllm-differential.md). + +### Open a ledger before the first change + +For anything spanning more than one session, keep an append-only ledger: a fixed +protocol table (hardware, model, batch, output length, parallelism, warmup/timed +counts), one row per experiment **including every rejection with its root cause**, +a running distance-to-target, and a next-levers list re-derived after each profile. +Never edit a recorded result; supersede it. See +[references/measuring.md](references/measuring.md). The negative results are the +higher-value half — they are what stops the next person re-deriving a dead end. + +## Step 2: CUDA-graph the largest region that is safe + +Full detail in [references/cuda-graphs.md](references/cuda-graphs.md). + +Set `cuda_graph_impl="local"` and `inference_cuda_graph_scope=block` to capture +the whole decoder block in one graph. For hybrid models, graph ownership lives on +`HybridModel`, not the stack, so the embedding and output layers land inside the +same capture — that widening was itself a measurable win. + +Then fix coverage. A step whose shape matches no captured bucket silently falls +back to eager, so a "CUDA graphs enabled" run can still be launch-bound. Check +`num_cuda_graphs` (`-1` auto-sizes), `cuda_graph_max_tokens` (512 by default, so +prefill and mixed steps up to 512 tokens get a graph), and +`cuda_graph_sizing_distribution`. + +Know what the wide capture costs, though — it is a trade, not a free win. Under +full-iteration inference capture, the flashinfer sampling backend cannot run, async +scheduling is guarded off for EP (and deadlocks if you open the guard), and +multi-stream comm/compute overlap is bounded by the graph's structure rather than by +hardware queues, so `CUDA_DEVICE_MAX_CONNECTIONS` does nothing. Details and error +signatures are under *What a wide capture costs you* in +[references/cuda-graphs.md](references/cuda-graphs.md). + +## Step 3: Enable the model-path stack + +**MoE** — see [references/moe-inference.md](references/moe-inference.md). +Use `--transformer-impl inference_optimized`, which swaps in +`InferenceTopKRouter` and `InferenceGroupedMLP` and picks a dispatcher from +`inference_moe_token_dispatcher_type`. Prefer `nvls` (the default): the NCCL +fallback requires equal token counts across EP ranks, which forces decode-only +graphs. Then check the grouped-GEMM backend, shared-expert overlap, and that +padding rows route to no expert. + +**Hybrid / Mamba** — see +[references/mamba-and-triton.md](references/mamba-and-triton.md). +Size the extraction scratch to the real per-step bound before the durable prefix +cache claims the rest, and use fused gather-plus-scatter kernels gated on a +runtime count instead of materializing intermediates. + +**Dense GPT** — Steps 2, 4, and 5 usually cover it. Still apply the padding and +dtype contracts to any fused kernel you add. + +## Step 4: Kernel and Triton hygiene + +Full detail in [references/mamba-and-triton.md](references/mamba-and-triton.md). +The two rules that cause the most damage when violated are the `tl.constexpr` +specialization rule and the no-autotune-in-production rule, both in the hard +rules below. + +## Step 5: Trim the host path + +Full detail in [references/host-path.md](references/host-path.md). The per-step +host critical section is the `bookkeeping` to `detokenization` to +`coordinator_communication` span in +[dynamic_engine.py](megatron/core/inference/engines/dynamic_engine.py). Anything +there runs once per step and blocks the next one. + +## Step 6: Re-measure, keep a kill switch, and feed the result back + +Re-run the same measurement from Step 1 with the identical config. Then confirm +correctness with `run-inference-functional-tests`, and regenerate golden values +if you changed bucketization (padding changes shift outputs — expected, not a +bug). + +Re-measure **in the same allocation, back to back, arms alternating** — identical +configs drifted by up to 1.6% between sessions on this workload, which is larger +than most individual wins, so a cross-session comparison will mislead you in both +directions. Accept on **distribution separation** (slowest ON beats fastest OFF), +not on mean delta. Protocol and worked examples: *The noise floor is bigger than +your win* in [references/measuring.md](references/measuring.md). + +Expect the end-to-end result to differ from the kernel-level result, often by +several times in either direction. Work on the serial dependency chain converts at +more than 1:1 (a ~1% microbench ceiling delivered +2.9%, because removing a launch +also removes a graph node and a dispatch gap ×48 layers); already-overlapped work +converts at a third to a half; work off the critical path converts at roughly zero +(a 1.25× kernel win delivered a wash). If the measured conversion is far from your +prediction, you mis-identified where the work sits. + +Every optimization here shipped with a way to turn it off +(`inference_disable_triton_nvls_kernels`, +`inference_moe_disable_fused_quant_kernels`, the backend enums). Add one. It is +how the next person A/Bs your change instead of reverting it. + +Then close the loop: append the result to the campaign ledger — **including +rejections, with their root cause** — and promote whatever generalizes into this +skill. The promotion test is whether it would change what someone does on a +*different* model; if it only describes this one, it stays in the ledger. Mechanics +in [references/updating-this-skill.md](references/updating-this-skill.md). + +--- + +## Hard rules + +These are non-negotiable. Each one exists because violating it broke something. + +### 1. No host sync on the per-step path + +No `.item()`, `.tolist()`, `.cpu()`, or data-dependent Python branch in code that +runs every step. To publish a per-step scalar, `fill_` it into a **preallocated, +fixed-address** GPU tensor and read it inside the kernel: + +```python +# megatron/core/inference/contexts/attention_context/mamba_metadata.py +# fill_ is async (no host sync) and keeps the tensor at the same address +# captured graphs reference. +self._intermediate_real_count_buffer.fill_(self.intermediate_count) +``` + +The fixed address is the whole point: a captured graph records pointers, so the +*value* may change between replays but the *address* may not. + +### 2. `tl.constexpr` only for values fixed for the process lifetime + +A `tl.constexpr` parameter is baked into the compiled kernel, so **every distinct +value triggers a fresh JIT compile.** Marking a per-step batch or token count +`constexpr` means recompiling on every step. That was the entirety of commit +`f29c747` — a four-line fix worth a large latency spike. + +Constexpr is correct for block and tile sizes, `tl.arange` bounds, and unroll +counts. It is wrong for anything that varies per step. Note that even plain +`int` arguments get specialized on divisibility-by-16 and `== 1`; if such a value +varies, use `@triton.jit(do_not_specialize=["batch"])`. + +### 3. Graph-safe kernel shape: worst-case grid, data-conditional body + +Size the launch grid at capture time to the maximum, then let padded programs +exit immediately: + +```python +# megatron/core/ssm/ops/intermediate_extraction.py +real_count = tl.load(real_count_ptr).to(tl.int32) +if pid_slot >= real_count: + return +``` + +This keeps the grid static for replay while making padded slots nearly free. + +### 4. Size per-step work to the matched bucket, not the global max + +Once a step matches a CUDA-graph bucket, all metadata updates and scratch writes +should be bounded by that bucket, never by global `max_tokens` or `max_requests`. +Commit `9b4074b` was largely this one change applied to Mamba prefill. + +### 5. Padding must not do real work + +Three concrete forms: + +- CUDA-graph pad rows get routing index `-1` (via `mask_routing_padding`) so they + activate no expert. +- Pad tokens point at a reserved `dummy_block_idx` so KV-append writes somewhere + valid and throwaway. +- Do not *zero* rows past `valid_tokens`. Downstream only reads the valid prefix, + so a zeroing pass is pure wasted bandwidth — just don't write them. + +### 6. Tune for the expected batch; never autotune in production + +Pick tile, warp, and stage counts from a host-side heuristic keyed on the typical +token count. Commit `20f09364` replaced 25 `@triton.autotune` configs with +vLLM's `_get_default_config` heuristic, which cut compile time and picked better +tiles for decode-sized batches. Route any remaining autotuning through +`autotune_configs` in +[determinism.py](megatron/core/ssm/ops/determinism.py). + +### 7. Preserve Transformer Engine `nn.Parameter` identity + +Grouped GEMM wants one stacked `[num_experts, out, in]` weight; TE stores +per-expert parameters. Redirect `param.data` to a view into the stacked buffer +rather than replacing the `Parameter` object, and build it **lazily on first +forward**, after checkpoint load: + +```python +# megatron/core/transformer/moe/experts.py +# Redirect param.data to view into contiguous buffer. +# The nn.Parameter object stays the same - TE's internal state is preserved. +fc1_param.data = _fc1_weight[i] +``` + +Replacing the object corrupts TE's FP8 and bookkeeping state; building eagerly at +`__init__` reads weights that do not exist yet. Commit `905c0e38` fixed both +after they broke the RL integration. + +### 8. Hold the dtype contract at every kernel boundary + +Token ids are `int64` throughout the pipeline. FlashInfer's sampling kernels +return `int32`, so their results are cast with `.long()` at the boundary. +Normalize dtype where the external kernel enters, not by letting it propagate — +the cast is free once, a mismatch inserts conversion kernels on every step. + +Related, and a sharper trap: **not everything belongs in a graph.** FlashInfer +sampling is deliberately left eager, because its kernel choice is data-dependent +*and* FlashInfer bakes the philox RNG state into a graph as a by-value constant at +capture — so a captured sampler replays identical random numbers every step. Check +for baked-in state before widening a graph over anything stateful or random. + +### 9. Guard correctness with a reference implementation + +The testing pattern throughout: compare the fast kernel against an obvious +PyTorch or plain-loop reference with `torch.testing.assert_close`; assert the +gating behavior explicitly by prefilling the output with a sentinel and checking +padded slots still hold it; and re-derive any sizing formula independently in the +test so a changed bound fails loudly. See `add-inference-unit-tests`. + +### 10. A sub-1% claim needs same-session, back-to-back, separated arms + +Session-to-session drift on identical configs reached 1.6% — larger than most +individual wins once the easy ones are gone. So a number compared against last +session's baseline is not evidence. + +Re-baseline the current best config **in the allocation you will test in**, run the +OFF and ON arms **back to back**, repeat the pair, and accept only if the arms do +not overlap: + +``` +min(ON iterations) > max(OFF iterations) +``` + +Report pairwise deltas rather than one average, and treat the first timed iteration +as a suspected cold outlier. A +1% mean with overlapping arms is not a result; a ++0.9% with separated arms is. + +### 11. If it is not bit-exact, say so and justify it differently + +Bit-exactness is achievable more often than assumed — retuning tiles while holding +`BLOCK_SIZE_K` fixed preserves the fp32 reduction order, and a masked slot adding an +exact 0.0 keeps a fused reduction bit-exact. **Check whether a bit-exact formulation +exists before accepting drift.** + +When none does — norm and reduction fusions frequently land one ulp off, since TE's +internal rsqrt and reduction order differ — the acceptance argument changes shape and +must be made explicitly: bound the deviation in ulps against the reference across +several token counts and seeds (`max_rel ≤ 7.9e-3` is about one bf16 ulp); diff +fixed temperature-0 coherence output against the gate-OFF arm; and for any prompt +that diverges, **inspect where**. A divergence at a genuinely low-confidence branch +with both continuations fluent and factually correct is acceptable; one that degrades +fluency or correctness is not. Record which prompts diverged — "not bit-exact" must +be lookup-able, not rediscovered. + +Never let an ulp-level deviation ride on `assert_close` alone: the tolerance that +passes it also passes a real bug. + +--- + +## Code entry points + +| Area | Paths | +|---|---| +| MoE dispatchers | [token_dispatcher_inference.py](megatron/core/transformer/moe/token_dispatcher_inference.py) — `NCCLAllGatherDispatcher`, `NVLSAllGatherVDispatcher` | +| MoE experts / router | [experts.py](megatron/core/transformer/moe/experts.py) `InferenceGroupedMLP`, [router.py](megatron/core/transformer/moe/router.py) `InferenceTopKRouter`, [moe_layer.py](megatron/core/transformer/moe/moe_layer.py) | +| Fused MoE kernels | [megatron/core/inference/moe/](megatron/core/inference/moe/) — `fused_moe.py`, `vllm_fused_moe.py`, `permute.py`, `activations.py`, `metadata.py` | +| Pad-row masking | [inference_routing_mask_kernel.py](megatron/core/transformer/moe/inference_routing_mask_kernel.py) | +| Backend selection | [backends.py](megatron/core/models/backends.py), [moe_module_specs.py](megatron/core/models/gpt/moe_module_specs.py) | +| CUDA-graph buckets | [batch_dimensions_utils.py](megatron/core/inference/batch_dimensions_utils.py) | +| Graph scope hooks | [enums.py](megatron/core/transformer/enums.py), [transformer_block.py](megatron/core/transformer/transformer_block.py), [hybrid_model.py](megatron/core/models/hybrid/hybrid_model.py) | +| Per-step context state | [dynamic_context.py](megatron/core/inference/contexts/dynamic_context.py), [gpu_view.py](megatron/core/inference/contexts/gpu_view.py) | +| Mamba / SSM | [mamba_mixer.py](megatron/core/ssm/mamba_mixer.py), [intermediate_extraction.py](megatron/core/ssm/ops/intermediate_extraction.py), [mamba_metadata.py](megatron/core/inference/contexts/attention_context/mamba_metadata.py), [mamba_slot_allocator.py](megatron/core/inference/contexts/mamba_slot_allocator.py) | +| Symmetric memory / NVLS | [symmetric_memory.py](megatron/core/inference/symmetric_memory.py), [torch_symm_triton/](megatron/core/inference/communication/torch_symm_triton/), [inference_layers.py](megatron/core/tensor_parallel/inference_layers.py) | +| Engine / host path | [dynamic_engine.py](megatron/core/inference/engines/dynamic_engine.py), [inference_request.py](megatron/core/inference/inference_request.py), [data_parallel_inference_coordinator/](megatron/core/inference/data_parallel_inference_coordinator/) | + +## Flags + +Verified against the current tree. Defaults matter — several are already the +tuned value, so the useful move is often *checking* rather than changing them. + +| Flag | Default | Effect | +|---|---|---| +| `transformer_impl` | `transformer_engine` | `inference_optimized` swaps in the whole inference MoE stack | +| `inference_moe_token_dispatcher_type` | `'nvls'` | `nccl` fallback forces equal EP token counts and decode-only graphs | +| `inference_grouped_gemm_backend` | `"vllm"` | `vllm`, `flashinfer`, or `torch`; MXFP8 needs `torch` | +| `inference_disable_triton_nvls_kernels` | `False` | Kill switch: fall back to NCCL collectives | +| `inference_moe_disable_fused_quant_kernels` | `False` | Kill switch: unfuse activation + quantize | +| `moe_router_dtype` | — | Must be `fp32` for `inference_optimized`, to avoid per-decode dtype conversion | +| `cuda_graph_impl` | `"none"` | Inference graphs need `local` | +| `inference_cuda_graph_scope` | derived | `none`, `layer`, or `block`; `local` derives `layer`, so set `block` explicitly | +| `cuda_graph_max_tokens` | `512` | Token ceiling for prefill and mixed graphs | +| `num_cuda_graphs` | `16` | `-1` auto-sizes from `log2(max_tokens)` | +| `cuda_graph_sizing_distribution` | `EXPONENTIAL` | `LINEAR` gives the dense small-batch ladder | +| `ep_consensus_interval` | `20` | Skip EP consensus all-reduces while busy | +| `prefix_caching_routing_alpha` | `0.5` | 0 is pure load balance, 1 is pure cache affinity | +| `SamplingParams.return_prompt_tokens` | `False` | Opt in to echoing prompt ids over the wire | +| `moe_enable_routing_replay` | `False` | Record per-token expert choices for imbalance analysis | + +Three further gates exist as environment variables in PR #6064 and are **not merged**, +so do not expect them in the tree: `MCORE_NVLS_RS_BF16` (bf16 reduce-scatter buffer, ++2.5%), `MCORE_FLASHINFER_DECODE` (flashinfer `trtllm-gen` decode attention, +2.6%), and +`MCORE_FLASHINFER_PDL` (Programmatic Dependent Launch for that kernel, +0.3%). If they +have since landed as config flags, correct this section. + +### Flags that look like free wins and are not + +Measured and rejected on Qwen3-30B-A3B EP4 under `inference_optimized` + block +scope. Each is the sort of thing you would reasonably try first. + +**Record the *reason*, not just the verdict, and re-test when the reason changes.** A +rejection is conditional on the state of the tree and on the rest of your config, and +two entries here already moved: + +- `CUDA_DEVICE_MAX_CONNECTIONS=8` measured **+0.9%** at one point in the campaign and + went flat only *after* the bf16 reduce-scatter removed the contention it was + relieving. "Flat" is true of the current config, not of the flag. +- FA4's rejection rested on a hardcoded `num_splits=1`, which upstream later fixed. It + had to be re-measured after the rebase — it still lost, but the original reason no + longer existed, so the old entry was no longer evidence. + +So a rejection is only reusable if it carries its mechanism and a date. One that says +only "flat" has to be re-run by the next person anyway. + +| Flag | Result | +|---|---| +| `--moe-router-fusion`, `--moe-permute-fusion` | **Crash**: `AssertionError: hidden_size mismatch: 128 vs 8`. Wired to the training MoE path; TE's fused router emits a dense `num_experts` map, not the dense top-k the inference dispatcher needs. | +| `--inference-dynamic-batching-sampling-backend flashinfer` | **Crash** under graph capture: `Generator not registered with the capturing graph`. | +| `--inference-dynamic-batching-async-sched-mode serial` | Guarded off for EP; opening the guard deadlocks the first single-request decode step. Gains +0.85% at batch — the post-sampling chain is a genuine data dependency, not a scheduling artifact. | +| `CUDA_DEVICE_MAX_CONNECTIONS=8` | Flat *as of 2026-08*, on a config whose collectives are bf16. Overlap is bounded by graph structure, not queue count. Was +0.9% when the reduce-scatter was fp32 — re-test if you change collective volume. | +| Flash Attention 4 (`num_splits` fix rebased in) | Still slower than FA2 at decode shapes, re-measured 2026-08 after the upstream fix removed the original reason. The win was in a different package entirely — see the flashinfer note in [vllm-differential.md](references/vllm-differential.md). | +| `--inference-moe-token-dispatcher-type nccl` | **0.66×**. Pads to the worst-case per-rank count, roughly doubling comm volume. Correctness fallback only. | +| `--inference-grouped-gemm-backend torch` | **0.82×** for BF16. Use only when MXFP8 forces it. | + +## References + +- [references/decision-gates.md](references/decision-gates.md) — quantify a lever's ceiling before building it; per-launch fixed costs; three gates that killed multi-week efforts +- [references/vllm-differential.md](references/vllm-differential.md) — kernel-level comparison against a competitor trace; slower kernels vs more kernels +- [references/cuda-graphs.md](references/cuda-graphs.md) — scope, bucketing, graph-safety invariants, idle EP ranks, what a wide capture forecloses +- [references/moe-inference.md](references/moe-inference.md) — dispatchers, grouped GEMM, overlap, pad masking +- [references/mamba-and-triton.md](references/mamba-and-triton.md) — Triton rules, fused extraction, scratch sizing +- [references/host-path.md](references/host-path.md) — serialization, IPC payloads, DP routing +- [references/measuring.md](references/measuring.md) — profiler endpoints, NVTX, built-in counters, nsys flags that deadlock, union-busy idle accounting, A/B protocol, kernel-to-e2e conversion, ledger +- [references/commit-log.md](references/commit-log.md) — every commit mapped to its pattern, plus superseded APIs +- [references/updating-this-skill.md](references/updating-this-skill.md) — how to extend, correct, and prune this skill as you learn +- [assets/review-checklist.md](assets/review-checklist.md) — pre-PR checklist +- [scripts/](scripts/) — runnable trace analysis: `steady_window.py` (steady-state + per-kernel budget), `union_window.py` (union-busy, overlap ratio, gap bands), + `compare_budget.py` (two engines, matched taxonomy), `kernel_neighbors.py` (attribute + a generically-named kernel). All take a one-per-layer-per-step anchor kernel and count + back from the end of the trace. Prefer these over hand-written SQL: the windowing + arithmetic is where the analysis errors happen, not the reasoning. + +## Out of scope + +- Authoring perf or functional test recipes — use `add-inference-performance-test` +- Cluster and container bootstrap — use `cog-setup-and-help` +- Non-hot-path product features (reasoning parsers, chat-template retention) + +## Keeping this skill current + +Edit this skill at the end of a piece of work, once you have a number and a root +cause — never mid-experiment on a hypothesis. Triggers: an accepted optimization, a +**rejection** (the highest-value entry type), a negative decision gate, a tooling +failure that cost over an hour, a measurement that contradicts something here, a +flag or default that differs from what is documented, or an invariant whose +violation broke something. + +Route by subject into the reference that owns it — only new invariants, flag +behavior, and routing lines belong in this file, because `SKILL.md` is the router +and it stops working once it needs its own router. An addition must be measured, +root-caused, scoped, and actionable. Delete what is false; never delete a measured +negative result. + +Triggers, routing table, house style, deletion policy, size budgets, and the +revision log: [references/updating-this-skill.md](references/updating-this-skill.md). diff --git a/skills/optimize-inference-siddharth/assets/review-checklist.md b/skills/optimize-inference-siddharth/assets/review-checklist.md new file mode 100644 index 00000000000..ec43bcfd565 --- /dev/null +++ b/skills/optimize-inference-siddharth/assets/review-checklist.md @@ -0,0 +1,167 @@ +# Inference Optimization Pre-PR Checklist + +Copy this into the working notes and walk it before opening the PR. Most items map +to a hard rule in [SKILL.md](../SKILL.md); the rest are review feedback these +commits actually received. + +## Decision gate (before writing the code, not before the PR) + +``` +- [ ] Floor established from MEASURED machine constants, not datasheet peaks +- [ ] Current cost measured under graph replay, at the real decode token count +- [ ] Gross ceiling stated as a ratio AND as a % of the step +- [ ] Net ceiling subtracts what the fix costs (added launches, grid syncs, + atomics, extra passes) +- [ ] Mechanism identified (bytes / FLOPs / occupancy / CTA utilization / + dependent load / pure launch overhead) — not just the magnitude +- [ ] Target located on the serial chain vs already-overlapped vs off-path, so the + expected kernel-to-e2e conversion is predicted before measuring +- [ ] Verdict recorded in the ledger, including gates that came out negative +``` + +## Measurement + +``` +- [ ] Baseline captured before any change, with the config recorded +- [ ] Baseline re-run in the SAME allocation as the test arm (cross-session drift + reached 1.6%, larger than most wins) +- [ ] OFF and ON arms run back to back, and the pair repeated at least once +- [ ] Arms do not overlap: slowest ON beats fastest OFF (report pairwise deltas, + not one average) +- [ ] First timed iteration checked for cold-start outlier before averaging +- [ ] After-measurement uses the identical batch size, sequence length, parallelism +- [ ] Warmup ran long enough that CUDA-graph capture is excluded from both numbers +- [ ] The win is attributed to a specific kernel or host span, not just end-to-end +- [ ] Mechanism confirmed by a LAUNCH COUNT change, not only by a bucket time — + counts are sequence-length invariant, per-bucket times are not +- [ ] If two traces are compared: buckets the change cannot touch held still. If they + moved, the windows sit at different average KV length and no time is comparable +- [ ] Measured conversion is consistent with the predicted one; if not, the target + was mis-located +- [ ] One mechanism per PR, so the measurement attributes cleanly +- [ ] Profiler flags known-good (no osrt, no process-tree sampling, no NVTX under + graph capture) or the capture will not finalize +``` + +## Host path + +``` +- [ ] No .item() / .tolist() / .cpu() added to per-step code +- [ ] No dataclasses.asdict() on any object that can hold a tensor +- [ ] No torch.save / pickle on the IPC path +- [ ] New per-request fields on the wire are scalars, or opt-in behind a flag +- [ ] Work not needed for the next step is off the step loop +- [ ] New host work on the step loop has an NVTX range around it +``` + +## CUDA graph safety + +``` +- [ ] Per-step scalars go into preallocated fixed-address GPU tensors via fill_/copy_ +- [ ] No new buffer allocation inside a captured region +- [ ] Grow-only buffers pre-sized to the worst case BEFORE capture +- [ ] No Python-level assignment in forward() relied on at replay time (use copy_) +- [ ] Nothing stateful or RNG-bearing newly captured without checking what gets + frozen by value +- [ ] Per-step work bounded by the matched bucket, not global max_tokens/max_requests +- [ ] Bucket coverage verified: confirm real steps match a graph rather than + silently falling back to eager +``` + +## Padding + +``` +- [ ] Pad rows route to expert -1 (no expert activated) +- [ ] Pad tokens point at reserved dummy storage, not real blocks +- [ ] Kernels skip rows where permutation_map == -1 +- [ ] No zeroing pass over rows past valid_tokens (don't write them at all) +- [ ] Idle EP ranks use the lightweight dummy path, never add_request +``` + +## Triton kernels + +``` +- [ ] tl.constexpr only on values fixed for the process lifetime +- [ ] Per-step ints either do_not_specialize or passed as a 0-d GPU tensor +- [ ] Grid sized to worst case, body gated by `if pid >= real_count: return` +- [ ] Any autotune goes through autotune_configs(), with a short config list +- [ ] Tile/warp/stage choice driven by a typical-batch hint, not buffer capacity +- [ ] torch.empty rather than torch.zeros where the kernel overwrites every row +``` + +## MoE specifics + +``` +- [ ] moe_router_dtype is fp32 (hard requirement for inference_optimized) +- [ ] NVLS eligibility checked via are_tensors_nvls_eligible, not re-derived +- [ ] NCCL fallback path still correct for non-NVLink / non-bf16 / unaligned shapes +- [ ] TE nn.Parameter identity preserved (param.data views, built lazily after + checkpoint load) +- [ ] Shared-expert overlap re-measured on THIS model; CTA caps not copied forward +- [ ] Unsupported config combinations rejected with a clear error, not silently wrong +``` + +## Buffer sizing + +``` +- [ ] Sized from the true per-step bound, not a loose upper bound +- [ ] DTYPE checked, not just length: does any buffer carry a tensor narrower than + itself? A collective buffer wider than its payload doubles wire bytes and adds + a cast per layer (this was worth 2.5%) +- [ ] The bound is a single shared value, not duplicated across modules +- [ ] Scratch reserved before durable pools +- [ ] Over-budget raises a config-time error naming the knob to reduce +- [ ] Assert on overrun rather than trusting the caller +``` + +## Correctness + +``` +- [ ] Reference implementation test (plain PyTorch or loop) with assert_close +- [ ] Gating asserted with a sentinel: padded slots still hold the sentinel after + the kernel runs +- [ ] Sizing formulas re-derived independently in the test, covering BOTH regimes + of any min()/max() +- [ ] Boundary cases covered (window before position 0, sub-block prefills, + non-default chunk sizes) +- [ ] Error paths tested (too-small budget raises, not OOM later) +- [ ] Golden values regenerated if bucketization or padding changed, and the + generated text sanity-checked for coherence +- [ ] Inference and training paths still agree where they should (e.g. router + picks the same experts) +- [ ] Checked whether a BIT-EXACT formulation exists (hold BLOCK_SIZE_K fixed; + masked slots adding exact 0.0) before accepting any deviation +- [ ] If not bit-exact: deviation bounded in ulps across several token counts and + seeds, temperature-0 coherence diffed against gate-OFF, every divergence + inspected for low-confidence branch + fluent + factually correct, and the + divergence recorded in the ledger +``` + +## Shipping + +``` +- [ ] Kill switch exists so the next person can A/B without reverting +- [ ] New flags documented with defaults in arguments.py and the config docstring +- [ ] Imports run through `uv run isort` +- [ ] Ledger entry appended (hypothesis, changed files/flags, throughput, delta, + correctness, run id, conclusion) — including for rejected attempts +- [ ] PR opened as a draft, commits signed with both -s and -S +``` + +## Skill maintenance (after the work, not before the PR) + +``` +- [ ] Anything that generalizes beyond this model promoted into the skill, routed + to the right reference rather than piled into SKILL.md +- [ ] Rejections and negative gate verdicts recorded with their root cause +- [ ] Anything this work proved wrong in the skill corrected in place; contradictions + scoped by hardware/version rather than overwritten +- [ ] Content that is now false deleted; measured negative results kept +- [ ] Links resolve, frontmatter parses, revision log appended +``` + +See [../references/updating-this-skill.md](../references/updating-this-skill.md). + +See `add-inference-unit-tests` for where new tests belong, `mcore-testing` for +recipe structure, and `run-inference-functional-tests` for verification on the +cluster. diff --git a/skills/optimize-inference-siddharth/references/commit-log.md b/skills/optimize-inference-siddharth/references/commit-log.md new file mode 100644 index 00000000000..b6fa11ebdee --- /dev/null +++ b/skills/optimize-inference-siddharth/references/commit-log.md @@ -0,0 +1,121 @@ +# Commit Provenance + +Every commit by `sidsingh@nvidia.com` on `main` since 2026-01-01, grouped by the +pattern it demonstrates. Use this to read the primary example of any technique: + +```bash +git show +git show --stat # scope first +git log -1 --format=%B # message and co-authors +``` + +Regenerate this list: + +```bash +git log --format='%h|%ad|%s' --date=short --since=2026-01-01 --author='sidsingh' main +``` + +## MoE inference stack + +| SHA | PR | Pattern | +|---|---|---| +| `7d1c01685` | #3496 | **Foundational.** Forks a parallel inference MoE hierarchy: `InferenceTopKRouter`, `InferenceGroupedMLP`, graph-safe AllGather dispatcher, GPU-resident offsets, fused NVLS collectives, centralized `are_tensors_nvls_eligible`. | +| `905c0e386` | #3851 | Lazy `SymmetricMemoryManager`; lazy non-destructive weight concatenation via `param.data` views to preserve TE `Parameter` identity. Fixes the RL integration. | +| `589cd9e12` | #3858 | `megatron/core/inference/moe/` package: `mcore_fused_moe`, Triton permute, padding-aware activations, fused activation+MXFP8-quantize+swizzle. Backend enum replaces a boolean. | +| `bfd45740c` | #4258 | **Architectural rewrite.** Variable-count AllGather-V / ReduceScatter-V so EP ranks carry different token counts; `fused_metadata_update` collapses 5 kernels into 1; buffers allocated once at init; dispatcher swap moves into `MoELayer.train()`. | +| `442a936a1` | #4570 | Shared-expert overlap on `SharedExpertMLP.stream`, with the AGV capped to 16 CTAs so it does not starve the side stream. | +| `c817dad28` | #4587 | Cuts the EP graph-size sync back to the minimum (max token count + is-anyone-non-decode) after it accumulated dead complexity. | +| `20f09364e` | #4603 | Replaces 25 `@triton.autotune` configs with vLLM's host-side `_get_default_config` heuristic keyed on a token-count *hint*; persistent-grid `moe_sum` that stops zeroing rows past `valid_tokens`; latent-MoE shared-expert overlap; `ep_consensus_interval`. | +| `3a253ac5c` | #4922 | `mask_routing_padding` writes `-1` into padded rows' topk slots so CUDA-graph padding activates no expert. Fixed-address `real_token_count` scalar. | + +Start with #3496 for the structure and #4258 for the current design. Read them in +that order — #4258 only makes sense as a response to #3496's equal-token-count +limitation. + +## CUDA graphs + +| SHA | PR | Pattern | +|---|---|---| +| `32efeffd2` | #3250 | Splits inference graph scope from the training `full_iteration` scope; adds block-scope graphs for the Mamba block. Establishes the `_should_call_local_cudagraph` predicate. | +| `fde3b90a8` | #3527 | `num_cuda_graphs=-1` auto-sizing with a dense small-batch ladder, `[1,2,4] + range(8,256,8) + range(256,max,16)` — fine granularity where decode batches actually live. | +| `60a25aa67` | #3525 | `add_dummy_requests_for_expert_parallel_step`: idle EP ranks `fill_` preallocated tensors instead of constructing request objects. Rewrote 1448 lines of golden values because bucketization changed. | +| `35f76df3f` | #4440 | Moves hybrid graph ownership from `HybridStack` up to `HybridModel` so embedding and output layers fall inside the capture. Widening scope as an optimization. | +| `740c16e6b` | #5797 | `cuda_graph_max_tokens` default 512, so prefill and mixed steps up to 512 tokens get a graph instead of falling back to eager. | + +## Mamba / SSM and Triton + +| SHA | PR | Pattern | +|---|---|---| +| `ab2b33d54` | #4397 | `do_not_specialize=["batch"]`; `fast_exp` via `exp2`; capability-specialized `BLOCK_SIZE_M`/`num_warps` for Blackwell; `torch.zeros` to `torch.empty` where the kernel overwrites every row. | +| `9b4074b51` | #4764 | Bounds per-step Mamba work by the padded bucket rather than the global max; retunes the varlen conv autotune list to two regimes; removes the AGV CTA cap from #4570 after measuring it net-negative. | +| `f29c747f2` | #5608 | **Four lines, high value.** Demotes per-step batch sizes from `tl.constexpr` to runtime args, ending per-step JIT recompilation. | +| `411a5d8b2` | #5863 | Scratch sized `min(ceil(max_tokens / block_size_tokens), 3 * max_requests)` — an order of magnitude at low concurrency. Also fixes a hardcoded `mamba_chunk_size = 128`. | +| `648bc011f` | #5866 | Fused gather-plus-conditional-scatter Triton kernels replacing dense gather + `copy_`; runtime `real_count` gate; transpose folded into the write address. | + +## Host path, scheduling, observability + +| SHA | PR | Pattern | +|---|---|---| +| `53a2b19a2` | #2920 | Drops `dataclasses.asdict` (recursive deepcopy) and `torch.save` from request IPC; moves detokenization to the coordinator so it overlaps the next engine step; adds the NVTX ranges that make the host critical path visible. | +| `faced5128` | #3034 | MoE routing replay: records per-token expert choices into a static graph-safe buffer for load-imbalance analysis. | +| `eadbaa618` | #5611 | `/start_profile` and `/stop_profile` endpoints relaying `cudaProfilerStart/Stop` to every engine, to pair with `nsys --capture-range=cudaProfilerApi`. | +| `edd45620b` | #5609 | Prefix-cache hit reporting as `usage.prompt_tokens_details.cached_tokens`. | +| `bcf4c8fb5` | #5607 | Load-aware DP routing: `alpha * match + (1-alpha) * free_capacity`, replacing round-robin. Vectorized numpy because it runs per request. | +| `602fad039` | #5918 | Drops `prompt_tokens` from the wire by default, keeping `prompt_length` for the usage contract. | + +## Test and CI maintenance + +Not optimization patterns, but useful for seeing what breaks when the above lands. + +| SHA | PR | Note | +|---|---|---| +| `369e0eba7` | #3071 | Fixes functional tests broken by #2920 | +| `66ec17eac` | #3357 | Inference functional test fixes | +| `a1165fabc` | #4454 | Fixes GitLab functional tests | + +## Non-performance work + +Included for completeness; these are product features on the same subsystem, not +hot-path optimizations. + +| SHA | PR | Note | +|---|---|---| +| `a79f49d37` | #5276 | Reasoning-token retention delegated to the chat template | +| `82e9dc69c` | #5634 | `nemotron_v3` reasoning parser | + +## Superseded APIs — do not copy from these + +Several commits above contain code that has since been replaced. If you read them +directly, know what no longer exists. + +**Dispatchers.** `InferenceCUDAGraphTokenDispatcher` (#3496) and the manual +`set_inference_cuda_graphed_iteration` / `unset_...` plumbing were replaced in #4258 +by `NCCLAllGatherDispatcher` / `NVLSAllGatherVDispatcher` plus the automatic +`MoELayer.train()` swap. + +**Grouped GEMM backend.** `inference_disable_torch_grouped_mm` (#3496) → +`inference_grouped_gemm_backend: 'auto'|'torch'|'te'` (#3858) → +`'flashinfer'|'torch'` (#4258) → `'flashinfer'|'torch'|'vllm'` with `vllm` as the +default (#4603). + +**Deleted modules.** `megatron/core/inference/moe/pad.py` and the `skip_permute` +branches (added #3858, removed #4258). `smallest_non_decode_cuda_graph_size` +(removed #4587). + +**Symmetric memory globals.** `_GLOBAL_SYMMETRIC_MEMORY_BUFFER_TP/_EP` in +`parallel_state.py` (#3496) → `SymmetricMemoryManager` in +`megatron/core/inference/symmetric_memory.py` (#3851). + +**CUDA graph scope.** The `cuda_graph_scope` list holding +`CudaGraphScope.full_iteration_inference` (#3250, #4440) → `cuda_graph_impl` plus +`inference_cuda_graph_scope` (`none`/`layer`/`block`). `CudaGraphScope` survives +only for checkpoint deserialization; its docstring carries the migration table. + +## Related work on branches, not `main` + +`44334edad` "Fix flashinfer sampling kernels to return int64 instead of int32" +lives on the unmerged `fix_flashinfer_sampling` branch. The same `.long()` casts +reached `main` via `650b7838` (#5791, Keshav Santhanam, co-authored by Siddharth), +which also concluded that FlashInfer sampling must run **eagerly** — its kernel +choice is data-dependent and FlashInfer bakes the philox RNG state into a graph as +a by-value constant. See [host-path.md](host-path.md). diff --git a/skills/optimize-inference-siddharth/references/cuda-graphs.md b/skills/optimize-inference-siddharth/references/cuda-graphs.md new file mode 100644 index 00000000000..32a5ee25554 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/cuda-graphs.md @@ -0,0 +1,302 @@ +# CUDA Graphs for Inference + +Decode is launch-bound. A hybrid or MoE decode step issues hundreds of small +kernels whose combined GPU time is less than the CPU time to launch them, so the +device idles between kernels. CUDA graphs replace per-step launches with one +replay. Getting them to actually engage is most of the work. + +Source commits: `32efeffd` (#3250), `fde3b90a` (#3527), `60a25aa6` (#3525), +`35f76df3` (#4440), `740c16e6` (#5797). + +## Scopes + +Two config fields control this today: + +- `cuda_graph_impl`: `none` | `local` | `transformer_engine` | `full_iteration`. + Inference graphs require **`local`**. +- `inference_cuda_graph_scope`: `none` | `layer` | `block`. + +| Scope | Graphs per step | When to use | +|---|---|---| +| `block` | One, wrapping the whole decoder block | The latency target. Requires the entire block to be graph-safe. | +| `layer` | One per transformer or Mamba layer | Fallback when block capture is not safe. Still cuts launches substantially. | +| `none` | Eager | Baseline / debugging. | + +`local` derives `layer` by default, so **set `block` explicitly** — a config that +merely says `cuda_graph_impl=local` is leaving the biggest win on the table. + +`CudaGraphModule` (`attn`, `mlp`, `moe`, `moe_router`, …) is a separate, +training-oriented axis for capturing sub-layer regions. Not used for inference. + +
+Deprecated API you will see in older commits and docs + +Commits before the refactor use a `cuda_graph_scope` *list* holding +`CudaGraphScope.full_iteration_inference`. The migration guide lives in +[enums.py](megatron/core/transformer/enums.py): + +- `full_iteration` → `cuda_graph_impl="full_iteration"` (training) +- `full_iteration_inference` → `inference_cuda_graph_scope=block` +- everything else → the equivalent `CudaGraphModule` member + +`CudaGraphScope` is retained only for checkpoint deserialization. Do not write +new code against it. +
+ +## Who owns the graph + +Two hooks decide, and they must agree — if both a block and its layers think they +own the graph you get nested capture. + +**`create_mcore_cudagraph_manager(config)`** attaches the manager. Each module +checks whether the scope selects it: + +```python +# megatron/core/models/hybrid/hybrid_model.py +def create_mcore_cudagraph_manager(self, config): + if config.inference_cuda_graph_scope == InferenceCudaGraphScope.block: + from megatron.core.transformer.cuda_graphs import CudaGraphManager + self.cudagraph_manager = CudaGraphManager(config) +``` + +**`_should_call_local_cudagraph(*args, **kwargs)`** decides per call whether to +replay or run eager. It checks the scope, that a real inference context is +present, that `attention_mask is None`, and finally +`inference_context.using_cuda_graph_this_step()`. + +Ownership by model family: + +| Model | Owner under `block` scope | +|---|---| +| Hybrid | `HybridModel` — the *model*, so embedding + stack + output layer are one capture | +| GPT | `model.decoder`, the `TransformerBlock` | +| Mamba | The `MambaStack` | + +### Widening the scope is itself an optimization + +Commit `35f76df3` moved hybrid graph ownership up from `HybridStack` to +`HybridModel`, making `HybridStack` a plain `MegatronModule` and adding +`GraphableMegatronModule` to `HybridModel`. The embedding lookup and the final +projection are individually small kernels, but their launch overhead is on the +critical path of every decode step. Folding them into the existing capture cost +nothing and removed those launches. + +**Generalizable:** when you add a new model class, ask what is still outside the +graph. Small kernels bracketing a graphed region are pure launch overhead. + +### What a wide capture costs you + +Widening is the right default, but it is a trade, not a free win, and the things it +forecloses are not obvious until you try them. On Qwen3-30B under full-iteration +inference capture, three separate levers turned out to be unreachable *because* the +capture was wide: + +| Lever | Outcome under wide capture | +|---|---| +| flashinfer sampling backend | Server crash: `RuntimeError: Generator not registered with the capturing graph`. The sampler cannot be inside the capture — see the RNG hazard below. | +| Async scheduling (`async-sched-mode=serial`) | Guarded off for EP (`Async scheduling does not support expert parallelism`); opening the guard ran at batch, but deadlocked on the first single-request decode step. The guard encodes a real limitation. | +| Comm/compute overlap on separate streams | `CUDA_DEVICE_MAX_CONNECTIONS=8` measured flat. Overlap is bounded by the captured graph's structure and data dependencies, not by hardware queue count. Chunked collective pipelining needs concurrent streams under capture and hits the same wall. | + +The async-scheduling result is worth internalizing beyond its own flag: async +overlap cannot hide the post-sampling host chain, because that chain is +**data-dependent on the current step's sampled tokens**. Measured end-to-end gain +from opening the guard was +0.85% — consistent with a genuine serial dependency +rather than a scheduling artifact. + +**Generalizable:** before proposing overlap, streams, or asynchrony as the fix for +GPU idle, check whether the idle is a *data dependency*. If step N+1's input needs +step N's sampled token, no scheduler will overlap them, and the only fix is making +that host chain cheaper. + +## Bucket coverage + +A graph replays at a fixed shape, so every step is padded up to the smallest +captured bucket that fits. Two failure modes: + +- **No bucket fits** → silent fallback to eager. A run with graphs "enabled" can + still be launch-bound. This is the single most common reason a CUDA-graph + change shows no improvement. +- **Nearest bucket is far too large** → wasted compute on padding. + +All the logic is in +[batch_dimensions_utils.py](megatron/core/inference/batch_dimensions_utils.py). +The bucket unit is `InferenceBatchDimensions(token_count, prefill_req_count, +decode_req_count)`. + +### Sizing distributions + +`EXPONENTIAL` (default) halves down from the ceiling. This bounds the graph count +at roughly `log2(max_tokens)` and the relative padding at about 2x: + +```python +sizes = set() +val = cuda_graph_max_tokens +for _ in range(num_cuda_graphs): + rounded = max(rounder, (val // rounder) * rounder) + rounded = math.ceil(rounded / tp_size) * tp_size + sizes.add(rounded) + val //= 2 + if val < 1: + break +``` + +`LINEAR` gives the dense small-batch ladder introduced by `fde3b90a`, which +mirrors vLLM's: + +```python +sizes = ( + [1, 2, 4] + list(range(8, 256, 8)) + list(range(256, cuda_graph_max_tokens + 1, 16)) +) +``` + +The reasoning: decode batches are usually small, so spend granularity at the low +end where relative padding hurts, and be coarse above 256 to bound the graph +count. Note the hygiene steps that follow in both paths — TP-align each entry, +clamp to the ceiling, and force the endpoints to be present. + +`num_cuda_graphs=-1` auto-sizes. Under `EXPONENTIAL` it derives the count from +`log2(cuda_graph_max_tokens)` plus headroom, floored at 4. + +### Two ceilings + +Decode and prefill graphs are sized separately: + +- **Decode** is always capped at `max_requests * (num_speculative_tokens + 1)` — + one token per request, or 1 + speculative. +- **Prefill and mixed** are capped at `cuda_graph_max_tokens`, default **512** + (commit `740c16e6`). Set `cuda_graph_all_prefills` to extend to `max_tokens`, + at the cost of many large graphs. + +At runtime `match_graph_config` picks `min()` over the applicable buckets and +returns `None` to signal eager. + +### Cost of more graphs + +Each bucket is a separate capture: its own activation memory in the graph mempool +plus capture time. `create_cuda_graphs` in +[dynamic_engine.py](megatron/core/inference/engines/dynamic_engine.py) runs a +full forward per bucket and logs elapsed time and memory deltas, so capture cost +is directly observable. `cuda_graph_max_tokens=512` exists to keep this bounded +by default. + +## Making code graph-safe + +### What breaks capture or replay + +1. **Dynamic shapes** — handled by bucketing plus padding. +2. **Reallocated buffers.** A graph records raw pointers. If a buffer is freed and + reallocated between capture and replay, replay writes into freed memory. Any + grow-only buffer must be pre-sized to the worst case *before* capture: + +```python +# megatron/core/inference/engines/dynamic_engine.py +# A forward larger than the capture-time size would reallocate (and free) the buffer +# whose address a captured graph still writes to on replay, corrupting whatever later +# reuses that freed block. +if getattr(model_config, "sequence_parallel", False): + max_ag_numel = self.context.max_tokens * model_config.hidden_size + get_global_memory_buffer().get_tensor((max_ag_numel,), model_config.params_dtype, "mpu") +``` + +3. **Host syncs and data-dependent Python** inside the captured region. +4. **Python-level assignment inside `forward`.** Under `block` scope it executes + only during capture, so replays never see it. Use `copy_()` into a + preallocated buffer instead — this is exactly why + `mtp_decoder_hidden_states` exists. +5. **State baked in by value at capture.** The subtlest failure, because it + produces no error. FlashInfer's sampling kernels bake the philox RNG state into + the graph as a by-value constant, so a captured sampler replays identical + "random" numbers every step. Its kernel choice is also data-dependent. For both + reasons FlashInfer sampling is deliberately left eager — see + [host-path.md](host-path.md). Before widening a graph over anything holding RNG + or counter state, check what gets frozen. + +### The toolkit + +**One contiguous buffer, one memcpy.** +[ContextGPUView](megatron/core/inference/contexts/gpu_view.py) is the only +interface GPU code uses to read context state. Every field is a `view(dtype)` +onto a slice of a single `uint8` buffer, mirroring a pinned CPU buffer with the +identical layout, so publishing a step's bookkeeping is one `cudaMemcpyAsync` +rather than one per field. The convention is worth internalizing: + +``` +context.foo -> CPU (source of truth, used by bookkeeping) +context.gpu_view.foo -> GPU (snapshot, used by forward pass) +``` + +**Cache views instead of reconstructing them.** Slicing and unsqueezing per step +constructs new `TensorImpl`s at 30-60us each: + +```python +# megatron/core/inference/contexts/dynamic_context.py +# Instead of slicing and unsqueezing on every new inference step (constructing +# new TensorImpls at 30-60 us), we fix the underlying storage so views are +# reusable across steps. +self._input_position_views: Dict[int, Tuple[Tensor, Tensor]] = {} +``` + +**Point padding at reserved storage.** Pad tokens get `dummy_block_idx` so the +KV-append kernel writes to a valid throwaway block. Pad rows get routing index +`-1` so they activate no expert. + +**Keep EP sync off the compute stream.** `adjust_batch_dims_for_expert_parallelism` +can do the cross-rank max over ZMQ on the CPU, avoiding both a per-step NCCL +all-reduce on the compute stream and the H2D/D2H pair around it. + +## Expert parallelism: idle ranks still have to run + +Under EP, every rank must issue the same collectives in lockstep, and all ranks +must select the *same* graph bucket. So a rank with no real requests still runs a +forward. Commit `60a25aa6` found that idle rank was building its dummy batch +through the heavyweight graph-capture warmup path, constructing real request +objects and allocating KV blocks every step. + +The fix, `add_dummy_requests_for_expert_parallel_step`, pokes only the +preallocated tensors that `initialize_attention_state` and the forward actually +read: + +```python +N = smallest_cuda_graph_dimensions.decode_req_count +dummy_block_idx = self.block_allocator.dummy_block_idx +self.total_request_count = N +self.active_token_count = N +self.num_prefill_requests = 0 +self.request_query_lengths[0:N].fill_(1) +self.request_kv_length_offsets[0:N].fill_(0) +self.request_to_kv_block_ids[0:N, 0] = dummy_block_idx +self.token_to_block_idx[0:N] = dummy_block_idx +``` + +Hybrid models also need Mamba slots, so `MambaMetadata.batch_allocate_slots(n)` +grabs them in a batch without allocating. + +**Generalizable:** for any lockstep collective, idle participants must still run, +but give them the cheapest valid inputs. Reserve the dummy block and slots once, +then only `fill_` preallocated tensors. Never let the idle path allocate or +construct objects. + +The EP token-count sync all-reduces the max so every rank agrees on the bucket, +and returns `None` (eager everywhere) if any rank is non-decode. Commit `c817dad2` +cut this back to the minimum after it had accumulated dead complexity: + +```python +(max_token_count, max_is_non_decode) = ep_zmq_communicator.sync_all_reduce_max( + local_batch_dims.token_count, int(is_non_decode) +) +``` + +## Expect golden values to change + +Changing bucketization changes which padded shape each real step maps onto, which +changes padding tokens and MoE routing masks, which shifts logits. Commit +`60a25aa6` rewrote a 1448-line golden values file for exactly this reason. That is +expected — regenerate, and sanity-check that the *generated text* is still +coherent rather than assuming a diff means a bug. + +## Validation gates + +Under `block` scope, fp8 requires `--transformer-impl=inference_optimized` and +`--fp8-recipe=mxfp8`. `cuda_graph_impl=local` requires +`inference_dynamic_batching_num_cuda_graphs` to be positive or `-1`. Both are +asserted in [arguments.py](megatron/training/arguments.py). diff --git a/skills/optimize-inference-siddharth/references/decision-gates.md b/skills/optimize-inference-siddharth/references/decision-gates.md new file mode 100644 index 00000000000..c16f6a6052d --- /dev/null +++ b/skills/optimize-inference-siddharth/references/decision-gates.md @@ -0,0 +1,237 @@ +# Decision Gates: Prove the Ceiling Before You Build + +Profiling tells you where time is spent. It does not tell you whether that time is +**recoverable**. Those are different questions, and the gap between them is where +whole sessions get lost — writing a kernel whose best possible outcome was never +worth the week. + +A decision gate is a short, measurement-backed answer to one question: *if this +optimization worked perfectly, how much would the step actually get faster?* If +the answer is small, or if the mechanism you plan to attack is not the mechanism +that is slow, you stop before writing production code. + +Source: the Qwen3-30B-A3B EP4 campaign ledger +(`skills/run-qwen-model/EXPERIMENTS.md`), where three gates each killed a +multi-session effort, and a fourth picked the right candidate out of six. + +## Why share is not headroom + +The instinct is to rank categories by their share of device time and attack the +top one. That instinct is wrong twice over on this stack. + +**Some categories are already at a hardware floor.** A category can be 33% of +device time and have 2% of headroom, because it is moving the bytes it has to +move. Ranking by share puts it first; ranking by headroom puts it last. + +**Device time is not wall time.** Kernels overlap across streams, so the step's +wall clock is the *critical path* through the per-layer dependency chain, not the +sum of kernel durations. A category with a large device-time share that sits on a +side stream, already overlapped with compute, contributes far less to wall time +than its share suggests. Conversely, a tiny kernel on the serial chain costs its +duration *plus* the dispatch gap behind it, every layer. + +So the gate has to produce a **wall-time ceiling**, not a device-time share. + +## The gate, in four steps + +### 1. Establish the floor + +Compute what the operation could not go faster than, from first principles, using +*measured* machine constants rather than datasheet peaks. + +For a memory-bound op, the floor is bytes / achievable bandwidth. Measure the +achievable bandwidth yourself with a streaming-read microbenchmark on the same +device — datasheet HBM numbers overstate it, and being wrong here invalidates the +whole gate. + +For a latency-bound op, decompose the measured time into launch, synchronization, +and transfer, and identify which term dominates. An empty-kernel launch and a +barrier-only kernel are the two calibration points you need. + +### 2. Measure the current cost the same way + +Under **CUDA-graph replay**, not eager, and at the token count decode actually +runs at. Eager per-kernel timings and replay timings differ by more than the +effect sizes you are chasing. + +### 3. Ceiling = current / floor, then subtract what the fix costs + +The gross ceiling is the ratio. The net ceiling subtracts the machinery the fix +introduces: an added zeroing launch, a grid sync, extra atomics, a second +streaming pass. Fusions frequently give back most of the gross win this way, and +occasionally more than all of it. + +Then convert to a fraction of the step, because that is the number that decides +whether to proceed. + +### 4. Write down a verdict with a threshold + +State the gross ceiling, the net ceiling, the fraction of the step, and either +**proceed** or **gated out**, with the reason. Record it in the ledger even when +the answer is "don't build this" — especially then. A recorded negative gate is +what stops the next person re-deriving it. + +## Calibrate the fixed cost of a launch + +Several gates reduce to "is removing this launch worth it," so measure the two +constants once per platform and reuse them. + +On GB200 under full-iteration graph capture, the campaign measured: + +| Constant | Value | How | +|---|---|---| +| Empty-kernel launch | 0.72 µs | Time a no-op kernel under replay | +| Inter-kernel dispatch gap | 0.55 µs | Median gap between consecutive kernels in the trace | +| Floor per launch | 1.27 µs | Sum of the two | +| Graph node cost | ~0.17 µs | `cudaGraphLaunch` 199.1 µs / 1158 nodes | + +The consequence: **a removed launch is worth its own duration plus the dispatch +gap behind it**, so eliminating a 1.5 µs kernel buys ~2.05 µs of wall time, not +1.5. And no kernel can be optimized below 1.27 µs — a 1.33 µs kernel has no +headroom left, however inefficient its body looks. + +## Diagnose the mechanism, not just the magnitude + +A gate that identifies the wrong *cause* is as expensive as no gate. The routing +token-count kernel measured 7.52 µs per layer — clearly anomalous for its byte +count — and the first attempt rewrote its reduction, replacing per-pair +`atomic_add` with a `tl.histogram` variant. It measured **0.96×, a wash**, and +the reason was that atomic contention was never the problem: at `BLOCK_SIZE=1024` +only **2 of the kernel's 152 CTAs received any work**. The cost was launch +overhead and CTA underutilization, which an in-kernel reduction rewrite cannot +touch. + +Before attacking a slow kernel, establish *why* it is slow: bytes, FLOPs, +occupancy, CTA utilization, a dependent load serializing a branch, or pure launch +overhead. The fix follows from the mechanism, and the wrong mechanism produces a +technically correct kernel that changes nothing. + +## Case study 1: the grouped GEMM that could not win + +**The proposal.** MoE expert GEMM was the largest single category at ~33-40% of +decode device time. The obvious plan was a hand-written CUTLASS/CuTe grouped GEMM +to replace the Triton path — a multi-week effort. + +**The gate.** Weight traffic is 302 MB per layer per rank. A measured streaming-read +ceiling of 6.081 TB/s gives a **49.66 µs floor**. Production FC1+FC2 measured +**72.13 µs**. So the entire ceiling for *any* implementation is **1.45×** — and +achieved throughput on valid FLOPs was 63-83 TFLOP/s, about 3% of the device's +BF16 peak, confirming the op is nowhere near compute-bound. + +Padding waste was real but nearly free: 74.3% dead rows at `BLOCK_M=64`, yet +cutting the padding ratio from 3.89× to 1.49× buys only ~1.2×, because the dead +rows re-read weights that are already resident. + +**The verdict.** A hand-written GEMM's whole ceiling was 1.45×, and **1.26× of it +was reachable by retuning Triton tiles alone**. The kernel effort was gated out; +the tile retune shipped instead and delivered **+4.32% end-to-end** for a +day's work. A later profile confirmed the gate independently: expert GEMM ended up +only 229 µs/step — 2.3% of the step — above the weight-bandwidth floor. + +**Generalizable:** roofline the category *before* proposing an implementation +swap. A memory-bound op with a 1.45× ceiling does not deserve a new kernel; it +deserves better tiles. + +## Case study 2: the barrier that was a hardware floor + +**The proposal.** Exposed NVLS EP communication was ~11.7% of device time and +100% exposed (the interval union equalled the sum). Plans on the table: pipeline +the collective in chunks, or fuse ReduceScatter into the FC2 epilogue. + +**The gate.** Decomposing the collectives showed they are **latency-bound, not +bandwidth-bound**: + +``` +AllGather-V 6.57 µs = 0.72 launch + 5.08 barrier + 0.77 transfer +ReduceScatter 7.88 µs = 0.72 launch + 5.03 barrier + 2.13 transfer +``` + +Bytes accounted for only 127 µs of a ~693 µs/step total, and the +ReduceScatter transfer was already at 82% of the NVLink floor — so batching bytes +cannot help. A follow-up microbenchmark isolated the barrier further: a +barrier-only kernel costs 5.75-5.83 µs against a 0.72 µs empty kernel, so the +5.05 µs **is the four-way system-scope flag round trip itself**, not polling +granularity. Rewriting the spin (`atom.cas` → `ld.acquire.sys` poll) measured a +**2.3% regression**, exactly as the gate predicted. + +Of 1060 µs/step, 632 µs was intrinsic and 428 µs was **inter-rank skew** — ranks +waiting on the slowest rank's expert GEMM, which is a routing-balance problem, not +a communication problem. + +**The verdict.** Recoverable critical path: **6.4-7.0% of the step**, and each +candidate gave most of it back. Epilogue fusion keeps the barrier (≈1.4% left); +2-chunk pipelining *adds* a barrier per chunk and needs concurrent streams under +graph capture, which full-iteration capture forbids. CTA count was already at its +optimum. Lever abandoned. + +**Generalizable:** decompose a collective into launch + barrier + transfer before +optimizing it. If the barrier dominates, you are looking at a fabric round trip +and no kernel rewrite will move it. Check whether the residual is really *skew*, +in which case the fix is load balance somewhere else entirely. + +## Case study 3: pricing six fusions before writing one + +**The proposal.** The routing/permute category was 1251 µs/step across 242 +kernels — visibly a launch storm, with several plausible fusions. + +**The gate.** Break the category down per kernel name — launches, device µs, +dispatch gap — then price every candidate against the 1.27 µs per-launch floor: + +| Kernel | µs | × per step | +|---|---:|---:| +| `_moe_sum` | 7.79 | 48 | +| `_count_local_tokens` | 7.52 | 48 | +| `gatherTopK` | 6.05 | 48 | +| `_scatter_token_indices` | 2.66 | 48 | +| router softmax | 1.94 | 48 | +| `_prefix_fill_init` | 1.33 | 48 | + +Candidate ceilings, gross and net: + +| Candidate | Gross | Net after costs | Verdict | +|---|---:|---|---| +| `_moe_sum` → FC2 epilogue | 400 µs (4.03%) | ~3.3%, possibly negative — needs cross-CTA fp32 atomics **plus** a zeroing launch | gated out | +| Cooperative-grid merge of count→fill→scatter | 122 µs (1.23%) | less, after two grid syncs | gated out | +| Fold count + its `zeros` fill into `_prefix_fill_init` | 449 µs (4.52%) | no grid sync, integer-exact | **chosen** | + +**The verdict.** The chosen fusion shipped at **+3.01% end-to-end**. The +`_moe_sum` epilogue candidate — superficially the most attractive, since +`_moe_sum` was the single most expensive routing kernel — had the exact failure +shape of an earlier rejected mega-fusion: cross-CTA atomics plus a zeroing pass +eating the gain. It was never built. + +Notably, `_moe_sum` was later improved anyway, but by a different mechanism the +gate exposed: its per-topk-slot locality test was a uniform scalar branch on a +*dependent* load, serializing the topk walk. Predicating it into the load mask +gave 1.36× and **+1.83%** end-to-end, bit-exact, with no fusion at all. + +**Generalizable:** when a category has several fusion candidates, price all of +them before building any. Rank by *net* ceiling, and prefer candidates that need +no grid sync and no atomics — those two costs are what turn a 4% gross win into a +0% net win. + +## Verdict template + +Record this in the ledger for every gate, including the negative ones: + +``` +GATE — +Question: +Floor: +Current: +Gross ceiling: +Fix costs: +Net ceiling: <% of step after costs> +Mechanism: +Verdict: PROCEED | GATED OUT — +``` + +## When to skip the gate + +Gates cost time too. Skip straight to measurement when the change is cheap and +reversible: flipping an existing flag, retuning tile sizes, swapping a backend +enum. Those are their own experiments — a gate on a config flip costs more than +the flip. + +Gate anything that would take more than about a day to build, anything requiring +a new kernel, and anything whose category you have not yet rooflined. diff --git a/skills/optimize-inference-siddharth/references/host-path.md b/skills/optimize-inference-siddharth/references/host-path.md new file mode 100644 index 00000000000..14b27be5da8 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/host-path.md @@ -0,0 +1,225 @@ +# The Host Path: Per-Step CPU Overhead and IPC + +Once the GPU forward is graphed and fused, the bottleneck moves to the CPU. Every +engine step ends with bookkeeping, detokenization, serialization, and a ZMQ send, +all of which block the next step. This is where several of the largest wins came +from, and it is the part people forget to profile. + +Source commits: `53a2b19a` (#2920), `602fad03` (#5918), `bcf4c8fb` (#5607), +`edd45620` (#5609), `650b7838` (#5791, co-authored). + +## Where per-step host work lives + +Always inspect these four places. + +**The engine step tail** in +[dynamic_engine.py](megatron/core/inference/engines/dynamic_engine.py). The +critical section is bracketed by NVTX ranges so it is directly visible in a trace: +`bookkeeping` → `detokenization` → `coordinator_communication`. Anything added here +runs once per step. + +**The controller**, `text_generation_controller.py`: routing-record bookkeeping, +log-prob detokenization, and building the step-result dict. CPU work here overlaps +already-enqueued GPU kernels, which makes it the right place to *put* unavoidable +work. + +**Serialization** in +[inference_request.py](megatron/core/inference/inference_request.py): +`InferenceRequest.serialize`, `DynamicInferenceRequest.serialize`, +`DynamicInferenceRequestRecord.serialize` and `.merge()`. Runs once per finished +request per step. + +**The coordinator** in +[data_parallel_inference_coordinator/](megatron/core/inference/data_parallel_inference_coordinator/): +per-request routing scoring in `coordinator.py`, pending-count updates in +`handlers.py`. + +## Serialization: the two rules + +### Never `dataclasses.asdict()` on anything holding tensors + +`asdict()` recursively deepcopies every field. With CUDA tensors in the object +that is catastrophic. Use a shallow dict copy and handle fields explicitly: + +```python +# megatron/core/inference/inference_request.py +# Dataclass to dict. +# do not use asdict(self) - it has very high CPU overheads +# and if there are tensors, it will try to deepcopy them +obj = self.__dict__.copy() # shallow dict copy +obj["status"] = self.status.name if self.status else None +obj["sampling_params"] = self.sampling_params.serialize() if self.sampling_params else None +``` + +### Never `torch.save` for IPC + +`torch.save` into a `BytesIO` pickles and writes the full tensor blob. For token +ids, a list is dramatically cheaper: + +```python +def serialize_tensor(tensor: torch.Tensor) -> List: + nvtx_range_push("serialize_tensor") + # simply convert tensor into a list + tensor = tensor.cpu().tolist() + nvtx_range_pop("serialize_tensor") + return tensor +``` + +Note the NVTX range on a function this small. That is deliberate — you cannot +attribute host cost you cannot see. + +## Don't put large tensors on the wire by default + +The engine was serializing and shipping the entire `prompt_tokens` tensor +engine → coordinator → API for every finished request. For long agentic or RL +prompts that dominates wire cost, and the client usually does not want the ids +echoed back. + +Commit `602fad03` made it opt-in: + +```python +# megatron/core/inference/sampling_params.py +# Echo prompt token ids back in the response. When False (default), the engine +# drops prompt_tokens before serializing the finished request, saving the ZMQ +# transmission cost for long prompts. Opt in when the client needs them. +return_prompt_tokens: bool = False +``` + +The API contract still needs `usage.prompt_tokens`, so the technique is to **keep +the scalar and drop the payload**: a `prompt_length: Optional[int]` field is always +populated during serialize even when the tensor is omitted. The drop is wire-only +— `serialize` nulls `self.prompt_tokens` around the `super().serialize()` call and +restores it, so the local object is unchanged. + +Endpoints opt in where they must: `/v1/completions` always echoes +`prompt_token_ids`, so it sets `return_prompt_tokens=True` unconditionally. + +**Generalizable:** for any per-request field on the IPC path, ask what the consumer +actually needs. Usually it is a count or a length, not the tensor. + +## Move CPU work off the engine, don't just make it faster + +The engine was detokenizing the full prompt plus generated sequence for every +finished request, inline in the step loop. Making detokenization faster would have +helped a little; moving it to a different process helped much more, because it then +overlaps the next engine step: + +```python +# Detokenize all finished requests if not using +# the coordinator. Otherwise, the coordinator will +# overlap detokenization with the engine. +if not self.use_coordinator: + nvtx_range_push("detokenization") + ... +``` + +The coordinator gained a tokenizer and a `detokenize()` method; the engine ships +raw token ids. When a coordinator is present the engine does no detokenization at +all. + +**Generalizable:** work that is not needed to compute the *next* step does not +belong in the step loop. Either defer it or move it to a process that runs +concurrently. + +## Load-aware DP routing + +With data-parallel engines, the coordinator decides which rank gets each request. +Round-robin ignores both prefix-cache affinity and current load, so it can send a +request to a busy rank while an idle rank holds exactly the KV blocks that request +needs. + +`bcf4c8fb` replaced round-robin with `LOAD_BALANCED`, which is now the default +policy, and unified everything under one score: + +```python +# megatron/core/inference/data_parallel_inference_coordinator/coordinator.py +if self.prefix_caching_coordinator_policy == PrefixCachingCoordinatorPolicy.LOAD_BALANCED: + return self.get_least_loaded_data_parallel_rank() + +# Without prefix caching (or when the request has no hashes to match on) +# fall back to load-balanced routing. +if not self.enable_prefix_caching or not request_hashes: + return self.get_least_loaded_data_parallel_rank() + +match, recency = self._match_vector(request_hashes) +alpha = self.prefix_caching_routing_alpha + +# Vectorized score: alpha * match + (1-alpha) * free_capacity_fraction. +free_slots = np.maximum(0, self.max_requests - self._pending_counts).astype(np.float64) +scores = alpha * match + (1.0 - alpha) * (free_slots / self.max_requests) + +# Tiebreak: highest score, then highest recency, then lowest rank index. +n_ranks = len(self._identities_list) +order = np.lexsort((np.arange(n_ranks), -recency, -scores)) +``` + +`prefix_caching_routing_alpha` (default 0.5) is the tradeoff knob: + +| alpha | Behavior | Risk | +|---|---|---| +| 0 | Pure load balance | Ignores cache; re-prefills work another rank already has | +| 0.5 | Balanced (default) | — | +| 1 | Pure cache affinity | Piles affine requests on one rank, starves the others | + +Two implementation details worth copying. `match` is policy-dependent and +normalized to `[0, 1]` — binary for `first_prefix_block`, normalized prefix depth +for `longest_prefix` — so it is commensurable with the load term. And the whole +computation is vectorized numpy, because it runs per request on the coordinator's +hot path; a Python loop over ranks here would be its own bottleneck. + +Pending counts are maintained incrementally in `handlers.py` (incremented on +dispatch, decremented on finish), and `_remove_engine` rebuilds them when an engine +disconnects mid-flight. + +## Hold the dtype contract at kernel boundaries + +FlashInfer's sampling kernels return sampled token ids as `int32`. The rest of the +pipeline — token buffers, `input_ids`, scatter into KV and embedding gather — is +`int64`. Every return in +[flashinfer_sampling.py](megatron/core/inference/sampling/flashinfer_sampling.py) +is cast with `.long()`. + +A mismatch inserts implicit conversion kernels on the per-step path. Normalize +dtype where the external kernel enters rather than letting it propagate: the cast +is free once, the mismatch costs every step. + +> Provenance note: Siddharth's `44334eda` ("Fix flashinfer sampling kernels to +> return int64 instead of int32") sits on the unmerged `fix_flashinfer_sampling` +> branch. The same casts reached `main` through #5791 (`650b7838`), which he +> co-authored. + +### Corollary: not everything belongs in a graph + +That same PR concluded FlashInfer sampling should run **eagerly**, and the reasons +generalize to any stateful or random kernel you are tempted to capture: + +``` +The sampler runs eagerly. Its kernel choice is data-dependent (it varies with +which filters the batch uses), so it cannot be captured in a CUDA graph; running +eagerly also lets the controller's seeded RNG generator advance its philox offset +normally between steps -- fresh randomness per step, reproducible from the seed. +(FlashInfer bakes the philox state into a graph as a by-value constant at capture, +so a captured sampler replays identical random numbers ...) +``` + +Two distinct hazards: **data-dependent kernel selection** cannot be captured at +all, and **state baked in by value at capture** produces a graph that replays +silently wrong — identical "random" numbers every step, with no error. Before +widening a graph over anything holding RNG or counter state, check what gets frozen. + +Note also how the dispatch flags are read: from the pinned CPU sampling metadata, +so evaluating them costs no GPU sync. Same rule as everywhere else in this skill. + +## Host-side rules, condensed + +1. No `dataclasses.asdict()` on objects holding tensors. +2. No `torch.save` / pickle for IPC; use lists or ndarrays. +3. Drop large tensors from the wire by default; ship the scalar the contract needs. +4. Work not needed for the next step goes off the step loop, ideally to a + concurrent process. +5. Route with load awareness, not blind round-robin; keep the scoring vectorized. +6. Normalize dtypes at external kernel boundaries. +7. NVTX-annotate anything you suspect, including small functions. Unmeasured host + cost is invisible cost. +8. Any per-step recording must use preallocated static buffers and respect + `using_cuda_graph_this_step()`. diff --git a/skills/optimize-inference-siddharth/references/mamba-and-triton.md b/skills/optimize-inference-siddharth/references/mamba-and-triton.md new file mode 100644 index 00000000000..d7f24f63015 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/mamba-and-triton.md @@ -0,0 +1,321 @@ +# Mamba / SSM Inference and Triton Production Rules + +Hybrid Mamba models (Nemotron-H, Nano) push most inference time through hand- +written Triton kernels, so kernel hygiene matters more here than anywhere else in +the stack. The Triton rules in this document are not Mamba-specific — they apply +to any kernel on the inference path. + +Source commits: `ab2b33d5` (#4397), `f29c747f` (#5608), `9b4074b5` (#4764), +`411a5d8b` (#5863), `648bc011` (#5866). + +## The Triton specialization trap + +This is the highest-value rule in the skill relative to how easy it is to get +wrong. Commit `f29c747f` is four lines: + +```diff + def _tensor_get_slice_after_kernel( +- INPUT_BATCH_SIZE: tl.constexpr, +- OUTPUT_BATCH_SIZE: tl.constexpr, ++ INPUT_BATCH_SIZE, ++ OUTPUT_BATCH_SIZE, + ROW_SIZE: tl.constexpr, + BLOCK_SIZE: tl.constexpr, +``` + +A `tl.constexpr` parameter is **baked into the compiled kernel**, so every +distinct value produces a separate compilation. `INPUT_BATCH_SIZE` is the number +of active requests, which changes essentially every step. The kernel was +JIT-compiling on the hot path, on nearly every step. + +`ROW_SIZE` and `BLOCK_SIZE` correctly stay `constexpr` — they depend only on the +fixed state shape and are needed for `tl.arange` bounds and loop unrolling. + +Current state in +[tensor_ops.py](megatron/core/inference/contexts/attention_context/triton/tensor_ops.py): + +```python +@triton.jit +def _tensor_get_slice_after_kernel( + INPUT_TENSOR, + OUTPUT_TENSOR, + POS_ON_DEVICE, + INPUT_BATCH_SIZE, + OUTPUT_BATCH_SIZE, + ROW_SIZE: tl.constexpr, + BLOCK_SIZE: tl.constexpr, +): +``` + +### The full decision rule + +| Value | How to pass it | +|---|---| +| Block/tile size, `tl.arange` bound, unroll count, `d_conv` | `tl.constexpr` | +| Per-step count, but you want one compilation | plain arg + `do_not_specialize` | +| Per-step count, and the kernel is CUDA-graphed | fixed-address GPU tensor, `tl.load` it | + +The subtle middle case: **even plain `int` arguments get specialized** by +`@triton.jit` on divisibility-by-16 and `== 1`. So demoting from `constexpr` is +not always sufficient. If the value varies per step, be explicit: + +```python +# megatron/core/ssm/ops/mamba_ssm.py +@triton.jit(do_not_specialize=["batch"]) +def _selective_scan_update_kernel( +``` + +The third case is strictly better when applicable: passing the value as a 0-d +tensor and loading it inside the kernel avoids specialization entirely *and* keeps +CUDA graphs valid, since the address is fixed while the value may change between +replays. + +## Never autotune on the hot path + +`@triton.autotune` times every config on first sight of each new key. In +production that is a multi-second stall and a source of nondeterminism. + +Two mitigations, both in use: + +**Keep config lists short.** Commit `9b4074b5` trimmed the varlen causal conv from +four configs to two, keeping one from each regime. The comment explains why both +regimes exist, which is the useful part: + +```python +# megatron/core/ssm/ops/causal_conv1d_varlen.py +# Two block-dim regimes: +# 1. vLLM-style: small BLOCK_T, large BLOCK_C, pipelined. Many small programs maximize +# GPU occupancy, BLOCK_C=256 fully fills HBM transactions, sequence-pure programs +# avoid in-block boundary branching. Usually wins at moderate-to-large conv_dim. +# 2. Large-block fallback: bigger tiles, fewer programs. Can win for small conv_dim +# where vLLM's regime over-parallelizes, or when launch overhead dominates. +``` + +**Route everything through `autotune_configs`.** In +[determinism.py](megatron/core/ssm/ops/determinism.py), this is the single funnel: +in deterministic mode it either enables cached autotuning +(`TRITON_CACHE_AUTOTUNING=1`, Triton >= 3.4.0) or picks the single cheapest config +by `block_product * stages`, tie-broken on warp count — so exactly one +compilation and no timing sweep. + +```python +def autotune_configs(configs): + if not configs or not use_deterministic_mode(): + return configs + if TRITON_HAS_CACHE_RESULTS and os.environ.get("TRITON_CACHE_AUTOTUNING") == "1": + return configs + ... + return [min(configs, key=_estimate_config_cost)] +``` + +Wrap any new autotuned kernel in `autotune_configs([...])`, not a bare +`configs=[...]`. + +## Fused gather-plus-scatter instead of materialized intermediates + +Prefix caching needs to extract Mamba states at block boundaries. The original +implementation was a three-step tensor dance: the chunk scan gathered requested +chunks internally, `mamba_mixer` copied the result into scratch, and conv windows +were extracted with a separate PyTorch gather plus `clamp_` plus `transpose` plus +`copy_`. That materialized an intermediate tensor and made two passes over HBM. + +Commit `648bc011` replaced it with two kernels in +[intermediate_extraction.py](megatron/core/ssm/ops/intermediate_extraction.py). +The module docstring states the design contract: + +``` +These replace the two-step ``states[indices]`` (dense gather) + ``.copy_()`` +(scratch write) pattern with a single kernel that: + +1. Reads a runtime ``real_count`` from a fixed-address GPU tensor. +2. For each slot ``i < real_count``, gathers the source row indexed by the + per-slot index/position and writes it directly into the destination scratch. +3. For each slot ``i >= real_count``, returns immediately (no work, no write). + +This is CUDA-graph safe: the launch grid is sized at capture time to the maximum +possible slot count, but per-program execution is data-conditional on the +runtime ``real_count``, so padded slots cost almost nothing. +``` + +The gating idiom, which generalizes to any graphed kernel with a variable count: + +```python +pid_slot = tl.program_id(0) +pid_col = tl.program_id(1) + +real_count = tl.load(real_count_ptr).to(tl.int32) +if pid_slot >= real_count: + return +``` + +The conv kernel adds two more techniques worth stealing: + +**Fold the transpose into the write address.** The old code did +`.transpose(1, 2).copy_()`; the kernel just indexes the destination differently: + +```python +for j in tl.static_range(D_CONV): + p_raw = abs_pos - D_CONV + j + p = tl.maximum(0, tl.minimum(p_raw, seq_len - 1)) + src = src_ptr + p.to(tl.int64) * src_stride_s + c_idxs.to(tl.int64) * src_stride_c + dst = out_ptr + slot_base + c_idxs.to(tl.int64) * D_CONV + j + val = tl.load(src, mask=c_mask) + tl.store(dst, val, mask=c_mask) +``` + +**`D_CONV` is `constexpr` so the window loop is `tl.static_range`** — fully +unrolled, no dynamic loop. This is the legitimate use of `constexpr`: `d_conv` is +a model constant. + +The scan correspondingly got simpler. `ssd_combined.py` no longer gathers +internally; `return_raw_states` hands the caller the full state tensor and the +caller extracts what it needs. + +### Publishing the count + +```python +# megatron/core/inference/contexts/attention_context/mamba_metadata.py +# Publish real_count to the fixed-address GPU tensor the scatter +# kernels consult. fill_ is async (no host sync) and keeps the tensor +# at the same address captured graphs reference. +self._intermediate_real_count_buffer.fill_(self.intermediate_count) +self.intermediate_real_count = self._intermediate_real_count_buffer +``` + +## Right-size scratch to the real per-step bound + +The extraction scratch was sized `MAX_INTERMEDIATE_OFFSETS_PER_REQUEST (=3) * +max_requests`. But a step processes at most `max_tokens` tokens, and a state can +only be extracted at a block boundary — one per `block_size_tokens`. So there are +two independent bounds and the truth is the tighter one: + +```python +# megatron/core/inference/contexts/dynamic_context.py +# Per-step upper bound on Mamba intermediate-state extractions, shared with +# MambaMetadata and MambaSlotAllocator so scratch/metadata buffers and the +# budget accounting agree. Bounded both by the token budget (one block +# boundary per block_size_tokens) and by the request budget +# (MAX_INTERMEDIATE_OFFSETS_PER_REQUEST per request); +token_based_count = math.ceil(self.max_tokens / self.block_size_tokens) +request_based_count = MAX_INTERMEDIATE_OFFSETS_PER_REQUEST * self.max_requests +self.max_mamba_intermediate_states_per_step = min(token_based_count, request_based_count) +``` + +For low-concurrency long-context configs this is an order of magnitude: 1 request +with 16384 tokens once reserved 65 scratch slots a single request could never +fill. Since scratch is reserved *before* the durable prefix cache, every wasted +slot directly shrinks the cache. + +Two structural points that make this pattern work: + +**One value, shared everywhere.** `max_mamba_intermediate_states_per_step` flows +into `MambaMetadata(max_intermediate_count=...)`, +`MambaSlotAllocator.max_intermediate_count`, and the byte budget. When a bound is +duplicated it drifts; when it is shared, a wrong bound fails loudly. + +**Fail at config time, not at OOM.** If reserving scratch leaves fewer than one +durable slot, the context raises `ValueError` naming the budget and advising which +knob to reduce. + +The same commit fixed a hardcoded `mamba_chunk_size = 128`, now read from config. +States can only be extracted at multiples of the chunk size the kernel actually +runs with, so a hardcoded value silently skipped valid boundaries for any other +chunk size. + +## Bound per-step work by the bucket, not the global max + +Commit `9b4074b5` is essentially this rule applied to Nemotron prefill. The +extraction bookkeeping looped over the *global* `max_intermediate_count`, so every +prefill step did work proportional to the largest possible batch even in a tiny +graph bucket: + +```diff +- max_count = self.max_intermediate_count ++ max_count = padded_prefill_count * MAX_INTERMEDIATE_OFFSETS_PER_REQUEST +... +- self._intermediate_chunk_indices_buffer[real_count:].fill_(0) ++ self._intermediate_chunk_indices_buffer[real_count:max_count].fill_(0) +``` + +Same for the scratch writes: `intermediate_ssm_out.copy_(...)` became +`intermediate_ssm_out[:n].copy_(...)`. Whenever you see a `fill_` or `copy_` over +a whole preallocated buffer on the per-step path, ask what the real bound is. + +## Inner-loop micro-optimizations + +From `ab2b33d5`, in the single-token decode SSM update: + +**Use `exp2` for `exp`.** Hardware has a native `exp2`: + +```python +@triton.jit +def fast_exp(x): + """ + Fast calculation of exponent via exponent of 2. + """ + LOG2E = tl.constexpr(1.4426950408889634) + return tl.math.exp2(LOG2E * x) +``` + +**`torch.empty` when the kernel overwrites every row.** Zero-init is wasted HBM +bandwidth. Applied to the squared-ReLU output in +[activations.py](megatron/core/inference/moe/activations.py). + +**Specialize the launch shape on compute capability.** The tile and warp choice +was derived from `dstate` alone; Blackwell wants a different point: + +```python +# megatron/core/ssm/ops/mamba_ssm.py +is_blackwell = torch.cuda.get_device_capability(x.device)[0] >= 10 +... +else: + # dstate > 64 + if is_blackwell: + # Optimized for B200 with dstate>64 + BLOCK_SIZE_M, num_warps = 32, 8 + elif dstate <= 128: + BLOCK_SIZE_M, num_warps = 4, 4 +``` + +## File map + +| File | Role | +|---|---| +| [mamba_mixer.py](megatron/core/ssm/mamba_mixer.py) | The mixer. Prefill: conv, chunk scan, state update, extraction. Decode: `selective_state_update` + `causal_conv1d_update`. | +| [ssd_combined.py](megatron/core/ssm/ops/ssd_combined.py) | `mamba_chunk_scan_combined_varlen`, the packed varlen chunk scan. `return_raw_states` exposes intermediates. | +| [mamba_ssm.py](megatron/core/ssm/ops/mamba_ssm.py) | `selective_state_update`, the decode-path kernel. | +| [intermediate_extraction.py](megatron/core/ssm/ops/intermediate_extraction.py) | Fused conditional gather+scatter for prefix-cache extraction. | +| [causal_conv1d_varlen.py](megatron/core/ssm/ops/causal_conv1d_varlen.py) | Varlen depthwise causal conv with fused SiLU (prefill). | +| [determinism.py](megatron/core/ssm/ops/determinism.py) | `autotune_configs` — the single autotuning funnel. | +| [mamba_metadata.py](megatron/core/inference/contexts/attention_context/mamba_metadata.py) | Per-step Mamba metadata in fixed-address buffers. | +| [mamba_slot_allocator.py](megatron/core/inference/contexts/mamba_slot_allocator.py) | Durable state cache plus extraction scratch. | + +## How these are tested + +The pattern is consistent and worth imitating: + +**Compare against an obvious reference.** `_ref_conv` is a plain nested-loop +PyTorch implementation; the kernel is checked against it with +`torch.testing.assert_close`. A fast kernel with no slow reference has no test. + +**Assert the gating with a sentinel.** Fill the output with `12345.0`, run with +`real_count < max_count` where the trailing indices are *valid but must not be +gathered*, then assert both that `out[:real_count]` matches the reference and that +`torch.all(out[real_count:] == sentinel)`. The second half proves padded slots +produced no HBM writes — the actual performance claim. + +**Re-derive formulas independently.** `test_max_intermediate_states_per_step_formula` +recomputes `min(token_based, request_based)` in the test, asserts it matches, and +asserts the same value is shared by the allocator, the metadata, and the buffer +shape. It covers both regimes explicitly, including one where the request bound is +strictly tighter. + +**Cover the boundary.** `test_scatter_conv_sub_dconv_clamp` uses `pos=2` with +`d_conv=4` so the window start goes negative, checking all three out-of-range +positions clamp to token 0. + +**Test the error path.** A too-small budget must raise `ValueError`, not silently +OOM later. + +Tests: `tests/unit_tests/ssm/ops/test_ssd_combined.py`, +`tests/unit_tests/inference/contexts/test_dynamic_prefix_caching.py`. diff --git a/skills/optimize-inference-siddharth/references/measuring.md b/skills/optimize-inference-siddharth/references/measuring.md new file mode 100644 index 00000000000..5270f206346 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/measuring.md @@ -0,0 +1,440 @@ +# Measuring Inference Performance + +Measure first, then change one thing, then measure again. Every commit in this +history started from a profile, and the bottleneck was frequently not where it +seemed. This document covers the observability that exists specifically so +inference changes can be attributed. + +Source commits: `eadbaa61` (#5611, profiler endpoints), `edd45620` (#5609, +cache-hit reporting), `faced5128` (#3034, routing instrumentation), `53a2b19a` +(#2920, NVTX ranges). + +## Bracketing an nsys capture with the profile endpoints + +Profiling a server is awkward: you want the steady state, not startup, warmup, or +CUDA-graph capture. The profile endpoints solve this by relaying +`cudaProfilerStart` / `cudaProfilerStop` to every engine, so an outer nsys process +records only the window you ask for. + +From +[profile.py](megatron/core/inference/text_generation_server/dynamic_text_gen_server/endpoints/profile.py): + +``` +POST /start_profile and /stop_profile relay a control signal through the +InferenceClient -> data-parallel coordinator -> every connected EP/DP engine, +which calls cudaProfilerStart()/cudaProfilerStop(). Pair with an outer +`nsys profile --capture-range=cudaProfilerApi` to bracket a capture window. +``` + +Workflow: + +1. Launch the server under nsys with the capture range armed: + +```bash +nsys profile --capture-range=cudaProfilerApi --capture-range-end=stop \ + -o inference_profile \ + +``` + +2. Send warmup load until the engine reaches steady state, including enough steps + that all CUDA graphs are captured. +3. Start recording, drive the load you care about, stop recording: + +```bash +curl -X POST http://localhost:5000/start_profile +# ... drive the workload ... +curl -X POST http://localhost:5000/stop_profile +``` + +Both `/start_profile` and `/v1/start_profile` are routed, likewise for stop. The +signal reaches **every rank**, so a multi-GPU trace covers all of them. + +Then analyze the trace with the `nsight-system-analysis` skill. + +### Flags that deadlock finalization — read before adding host visibility + +On a MoE decode workload under full-iteration CUDA graphs, nsys **hangs in +finalization** and `QdstrmImporter` then rejects the `.qdstrm` if any of the +following is enabled. This was bisected one flag at a time over several sessions +on Qwen3-30B / 4×GB200, and it is independent of `--cuda-graph-trace` level: + +| Trigger | Symptom | +|---|---| +| `osrt` in `--trace` | Finalization deadlock; unusable qdstrm | +| `--sample=process-tree` (CPU / Python sampling) | Same | +| NVTX ranges active under graph capture | Same — capture inflates NVTX event volume enormously | + +The only set that reliably finalizes on this workload: + +```bash +nsys profile --trace=cuda,nvtx --sample=none \ + --capture-range=cudaProfilerApi --capture-range-end=stop \ + -o inference_profile +``` + +This is a real constraint on method, not a nuisance: **the documented way to +attribute host time is exactly the way that breaks.** Two consequences, both with +working substitutes below — attribute host-side idle from the CUDA-API rows in a +clean GPU-only trace (see *Idle is not one thing*), and name host phases with +`perf_counter` instead of NVTX (see *When NVTX is unavailable*). + +### Validate that your instrumentation is not the measurement + +`--cuda-graph-trace=node` resolves individual kernels inside a graph, which is +what makes per-kernel analysis possible, but it instruments every node — 1158 of +them on this model — so it is fair to suspect it inflates the step. + +Check rather than assume. Capturing the same steady-state window both ways gave a +step period of **8882.5 µs (node) against 8945.1 µs (graph)** — 0.71% apart and in +the *wrong direction* for an instrumentation artifact — and host `cudaGraphLaunch` +of 190.0 against 184.7 µs. Node mode also reported 1169 kernels/step, matching the +independent count. Only after that control did a 900 µs/step graph-machinery +attribution become safe to act on. + +Run this control once per workload before trusting node-mode wall-time +attributions. Use `--cuda-graph-trace=graph` when you only need step periods, since +it produces far fewer events. + +## NVTX ranges label the host critical path + +Device timelines show kernels; they do not show why the host fell behind. The +per-step host critical section is annotated, so the CPU side of the trace is +readable without guessing: + +| Range | What it covers | +|---|---| +| `bookkeeping` | Per-step request state updates | +| `detokenization` | Only when no coordinator is present | +| `coordinator_communication` | msgpack serialize + ZMQ send | +| `serialize_tensor` | Individual tensor to list conversion | +| `drain_zmq_socket`, `add_request` | Inbound request handling | + +If you add host work to the step loop, add an NVTX range around it. `nvtx_range_push` +/ `nvtx_range_pop` come from `megatron/core/utils.py`. The precedent is that even +`serialize_tensor` — a two-line function — is annotated, because unmeasured host +cost is invisible cost. + +Two gotchas before you rely on these ranges. The engine's own helpers are **inert +unless NVTX profiling is enabled**, which only the training loop does by default — +a whole 187.9 µs/step host phase (`post_process_requests` via `async_bookkeep`) was +invisible in one host trace for exactly this reason, and had to be found by +elimination instead. And enabling NVTX under graph capture is one of the +finalization-deadlock triggers above, so on this workload the ranges you most want +are the ones you cannot record. + +### When NVTX is unavailable: `perf_counter` phase timing + +Since the range names are the useful part and nsys is the broken part, keep the +names and drop nsys. Monkeypatch the same push/pop symbols with +`time.perf_counter_ns` self-timers over a shared nesting stack — each label then +reports self-time excluding its children — and print a per-step breakdown every N +steps. No profiler, no CUDA-graph interaction, and it names precisely the +between-step Python phases a GPU-only trace can only localize. + +Gate it behind an env var and make `install()` a no-op by default, the same way +every other diagnostic here ships. + +## Built-in counters, no profiler needed + +Two things are exposed through the API and are often enough to confirm or reject a +hypothesis in seconds. + +### Prefix cache hit rate + +`DynamicInferenceRequest.num_cached_tokens` accumulates +`num_matched_blocks * block_size_tokens` across prefill chunks, and surfaces in the +chat-completions response as `usage.prompt_tokens_details.cached_tokens` +(OpenAI-compatible). + +Use it to check that a prefix-caching or routing change actually improved cache +hits, rather than inferring it from end-to-end latency. Particularly relevant when +tuning `prefix_caching_routing_alpha` — see +[host-path.md](host-path.md). + +### MoE expert load imbalance + +Set `moe_enable_routing_replay` (requires `num_moe_experts`). The context then +records, per request, which experts every token was routed to, in a static +CUDA-graph-safe buffer of shape +`[max_tokens, num_moe_layers, moe_router_topk]`, exposed as +`DynamicInferenceRequest.routing_indices` with shape +`[total_tokens, num_layers, topk]`. + +Aggregate a histogram across many requests. A few hot experts with the rest idle +means wasted expert capacity and straggler experts gating every step — a +throughput problem no kernel tuning will fix. + +Implementation: [routing_metadata.py](megatron/core/inference/contexts/routing_metadata.py), +[router_replay.py](megatron/core/transformer/moe/router_replay.py), collected in +`TextGenerationController._router_record_bookkeeping()`. Note that the recording +itself follows the rules in this skill: a preallocated static buffer, toggled by +`using_cuda_graph_this_step()`, sliced to `active_token_count` so padding is +excluded. + +## Benchmark harnesses + +| Tool | Use | +|---|---| +| [tools/run_inference_performance_test.py](tools/run_inference_performance_test.py) | Standalone driver: builds the model and engine, sweeps requests, times steps. Fastest local loop. | +| `tests/performance_tests/shell_test_utils/run_perf_test.sh` | The harness CI perf recipes invoke. | +| `tests/performance_tests/shell_test_utils/compare_to_baseline.py` | Compares a run against checked-in `baseline_values.json`. | +| `tests/test_utils/recipes/h100/*inference*.yaml` | Registered perf and functional recipes. | +| [examples/inference/](examples/inference/) | Server launch and offline inference examples. | + +For cluster runs use the `run-inference-performance-tests` skill (throughput and +latency versus baseline) and `run-inference-functional-tests` (correctness, +including the CUDA-graph paths). Neither is duplicated here. + +## Interpreting the profile + +The four scripts in [scripts/](../scripts/) implement this section, because the +windowing arithmetic below is easy to get wrong by hand and both of the worst analysis +errors in this campaign were windowing errors, not reasoning errors: + +| Script | Answers | +|---|---| +| `steady_window.py` | Per-kernel budget over an auto-located steady-state decode window | +| `union_window.py` | Union-busy, overlap ratio, and the gap histogram by band | +| `compare_budget.py` | Two engines, matched functional taxonomy, per-bucket ms and launch counts | +| `kernel_neighbors.py` | Which call site an unattributed kernel belongs to | + +All four take a **one-per-layer-per-step anchor kernel** as a `LIKE` pattern and count +back from the end of the trace, which is what makes the window steady-state by +construction. + +### Pick a clean window, and measure busy as a union + +Two arithmetic mistakes will make every category number wrong. + +**Do not take a fraction of the capture span.** A span fraction mixes prefill, +graph capture, and ramp into the "steady state." Instead search for the **densest +window** of the length you want — kernel starts per unit time — and verify the +kernel sequence repeats with a stable period inside it. A 1-2 s window is enough. + +**Do not sum per-stream kernel time.** Kernels overlap across streams, so summing +double-counts and can exceed the wall clock. Compute the **union of kernel +intervals** across all streams on the device; that is busy time, and +`window − union` is real idle. + +On the twelve-gate Qwen config a clean 2 s window gave union-busy **81.6%** and +idle **18.4%**, with device-time by category of expert GEMM 587 ms, NVLS comm +247 ms, dense GEMM 236 ms, attention 211 ms, and then routing 69, norm 67, +elementwise 66, `moe_sum` 65, rope/KV 15. Those small buckets being ~65-70 ms each +*and partly overlapped on other streams* is what established that further +small-kernel fusion was worth only ~0.5-1% apiece — a conclusion that a +per-stream sum would have hidden. + +Keep **both** numbers, because their ratio is itself a lever. Sum-of-durations ÷ union +is the overlap ratio: how much work the engine is hiding behind other work. Across the +Qwen campaign mcore measured **1.003-1.008** against vLLM's **1.050** — mcore ran +essentially everything serially while vLLM hid ~5% of its device work. That ~0.3 ms/step +was the largest single remaining component of the gap, and it is invisible in either +number alone. (Unlike union-as-a-share-of-wall, this ratio stays meaningful even in a +window that contains an iteration boundary, since it compares device work to device +span.) A category whose ratio is 1.00 and whose kernels are already efficient has +one lever left, and it is scheduling, not tiling. + +### Two traces of the same config are not the same workload + +This one produced a confident wrong conclusion before it was caught, and it is specific +to decode: **the KV cache grows as the run proceeds, so per-token cost grows with it.** +Two windows taken as "the last 60 steps" of two different runs therefore sit at +different average sequence lengths, and then *every* bucket differs at once — attention +and the dense GEMMs most, but nothing is exempt. Two mcore traces of the same config +differed by **27% on the expert GEMM**, which nearly credited a reduce-scatter change +with speeding up expert compute. + +Three defenses, in order of cost: + +**Compare launch counts, not times.** Counts are invariant to sequence length; times +are not. A fusion or a removed cast shows up as an exact integer change per layer +(23.9/step, 48/step), which is both unambiguous and immune to drift. This is the primary +evidence for "did the mechanism fire?" — the +2.5% bf16 win was confirmed by +`bfloat16_copy_kernel` instances dropping 98% (258,564 → 4,572), not by any bucket time. + +**Check that untouched buckets held still.** Before attributing a bucket's change to +your change, verify the buckets your change *cannot* touch are within noise. If they +moved, the windows are not comparable and no per-bucket time in that pair means anything. + +**Infer the window's sequence length.** Attention time per launch is a near-monotone +function of KV length, so a microbenchmark sweep of the attention kernel doubles as a +ruler: match the trace's µs/launch against the curve to recover the window's effective +seqlen. That is how the final trace was found to sit at ~260 tokens rather than the 512 +the harness had used, which explained a per-launch time that otherwise looked impossibly +fast. + +The corollary for the whole method: **use throughput for the verdict and the trace for +the mechanism.** The trace tells you *what changed structurally*; only the same-session +A/B (below) tells you whether it was worth anything. + +### Name a kernel by its neighbours, not by instrumenting the framework + +When the profile shows an expensive kernel with a generic name — +`elementwise_kernel`, `vectorized_elementwise_kernel`, `CatArrayBatchedCopy` — the +instinct is to go find the launcher with a `TorchDispatchMode` hook, a monkeypatched +`copy_`, or a stack capture. **Under CUDA-graph replay that instinct costs job launches +and usually returns nothing.** Four consecutive attempts failed here for two independent +reasons: the hook fires at graph *capture*, not replay, so anything armed on per-step +state never becomes true; and an arming predicate that reads the inference context does +not see it at all, because the context does not reach the layer as a keyword argument +under capture. Each failure cost a full allocation slot to discover. + +The trace already contains the answer. A kernel's immediate **predecessor and successor** +on the device are stable under graph replay, and the pair usually names the call site +outright. Aggregating them across a window took one query and named both unattributed +copies immediately: 23.9/step between the reduce-scatter and the residual add (the +collective's output cast) and 23.9/step between the QKV GEMM and the QK-norm (the split). +`kernel_neighbors.py` in [scripts/](../scripts/) does this. + +Reach for framework instrumentation only when the neighbours are genuinely ambiguous, +and when you do, **arm it on a call counter rather than on any per-step state**, so it +cannot silently never fire. + +### Idle is not one thing + +Total idle is not an actionable number, because a large fraction of it is not +attackable. Split it by gap size first: + +| Gap size | Share (Qwen, 178 ms idle in 1 s) | What it is | +|---|---:|---| +| < 10 µs | 59.6 ms (33%) | Intra-graph kernel scheduling — unavoidable | +| ≥ 10 µs | 118.8 ms (67%) | Host chain between steps — attackable, median 37 µs | + +Then attribute the large gaps **without a host trace**, using the CUDA-API +(`CUPTI_ACTIVITY_KIND_RUNTIME`) rows that a clean `cuda,nvtx` capture already +contains. For each gap, ask which host API calls fall inside it: + +| Inside the gap | Share of large-gap idle | +|---|---:| +| Nothing — no CUDA call at all (Python/CPU compute) | **73.9%** | +| Token D2H memcpy | 8.9% | +| Graph launch | 7.5% | +| Kernel launch | 6.9% | +| Sync wait | 1.6% | + +So ~74% of attackable idle — about 8-9% of wall time — was **Python on the +critical path**, and only 1.6% was waiting on the GPU. Localize it further by +bracketing each uncovered region with the CUDA APIs immediately before and after +it: 47.6% sat between `cudaMemcpyAsync` and `cuKernelGetName`, another 15.7% +between `cudaLaunchKernel` and `cudaMemcpyAsync`. That pins the cost to +**between-step host orchestration after the sampled-token D2H copy** — post-sampling +bookkeeping, scheduling, attention-state prep — which is where the host-path wins +in [host-path.md](host-path.md) came from. + +This RUNTIME-row method is the recommended one, not a fallback. It works on the +only capture configuration that finalizes. + +### Map the signal to a section + +Map the dominant signal to the section of this skill that addresses it: + +| Signal | Reading | Go to | +|---|---|---| +| GPU idle gaps, host running ahead | Launch-bound: too many launches, or no matching graph bucket | [cuda-graphs.md](cuda-graphs.md) | +| Long host span after the last kernel | Post-processing on the critical path | [host-path.md](host-path.md) | +| NCCL kernels not overlapped | Exposed collective, or EP ranks out of step | [moe-inference.md](moe-inference.md) | +| Multi-ms stalls, or slow early steps | Triton recompilation or autotune | [mamba-and-triton.md](mamba-and-triton.md) | +| Grouped GEMM wide but shallow | Tiles tuned for the wrong batch size | [moe-inference.md](moe-inference.md) | +| Many sub-microsecond kernels in a row | Fusion candidate | [moe-inference.md](moe-inference.md) | + +## Measurement discipline + +- **A/B against a matched config.** Same batch size, sequence length, and + parallelism as the vLLM or baseline number you are comparing to. A latency + comparison across different `max_requests` values means nothing. +- **Exclude capture from the measurement.** Graph capture runs a full forward per + bucket. Warm up until capture is done before you start recording. +- **Change one thing.** Most of these commits are single-mechanism precisely so the + measurement attributes cleanly. +- **Keep the kill switch and measure both ways.** `inference_disable_triton_nvls_kernels`, + `inference_moe_disable_fused_quant_kernels`, and the backend enums exist so a + regression can be bisected by config rather than by revert. +- **Re-measure per model.** The shared-expert CTA cap was a win on one model and a + loss on another (see [moe-inference.md](moe-inference.md)). Overlap and tile + tuning do not transfer. + +## The noise floor is bigger than your win + +Once the easy multi-percent wins are gone, the remaining levers are worth well +under 1% each, and at that point the above discipline is necessary but **not +sufficient**. Across the Qwen campaign, throughput for an *identical* config +measured in different allocations drifted by **−1.61%, −0.57%, and +1.22%** on +three separate occasions — larger than most of the individual wins that followed +(+0.94%, +0.90%, +0.98%, ~+0.5%). Comparing against last session's number will +produce confident, wrong conclusions in both directions. + +Four rules make a sub-1% claim falsifiable. + +**1. Re-baseline in the same session.** Before each lever, run the current best +config in the *same allocation* you will test in, and use that as the reference. +Record it as its own ledger entry. This is the single highest-value habit here. + +**2. Run the arms back to back.** Same allocation, same warm process where +possible, alternating OFF and ON. An OFF arm collected an hour earlier is a +different measurement. + +**3. Judge by distribution separation, not mean delta.** With N timed iterations +per arm, the acceptance criterion is that the **arms do not overlap** — the +slowest ON iteration beats the fastest OFF iteration. Examples that met it: + +``` +QWEN-026: min ON 27,102.9 > max OFF 26,425.7 (fully separated) +QWEN-025: 13 of 15 ON iterations beat all 15 OFF (near-separated) +QWEN-023: slowest ON 25,992.1 > fastest OFF 25,893.9 (fully separated) +``` + +A +1% mean delta with overlapping arms is not a result. A +0.9% delta with +separated arms is, even though it is smaller than the session drift — because +drift is between sessions and separation is measured within one. + +**4. Repeat the pair, and report pairwise deltas.** Two or three OFF/ON pairs, each +delta stated separately (`+0.99% / +1.26% / +0.68%`), tells the reader the spread. +A single averaged number hides it. + +Also discard the first timed iteration or check it explicitly — cold-start outliers +of ~1.3% appeared repeatedly while iterations 2-5 spanned 0.23%. + +## Predict how much of a kernel win will convert + +The most common surprise in this campaign was the ratio between a kernel-level +speedup and the end-to-end result. It went both ways, by large factors, and it is +predictable enough to estimate before you build. + +| Where the work sits | Conversion | Evidence | +|---|---|---| +| Serial per-layer dependency chain | **more than 1:1** | Fused QK-norm: ~1% microbench ceiling → **+2.9% e2e**, ~3×. Removing a launch also removes a graph node and the host dispatch gap behind it, ×48 layers. | +| Overlapped with other work | **well under 1:1** | A 183.9 µs/step host saving converted at only **1/3 to 1/2**, because that phase was already partly overlapped with GPU work. | +| Off the critical path | **~0** | A 1.25× speedup of the MoE activation path delivered **+0.55%, then +0.13%** — a wash. The op was never the wall. | + +The rule: **estimate against the critical path, never against device-time share.** +Before building, establish whether the target is (a) on the serial chain, (b) +already overlapped, or (c) off it — the same determination the +[decision-gates.md](decision-gates.md) ceiling calculation needs. + +The corollary is that a small kernel on the serial chain can be worth more than a +large kernel off it, which inverts the usual ranking. Two ~1.5-4 µs norm fusions +delivered +2.9% and +1.37%, while the largest device-time category in the profile +had 2.3% of headroom left in total. + +## Keep an append-only ledger + +For anything longer than a single session, the ledger *is* the method. Sessions get +preempted, context is lost, and without a durable record the same dead end gets +re-explored — a routing-kernel rewrite that measured a wash was nearly repeated for +exactly this reason. + +Requirements that earned their place: + +- **A fixed protocol table at the top** — cluster, hardware, model, batch, output + length, parallelism, warmup/timed counts, correctness gate, primary metric — plus + the explicit rule that results are not comparable when any of them differ. +- **One row per experiment, including every rejection**, with its root cause. The + negative results are the higher-value half: they are what stops the next person, + and three of them killed multi-week efforts here. +- **Append-only. Never edit a recorded result.** Supersede it with a new entry. +- **A running distance-to-target number** so priorities stay honest. +- **A ranked next-levers list**, re-derived after each profile rather than carried + forward. + +Record the baselines in a fixed order — competitor first, then yours, then the +gap — and only then change code. diff --git a/skills/optimize-inference-siddharth/references/moe-inference.md b/skills/optimize-inference-siddharth/references/moe-inference.md new file mode 100644 index 00000000000..167538e7404 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/moe-inference.md @@ -0,0 +1,407 @@ +# MoE Inference: the `inference_optimized` Stack + +The training MoE path is unusable for low-latency decode. Its AlltoAll dispatcher +does host synchronizations (`.item()`, `.tolist()`, `.cpu()`) that make CUDA-graph +capture impossible; the router computes z-loss, aux-loss, token-drop, and +expert-bias updates that are pure training overhead; and TE GroupedGEMM needs a +host-resident `tokens_per_expert`. + +Rather than branch the training code, Siddharth forked a parallel inference +module hierarchy. Understanding that structure is most of what you need. + +Source commits: `7d1c0168` (#3496, foundational), `905c0e38` (#3851), +`589cd9e1` (#3858), `bfd45740` (#4258, architectural rewrite), `442a936a` (#4570), +`20f09364` (#4603), `c817dad2` (#4587), `3a253ac5` (#4922). + +## How the backend is selected + +```mermaid +flowchart TD + A["--transformer-impl inference_optimized"] --> B["gpt_layer_specs / mamba_layer_specs
take the inference branch"] + B --> C["get_inference_optimized_moe_spec()"] + C --> D["MoELayer with MoESubmodules"] + D --> E["InferenceTopKRouter"] + D --> F["InferenceGroupedMLP
(via InferenceSpecProvider)"] + D --> G["SharedExpertMLP"] + H["MoELayer.__init__ -> _setup_inference_mode"] --> I["instantiate dispatcher from
inference_moe_token_dispatcher_type"] + J["MoELayer.train(mode)"] --> K["eval -> inference dispatcher
train -> standard dispatcher"] +``` + +The load-bearing detail: **`MoELayer.train()` performs the swap.** Eval mode means +the inference stack is live. There is no per-iteration flag to set — the original +`set_inference_cuda_graphed_iteration` mechanism from #3496 was removed in #4258 +precisely because manual plumbing was error-prone. + +Relevant files: [backends.py](megatron/core/models/backends.py) +(`InferenceSpecProvider.grouped_mlp_modules`), +[moe_module_specs.py](megatron/core/models/gpt/moe_module_specs.py), +[moe_layer.py](megatron/core/transformer/moe/moe_layer.py). + +## What each replaced component does differently + +### `InferenceTopKRouter` + +In [router.py](megatron/core/transformer/moe/router.py). Strips z-loss, +aux/load-balance loss, token drop, and expert-bias updates. It is +`@torch.compile`d, and it returns **dense** `[num_tokens, topk]` tensors rather +than sparse `[num_tokens, num_experts]`, which is what the fused GEMM kernels +want. + +`moe_router_dtype` must be `fp32` — enforced as a hard config error specifically +to avoid a dtype conversion on every decode step. + +### `InferenceGroupedMLP` + +In [experts.py](megatron/core/transformer/moe/experts.py). Three things matter: + +**Stacked contiguous weights.** TE stores per-expert `weight0..weightN`; grouped +GEMM wants one `[num_experts, out, in]` tensor. Built lazily on first forward and +shared by pointer, never by replacing the `Parameter` — see hard rule 7 in +[SKILL.md](../SKILL.md). + +**GPU-resident expert offsets.** Grouped GEMM needs to know where each expert's +token block starts and ends. The training path gets that from a host-resident +`tokens_per_expert`; the inference path computes it entirely on device. Commit +#3496 used `tokens_per_expert.cumsum(0)`; it is now a Triton prefix-sum kernel in +[permute.py](megatron/core/inference/moe/permute.py) that also folds in the +per-expert alignment rounding: + +```python +h = tl.load(tokens_per_expert_ptr + r, mask=mask, other=0) +# Round up non-zero counts to alignment boundary +if alignment > 1: + h = tl.where(h > 0, ((h + alignment - 1) // alignment) * alignment, h) +inc = tl.cumsum(h, axis=0) +tl.store(exclusive_offsets_ptr + r, inc - h, mask=mask) +tl.store(inclusive_offsets_ptr + r, inc, mask=mask) +``` + +The inclusive offsets are passed straight through as `offs` to `grouped_mm` / +`scaled_grouped_mm`, so no expert boundary ever touches the host. + +**Backend dispatch** on `inference_grouped_gemm_backend`: `vllm` (default), +`flashinfer`, or `torch`. MXFP8 requires `torch`, because +`scaled_grouped_mm` is the only path supporting blockwise 1x32 scaling. Note the +alignment constraint that falls out of combining requirements: + +```python +# scaled_grouped_mm requires each expert's token count aligned to 32, +# but swizzled MXFP8 scales require alignment to 128. Use 128 to +# satisfy both constraints. +expert_alignment = 128 +``` + +For BF16 decode on GB200, all three backends were measured against each other and +**`vllm` won**, so the default is also the fast path — but the two rejections are +informative rather than obvious: + +- **`torch`** (`grouped_mm`, cuBLAS) measured **0.82×**. Reach for it only when + MXFP8 forces it. +- **`flashinfer`** (`cutlass_fused_moe`) measured **90.81 µs against 83.63 µs** for + the retuned Triton path — 8% slower. Getting to that number first required fixing + two independent wiring bugs: an `ActivationType` mis-map that hard-fails, and a + gate/up ordering mismatch that **silently corrupts numerics**. Note the shape of + that second one — a backend can be wired wrongly and still run. +- **`trtllm_bf16_routed_moe`**, the TRT-LLM-Gen kernel vLLM actually wins with, + cannot currently be dropped in at all: `weight_layout=MajorK` is rejected + (`BF16 Moe: weight_layout must be BlockMajorK`) and `BlockMajorK` indexes + `size(3)`, i.e. it requires **4-D pre-shuffled block-major** weights against + mcore's 3-D `[E, 2*ffn, H]`. The blocker is the weight-preparation pipeline, not + the kernel or the EP contract. + +## Dispatchers: NCCL AllGather vs NVLS AllGather-V + +Both live in +[token_dispatcher_inference.py](megatron/core/transformer/moe/token_dispatcher_inference.py) +under `InferenceAllGatherDispatcherBase`. Selected by +`inference_moe_token_dispatcher_type`. + +The progression is worth understanding because it explains why the default is +what it is. AlltoAll was replaced by AllGather/ReduceScatter for graph +friendliness. But a fixed AllGather requires every EP rank to contribute the same +token count, so ranks pad to the global max — wasted compute and communication — +and the original version only worked with decode-only graphs, forcing prefill to +run eager. AllGather-V removes the equal-count requirement. + +| | `nccl` | `nvls` (default) | +|---|---|---| +| Token counts across EP | Must be equal on the graph path | Variable per rank | +| Non-decode CUDA graphs | Disabled | Supported | +| Padding | Pads to max, then compacts | None | +| Requires | Nothing special | Hopper+, NVLink, bf16, 16-byte-aligned | + +Choose `nccl` only when NVLS is unavailable. It is the correctness fallback, not a +tuning option. + +### NVLS eligibility + +Checked centrally so the conditions live in one place: + +```python +# megatron/core/inference/communication/torch_symm_triton/utils.py +def are_tensors_nvls_eligible(*tensors: torch.Tensor) -> bool: + """ + Requirements: + - Hopper+ GPU (SM >= 9) + - All tensor byte sizes are divisible by 16 (128-bit), since NVLS + kernels process data in 128-bit chunks. + """ + if not tensors: + return False + return is_device_nvls_capable(tensors[0].device) and all( + t.element_size() * t.numel() % 16 == 0 for t in tensors + ) +``` + +If you add a collective, call this rather than re-deriving +`dtype == bf16 and device.major >= 9`. Centralizing it was part of #3496. + +### Why symmetric memory and multimem win + +`multimem.st` / `.ld` / `.red` operate on a symmetric-memory buffer mapped across +NVLink. For the small, latency-bound messages of decode: + +- A multicast store *is* the all-gather — one kernel, one end barrier, versus + NCCL's multi-step algorithm. +- Multiple tensors fuse into one kernel and one barrier. The dispatch of + `routing_map`, `probs`, and `hidden` went from 3 collectives to 1. +- Reduce-scatter accumulates in fp32 then casts, so results are *exactly* + reproducible — the unit tests assert `atol=0, rtol=0`. **But the fp32 *buffer* is + not part of that guarantee, and it is expensive** — see below. +- Buffers are fixed-address, hence graph-safe. +- CTA count is tunable, so the collective can be throttled to share the GPU with + overlapped compute. + +Buffers come from `SymmetricMemoryManager` in +[symmetric_memory.py](megatron/core/inference/symmetric_memory.py), a lazy +registry keyed by string (`"tp"`, `"ep"`). Lazy specifically so there is no +init-ordering coupling with the context — the eager version broke the RL +integration. + +### Audit the dtype of every symmetric buffer: the combine buffer was fp32 + +The `ep_rsv` reduce-scatter buffer was allocated fp32 while the tensor it carries is +bf16 everywhere on both sides of the collective. That single dtype **cost 2.5% +end-to-end** (measured, `MCORE_NVLS_RS_BF16=1`, Qwen3-30B-A3B EP4 / 4×GB200, +2026-08), in two ways at once: + +- It doubled the NVLink bytes of the largest collective in the step. +- It forced a bf16 cast on the output, one per layer — 23.9 launches/step of + otherwise unattributed `elementwise_kernel`, on the serial chain. + +The precision argument for fp32 here is weaker than it looks: `multimem.ld_reduce` +**accumulates in f32 internally even with bf16 operands**, so the reduction order and +accumulator width are unchanged by the buffer dtype. What changes is only the width of +the operands crossing the wire, and vLLM crosses that same wire in bf16. The +`atol=0, rtol=0` unit-test guarantee is about the fp32 *accumulator*, and it survives. + +**Generalizable:** a collective's buffer dtype is a per-step bandwidth decision, not a +numerics decision, and it is easy to over-specify at allocation time where nobody +re-reads it. Audit every symmetric and collective buffer against the dtype of the +tensor actually being moved, and against what the reference implementation moves. A +cast kernel sitting immediately after a collective is the tell — see *Name a kernel by +its neighbours* in [measuring.md](measuring.md). + +## Fusing the metadata update + +Per-step metadata (`valid_tokens`, `rank_token_offset`, `ep_max_tokens`) used to +take five kernels plus a NCCL all-gather. `fused_metadata_update` in +[metadata.py](megatron/core/inference/moe/metadata.py) does it in one Triton +kernel that multicast-stores the local count, barriers, then computes all three +in place. From its docstring: + +``` +Replaces the multi-kernel sequence: + dist.all_gather_into_tensor(...) # NCCL + local_tokens_per_rank.sum() # kernel + local_tokens_per_rank[:rank].sum() # kernel + local_tokens_per_rank.max() # kernel + _step_metadata.copy_(...) # kernel +with a single Triton kernel ... +``` + +**Generalizable:** a chain of tiny reductions over a handful of values is all +launch overhead. If you see several sub-microsecond kernels in sequence, they are +one kernel. + +## Buffers are allocated once, at model init + +`allocate_buffers()` is a classmethod called at model init from the dynamic +context, never inside a captured graph. Everything downstream operates on +max-sized buffers and gates work by a single fixed-address scalar +`_valid_tokens_tensor`. Expert output is written *directly* into the +reduce-scatter symmetric buffer via `out=`, avoiding a copy before the collective. + +## Padding must not reach an expert + +CUDA-graph replay pads the token dimension to the captured size. Those rows carry +garbage routing indices, and before `3a253ac5` they were dispatched to real +experts — wasted GEMM work and potentially corrupted reductions. + +The fix publishes a fixed-address GPU scalar `real_token_count` each step, and a +Triton kernel writes `-1` into every topk slot of the padded rows: + +```python +# Mask out CUDA-graph padding rows of the local routing map so the AGV +# propagates -1 into agv_r for those slots; padding tokens then route +# to no expert. +if self.__class__._real_token_count_tensor is not None: + mask_routing_padding( + self.routing_map, self.__class__._real_token_count_tensor, self.sp_rank + ) +``` + +Two subtleties: the count is in the *global* frame, so `sp_rank` shifts local rows +for the comparison; and on capture or dummy steps `real_token_count` is set to 0 +so *all* rows are masked. + +The complementary rule is in the activation and reduction kernels — skip rows +where `permutation_map == -1`, and do not *zero* rows past `valid_tokens`, since +downstream only reads the valid prefix. Commit `20f09364` removed exactly such a +zeroing pass from `moe_sum`. + +## Tune for the typical batch, not the buffer + +The vLLM-derived grouped GEMM originally used `@triton.autotune` over 25 configs. +That meant long compile times and, worse, tile choices made for the wrong batch +size. It was replaced with vLLM's host-side heuristic: + +```python +# megatron/core/inference/moe/vllm_fused_moe.py +def _get_default_config(M: int, E: int, top_k: int) -> dict: + """ + M here is the host-side token-count hint (``num_tokens_hint`` in + ``vllm_fused_moe``), NOT ``hidden_states.size(0)``. The hint is the + expected per-step token count; the worst-case buffer size would over-tune + for prefill on every decode step. + """ + # BLOCK_SIZE_M: shrink at small M to limit per-expert padding waste. + if M <= 32: + block_m = 16 + elif M <= 96: + block_m = 32 + elif M <= 512: + block_m = 64 + else: + block_m = 128 +``` + +The grid is sized from the same hint; on rare prefill spikes each CTA strides over +extra tiles via `tl.range` — correct, with reduced parallelism, which is the right +trade for a rare case. + +**Generalizable:** buffers are sized for the worst case, but kernels should be +*tuned* for the common case. Pass an explicit `num_tokens_hint` rather than +inferring from `tensor.size(0)`. + +### One shared config cannot serve both expert GEMMs + +vLLM's heuristic is the right starting point, not the end of the tuning. It returns +**one** config, but FC1 and FC2 have different output widths — on Qwen3-30B-A3B, +N=768 for FC1 against N=2048 for FC2 — and they want *opposite* `BLOCK_SIZE_N`. A +single shared config cannot express that, so one of the two GEMMs is always +mistiled. + +Splitting into per-GEMM decode-tuned configs (`config_fc1` / `config_fc2`) took +expert-GEMM device time from 77.65 to 61.89 µs (**1.255×**), the whole MoE call from +100.64 to 84.31 µs, and delivered **+4.32% end-to-end** — the largest single win of +that campaign, from tuning alone, after a roofline gate had shown a hand-written +CUTLASS kernel could not beat it (see +[decision-gates.md](decision-gates.md)). + +Two details worth copying: + +**It can be bit-exact.** Changing `BLOCK_SIZE_M` / `BLOCK_SIZE_N` while leaving +`BLOCK_SIZE_K` **unchanged** preserves the fp32 K-reduction order, so the output is +bit-identical (max abs diff 0.0). Retuning K would not be. This makes tile tuning +the cheapest possible optimization to validate — reach for it first. + +**Tuned tiles have a range, so gate them.** The tuned configs measured +1.289× / 1.197× / 1.120× / **0.952×** at 128 / 256 / 384 / 512 tokens — a 5% +*regression* at the top of the intended range. Fall back to the default heuristic +above the token count you actually tuned for (384 here, not 512). Always measure the +upper edge of a tuned range; a win at the decode point can hide a loss just past it. + +### Two built-in fusion flags do not work on this path + +`--moe-router-fusion` and `--moe-permute-fusion` look like free wins and are the +first thing anyone tries. Under `--transformer-impl inference_optimized` both crash: + +``` +AssertionError: hidden_size mismatch: 128 vs 8 +``` + +TE's fused router emits a **dense `num_experts`-wide routing map** (128), while +`InferenceTopKRouter` produces the **dense top-k** contract (8) that the vLLM/NVLS +dispatcher consumes. The built-in fusions are wired to the training MoE path only. A +hand-written fused softmax+top-k that honors the top-k contract *is* worth building +— doing so collapsed four kernels into one, 16.41 → 4.10 µs under replay, for +**+3.89% end-to-end**, with bit-exact probabilities and identical expert sets. + +## Overlapping the shared expert + +The shared expert (a dense FFN over every token) is independent of the routed +path, so it can run concurrently with dispatch, expert GEMMs, and combine. Launch +it on a side stream in `dispatch_preprocess` and join in `combine_postprocess`: + +```python +if self.shared_experts is not None and not self._external_shared_expert_launch: + stream = SharedExpertMLP.stream + stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(stream): + self._shared_expert_output = apply_module(self.shared_experts)(hidden_states) +``` + +Only enabled for `nvls` plus `use_shared_expert` plus `moe_shared_expert_overlap`. + +**Latent MoEs** (SuperV3, UltraV3) run experts in a projected latent dimension, +but the shared expert needs full hidden dim, so it sits outside the dispatcher's +view. There, `MoELayer.preprocess`/`postprocess` own the launch and the add, and +the dispatcher is told `_external_shared_expert_launch = True`. + +### A cautionary tale on overlap tuning + +Commit `442a936a` capped the AllGather-V to 16 CTAs when overlapping, so the +collective would not starve the shared-expert GEMMs on the side stream. Commit +`9b4074b` **removed that cap** — on Nemotron it measured net worse. + +Overlap tuning is model- and shape-specific. Re-measure per model; do not copy a +CTA cap forward on the assumption it generalizes. + +## Reduce how often you synchronize + +`ep_consensus_interval` (default 20) skips EP consensus all-reduces while the +engine is busy, running them immediately only when idle so new arrivals are still +picked up promptly. Same instinct as everything else here: the collective was not +slow, it was just happening more often than the semantics required. + +## Exposed comm is partly skew, and skew is a routing problem + +Before optimizing an exposed collective, split its cost into what the collective +intrinsically costs and what is **other ranks arriving late**. On Qwen3-30B, of +1060 µs/step of exposed EP communication, 632 µs was intrinsic and **428 µs was +inter-rank skew** — ranks blocked at the barrier waiting on the slowest rank's +expert GEMM. + +That 428 µs is not addressable by any change to the collective. It is expert load +imbalance, surfacing at the one place in the step where ranks must agree. The +diagnostic is `moe_enable_routing_replay` plus a histogram over +`DynamicInferenceRequest.routing_indices` (see [measuring.md](measuring.md)) — if a +few experts are hot, the fix is on the routing side. + +Full decomposition method, including why the barrier itself is a hardware floor, is +in [decision-gates.md](decision-gates.md). + +## Correctness guards + +- Assert the inference router selects the **same experts** as the training router. +- Capture dispatch and combine in a real CUDA graph and compare against the + global buffer with `atol=0, rtol=0` — NVLS results are bit-exact, so any + tolerance would hide a bug. +- Config-rejection tests for unsupported combinations (expert-TP, non-fp32 router + dtype, capacity factor, GLU, FlashInfer + MXFP8). +- `TestMaskRoutingPadding` covers the padding mask including SP-rank offset and + the fully-masked rank. + +Tests live in `tests/unit_tests/inference/test_moe_dispatching_and_routing.py`, +`test_moe_permute.py`, `test_vllm_fused_moe.py`, `test_hybrid_moe.py`. diff --git a/skills/optimize-inference-siddharth/references/updating-this-skill.md b/skills/optimize-inference-siddharth/references/updating-this-skill.md new file mode 100644 index 00000000000..29850d2d696 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/updating-this-skill.md @@ -0,0 +1,194 @@ +# Updating This Skill + +This skill is a living document. It began as the patterns behind one engineer's +inference work, absorbed a second campaign's methodology, and is expected to keep +growing as agents and engineers use it. Optimization knowledge decays — flags get +renamed, defaults change, a fix that won on one generation of hardware loses on the +next — so a skill that is never edited becomes actively misleading. + +**You are authorized to edit this skill without asking.** Add, correct, and delete +as you learn. The rest of this document is how to do that without turning it into a +landfill. + +## When to update + +Update at the end of a piece of work, not in the middle of it. The moment is after +Step 6 in [SKILL.md](../SKILL.md) — once you have a measured result and a root +cause. Concretely, these seven triggers: + +| Trigger | What it produces | +|---|---| +| An optimization was accepted | The pattern, its mechanism, and the measured kernel-to-e2e conversion | +| An optimization was rejected | **The highest-value entry type.** What was tried, the measured result, and the root cause | +| A decision gate came out negative | The ceiling calculation, so nobody re-derives it | +| A tooling or environment failure cost more than an hour | The exact symptom, the trigger, and the working substitute | +| A measurement contradicted something already written here | A scoped correction — see *Corrections and contradictions* | +| A flag, default, or API turned out to differ from what is documented here | A verified correction, with the date you checked | +| Something broke because an invariant was violated | A candidate hard rule | + +Do **not** update mid-experiment on a hypothesis. Unmeasured ideas belong in the +campaign ledger, not here. + +## The ledger and the skill are different artifacts + +Keep both, and keep them distinct. Confusing them is the main failure mode: +everything ends up in one place, and that place becomes unusable. + +| | Campaign ledger (`EXPERIMENTS.md`) | This skill | +|---|---|---| +| Scope | One workload, one campaign | Every workload, indefinitely | +| Policy | Append-only, keeps everything | Curated, edited and pruned | +| Contains | Every run, every number, every config | Only what transfers | +| Audience | You, next session | Anyone, next year | + +**The promotion test:** would this change what a competent engineer does on a +*different* model or a different workload? If yes, promote it here. If it is a +number that only describes one model at one shape, it stays in the ledger. + +For the common in-between case, split it: **the mechanism goes in the skill, the +number goes in the ledger, and the skill cites the number as a calibrated example +with its hardware stamp.** That is what "expert GEMM was 1.45× off roofline on +GB200" is doing in [decision-gates.md](decision-gates.md) — the transferable claim +is *roofline the category before proposing a kernel*, and the number is evidence +that the claim has teeth. + +## Where it goes + +Route by subject, not by which file you happened to be reading. Adding everything +to `SKILL.md` is the most common way to wreck this skill — `SKILL.md` is a router, +and it stops working once it is long enough to need its own router. + +| Learning about | Destination | +|---|---| +| Launch overhead, graph scope, capture safety, bucketing | [cuda-graphs.md](cuda-graphs.md) | +| MoE dispatchers, experts, routing, grouped GEMM, EP comm | [moe-inference.md](moe-inference.md) | +| Per-step CPU work, serialization, IPC, coordinator, DP routing | [host-path.md](host-path.md) | +| Triton rules, kernel hygiene, Mamba/SSM, scratch sizing | [mamba-and-triton.md](mamba-and-triton.md) | +| Profiling tooling, trace analysis, A/B protocol, idle accounting | [measuring.md](measuring.md) | +| Prioritization, ceilings, whether a lever is worth building | [decision-gates.md](decision-gates.md) | +| Comparing against vLLM or another engine | [vllm-differential.md](vllm-differential.md) | +| A new invariant whose violation broke something | `SKILL.md` → Hard rules | +| Flag behavior, defaults, flags that look free and are not | `SKILL.md` → Flags | +| A check that belongs before every PR | [../assets/review-checklist.md](../assets/review-checklist.md) | +| A subject none of the above covers | A new `references/*.md`, linked from `SKILL.md` | + +When you add a new reference file, three things must follow or it is invisible: link +it from the `## References` list in `SKILL.md`, mention it from whichever step it +serves, and extend the frontmatter `description` if it represents a genuinely new +capability — that description is what causes this skill to be selected at all. + +## The bar for an addition + +All four, or it does not go in: + +1. **Measured.** A number, and enough method that someone could reproduce it. "Felt + faster" is not a finding. +2. **Root-caused.** The mechanism, not the symptom. A kernel that was slow because + 2 of 152 CTAs got work teaches something; "the count kernel was slow" teaches + nothing and invites the wrong fix. +3. **Scoped.** State the conditions under which it holds — hardware, model shape, + token count, config. An unscoped claim will be applied where it is false. +4. **Actionable.** It changes a decision. If a reader finishes the paragraph and + does nothing differently, cut it. + +Reject: restatements of content already here, cluster or queue incidents with no +methodological consequence, and single anomalous runs that were never repeated. + +## House style + +Match what is already here, so the file reads as one voice rather than a pile of +contributions: + +- Lead with the finding, then the number. Not the narrative of how you got there. +- Imperative and second person. "Roofline the category before proposing a kernel." +- Quote **exact error strings** for failure modes. `AssertionError: hidden_size + mismatch: 128 vs 8` is findable; "it crashes with a shape error" is not. +- Close a transferable pattern with a **`Generalizable:`** line. That convention is + load-bearing — it is how a reader tells a specific finding from a rule. +- Cite the source: a commit sha for merged work, a ledger entry id for campaign + work. Unattributed claims cannot be re-verified. +- Show real code with a file-path comment as the first line, as the existing + snippets do. +- Tables for enumerable facts, prose for reasoning. Never a table of paragraphs. + +## Corrections and contradictions + +**If something here is wrong, fix it in place.** If it was wrong in a way someone +may have already acted on, say so in one clause rather than silently rewriting. + +**If a new measurement contradicts an existing one, do not overwrite it.** Both are +facts; they differ in scope. Scope them: + +> Capping the AllGather-V to 16 CTAs when overlapping was a win on one model and +> measured net worse on Nemotron. + +That is the correct shape — it preserves both results and tells the reader the +outcome is model-dependent, which is itself the lesson. Overwriting would have +produced a confident claim that is wrong half the time. + +**Never delete a measured negative result** because it is old or inconvenient. Those +entries are the ones preventing repeated work; three of them here each killed a +multi-week effort. If new hardware changes the answer, add the condition, do not +replace the record. + +## What to delete + +Deletion is part of maintenance, not vandalism. Delete: + +- **Content that is now false** — a flag that no longer exists, a default that + changed, a code path that was removed. Verify against the tree before restating + any default, and note the date you checked. +- **Superseded APIs**, once nothing in-tree references them. Until then keep the + old→new mapping in a collapsed `
` block; the deprecated + `CudaGraphScope` block in [cuda-graphs.md](cuda-graphs.md) is the precedent. +- **Duplication.** The same lesson stated in three files means two of them are + drift waiting to happen. Keep the fullest treatment, link to it from the others. +- **Prose that changes no decision**, including your own earlier additions. Prefer + cutting recent additions over the original source-commit material, which carries + provenance you cannot reconstruct. + +When you delete something substantive, note it in the revision log so the removal +is discoverable rather than mysterious. + +## Size budget + +Progressive disclosure only works if the entry point stays navigable. + +| File | Soft target | Action when over | +|---|---|---| +| `SKILL.md` | ~500 lines (currently ~466) | Move detail into a reference; keep the routing line | +| Any reference | ~400 lines | Split by subject, link both from `SKILL.md` | + +These are soft targets, and the line count is a proxy, not the thing that matters. +The real test for `SKILL.md`: **can a reader scan it once and know which reference +answers their question?** If a section explains a mechanism rather than pointing at +one, the mechanism belongs in a reference regardless of the current line count. Do +not trim readable prose to hit a number — that trade makes the skill worse. + +For references the test is subject coherence. If one has grown to cover two +unrelated subjects, that is a split rather than a trim; `mamba-and-triton.md` +covering both Triton hygiene and SSM scratch sizing is already near that line. + +## After any edit + +``` +- [ ] Every intra-skill link still resolves +- [ ] Frontmatter still parses, and `name` still matches the directory +- [ ] `description` extended if a new capability area was added +- [ ] New reference linked from the References list AND from the step it serves +- [ ] Measured claims carry model + hardware + date, or a commit / ledger citation +- [ ] Revision log below appended +``` + +## Revision log + +One line per substantive change, newest last. This exists so a reader can tell what +is original source-commit material and what was added later, and so deletions are +discoverable rather than mysterious. + +| Date | Change | Source | +|---|---|---| +| 2026-07 | Initial skill: the five moves, hard rules 1-9, the CUDA-graph / MoE / Mamba-Triton / host-path / measuring references, commit log, review checklist | Siddharth Singh's 2026 inference work, 29 commits | +| 2026-07-28 | Added `decision-gates.md` (ceiling-before-building, per-launch fixed costs, three gates that each killed a multi-week effort) and `vllm-differential.md` (competitor-trace comparison). Hardened `measuring.md`: nsys flag combinations that deadlock finalization, the node-vs-graph trace control, `perf_counter` phase timing, union-busy idle accounting with gap-size decomposition, the same-session back-to-back A/B protocol, kernel-to-e2e conversion, ledger requirements. Added *What a wide capture costs you* to `cuda-graphs.md`. Added per-GEMM tile tuning, the measured backend comparison, the training-path-only fusion flags, and comm-vs-skew to `moe-inference.md`. Added hard rules 10 (A/B protocol) and 11 (non-bit-exact acceptance) plus the *flags that look like free wins* table | Qwen3-30B-A3B EP4 on 4×GB200, `skills/run-qwen-model/EXPERIMENTS.md` | +| 2026-07-28 | Made the skill self-maintaining: this file, the editing authorization at the top of `SKILL.md`, the Step 6 feedback loop, the checklist's skill-maintenance block, and this log | — | +| 2026-08-03 | Added `scripts/` (`steady_window.py`, `union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) so the trace method is runnable rather than only described. **Corrected** the fp32 reduce-scatter buffer in `moe-inference.md`, which was documented as a reproducibility feature and costs 2.5% (the `atol=0` guarantee is about the fp32 accumulator, and `multimem.ld_reduce` accumulates in f32 regardless of operand width). Added to `measuring.md`: cross-trace window drift at differing KV length (27% bucket error, and the launch-count / untouched-bucket / seqlen-inference defenses), the sum÷union overlap ratio, and naming a kernel from its trace neighbours instead of framework instrumentation that cannot fire under graph replay. Added to `vllm-differential.md`: resolve the competitor's kernel to a *package* and check whether it is already installed (+2.6% from flashinfer `trtllm-gen`, after two sessions tuning inside flash-attn), and the packing split. Added thesis moves 5-width and 6, three routing-table rows, and the convention that a rejection records its mechanism and date so it can be re-tested | Qwen3-30B-A3B EP4 on 4×GB200, PR #6064, `skills/run-qwen-model/EXPERIMENTS.md` sessions 16-18 | diff --git a/skills/optimize-inference-siddharth/references/vllm-differential.md b/skills/optimize-inference-siddharth/references/vllm-differential.md new file mode 100644 index 00000000000..c7a4963fea4 --- /dev/null +++ b/skills/optimize-inference-siddharth/references/vllm-differential.md @@ -0,0 +1,191 @@ +# Differential Analysis Against vLLM + +When the goal is stated as "match vLLM," the competitor's trace is not just a +scoreboard — it is a specification. It tells you what the same model on the same +hardware costs when someone else built the decode path, which is a far more useful +target than a roofline, because it is known to be achievable. + +Everything else in this skill looks inward at Megatron. This document is the one +workflow that looks at both. + +Source: the Qwen3-30B-A3B EP4 campaign, where a single kernel-level differential +simultaneously justified the fusion work and killed the grouped-GEMM work. + +## The one question worth asking + +Given matched configs and matched trace windows: + +> **Is any individual kernel slower in mcore, or does mcore just launch more +> kernels to do the same thing?** + +These have opposite fixes, and guessing wrong is expensive. Slower kernels mean +tiling, occupancy, or implementation work. More kernels mean fusion, graph-node +removal, and host-dispatch reduction. A third answer — mcore has *categories the +competitor does not have at all* — means the structure differs and no amount of +kernel work will close it. + +On Qwen3-30B the answer was unambiguous, and it was the second one. One forward +block: **~467 kernels in vLLM against ~1784 in mcore**, roughly 3.8×, for a step +of **4.4 ms against 9.5 ms**. Individual kernels were broadly competitive; mcore +was running three to four times as many of them, most on the serial per-layer +dependency chain. + +## Getting comparable windows + +The comparison is worthless unless both windows cover the same work in the same +regime. Two requirements beyond the usual matched-config discipline: + +**Both must be steady-state decode.** No prefill, no graph capture, no warmup. +Verify by checking that the window's kernel sequence repeats with a stable period. + +**Anchor on a once-per-block kernel, not on wall time.** Pick a kernel that fires +exactly once per transformer block with a stable name — the block's leading norm +is the natural choice — and take the window from one occurrence to the next +occurrence of the *closing* kernel. That gives exactly one forward block on each +side, independent of clock offsets between the two runs. + +For Qwen the anchors were: + +| Engine | Window start | Window end | +|---|---|---| +| vLLM | `triton_red_fused__to_copy_embedding_rms_norm_0` | `triton_red_fused__to_copy_add_mean_moe_forward_mul_pow_rsqrt_0` | +| mcore | `rmsnorm_fwd_tuned_kernel<...>` | `triton_poi_fused_add_copy__0` | + +Extract the ordered kernel list per window from the `.sqlite` export — kernel +rows joined to the string table, filtered to one device and the window, ordered by +start time. Keep start, duration, and name; you need the order, not just the +totals. `compare_budget.py` in [scripts/](../scripts/) does the whole comparison, +taxonomy included. + +**And match the sequence length, not just the config.** Decode cost grows with KV +length, so "the last N steps" of two runs can sit at different average seqlen and then +every bucket differs at once, by up to 27% on one pair here. Cross-engine windows are +especially exposed to this because the two engines reach a given step at different times. +Sanity-check with the buckets your change cannot touch, and prefer launch counts, which +are seqlen-invariant — full method in *Two traces of the same config are not the same +workload* in [measuring.md](measuring.md). + +## Decompose into three buckets + +Assign every kernel in both windows to a *role* — norm, QKV GEMM, attention, +router, dispatch, expert GEMM, combine, residual, sampling — and then compare +role by role. Do **not** match on kernel names: Inductor-generated Triton names +differ between engines and encode fusion structure rather than function, which is +exactly the thing you are trying to measure. + +Then split the total gap three ways: + +**Same role, more kernels.** The fusion signal. vLLM runs one +`triton_red_fused__to_copy_add_..._rms_norm` where mcore runs +`triton_poi_fused_add_copy` followed by `rmsnorm_fwd_tuned` — two launches, two +graph nodes, per boundary, per layer. That observation, repeated across the norm +and residual roles, is what produced the fused QK-norm and fused add-norm +optimizations, worth **+2.9%** and **+1.37%** respectively. + +**Same role, slower kernel.** The implementation signal. Rank by +`(mcore_time - competitor_time)` for the role, not by mcore time alone. When the two +sides show **the same launch count and different time**, do not go tune your kernel yet +— first find out *which library the competitor is calling*, then check whether it is +already installed. See below. + +**Same total work, different packing.** The scheduling signal, and the one the other +three miss entirely. Compare sum-of-durations ÷ union for both engines: mcore measured +**1.003-1.008** against vLLM's **1.050**, i.e. vLLM hides ~5% of its device work behind +other work and mcore hides essentially none. Worth ~0.3 ms/step here, and +it does not appear in any per-bucket comparison, because the buckets are identical — it +is the same work, packed differently. Note that a wide CUDA-graph capture bounds what +you can do about it: overlap becomes a property of the captured graph's structure rather +than of stream priorities or `CUDA_DEVICE_MAX_CONNECTIONS`, so the fix is a capture-time +structural change, not a runtime knob. See *Pick a clean window* in +[measuring.md](measuring.md) for the metric and +[cuda-graphs.md](cuda-graphs.md) for what the capture forecloses. + +**Role present on one side only.** The structural signal, and usually the largest. +mcore's decode carried **exposed NVLS AllGather-V dispatch and ReduceScatter-V +combine on the per-layer critical path**; vLLM, using a TRT-LLM fused MoE, had +**zero exposed communication** — dispatch, grouped GEMM, activation, and finalize +were one fused unit. No fusion or tile change on mcore's side reaches that; it is +a different decomposition of the same math. + +## Read the competitor's fusion boundaries as the target + +The most directly actionable output is the competitor's *fusion inventory*: which +adjacent operations it has merged that you have not. For Qwen decode: + +| Operation | vLLM | mcore (before) | +|---|---|---| +| Residual add + RMSNorm | one kernel | two kernels | +| Router softmax + top-k | one kernel | four kernels | +| Routing metadata / indirection table | folded in | up to five kernels | +| Dispatch + expert GEMM + activation + finalize | one fused MoE | six-plus kernels with exposed comm between | + +Each row is a candidate with a known-achievable target. Price them with +[decision-gates.md](decision-gates.md) before building — the routing chain +collapsed from five kernels to one across four separate experiments, but the +full MoE mega-fusion was attempted, measured **0.68-0.80×**, and rejected. That +the competitor fuses something does not prove *your* fusion of it will be faster; +it proves the fusion is legal and bounds what it is worth. + +## First ask which library the competitor is calling, and whether you already have it + +Before tuning anything, resolve the competitor's kernel to a **package and an entry +point**. Kernel names in the trace are usually enough (`fmhaSm100...`, `nvjet...`, +`trtllm_gen...`), and `pip list` in the competitor's environment finishes it. Then check +your own environment: on a shared image the answer is frequently that the faster kernel +is *already installed* and you are calling a different one. + +Attention was the clearest case. The differential showed the same launch count on both +sides with mcore **0.409 ms/step slower** — the largest remaining work-bucket deficit, +and unambiguously an implementation gap rather than a fusion gap. Two prior sessions had +been spent flag-flipping between FA2 and FA4 *inside flash-attn*, which is the wrong +package: vLLM was calling flashinfer's `trtllm-gen` Blackwell decode kernel, and +flashinfer was already in the same venv. A microbenchmark at mcore's exact shapes put it +**~24% under FA2**, it accepted mcore's existing paged KV layout unchanged, and wiring it +in behind `MCORE_FLASHINFER_DECODE` was worth **+2.6% end-to-end**. Enabling its +Programmatic Dependent Launch added a further +0.3%. + +The generalizable order is: identify the package, check whether it is installed, +microbenchmark it at *your* shapes under graph replay, verify the memory layout it +demands, and only then consider tuning your own. Steps 1 and 2 are minutes; the two +sessions spent skipping them were not. + +An important asymmetry: this works for **leaf kernels with a narrow contract** +(attention, norms, sampling) and fails for **fused mega-kernels**, which come with layout +and weight-preparation demands — the two traps below. + +## Do not conclude the competitor's kernel is better without checking why + +Two traps, both hit during this campaign. + +**The competitor may win by weight layout, not by kernel quality.** The TRT-LLM +BF16 fused MoE that vLLM uses was probed directly as a drop-in backend and turned +out to require **4-D pre-shuffled block-major weights**; mcore stores 3-D +`[E, 2*ffn, H]`. `weight_layout=MajorK` is rejected outright and `BlockMajorK` +indexes `size(3)`. The kernel was never the blocker — the weight preparation +pipeline was. + +**An existing backend may be mis-wired rather than slow.** The already-shipped +flashinfer `cutlass_fused_moe` path was documented as "blocked for SwiGLU." It is +not: the kernel supports BF16 gated SwiGLU, and mcore's wiring was broken in two +independent ways — an `ActivationType` mis-map that hard-fails, and a gate/up +ordering mismatch that *silently corrupts numerics*. Both were root-caused and +fixed in a harness. Only then was the comparison meaningful, and the answer was +**90.81 µs against 83.63 µs** for the retuned Triton path, i.e. 8% slower. The +ledger note had been wrong for the wrong reason, and the correct measurement still +rejected the backend. + +**Generalizable:** before adopting a competitor's kernel, verify what it demands +of the weight layout and whether your existing wiring to it is actually correct. +A hard failure and a silent numerics corruption look nothing alike in a profile. + +## What a differential cannot tell you + +It gives you counts and durations, not the critical path. A role with more kernels +in mcore may still cost nothing if those kernels sit on a side stream and overlap +with compute. Convert the differential's findings to wall-time impact with +union-busy analysis and the serial-chain reasoning in +[measuring.md](measuring.md) before ranking them. + +It also will not show host-side differences. Two engines with identical kernel +sequences can differ substantially in Python overhead between steps, which on this +workload was 18.4% of wall time. diff --git a/skills/optimize-inference-siddharth/scripts/compare_budget.py b/skills/optimize-inference-siddharth/scripts/compare_budget.py new file mode 100644 index 00000000000..03d17fd3e7d --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/compare_budget.py @@ -0,0 +1,146 @@ +"""Matched per-category decode budget for two engines, over steady-state windows. + +The point of this script is to answer "where does mcore spend device time that vLLM +does not" without any of the three traps that made earlier attempts wrong: + + 1. **Windowing.** Whole-trace sums include model load, capture, warmup and prefill, + where the collectives move 100x the bytes of a decode step. Both traces are cut to + the last N whole steps located from a one-per-layer-per-step anchor kernel. + 2. **Matched taxonomy.** The two engines share no kernel names, so a name-by-name diff + is unreadable. Every kernel is mapped into one of a dozen functional buckets that + mean the same thing on both sides, and anything unmatched is surfaced rather than + silently dropped. + 3. **Device time only.** Both traces were captured with `--cuda-graph-trace=node`, + which inflates host-side gaps and therefore wall time. Per-kernel *durations* are + unaffected, so only sum-of-durations per bucket is compared -- never wall or idle. + +One caveat that no windowing can remove: a device-side collective blocks until its peers +arrive, so its duration includes skew. Collective buckets are exposure, not cost. + +Usage: + python compare_budget.py A.sqlite '' B.sqlite '' [steps] [layers] +""" + +import re +import sqlite3 +import sys + +# Ordered: first match wins, so put specific patterns before general ones. +TAXONOMY = [ + ("expert GEMM", r"_fused_moe_kernel|^bmm_Bfloat16"), + ("attention", r"flash_?attn|fmhaSm100|FlashAttentionForward|splitkv|splitKV"), + ("dense GEMM", r"^nvjet|cutlass.*gemm|^sm100_xmma"), + ("collective", r"multimem|ncclDevKernel|nccl.*Kernel"), + ("MoE routing", r"routingIndices|_align_single|_mask_routing|softmax_topk|topk"), + ("MoE finalize", r"_moe_sum|finalizeKernel"), + ("norm", r"rmsnorm|layernorm|_fused_add_rmsnorm|rsqrt"), + ("rotary / KV", r"rotary|append_kv|reshape_and_cache|cache_kernel"), + ("elementwise / copy", r"elementwise_kernel|CatArrayBatched|copy|vectorized_elementwise"), + ("splitK reduce", r"splitKreduce|splitk"), + ("metadata / bookkeep", r"_fused_metadata|delayStream|memset|fill"), + ("sampling", r"sampl|argmax|multinomial|gumbel"), +] + + +def bucket(name): + for label, pat in TAXONOMY: + if re.search(pat, name, re.I): + return label + return "OTHER" + + +def window(con, anchor, layers, want): + ts = [ + r[0] + for r in con.execute( + """SELECT k.start FROM CUPTI_ACTIVITY_KIND_KERNEL k + JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND s.value LIKE ? ORDER BY k.start""", + (anchor,), + ) + ] + if not ts: + sys.exit(f"anchor {anchor!r} not found") + need = want * layers + sel = ts[-need:] if len(ts) > need else ts + lo, hi = sel[0], sel[-1] + steps = len(sel) / layers + gaps = [b - a for a, b in zip(sel, sel[1:])] + big = [g for g in gaps if g > 50_000_000] + return lo, hi, steps, big + + +def budget(db, anchor, layers, want): + con = sqlite3.connect(db) + lo, hi, steps, big = window(con, anchor, layers, want) + rows = con.execute( + """SELECT s.value, SUM(k.end-k.start), COUNT(*) + FROM CUPTI_ACTIVITY_KIND_KERNEL k JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND k.start>=? AND k.start<=? GROUP BY s.value""", + (lo, hi), + ).fetchall() + per = {} + for name, ns, n in rows: + b = bucket(name) + ms, cnt = per.get(b, (0.0, 0.0)) + per[b] = (ms + ns / 1e6 / steps, cnt + n / steps) + other = sorted( + ((ns / 1e6 / steps, n / steps, name) for name, ns, n in rows if bucket(name) == "OTHER"), + reverse=True, + ) + return { + "db": db, + "steps": steps, + "wall": (hi - lo) / 1e6 / steps, + "big": big, + "per": per, + "other": other[:8], + "total": sum(v[0] for v in per.values()), + "launches": sum(v[1] for v in per.values()), + } + + +def main(): + a_db, a_anchor, b_db, b_anchor = sys.argv[1:5] + want = int(sys.argv[5]) if len(sys.argv) > 5 else 60 + layers = int(sys.argv[6]) if len(sys.argv) > 6 else 48 + A = budget(a_db, a_anchor, layers, want) + B = budget(b_db, b_anchor, layers, want) + + for X in (A, B): + warn = f" WARNING {len(X['big'])} inter-iteration gaps in window" if X["big"] else "" + print( + f"{X['db']}: steps={X['steps']:.0f} wall/step={X['wall']:.3f} ms " + f"device={X['total']:.3f} ms launches/step={X['launches']:.0f}{warn}" + ) + + print( + f"\n{'bucket':<22}{'A ms':>8}{'A n':>7}{'B ms':>9}{'B n':>7}" + f"{'B-A ms':>9}{'B-A n':>8}" + ) + labels = sorted( + set(A["per"]) | set(B["per"]), + key=lambda l: -(B["per"].get(l, (0, 0))[0] - A["per"].get(l, (0, 0))[0]), + ) + for label in labels: + am, an = A["per"].get(label, (0.0, 0.0)) + bm, bn = B["per"].get(label, (0.0, 0.0)) + print( + f"{label:<22}{am:>8.3f}{an:>7.0f}{bm:>9.3f}{bn:>7.0f}" + f"{bm-am:>+9.3f}{bn-an:>+8.0f}" + ) + print( + f"{'TOTAL':<22}{A['total']:>8.3f}{A['launches']:>7.0f}" + f"{B['total']:>9.3f}{B['launches']:>7.0f}" + f"{B['total']-A['total']:>+9.3f}{B['launches']-A['launches']:>+8.0f}" + ) + + for X, tag in ((A, "A"), (B, "B")): + if X["other"]: + print(f"\nunclassified in {tag} ({X['db']}):") + for ms, n, name in X["other"]: + print(f" {ms:7.3f} ms/step {n:6.1f}/step {name[:78]}") + + +if __name__ == "__main__": + main() diff --git a/skills/optimize-inference-siddharth/scripts/kernel_neighbors.py b/skills/optimize-inference-siddharth/scripts/kernel_neighbors.py new file mode 100644 index 00000000000..f82f046391a --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/kernel_neighbors.py @@ -0,0 +1,92 @@ +"""Identify an unattributed kernel from its trace neighbours. + +A profile shows `elementwise_kernel` costing 0.16 ms/step and says nothing about which +line launched it. The instinct is to instrument the framework -- a `TorchDispatchMode` +hook, a monkeypatched `copy_`, a stack capture. Under CUDA-graph replay that instinct +is usually wrong: the hook fires at *capture* time, not replay, so an arming predicate +based on per-step state either never becomes true or becomes true during prefill, and +each attempt costs a full job launch to find out. + +The trace already knows. A kernel's immediate predecessor and successor on the same +device are stable under graph replay, and the pair usually names the call site outright: +a copy between a QKV GEMM and a QK-norm is the split; a cast between an expert GEMM and +a reduce-scatter is the collective's input buffer dtype. + +Reads the same steady-state window convention as the other scripts here. + +Worked example -- Qwen3-30B-A3B, 76.3 unattributed `elementwise_kernel` launches/step, +after four failed instrumentation attempts. Two neighbour pairs at exactly 23.9/step +(the 48-layer count halved, i.e. one per layer per boundary) named both sources: + + 23.9 _multimem_reduce_scatter_v_kernel -> triton_poi_fused_add_copy__0 + 23.9 nvjet_sm100_tst_64x8_... (QKV GEMM) -> transformer_engine::...rmsnorm_fwd + +The first is the reduce-scatter output cast -- fixed by making the collective's buffer +bf16, worth +2.5%. The second is the QKV split before QK-norm. + +Usage: + python kernel_neighbors.py trace.sqlite '' \ + '' [steps] [layers] +""" + +import sqlite3 +import sys +from collections import Counter + + +def main(): + db, target, anchor = sys.argv[1], sys.argv[2], sys.argv[3] + want = int(sys.argv[4]) if len(sys.argv) > 4 else 20 + layers = int(sys.argv[5]) if len(sys.argv) > 5 else 48 + + con = sqlite3.connect(db) + ts = [ + r[0] + for r in con.execute( + """SELECT k.start FROM CUPTI_ACTIVITY_KIND_KERNEL k + JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND s.value LIKE ? ORDER BY k.start""", + (anchor,), + ) + ] + if not ts: + sys.exit(f"anchor {anchor!r} not found") + sel = ts[-want * layers :] if len(ts) > want * layers else ts + lo, hi = sel[0], sel[-1] + steps = len(sel) / layers + + # Whole ordered kernel stream for the window; neighbours are positional, and a + # stream-aware version would need the correlation ids, which decode rarely needs + # because the per-layer chain is serial anyway. + stream = con.execute( + """SELECT k.start, k.end, s.value FROM CUPTI_ACTIVITY_KIND_KERNEL k + JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND k.start>=? AND k.start<=? ORDER BY k.start""", + (lo, hi), + ).fetchall() + + tgt = target.replace("%", "").lower() + hits = [i for i, (_, _, name) in enumerate(stream) if tgt in name.lower()] + if not hits: + sys.exit(f"target {target!r} not found in window") + + pairs = Counter() + dur = 0 + for i in hits: + dur += stream[i][1] - stream[i][0] + prev = stream[i - 1][2] if i else "" + nxt = stream[i + 1][2] if i + 1 < len(stream) else "" + pairs[(prev[:52], nxt[:52])] += 1 + + print(f"{db}") + print( + f" target {target!r}: {len(hits)/steps:.1f} launches/step, " + f"{dur/1e6/steps:.4f} ms/step over {steps:.0f} steps" + ) + print(f"\n {'n/step':>7} predecessor -> successor") + for (prev, nxt), n in pairs.most_common(12): + print(f" {n/steps:>7.1f} {prev}\n -> {nxt}") + + +if __name__ == "__main__": + main() diff --git a/skills/optimize-inference-siddharth/scripts/steady_window.py b/skills/optimize-inference-siddharth/scripts/steady_window.py new file mode 100644 index 00000000000..2a9a8d6ad18 --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/steady_window.py @@ -0,0 +1,59 @@ +"""Per-kernel budget over a *steady-state decode* window, auto-located. + +Taking MIN/MAX over a whole trace is wrong: the span then covers model load, graph +capture, warmup and the gaps between benchmark iterations, which both destroys wall/step +and pollutes every per-kernel average with prefill instances. This locates the window +from the cadence of the one-per-layer-per-step collective and reports only the last +`--steps` steps of it, which are steady-state decode by construction. +""" + +import sqlite3 +import sys + +DB = sys.argv[1] +ANCHOR = sys.argv[2] # LIKE pattern for the one-per-layer-per-step collective +LAYERS = int(sys.argv[3]) if len(sys.argv) > 3 else 48 +WANT = int(sys.argv[4]) if len(sys.argv) > 4 else 200 + +con = sqlite3.connect(DB) +rows = con.execute( + """SELECT k.start FROM CUPTI_ACTIVITY_KIND_KERNEL k JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND s.value LIKE ? ORDER BY k.start""", + (ANCHOR,), +).fetchall() +ts = [r[0] for r in rows] +if not ts: + sys.exit(f"anchor {ANCHOR!r} not found") + +# Step boundaries are the large gaps in the anchor sequence; a step contributes LAYERS +# launches close together. Walk back from the end to collect WANT whole steps. +need = WANT * LAYERS +sel = ts[-need:] if len(ts) > need else ts +lo, hi = sel[0], sel[-1] +steps = len(sel) / LAYERS + +# Guard: if the chosen span still contains a huge gap it is not steady state. +gaps = [b - a for a, b in zip(sel, sel[1:])] +big = [g for g in gaps if g > 50_000_000] # >50 ms => iteration boundary +print(f"{DB}") +print(f" window {(hi-lo)/1e9:.2f} s steps={steps:.0f} wall/step={(hi-lo)/1e6/steps:.3f} ms") +if big: + print(f" WARNING: {len(big)} inter-iteration gaps in window " + f"(max {max(big)/1e6:.0f} ms) -- narrow --steps") + +tot, cnt = con.execute( + """SELECT SUM(end-start), COUNT(*) FROM CUPTI_ACTIVITY_KIND_KERNEL + WHERE deviceId=0 AND start>=? AND start<=?""", + (lo, hi), +).fetchone() +print(f" sum-of-durations={tot/1e6/steps:.3f} ms/step launches/step={cnt/steps:.0f}") +print(f"\n{'ms/step':>9} {'n/step':>7} {'avg us':>7} kernel") +for ms, n, us, name in con.execute( + """SELECT SUM(k.end-k.start)/1e6/?, COUNT(*)/1.0/?, AVG(k.end-k.start)/1000.0, + substr(s.value,1,64) + FROM CUPTI_ACTIVITY_KIND_KERNEL k JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND k.start>=? AND k.start<=? GROUP BY s.value + HAVING SUM(k.end-k.start)/1e6/? > 0.02 ORDER BY 1 DESC""", + (steps, steps, lo, hi, steps), +): + print(f"{ms:9.4f} {n:7.1f} {us:7.2f} {name}") diff --git a/skills/optimize-inference-siddharth/scripts/union_window.py b/skills/optimize-inference-siddharth/scripts/union_window.py new file mode 100644 index 00000000000..944684270ca --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/union_window.py @@ -0,0 +1,95 @@ +"""GPU busy (interval union) and exposed gap per step, over a windowed decode span. + +Sum-of-durations answers "how much kernel work is there" but not "how well is it +packed": it double-counts any two kernels that overlap, so a workload that runs its +collectives concurrently with compute looks busier than it is. The union of kernel +intervals is the honest busy number, and `wall - union` is the exposed time. + +Both traces here were captured with `--cuda-graph-trace=node`, which adds host +overhead per launch, so the *absolute* gap is inflated -- and inflated unequally, +more for an engine that launches eagerly than for one replaying a whole-step graph. +Read the union as a lower bound on packing quality, and confirm any conclusion +against unprofiled step time before acting on it. + +A second caveat that bit once: the last N steps of two runs are not necessarily the +same *workload*. Decode cost grows with sequence length, so two windows sitting at +different average seqlen differ in every bucket at once -- one such pair differed by +27% on the expert GEMM. Before attributing a per-bucket time change to a code change, +check that the buckets the change cannot touch held still; if they did not, compare +launch *counts* (seqlen-invariant) instead. + +Usage: + python union_window.py trace.sqlite '' [steps] [layers] +""" + +import sqlite3 +import sys + + +def main(): + db, anchor = sys.argv[1], sys.argv[2] + want = int(sys.argv[3]) if len(sys.argv) > 3 else 60 + layers = int(sys.argv[4]) if len(sys.argv) > 4 else 48 + + con = sqlite3.connect(db) + ts = [ + r[0] + for r in con.execute( + """SELECT k.start FROM CUPTI_ACTIVITY_KIND_KERNEL k + JOIN StringIds s ON k.demangledName=s.id + WHERE k.deviceId=0 AND s.value LIKE ? ORDER BY k.start""", + (anchor,), + ) + ] + if not ts: + sys.exit(f"anchor {anchor!r} not found") + sel = ts[-want * layers :] if len(ts) > want * layers else ts + lo, hi = sel[0], sel[-1] + steps = len(sel) / layers + wall = hi - lo + + iv = con.execute( + """SELECT k.start, k.end FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.deviceId=0 AND k.end > ? AND k.start < ? ORDER BY k.start""", + (lo, hi), + ).fetchall() + + union = 0 + gaps = [] + cs, ce = iv[0] + for a, b in iv[1:]: + if a > ce: + union += ce - cs + gaps.append(a - ce) + cs, ce = a, b + else: + ce = max(ce, b) + union += ce - cs + + total_gap = wall - union + # sum/union is the overlap ratio: 1.00 means fully serialized, and it is the only + # place a packing difference shows up when the per-bucket budgets are at parity. + dur_sum = sum(b - a for a, b in iv) + print(f"{db}") + print(f" steps={steps:.0f} wall/step={wall/1e6/steps:.3f} ms") + print(f" busy(union)/step={union/1e6/steps:.3f} ms ({100*union/wall:.1f}% of wall)") + print( + f" sum-of-durations/step={dur_sum/1e6/steps:.3f} ms " + f"overlap ratio={dur_sum/union:.3f}x (1.000 = fully serialized)" + ) + print(f" gap/step={total_gap/1e6/steps:.3f} ms n_gaps/step={len(gaps)/steps:.0f}") + # Bands, because one 500 us stall and 400 sub-microsecond ones need different fixes. + bands = [(0, 1_000), (1_000, 5_000), (5_000, 20_000), (20_000, 100_000), (100_000, 1 << 62)] + print(f"\n {'gap band':>16} {'n/step':>8} {'ms/step':>9} {'share':>7}") + for a, b in bands: + sel_g = [g for g in gaps if a <= g < b] + ms = sum(sel_g) / 1e6 / steps + label = f"{a/1000:g}-{'inf' if b > 1e15 else f'{b/1000:g}'} us" + print( + f" {label:>16} {len(sel_g)/steps:>8.0f} {ms:>9.3f} " + f"{100*sum(sel_g)/total_gap if total_gap else 0:>6.1f}%" + ) + + +if __name__ == "__main__": + main() diff --git a/skills/run-qwen-model/EXPERIMENTS.md b/skills/run-qwen-model/EXPERIMENTS.md new file mode 100644 index 00000000000..b81ed871c6d --- /dev/null +++ b/skills/run-qwen-model/EXPERIMENTS.md @@ -0,0 +1,3590 @@ +# Qwen3-30B-A3B optimization ledger + +Goal: make Megatron-Core EP4 inference match or exceed the vLLM DP4+EP +throughput on one OCI 4×GB200 node without correctness regressions. + +This ledger starts from scratch. Append every experiment, including failures +and regressions. Never edit an earlier result after it is recorded. + +## Fixed protocol + +| Setting | Value | +|---|---| +| Cluster | OCI `oci-hsg` | +| Hardware | 1 node, 4×GB200 | +| Model | Qwen3-30B-A3B, BF16 | +| Dataset / batch | gsm8k / 256 | +| Throughput workload | OSL1024, 2 warmups, 5 timed iterations | +| Profile workload | OSL128, one short warmup request, one timed request | +| mcore layout | TP=1, PP=1, EP=4, ETP=1 | +| vLLM layout | TP=1, DP=4, `--enable-expert-parallel` | +| Correctness gate | Fixed temperature-0 coherence prompts plus benchmark success | +| Primary metric | Throughput (output tokens/s) | +| Secondary metrics | Average latency and TPOT | + +Do not compare results when hardware, checkpoint, batch size, output length, +parallelism, or warmup/timed counts differ. + +## Baselines + +| ID | Engine | Throughput | Avg latency | TPOT | Job / run | Nsight trace | Status | +|---|---|---:|---:|---:|---|---|---| +| VLLM-BASELINE | vLLM DP4+EP | 23,606.7 tok/s | 1,368.2 ms | 10.844 ms/tok | 5547673 / `vllm-qwen30b-nsys-20260722-093437` | `vllm_profile.nsys-rep`, `.sqlite` | Pass | +| MCORE-BASELINE | mcore EP4/TP1 | 12,346.1 tok/s | 2,590.1 ms | 20.735 ms/tok | 5553135 / `qwen-30b-nsys-20260722-161020` | `mcore_profile.nsys-rep`, `.sqlite` | Pass | + +Fresh profile gap: mcore delivers 52.30% of vLLM throughput and is 47.70% +below the target. vLLM is 1.912× faster on this profile workload. + +Baseline order is mandatory: + +1. Record `VLLM-BASELINE` with Nsight Systems. +2. Record `MCORE-BASELINE` with Nsight Systems. +3. Compute the absolute and percentage gap. +4. Only then modify Megatron-Core. + +## Experiment index + +| ID | Date | Hypothesis | Changed files / flags | Throughput | Delta vs mcore baseline | Correctness | Job / run | Conclusion | +|---|---|---|---|---:|---:|---|---|---| +| VLLM-BASELINE | 2026-07-22 | Establish the fixed competitor target | none | 23,606.7 | n/a | Benchmark pass | 5547673 | Target established | +| MCORE-BASELINE | 2026-07-22 | Establish the fixed EP4 starting point | `max_requests=256` | 12,346.1 | baseline | Benchmark pass | 5553135 | Starting point established | +| QWEN-001 | 2026-07-22 | Single-kernel FC1+SwiGLU+FC2+topk-reduce mega-fusion beats the 4-kernel vLLM MoE path | `megatron/core/inference/moe/fused_moe_decode.py` (new), `dev/moe_fused/harness.py` (new) | microbench only | n/a (0.68–0.80× kernel) | Numerics pass (max_abs 2.6e-5, allclose) | session `qwen-moe-kernel` | Rejected — fused kernel 20–50% slower than reference; not integrated | +| QWEN-002 | 2026-07-22 | Fusing SiLU(gate)*up into the FC1 GEMM epilogue (removing bounded_silu_mul + the 2N round-trip) speeds up the decode MoE path without hurting FC2 tiling | `vllm_fused_moe.py` (FUSE_SWIGLU), `experts.py` (`fuse_fc1_activation=True`), `dev/moe_fused/harness_fc1.py`, `dev/moe_fused/run_e2e_insession.sh` (new) | **22,741.9 tok/s** (OSL1024) | **+0.55%** vs same-env fusion-off (22,617.7) | Coherent + numerics (max_abs 3.9e-5) | session `qwen-moe-kernel` in-session A/B | Accepted — MoE path 1.25×, e2e +0.55% throughput / −0.55% TPOT, no regression | +| PROFILE-OSL1024 | 2026-07-22 | Re-profile at the real throughput regime (BS256/OSL1024) to find the true vLLM→mcore gap and the dominant decode bottleneck | none (profiling only); `dev/moe_fused/profile_insession.sh` (new), `dev/moe_fused/vllm_osl1024_tput.sbatch` (new) | vLLM **33,994.5** vs mcore **~22,700** tok/s | mcore = **66.8%** of vLLM (**vLLM 1.50×**) | Both coherent | vLLM job `5555787` (tput) + `5555868` (nsys); mcore in-session `prof256a` | Gap is real & large at OSL1024. mcore decode: GPU 79% busy; MoE grouped-GEMM 41%, attn 22%, MoE routing 12% (49k tiny kernels), **exposed EP comm 11.5%**, norm/elt 10%, GPU idle 21%. vLLM uses TRT-LLM fused MoE (1-kernel routing + cutlass bmm + fused finalize, ~~0 exposed comm~~ — **comm claim falsified, see VLLM-COMM-S16**). Next target: routing-kernel storm + exposed EP comm (pending approval) | +| SESSION2-BASE | 2026-07-23 | Re-establish clean OSL1024 baseline in fresh session `qwen-opt` (fusion on, histogram off) before autonomous optimization campaign | none (config = current best) | **22,398.9 tok/s** | baseline for session 2 (−1.5% vs QWEN-002 run, within variance) | Coherent | session `qwen-opt` run `e2ebase` | Clean reference; 65.9% of vLLM 33,994.5 | +| QWEN-003 | 2026-07-23 | Replacing per-pair `atomic_add` in the MoE local-token count kernel with a `tl.histogram` variant (one atomic/bin/CTA) cuts the routing-kernel cost | `permute.py` (`_count_local_tokens_kernel_histogram`, env `MCORE_MOE_HISTOGRAM_COUNT`), `dev/moe_fused/harness_count.py` (new) | microbench only | **0.96× (wash)** | **EXACT integer match** vs reference | session `qwen-opt` `harness_count.py` | Rejected — count-kernel cost is per-launch fixed overhead, not atomic contention; in-kernel rewrite can't help. Default OFF. Real lever = fewer launches + less host-scheduling idle | +| QWEN-004 | 2026-07-23 | flashinfer sampling backend is faster than torch sampling | `run_e2e_cfg.sh` (`--inference-dynamic-batching-sampling-backend flashinfer`) | server crash | n/a | n/a | session `qwen-opt` `flashinfer` | Rejected — incompatible with `full_iteration_inference` CUDA graph capture: `RuntimeError: Generator not registered with the capturing graph` in `flashinfer.sampling`. torch sampling stays. | +| QWEN-005a | 2026-07-23 | `async-sched-mode=serial` overlaps host resolve with next forward, hiding the ~2.3ms/step (21%) GPU idle | `run_e2e_cfg.sh` (`--inference-dynamic-batching-async-sched-mode serial`) | server crash | n/a | n/a | session `qwen-opt` `serial` | Blocked by explicit guards: `ValueError: Async scheduling does not support expert parallelism` (+ separate MoE guard). Guards env-gated for experiment → QWEN-005b | +| QWEN-005b | 2026-07-23 | Guards were merely conservative; async serial works for MoE+EP if opened | env-gated EP+MoE guards in `dynamic_engine.py` + `text_generation_controller.py` (`MCORE_ALLOW_ASYNC_MOE`) | **hang** | n/a | server init OK but **first decode request hangs** (>2min) | session `qwen-opt` `asyncmoe` (cancelled) | Rejected — the guard encodes a real limitation: async serial + nvls alltoall + EP deadlocks on the first decode step. Patch reverted (tree clean). Would need real engine work to support. | +| QWEN-006 | 2026-07-23 | `nccl` AllGather/ReduceScatter inference dispatcher overlaps/costs less than `nvls` | `run_e2e_cfg.sh` (`--inference-moe-token-dispatcher-type nccl`) | **14,677 tok/s** | **0.66× (much worse)** | Coherent | session `qwen-opt` `nccl` | Rejected — nccl pads to worst-case per-rank token count (fixed-count AllGather), inflating comm volume ~2×. `nvls` variable-count stays the best dispatcher. | +| QWEN-007 | 2026-07-23 | `async-sched-mode=serial` (guards opened) hides the between-graph idle at the real OSL1024 regime | env-gated guards + `run_e2e_cfg.sh` | **22,589 tok/s** | **+0.85% (marginal)** | Runs at BS256 (hang was single-request-only) | session `qwen-opt` `async1024` | Marginal. Proves decode at OSL1024 is **NOT idle-bound** (the earlier "21% idle" was an OSL256 prefill artifact). Not worth shipping (unsupported path + tiny gain). Guards left env-gated default-off. | +| QWEN-008 | 2026-07-23 | `CUDA_DEVICE_MAX_CONNECTIONS=8` (was hardcoded 1) lets comm & compute overlap on separate HW queues, hiding exposed NVLS comm | `run_e2e_cfg.sh` (env override) | **~22,556 tok/s** | **~flat (noise)** | Coherent | session `qwen-opt` `maxconn8` | Reject — overlap is bounded by the full-iteration CUDA-graph structure / data deps, not connection count. No effect. | +| QWEN-010 | 2026-07-23 | `torch` grouped-GEMM backend (`torch.nn.functional.grouped_mm`, cuBLAS) beats the vLLM Triton fused-MoE backend on GB200 | `run_e2e_cfg.sh` (`--inference-grouped-gemm-backend torch`) | **18,434 tok/s** | **0.82× (worse)** | Coherent | session `qwen-opt` `gemmtorch` | Rejected — vLLM Triton backend stays best. (`flashinfer` cutlass backend is blocked for SwiGLU; only torch/vllm allowed.) All three backends now evaluated → vLLM is optimal. | +| QWEN-009 | 2026-07-23 | Built-in `--moe-router-fusion` (TE fused softmax+topk) + `--moe-permute-fusion` cut the routing critical path (~18%) | `run_e2e_cfg.sh` EXTRA_SERVER_ARGS | server crash | n/a | n/a | session `qwen-opt` `routperm`/`routfus` | Rejected — both crash: `AssertionError: hidden_size mismatch: 128 vs 8`. TE fused router emits a dense **128-expert** routing map, but `InferenceTopKRouter` (transformer_impl=inference_optimized) uses a dense **top-8** contract for the vLLM/nvls dispatcher. The built-in fusions are wired to the training MoE path only. A hand-written fused softmax+topk would need to honor the top-8 inference contract. | +| PROFILE-DECODE | 2026-07-23 | Get the TRUE per-step decode bottleneck (prior OSL256 totals were prefill-contaminated) | analysis of archived `mcore_osl256.sqlite`, pure-decode window (t0+220s, big-dispatch-free) | n/a | n/a | n/a | local sqlite | **Corrected model** (decode GPU-time share): MoE grouped-GEMM (`_fused_moe_kernel`) **~40% #1**, routing (count/moe_sum/topk/scatter/softmax/meta) ~18%, exposed comm (dispatch 122k + combine 320k) ~16%, attention ~13%, norm/elt ~10%. Kernels overlap across streams → **wall = per-layer critical path** (attn→router→dispatch→GEMM→combine). vLLM wins via TRT-LLM fused MoE (fused routing+finalize; the ~~0 exposed comm~~ half is **falsified, see VLLM-COMM-S16**). Explains QWEN-002 1.25× kernel → +0.55% e2e. | +| CLEANBASE-S3 | 2026-07-24 | Fresh clean un-profiled OSL1024 baseline (SwiGLU fusion OFF) in session `qwen-fuse` before session-3 fusion campaign | none (`MCORE_FUSE_FC1_ACT=0`) | **22,241.5 tok/s** | baseline (session 3) | Coherent | session `qwen-fuse` `cleanbase` | Clean reference; 65.4% of vLLM un-profiled 33,994.5. avg_latency 11,494 ms | +| QWEN-002-CONFIRM | 2026-07-24 | Re-measure QWEN-002 SwiGLU FC1-epilogue fusion (lever #3) cleanly at OSL1024 | `experts.py` (`MCORE_FUSE_FC1_ACT=1`), `vllm_fused_moe.py` FUSE_SWIGLU | **22,269.5 tok/s** | **+0.13% vs OFF** (wash) | Coherent | session `qwen-fuse` `qwen002on` | Confirms QWEN-002: MoE-kernel 1.25× but e2e wash at OSL1024 — FC1 activation is not the decode wall bottleneck. Kept on (free, exact). | +| CLEANBASE-S4 | 2026-07-25 | Fresh clean un-profiled OSL1024 reference in session `qwen-cutlass2` at the session-3 best config, before the grouped-GEMM/cutlass campaign | none (`MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`, vllm backend) | **22,657.7 tok/s** | baseline (session 4); −0.17% vs QWEN-011 (within drift) | Coherent | session `qwen-cutlass2` `ref0-1785010433` | Clean reference; 66.65% of vLLM un-profiled 33,994.5. avg_latency 11,282 ms, TPOT 11.299 ms/tok | +| QWEN-012 | 2026-07-25 | **Decision gate**: quantify grouped-GEMM headroom before writing any kernel — is the ~40% decode share recoverable inefficiency or irreducible expert-weight traffic? | none (analysis); `dev/moe_fused/harness_roofline.py` (new) | n/a (microbench) | n/a | n/a | session `qwen-cutlass` `c05fd5f7` | **The decode grouped GEMM is memory-bound, not FLOP-bound.** Weight traffic 302 MB/layer/rank vs a *measured* 6.081 TB/s streaming-read ceiling ⇒ 49.66 µs floor; production FC1+FC2 = 72.13 µs = **1.45× off roofline**. Achieved 63–83 TFLOP/s on valid FLOPs (~3% of GB200 BF16 peak). Padding waste is real (74.3% dead rows at BLOCK_M=64) but nearly **free**: cutting it 3.89×→1.49× buys only ~1.2×. **⇒ Stage 4 (hand-written CUTLASS/CuTe grouped GEMM) cannot win**: its entire ceiling is 1.45×, and 1.26× of that is reachable by Triton tile retuning alone (→ QWEN-013). Remaining levers are elsewhere (exposed EP comm ~16%, routing ~18%) | +| QWEN-013 | 2026-07-25 | Tiling, not the GEMM implementation, is the gap QWEN-012 found: FC1 (N=768) and FC2 (N=2048) want *opposite* BLOCK_SIZE_N, which vLLM's single shared config cannot express | `vllm_fused_moe.py` (`_get_decode_tuned_configs`, per-GEMM `config_fc1`/`config_fc2`, env `MCORE_MOE_GEMM_TUNE`), `dev/moe_fused/harness_gemmtune.py` (new) | **23,636.0 tok/s** | **+4.32% vs CLEANBASE-S4** | **Bit-exact** (max_abs 0.0 at 128/256/384/512 tokens) + coherent | session `qwen-cutlass2` `gemmtune-1785010791` | **Accepted** — GEMM GPU time 77.65→61.89 µs (1.255×), whole MoE call 100.64→84.31 µs (1.194×), e2e +4.32% with tight variance (23.55–23.67k). Bit-exact because `BLOCK_SIZE_K` is unchanged, so the fp32 K-reduction order is identical. Default OFF; enable with `MCORE_MOE_GEMM_TUNE=1`. Now **69.53%** of vLLM | +| QWEN-013b | 2026-07-25 | The QWEN-013 tuned tiles regress past the decode point, so the fallback threshold should be 384, not 512 | `vllm_fused_moe.py` (`_get_decode_tuned_configs` gate `M > 512` → `M > 384`) | **23,646.0 tok/s** | **+4.36% vs CLEANBASE-S4** (+0.04% vs QWEN-013, i.e. same) | Bit-exact (unchanged kernel) + coherent | session `qwen-cutlass2` `gemmtune2-1785012743` | **Accepted** — CUDA-graph device time vs default tiles is 1.289×/1.197×/1.120×/**0.952×** at 128/256/384/512 tokens, so 512 was a 5% regression inside the tuned range. Decode runs at 256 so e2e is unchanged (23,646.0 vs 23,636.0, within variance); the guard just removes a latent regression for other batch shapes. **Current best config.** | +| QWEN-014 | 2026-07-25 | The already-wired `--inference-grouped-gemm-backend flashinfer` (`flashinfer.fused_moe.cutlass_fused_moe`) accepts BF16 and beats the vLLM Triton path; the ledger's "blocked for SwiGLU" note (QWEN-010) is unsupported by the code | none shipped; `dev/moe_fused/inspect_flashinfer.py`, `dev/moe_fused/harness_flashinfer.py` (new) | microbench only | **0.921× vs QWEN-013** (1.079× vs default tiles) | Numerics pass only with a weight-layout fix (see record) | session `qwen-cutlass2` exec `7590adcb`/`659d2633` | **Rejected on measurement, but the ledger note was wrong.** The kernel *does* support BF16 gated SwiGLU; mcore's wiring is broken in two independent ways (an `ActivationType` mis-map that hard-fails, and a gate/up ordering mismatch that silently corrupts numerics — both root-caused below). After fixing both in the harness, CUDA-graph device time is **90.81 µs vs 83.63 µs** for the retuned Triton path — 8% slower — so fixing the backend would not win. Consistent with QWEN-012: everything is pinned near the weight-bandwidth floor. | +| QWEN-015 | 2026-07-25 | `flashinfer.fused_moe.trtllm_bf16_routed_moe` — the actual TRT-LLM-Gen kernel vLLM wins with — can be dropped in as a new grouped-GEMM backend | none (blocked before any Megatron change); probe in `dev/moe_fused/harness_flashinfer.py` | not reached | n/a | n/a | session `qwen-cutlass2` exec `7590adcb` | **Blocked on weight layout** — not a contract or EP problem. `use_shuffled_weight=False, weight_layout=MajorK` is rejected outright (`BF16 Moe: weight_layout must be BlockMajorK`), and `BlockMajorK` reads `size(3)` of the weights (`IndexError: Index 3 out of bounds for tensor with 3 dimensions`): it requires **4-D pre-shuffled block-major** weights, not mcore's 3-D `[E, 2*ffn, H]`. Path forward and why it is now low priority are in the record. | +| PROFILE-TUNED | 2026-07-25 | Confirm the decode timing composition after QWEN-013 and re-rank the remaining levers | none (profiling only) | n/a | n/a | n/a | session `qwen-cutlass2` `prof/tuned-1785013555` (`mcore_profile.nsys-rep`, `.sqlite`) | Workflow C on one steady-state decode step (BS256/OSL128, device 3, 302-step window): forward pass **9.933 ms**, GPU-busy 7.884 ms, **idle 2.049 ms (20.6%)**, 1362 kernels/step. Share of GPU time: **MoE expert GEMM 2611 µs (32.9%, 96 kernels, 27.2 µs/kern)** — down from ~40% pre-retune — routing/permute 1251 µs (15.8%, 242 kernels), dense GEMM 1092 µs (13.8%), **exposed EP comm 929 µs (11.7%, 96 kernels, 9.7 µs/kern)**, attention 913 µs (11.5%), elementwise 528 µs (6.7%), norm 462 µs (5.8%). **Corroborates QWEN-012/013 independently**: 2611 µs / 48 layers = 54.4 µs per layer for FC1+FC2, matching the ~57 µs microbench, and leaving only 2611−48×49.66 = **229 µs/step (2.3%)** above the weight-bandwidth floor. The GEMM lever is spent; comm + routing + idle are what remain. | +| QWEN-011 | 2026-07-24 | Fuse the MoE indirection-table build (lever #1): merge `_init_sorted_ids` + `_prefix_sum` + `_fill_expert_block_ids` into one `_prefix_fill_init_kernel`, cutting the 5-kernel routing storm to 3 kernels/layer (−96 launches/step) | `vllm_fused_moe.py` (`_prefix_fill_init_kernel`, `_moe_align_block_size_fused`, env `MCORE_MOE_FUSED_ALIGN`), `dev/moe_fused/harness_align.py` (new) | **22,696.4 tok/s** | **+1.92% vs QWEN-002-CONFIRM** (+2.04% vs clean) | **EXACT** (max_abs_diff 0.0, allclose) + coherent | session `qwen-fuse` `fusedalign`, branch `perf/moe-fused-align` | **Accepted** — bit-exact, MoE microbench 1.029×, e2e +1.9% with tight variance (22.65–22.71k). Reduces routing/permute kernel count (lever #1). Now 66.8% of vLLM. | +| QWEN-016 | 2026-07-25 | **Decision gate**: is the 929 µs/step of exposed NVLS EP comm a worthwhile lever — how much is genuinely exposed, is it latency- or bandwidth-bound, and what is the floor under perfect overlap? | none (analysis); `dev/moe_fused/probe_comm.py`, `analyze_comm.py`, `analyze_comm_skew.py`, `harness_comm.py` (new) | n/a (analysis) | n/a | Microbench gate bit-exact (AGV max abs 0.0, RSV max rel 0.0 over 64 trials) | session `qwen-comm` job 5601961, execs `08e6a692`/`04bdfdd6` | **Redirect off this lever.** 100% exposed (interval union = sum), but **latency-bound**: AGV 6.57 µs = 0.72 launch + **5.08 barrier** + 0.77 transfer; RSV 7.88 = 0.72 + **5.03** + 2.13. Bytes are only 127 µs/step of ~693 µs and RSV transfer is already 82% of the NVLink floor, so batching bytes cannot win. Of 1060 µs/step, 632 µs is intrinsic and 428 µs is inter-rank skew (ranks waiting on the slowest rank's expert GEMM — a routing-balance problem). Recoverable critical path is **6.4–7.0% of the step**; RSV-into-FC2 fusion keeps the barrier (≈1.4%) and 2-chunk pipelining adds a barrier per chunk (≈4.3%, and needs concurrent streams under graph capture — the QWEN-008 blocker). CTA count already optimal at the shipped 128. | +| QWEN-018 | 2026-07-25 | **Decision gate**: break the 1251 µs/step routing/permute category down by kernel name — launches, device µs, dispatch gap, and fixed-vs-work split — and compute a wall-time ceiling per candidate fusion before building anything | none (analysis); `dev/moe_fused/analyze_routing.py` (new) | n/a (analysis) | n/a | n/a | session `qwen-comm` job 5601961 (login-node re-analysis of `qwen-cutlass2:prof/tuned-1785013555`) | **Two of the six routing kernels are pathological, and neither is launch-bound.** Per steady-state decode step (40-step mean, device 3): `_moe_sum` 7.79 µs × 48, **`_count_local_tokens` 7.52 µs × 48**, `gatherTopK` 6.05, `_scatter_token_indices` 2.66, router softmax 1.94, `_prefix_fill_init` 1.33, plus a 0.74 µs `torch.zeros` fill = 1531 µs/step of wall including gaps. Inter-kernel dispatch gap is **0.55 µs**, so a removed launch is worth `duration + 0.55`, and the fixed floor per launch is only 0.72+0.55 = **1.27 µs**. Ceilings: **(1) `_moe_sum`→FC2 epilogue = 400 µs = 4.03% gross**, but needs cross-CTA fp32 atomics into `out` plus a zeroing launch — QWEN-001's exact failure shape — so ~3.3% net and possibly negative; **(2) cooperative-grid merge of count→prefix_fill→scatter = 122 µs = 1.23%** before paying two grid syncs ⇒ **gated out**; **(3) new: fold the count + its zeros-fill into `_prefix_fill_init_kernel` = 449 µs = 4.52% gross**, no grid sync needed and integer-exact ⇒ **chosen**. `_count_local_tokens` is slow because at `BLOCK_SIZE=1024` only **2 of its 152 CTAs get work**; QWEN-003 changed the reduction but not that, which is why it measured a wash | +| CLEANBASE-S6 | 2026-07-25 | Fresh same-session OSL1024 reference at the session-4/5 best config, in session `qwen-comm`, before the routing campaign | none (`MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`) | **23,264.4 tok/s** | baseline (session 6); −1.61% vs QWEN-013b's 23,646.0 (session drift) | Benchmark pass, 5/5 iters | session `qwen-comm` job 5601961, run `e2e/ref-s6-1785024868` | Clean reference; 68.44% of vLLM 33,994.5. avg_latency 10,975 ms, TPOT 11.004 ms/tok. Synced snapshot `workspaces/megatron_lm/6b9355b187072223` verified to **not** contain the session-6 kernel | +| QWEN-019 | 2026-07-25 | Fold the per-expert token count and the `torch.zeros` fill of its counter buffer into `_prefix_fill_init_kernel`, taking the decode indirection-table build from 4 launches to 2. No grid sync is needed because that kernel already has every CTA redundantly recompute the cumsum in registers | `vllm_fused_moe.py` (`_count_prefix_fill_init_kernel`, `_moe_align_block_size_count_fused`, env `MCORE_MOE_FUSED_COUNT`, `_FUSED_COUNT_MAX_TOKENS`), `dev/moe_fused/harness_countfuse.py` (new) | **23,964.4 tok/s** | **+3.01% vs CLEANBASE-S6** | **Bit-exact** (max_abs 0.0 *and* max_rel 0.0 at 128/256/384/512) + tables integer-identical + coherent | session `qwen-comm` job 5601961, run `e2e/countfuse-1785025539` | **Accepted** — align call 16.38→10.24 µs under CUDA-graph replay (1.600×), whole MoE call 86.02→79.96 µs (1.076×), e2e +3.01% with tight variance (23.57–24.30k). TPOT 11.004→10.682 ms/tok. Default OFF; enable with `MCORE_MOE_FUSED_COUNT=1` (requires `MCORE_MOE_FUSED_ALIGN=1`). New best = **70.49%** of vLLM 33,994.5 | +| QWEN-017 | 2026-07-25 | `symm_mem_sync` spins on a system-scope `atom.cas` (a full uncached RMW per attempt); an `ld.acquire.sys` poll + clearing store should cut the 5.05 µs barrier and thus ~485 µs/step | `communication/torch_symm_triton/barrier.py` (`_wait_signal_ldpoll`, env `MCORE_SYMM_BARRIER_LDPOLL`) — **reverted**; `dev/moe_fused/harness_comm.py` | microbench only | **−2.3% (regression)** | Bit-exact both variants (AGV max abs 0.0, RSV max rel 0.0 over 64 trials) | session `qwen-comm` job 5601961, exec `04bdfdd6` | **Rejected** — per-step comm 693.7 µs baseline vs 709.8 µs ldpoll (3 alternating reps each, no overlap). The barrier-only kernel is unchanged at 5.75–5.83 µs vs a 0.72 µs empty kernel, so the 5.05 µs is the 4-way system-scope flag round trip itself, not polling granularity — a hardware/driver floor a Triton rewrite does not move. Patch reverted; `megatron/` back to the session-4 state. | +| QWEN-020 | 2026-07-25 | `_moe_sum_kernel` spends 7.79 µs/layer moving ~6.3 MB — 7× its own bandwidth floor — because the per-topk-slot locality test is a uniform scalar branch on a dependent `routing_map` load, serialising the topk walk. Predicate it into the load mask and widen the K tile to the full hidden size | `vllm_fused_moe.py` (`_moe_sum_kernel_fast`, env `MCORE_MOE_SUM_FAST`, `_FAST_MOE_SUM_MAX_BLOCK_K`) | **24,403.7 tok/s** | **+1.83% vs QWEN-019**, **+4.90% vs CLEANBASE-S6** | **Bit-exact** (max_abs 0.0 *and* max_rel 0.0 at 128/256/384/512) + coherent | session `qwen-comm` job 5601961, run `e2e/moesum-1785026515` | **Accepted** — `_moe_sum` 8.14 → 5.97 µs/layer (1.36×), whole MoE call 80.69 → 78.77 µs at 256 tokens. TPOT 10.682 → 10.490 ms/tok. Same reduction order and same fp32 arithmetic (masked slots add an exact 0.0), so bit-exactness is structural, not luck. Default OFF; enable with `MCORE_MOE_SUM_FAST=1`. New best = **71.79%** of vLLM 33,994.5 | +| QWEN-021 | 2026-07-25 | The router's `torch.softmax` + `torch.topk(sorted=False)` pair costs 436.8 µs/step (4.4%) for a `[256, 128]` fp32 reduction; one CTA per token can do softmax in registers and select the top-8 by max-then-mask in a single kernel | `megatron/core/inference/moe/router_topk.py` (new), `megatron/core/transformer/moe/router.py` (`InferenceTopKRouter._forward` fused-path branch), env `MCORE_ROUTER_FUSED_TOPK`; `dev/moe_fused/harness_routertopk.py` (new) | **25,352.9 tok/s** | **+3.89% vs QWEN-020**, **+8.98% vs CLEANBASE-S6** | **Bit-exact probs and identical expert sets** at 128/256/384/512 × 2 seeds | session `qwen-comm` job 5601961, run `e2e/routertopk-1785027029` | **Accepted, largest single win of the session.** 4 kernels → 1: 16.41 → 4.10 µs under graph replay at 256 tokens (4.00×), eager per-kernel 15.54 → 2.04 µs. TPOT 10.490 → 10.097 ms/tok (−393 µs/step vs a predicted ~312 µs). Default OFF; `MCORE_ROUTER_FUSED_TOPK=1`. New best = **74.58%** of vLLM 33,994.5 | +| QWEN-022 | 2026-07-25 | Now that every CTA streams the routing pairs to build its own histogram (QWEN-019), the scatter can be a second streaming pass in the same kernel — no atomics, no grid sync — taking the decode indirection-table build from 2 launches to 1 | `vllm_fused_moe.py` (`_align_single_kernel`, `_moe_align_block_size_single`, env `MCORE_MOE_FUSED_SCATTER`), `dev/moe_fused/harness_scatterfuse.py` (new) | **25,495.9 / 25,441.4 tok/s** (two runs) | **+0.46% vs QWEN-021** (mean of two), inside the noise band | **Bit-exact** whole-MoE output at 128/256/384/512 + table equality (16 cases: `npp`, `expert_ids`, row multiset, and every row placed under an expert id that owns it) | session `qwen-comm` job 5601961, runs `e2e/scatterfuse-1785027598`, `e2e/scatterfuse-rep2-1785027955` | **Accepted, marginally.** Align call 10.08 → 8.23 µs and whole MoE call 82.17 → 79.56 µs at 256 tokens (1.033×), predicting 1.24%; e2e delivered **+0.35% and +0.55%** in two runs. Both runs beat QWEN-021, so the sign is reliable, but the size is not — the honest attribution is ~0.5%. The candidate is a **wash at 384 tokens (1.007×) and a slight loss at 512 (0.997×)**, since the second streaming pass scales with pairs while the removed launch does not. Default OFF; `MCORE_MOE_FUSED_SCATTER=1` (requires `MCORE_MOE_FUSED_COUNT=1`). New best = **75.00%** of vLLM 33,994.5 | +| PROFILE-S6 | 2026-07-25 | Re-profile at the session-6 configuration (all six gates on) and re-rank, since the step lost ~10% of its wall time and ~820 µs of routing | none (`profile_insession.sh`) | n/a (BS256/OSL128 profile run) | n/a | n/a | session `qwen-comm` job 5601961, `prof/s6-all-1785028240/mcore_profile.{nsys-rep,sqlite}` | **Routing is done as a lever.** One decode step: wall **9.097 ms** (was 9.933), GPU-busy 7.183 (was 7.884), idle 1.914 (21.0%), **1170 kernels** (was 1362). Routing/permute is now **430 µs in 98 kernels** (was 1251 µs in 242). New ranking: MoE expert GEMM 2469 µs (27.1%, closed), **host-side idle 1914 µs (21.0%)**, dense GEMM 1084, comm 1063 (closed), attention 924, elementwise 488, norm 450, routing 430. The idle is concentrated in **three host gaps totalling ~1237 µs/step**: 548 µs after `index_elementwise_kernel`, 386 µs after `vectorized_elementwise_kernel`, 303 µs after `CatArrayBatchedCopy_vectorized` — sampling/detokenize/scheduling between graph replays | +| CLEANBASE-S7 | 2026-07-26 | Fresh same-session OSL1024 reference at the seven-gate session-6 best config, in session `qwen-attnstate`, before the host-side `initialize_attention_state` lever | none (all seven gates on) | **25,805.9 tok/s** | baseline (session 7); +1.22% vs QWEN-022's 25,495.9 (session drift) | Coherent, 5/5 iters | session `qwen-attnstate` job 5613090, run `e2e/ref-s7-1785084185` | Clean reference; **75.91%** of vLLM 33,994.5. Per-iter 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 (spread 0.23%, the tightest reference this campaign has had). avg_latency 9,867 ms, TPOT 9.920 ms/tok | +| QWEN-023 | 2026-07-26 | HOSTGAP-S6's ~311 µs/step of `initialize_attention_state` is recomputation, not computation: at fixed BS256 the request set, sampling metadata, KV block table and CUDA-graph selection are identical for 128 consecutive decode steps, and only the KV sequence lengths advance. Recompute those and reuse the rest, invalidated by a request-layout version counter | `megatron/core/inference/contexts/dynamic_context.py` (`_bump_request_layout_version`, `_incr_attn_state_cache_key/_store/_snapshot/_verify`, `_incremental_attention_state_update`, env `MCORE_INFER_INCR_ATTN_STATE`, `MCORE_INFER_INCR_ATTN_STATE_VERIFY`, profiling gate `MCORE_INFER_ATTN_PROF`), `megatron/core/inference/contexts/attention_context/mha_metadata.py` (`restore_state_data`), `dev/moe_fused/harness_attnstate.py` (new) | **26,032.1 / 26,092.5 tok/s** (two OFF/ON pairs) | **+0.94%** (pairwise +0.88% and +1.01%) | **130 consecutive decode steps bit-identical** (cached vs freshly recomputed, every host and GPU bookkeeping buffer) + temperature-0 coherence output byte-identical to gate-OFF | session `qwen-attnstate` job 5613090, runs `e2e/incr-attn-on-1785085393`, `e2e/incr-attn-on-rep2-*` | **Accepted.** Host CPU per call **174.1 → 51.7 µs (3.37×)**; the residual is the H2D bookkeeping copy, which is irreducible. TPOT 9.915 → 9.823 ms/tok (−92 µs/step). Across all twenty timed iterations the two gates are **fully separated** (slowest ON 25,992.1 > fastest OFF 25,893.9), so the sign is certain even though the size is under 1%. Delivered ~28% of the 3.4% ceiling: the removable phases are ~72% of the call, and the H2D bookkeeping copy that remains cannot be cached. Default OFF; `MCORE_INFER_INCR_ATTN_STATE=1`. New best = **76.67%** of vLLM 33,994.5 | +| CGTRACE-CONTROL | 2026-07-26 | HOSTGAP-S6 could not tell whether the ~900 µs/step of CUDA-graph machinery (a 199 µs `cudaGraphLaunch` plus GPU-side inter-node dispatch) is real or an artifact of `--cuda-graph-trace=node` instrumenting all 1158 graph nodes. Capture the same workload both ways and compare | none (control); `dev/moe_fused/profile_insession.sh` (env `CUDA_GRAPH_TRACE`, plus the hardened nsys stop/qdstrm recovery ported from `profile_host_insession.sh`), `dev/moe_fused/analyze_cgtrace.py` (new) | n/a (two BS256/OSL128 profile runs) | n/a | n/a | session `qwen-attnstate` job 5613090, `prof/cgt-node-1785086713` and `prof/cgt-graph-1785086964` | **The machinery is real; node-mode instrumentation is not inflating the step.** Same rank, same steady-state decode window: step period **8882.5 µs (node) vs 8945.1 µs (graph)** — 0.71% apart and in the *wrong* direction for an instrumentation artifact — and host `cudaGraphLaunch` **190.0 vs 184.7 µs** (2.8%). Node mode traces 1169 kernels/step, matching PROFILE-S6's 1170. **HOSTGAP-S6's wall-time attributions therefore all stand**, and lever 4 (reduce graph node count, ~0.78 µs/node) is live rather than chasing a measurement artifact. Caveat: `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` is too sparsely populated (417 rows total, n=6 in-window on the analyzed device) to independently confirm the *GPU-side* inter-node dispatch component; only the step period and host submit cost are settled | +| HOSTGAP-S6 | 2026-07-26 | Attribute PROFILE-S6's three inter-kernel host gaps to concrete host functions using the recovered host-visibility trace | none (analysis); `dev/moe_fused/analyze_hostgaps.py` (new) | n/a (analysis) | n/a | n/a | session `qwen-host` job 5607600, `prof/hosts6-1785048883` | **All three gaps attributed; the largest is a true serial data dependency, not a scheduling artifact.** Per step (GPU-only trace, mean of 4 ranks, 9.134 ms step / 1965.7 µs idle): **G1 `index_elementwise`→`vectorized_elementwise` 571.4 µs** = `transfer_samples_to_cpu`+`active_request_mask`+`update_requests`+engine window+`initialize_attention_state`, **CPU-bound Python** (100% of samples Running in `_PyEval_EvalFrameDefault`, only 8.6% of the gap in CUDA API, 0.5% in `cudaStreamSynchronize`); **G2 `vectorized_elementwise`→`vectorized_gather` 331.0 µs in 2 instances** (267.9 µs inside `initialize_attention_state`+`transfer_bookkeeping_to_gpu`+`forward_pass` head, **CPU-bound Python**; 63.1 µs inside `sampling`, launch+`cudaStreamSynchronize`); **G3 `CatArrayBatchedCopy`→`rmsnorm_fwd` 305.4 µs** = **one `cudaGraphLaunch` costing 199.1 µs median** for a 1158-node graph (0.172 µs/node), driver-side, host Running in `cuGraphLaunch`. No rank asymmetry (≤4%). Idle splits **1065 µs real host Python (11.7% of step)** vs **900 µs graph machinery (9.9%)** that is `--cuda-graph-trace=node`-artifact-suspect. Async overlap cannot fix G1: the chain is data-dependent on the current step's sampled tokens, which explains QWEN-007's +0.85% | +| QWEN-025 | 2026-07-26 | QWEN-024's decision-gate measurement showed 84% of `post_process_requests` is reducible by collapsing the per-request loop body — unvectorizable is not the same as irreducible. Add a decode fast path that is structurally unable to touch the request termination state machine | `dynamic_engine.py` (`_post_process_requests_decode_fast`, `_ppr_cache_epoch` invalidation at all three record-mutation sites, env `MCORE_INFER_FAST_POST_PROCESS`, `..._VERIFY`), `dev/moe_fused/harness_updreq.py` (`run_ppr_equivalence`, in-process gate A/B) | **26,361.7 tok/s** (three-pair ON mean); best run 26,430.7 | **+0.98%** vs same-session gate-OFF (pairwise +0.99% / +1.26% / +0.68%; +0.59% on a best-4-of-5 trim) | 400-step two-engine equivalence with 10 finisher steps and 200 steps past the token limit; 300 steps under VERIFY; coherence byte-identical | session `qwen-updreq` job 5616264, runs `e2e/ppr-{off,on}{,-r2,-r3}-*`, profiles `prof/g2-off-1785106875` and `prof/g2-on-1785107105` | **Accepted.** `post_process_requests` **211.0 → 27.2 µs/step (7.77×)**; whole post-sampling host chain **372.4 → 148.0 µs/step** across QWEN-024+025. The fast path **declines any step on which a request finishes**, so the termination state machine is unreachable from it — that is the safety argument, not extra testing. cProfile 261 calls/step vs the pre-registered <500 threshold. 13 of 15 ON iterations beat every one of the 15 OFF iterations. G1 −56.1 µs, idle −77.9 µs, kernels/step unchanged at 1169. **Only ~1/3–1/2 of the 183.9 µs host saving converts** (vs ~1:1 for QWEN-024) because `async_bookkeep` was already partly overlapped with GPU work. Default OFF. New best = **77.75%** of vLLM 33,994.5 | +| PROFILE-S10 | 2026-07-27 | Re-profile at the twelve-gate config (ten gates + QWEN-026 + QWEN-027) and, for the first time this campaign, take a **clean 2 s decode window** (densest-window search, not a span fraction) to get an honest category breakdown and decide whether more small-kernel fusion is worth it | none (`profile_insession.sh` BS256/OSL256; `prof/s10-addnorm-1785188237/window_breakdown.py`, new) | n/a | n/a | n/a | session `qwen-nodecut`, `prof/s10-addnorm-1785188237/mcore_profile.{nsys-rep,sqlite}` | **Both fusions confirmed live and each removed ~48 kernels/step**: `rmsnorm_fwd_tuned` 104→56/step, `triton_poi_fused_add_copy` 103→55/step; `_fused_qk_rmsnorm` and `_fused_add_rmsnorm` each 55/step (1/layer). Clean-window (union-busy **81.6%**, idle **18.4%** = 360 µs-scale host gaps) device-time by category: **moe_expert_gemm 587 ms**, **comm/NVLS 247 ms**, **dense_gemm 236 ms**, **attention(flash) 211 ms**, routing 69, norm 67, elementwise 66, moe_sum 65, rope_kv 15. **Verdict: small-kernel fusion is now in diminishing returns.** The fusion-addressable buckets (norm/elementwise/routing) are each ~65-70 ms and partly overlapped on other streams, so each remaining fusion is ~0.5-1%; the gap to vLLM is now dominated by expert GEMM (compute-bound), NVLS comm, attention, and the 18.4% host idle (blocked on the nsys host-capture failure). Next levers should target those, not more small-kernel merges | +| HOSTGAP-S10 | 2026-07-27 | With small-kernel fusion in diminishing returns (PROFILE-S10), attack the biggest remaining prize: the **18.4% GPU-idle**. Attribute it to host call sites **without** the deadlock-prone host-visibility capture, by using the CUDA-API host rows already in the clean `cuda,nvtx` trace | `idle_attrib2.py`, `bracket.py` (new, in `prof/s10-addnorm-*/`); env-gated `MCORE_INFER_NVTX` toggle added to `dynamic_engine.py`; `CUDA_GRAPH_TRACE` + `TRITON_LIBCUDA_PATH` + `HOST_LEVEL`/`OSRT`/`SAMPLE` knobs added to `profile_host_insession.sh` | n/a | n/a | n/a | traces on lustre `sessions/qwen-nodecut/prof/s10-addnorm-1785188237/`; session `qwen-host` (job 5638180) | **Idle decomposition (clean 1 s decode window, 178 ms idle):** small gaps <10 µs = 59.6 ms (33%, intra-graph kernel scheduling, unavoidable); **large gaps ≥10 µs = 118.8 ms (67%) = host chain between steps** (median 37 µs). Large-gap attribution via the decode thread's CUDA-API rows: **UNCOVERED (Python/CPU compute, no CUDA call) = 87.8 ms = 73.9%**, memcpy(token D2H) 8.9%, graph_launch 7.5%, launch 6.9%, sync_wait 1.6%. So ~74% of attackable idle (~8-9% of wall) is Python on the critical path. **Localized** by bracketing each Python gap with its surrounding CUDA APIs: 47.6% sits in `cudaMemcpyAsync -> cuKernelGetName`, +15.7% `cudaLaunchKernel -> cudaMemcpyAsync`, +11% `cuKernelGetName -> cudaMemcpyAsync` — i.e. the **between-steps host orchestration** after the sampled-token D2H copy (post-sampling bookkeeping / scheduling / attention-state prep), not the GPU-sync wait. **nsys dead-end recorded:** the host-visibility capture deadlocks nsys finalization and `QdstrmImporter` rejects the qdstrm whenever **osrt tracing**, **`--sample=process-tree`**, OR **`MCORE_INFER_NVTX` (NVTX ranges under CUDA graphs)** is on — independent of `--cuda-graph-trace` level. Only the plain `cuda,nvtx --sample=none` set finalizes. **Actionable conclusion:** future idle work should (a) use clean-trace RUNTIME attribution (this method), not host capture; (b) target the between-step Python orchestration via `perf_counter` phase timing (no nsys), since NVTX-based naming is blocked | +| QWEN-027 | 2026-07-27 | PROFILE-S9 showed the two dominant small kernels are `triton_poi_fused_add_copy` (bias-dropout-add residual, 103/step) and `rmsnorm_fwd_tuned` (104/step), adjacent on the serial path. Fuse the residual-add with the *standalone* `pre_mlp_layernorm` (MoE) into one kernel (vLLM's `add+rms_norm`), returning both the updated residual and the normed MLP input — the input norm is inside the TE `LayerNormLinear` QKV GEMM and not interceptable, so only the pre-MLP boundary (48/step) is fusible | `megatron/core/inference/fused_add_rmsnorm.py` (new: `_fused_add_rmsnorm_kernel`, `fused_add_rmsnorm`, `can_use_fused_add_rmsnorm`), `megatron/core/transformer/transformer_layer.py` (guarded fused path at the `self_attn_bda` site + `_forward_pre_mlp_layernorm` consumes the stashed normed output), env `MCORE_FUSED_ADD_NORM`, `MCORE_FUSED_ADD_NORM_MAX_TOKENS` (default 256); `dev/moe_fused/harness_addnorm.py` (new) | **27,525.2 / 27,532.1 tok/s** (two runs, mean 27,528.6) | **+1.37%** vs QWEN-026-on 27,151.9; arms **fully separated** (min ON 27,503.6 > max ref 27,206.4) | **residual BIT-EXACT** (fp32-accumulate add reproduces torch's bf16 add), **norm bf16-ulp** (max_rel ≤7.9e-3 ≈ one bf16 ulp). Coherence: **all 3** temperature-0 prompts identical to the QWEN-026-on arm (incl. the France prompt that QWEN-026 flipped), fluent + factually correct | session `qwen-nodecut`, runs `e2e/addnorm-on-1785184454`, `e2e/addnorm-on-2-1785185287`; microbench, profile `prof/s9-qknorm-1785182903` | **Accepted.** Microbench CUDA-graph replay at 256 tokens: (add + TE norm) 6.16 → fused 4.11 µs (**1.50×**), and **2.0× at 384/512** — unlike QWEN-026 this kernel does not regress at higher token counts (flat ~4.1 µs). Removes 48 add launches + 48 graph nodes/step. Heavily gated: decode-only, MoE standalone RMSNorm, no bias/dropout/recompute/offload, no cross-attn, with a reference fallback so training and all non-matching cases are byte-identical. TPOT 9.428 → 9.299 ms/tok. Default OFF; `MCORE_FUSED_ADD_NORM=1`. New best = **80.98%** of vLLM 33,994.5 | +| PROFILE-S9 | 2026-07-27 | Re-profile at the eleven-gate config (ten gates + QWEN-026) and re-rank the remaining small-kernel serial chain | none (`profile_insession.sh`, BS256/OSL256) | n/a | n/a | n/a | session `qwen-nodecut`, `prof/s9-qknorm-1785182903/mcore_profile.{nsys-rep,sqlite}` | **QWEN-026 confirmed live** (`_fused_qk_rmsnorm_kernel` 51.7/step = 1/layer; the old two `rmsnorm_fwd_general` are gone). The two dominant small kernels are now **`rmsnorm_fwd_tuned` 104/step @ 2.93 µs** (input + pre-MLP norm, 2/layer) and **`triton_poi_fused_add_copy` 103/step @ 1.51 µs** (the bias-dropout-add residual, 2/layer), and they are adjacent: the top consecutive pair by count is `triton_poi_fused_add_copy → rmsnorm_fwd_tuned` at 103/step. vLLM runs these as one `triton_red_fused__to_copy_add_..._rms_norm`; mcore runs them as two. **Next lever = fused add+RMSNorm** (~96 add-launches + 96 graph nodes/step removed, 4.44 → ~2.5 µs/boundary), the same convert-at-3× serial-chain + node-removal pattern QWEN-026 proved. Ranked #1 | +| QWEN-026 | 2026-07-27 | Qwen attention applies `q_layernorm` and `k_layernorm` as two separate per-head RMSNorm module calls → two `rmsnorm_fwd_general` launches (and two graph nodes) per layer, each normalizing a tiny 128-wide row (launch/latency bound, not bandwidth). Fuse into one Triton kernel: one CTA per row over the concatenated `[q_rows; k_rows]` space, selecting the q or k gamma per row; decode-gated to ≤256 tokens since the 1-row-per-CTA design loses to TE above that | `megatron/core/inference/attention/fused_qk_norm.py` (new: `_fused_qk_rmsnorm_kernel`, `fused_qk_rmsnorm`, `can_use_fused_qk_norm`), `megatron/core/inference/attention/__init__.py` (new), `megatron/core/transformer/attention.py` (guarded fused path in `get_query_key_value_tensors`), env `MCORE_FUSED_QK_NORM`, `MCORE_FUSED_QK_NORM_MAX_TOKENS` (default 256); `dev/moe_fused/harness_qknorm.py` (new) | **27,143.5 / 27,160.3 tok/s** (two runs, mean 27,151.9) | **+2.9%** vs back-to-back gate-OFF 26,363.0 (+2.83% vs CLEANBASE-S9 26,404.4); arms **fully separated** (min ON 27,102.9 > max OFF 26,425.7) | **bf16-ulp, not bit-exact** (max_rel ≤7.6e-3 ≈ one bf16 ulp; TE's internal rsqrt/reduction differs). Microbench MATCH at 128/256/384/512 × 2 seeds. Coherence: **2 of 3** temperature-0 prompts byte-identical to gate-OFF; the France prompt diverges at a low-confidence branch (" Paris, and…" vs " Paris. The…"), both fluent and factually correct | session `qwen-nodecut`, runs `e2e/qknorm-on-1785180416`, `e2e/qknorm-on-2-1785181060`, OFF `e2e/qknorm-off-bb-1785181867` | **Accepted, largest win since QWEN-021.** Microbench CUDA-graph replay at 256 tokens (BS256 decode): 10.25 → 8.20 µs (1.25×), but 0.83×/0.86× at 384/512 → gated decode-only. e2e delivered **+2.9%**, ~3× the ~1% microbench ceiling, because the two norms sit in the serial attention dependency chain and removing a launch + graph node per layer (×48) also cuts host dispatch gaps. TPOT 9.711 → 9.428 ms/tok. Default OFF; `MCORE_FUSED_QK_NORM=1`. Not bit-exact, so acceptance rests on coherence staying fluent+correct, not byte-identity. New best = **79.87%** of vLLM 33,994.5 | +| CLEANBASE-S9 | 2026-07-27 | Fresh same-session OSL1024 reference at the ten-gate best config, in session `qwen-nodecut`, before the fused QK-RMSNorm lever; re-measured back-to-back with the QWEN-026 ON arm to rule out session drift | none (all ten gates on) | **26,404.4 tok/s** (initial); **26,363.0 tok/s** back-to-back (`qknorm-off-bb`) | baseline (session 9) | Coherent, 5/5 iters | session `qwen-nodecut`, run `e2e/qknorm-off-bb-1785181867` | Clean reference; **77.67%** of vLLM 33,994.5. TPOT 9.695 / 9.711 ms/tok. The two OFF measurements agree to 0.16%, so the +2.9% QWEN-026 delta is not drift | +| QWEN-024 | 2026-07-26 | HOSTGAP-S6 lever 2. Measure the three parts of the post-sampling chain separately, name the unlabelled window, and attack only the part that is not per-request Python object churn: `update_requests` and `active_request_mask` are already vectorized, and roughly half their whole-tensor op cost is provably dead work at one generated token per request | `dynamic_context.py` (`_write_decode_token_bookkeeping_fast`, `_write_token_bookkeeping_reference`, `_verify_decode_token_bookkeeping`, env `MCORE_INFER_VEC_UPDATE_REQS`, `..._VERIFY`), `text_generation_controller.py` (no-finisher short-circuit), `dev/moe_fused/harness_updreq.py` (new) | **26,128.2 / 26,270.3 tok/s** (two pairs) | **+0.90% vs CLEANBASE-S8** (pairwise +0.71% and +1.10%) | 400-step two-context equivalence with 30 mid-batch terminations + block-boundary crossing; 300 steps under VERIFY; coherence byte-identical | session `qwen-updreq` job 5616264, runs `e2e/vecupd-on-1785100576`, `e2e/vecupd-on-rep2-1785101960` | **Accepted.** Named the unlabelled engine window: it is `post_process_requests` via `async_bookkeep`, invisible to HOSTGAP-S6 only because the engine's NVTX helper is inert unless `_nvtx_enabled` is set. It is **187.9 µs/step of pure per-request Python object churn** (1,719 `len()`, 256 dict lookups, 256 appends, 3 tensor ops per step) and was **deliberately left alone** per the falsification rule. The other half was op-eliminated: host chain 372.4 → 293.7 µs/step, `update_requests` 137.6 → 74.0 µs (1.86×). TPOT −88 µs/step. Default OFF; `MCORE_INFER_VEC_UPDATE_REQS=1`. New best = **77.28%** of vLLM 33,994.5 | +| CLEANBASE-S8 | 2026-07-26 | Fresh same-session OSL1024 reference at the eight-gate best config, in session `qwen-updreq`, before the `update_requests` host lever | none (all eight gates on) | **25,944.9 tok/s** | baseline (session 8); −0.57% vs QWEN-023's 26,092.5 (session drift) | Coherent, 5/5 iters | session `qwen-updreq` job 5616264, run `e2e/ref-s8-1785099726` | Clean reference; **76.32%** of vLLM 33,994.5. Per-iter 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7 — iteration 1 is a 1.3% cold outlier, iterations 2–5 span 0.23%. avg_latency 9,820.6 ms, TPOT 9.867 ms/tok | +| PROFILE-HOST-S6 | 2026-07-25 | Capture a host-visible trace (`osrt` + CPU sampling + Python sampling) so the ~1237 µs/step of host gaps PROFILE-S6 found between CUDA-graph replays can be attributed to concrete host call sites | none (`dev/moe_fused/profile_host_insession.sh`, new; `dev/moe_fused/dispatch_host_profile.sh`, new) | n/a (BS256/OSL128 profile run) | n/a | n/a | session `qwen-comm` job 5601961 **preempted mid-task**; replacement session `qwen-host` job 5607600 on `nvl72166-T15`, exec `d920241ea00a41618c63a75fefcb9ea2` | **Capture dispatched and running, not confirmed complete within the time box.** Job 5601961 was PREEMPTED (not expired) before any capture could run, so the whole capture had to be re-queued on a fresh allocation that only scheduled at the deadline. Artifact target `sessions/qwen-host/prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}`. **No gap attribution was performed.** Next session must first check whether the detached exec produced the artifact | +| REBASE-S11 | 2026-07-29 | Rebase the twelve-gate branch onto current `main` (141 commits ahead of the old base) and re-measure, since the PR was unreviewable on a stale base | rebase of 14 commits onto `upstream/main` `3ff70c006`; conflict resolutions in `dynamic_context.py`, `text_generation_controller.py`, `router.py`; `dev/moe_fused/run_rebased.sh` (new) | **27,277.4 tok/s** (tpot 9.385 ms, avg_latency 9,340.6 ms, p99 9,589.9 ms) | **−0.91% vs pre-rebase 27,528.6**; 80.24% of vLLM 33,994.5 (was 80.98%) | coherence passed (2+2→4, capital of France→Paris, 3+2 cows→5); all 12 gates present, 3 kernel files present, py_compile clean | job 5694153 on `nvl72159-T08` (5 timed iters, spread 27,214–27,314 tok/s = 0.36%) | **Rebase completed and statically verified; measurement blocked by a node-level Lustre fault, not by code.** Only 1 of 14 commits conflicted (`b9165d87b`), in 3 files, 1 hunk each. (1) `router.py`: upstream added a `qb_beta` selection path (0→26 occurrences); resolved by keeping it and gating the fused topk on `self.qb_beta is None`, since the fused kernel selects on raw logits and cannot honor a qb_beta-shifted selection. (2) `text_generation_controller.py`: purely additive (upstream dummy-forward helpers vs our cached-empty-tensor helpers); kept both. (3) `dynamic_context.py`: upstream made the bookkeeping H2D **conditional on a new `transfer_bookkeeping_to_gpu` param** and returns an `Optional[Event]`; resolved by adopting that contract while keeping the `_INCR_ATTN_STATE` store and profiling epilogue reachable before the return, making the QWEN-023 fast path **decline** whenever `transfer_bookkeeping_to_gpu=False` or `record_bookkeeping_done_event=True` (it publishes itself and records no event), and clearing upstream's new `_bookkeeping_no_real_work` inside the fast path so a stale `True` from an earlier capture step cannot publish `real_token_count=0`. **Note the PR's claimed blocker was wrong**: `vllm_fused_moe.py` is byte-identical between the old merge-base and current `main`, so the SwiGLU region needed no port and applied clean. **Infra blocker:** every run attempt died in `import megatron.core` on a different "missing" module (`sympy._trigonometric_special`, `networkx.preflowpush`, `jaraco.functools`, `nvidia_cutlass_dsl.static_persistent_tile_scheduler`, our own `transformer_layer`). All are the same fault: on `nvl72169-T17` those files **appear in directory listings but `os.path.isfile()` returns False**, while the login node reads all of them fine — a broken Lustre client that serves MDS metadata but cannot stat/open inodes. Do not chase these as package corruption or venv rot — job 5694153 on a healthy node imported everything cleanly with the **unmodified** venv and no overlay, which proves it. **Re-measured via `dev/moe_fused/rebased12.sbatch`** (submitted with plain `sbatch` because cog's repo sync hung three times from this client; runs the unmodified `run_e2e_cfg.sh` for comparability, excludes the bad node, and gates on an `os.path.isfile` health check that aborts in seconds rather than after the model load — that gate is the reusable lesson here). **Verdict: the rebase preserves the gains.** The −0.91% is small enough to be node-to-node variation (the pre-rebase number came from a different node), so 141 upstream commits cost at most ~1%. Log: `agents-space/auto_rebased/rebased12-5694153.out` | +| CLEANBASE-S12 | 2026-07-30 | Fresh same-node twelve-gate reference before testing upstream's new async scheduling, with no instrumentation loaded | none (all twelve gates on, `ASYNC_SCHED=legacy`) | **26,538.2 tok/s** (tpot 9.646 ms, avg_latency 9,597.3 ms, p99 9,892.8 ms) | baseline (session 12); **−2.71% vs REBASE-S11's 27,277.4 on a different node** | Coherent, 3/3 prompts | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/legacy-clean-1785379...` | Clean reference; 78.06% of vLLM 33,994.5. Per-iter 26,373.8 / 26,533.5 / 26,536.8 / 26,535.3 / 26,713.9 — spread 1.29%. **The −2.71% against REBASE-S11 is node-to-node, not a regression**, and is the reason every arm in this session is compared only against same-node runs: node variance here is larger than several accepted gates' individual effects | +| QWEN-029 | 2026-07-30 | HOSTGAP-S6 attributed the largest remaining lever to G1+G2B — 839 µs/step of resolve-before-prepare bookkeeping that is CPU-bound Python on the serial critical path — and concluded it had to be made cheaper rather than hidden, because `async-sched-mode=serial` (QWEN-005a/b, QWEN-007) refused EP and MoE outright, hung on a single request, and then delivered only +0.85%. **The 141 commits REBASE-S11 landed on replaced that implementation.** `serial` is gone; the mode is now `async` ("prepare-before-resolve"), the EP and MoE guards are gone (only MTP-depth and routing-replay remain), and it has a dedicated copy stream, CPU-ready events, and per-step overlap eligibility. Re-test it as a new feature, not a re-litigation | none — one flag: `--inference-dynamic-batching-async-sched-mode async` (`ASYNC_SCHED=async` in `run_e2e_cfg.sh`) | **27,515.0 tok/s** (tpot 9.304 ms, avg_latency 9,243.0 ms, p99 9,553.6 ms) | **+3.68% vs same-node CLEANBASE-S12 26,538.2**; arms **fully separated** (min ON 27,346.2 > max OFF 26,713.9) | Coherence **byte-identical to the legacy arm on all 3** temperature-0 prompts; no warnings, fallbacks or guard messages in the server log | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/async-on-1785379854` | **Accepted — largest single win of the campaign, and it is a flag, not a kernel.** Per-iter 27,346.2 / 27,369.3 / 27,451.1 / 27,569.2 / 27,845.4. **The ledger's "do not re-litigate async scheduling" guidance was correct about the old code and wrong about the new** — worth remembering as a general rule: a rebase that lands 141 upstream commits can invalidate a rejection, so re-check rejected levers whose blocking reason was an explicit upstream guard. **Caveat that makes this an underestimate:** REBASE-S11's conflict resolution makes the QWEN-023 incremental-attention fast path decline whenever the caller defers the bookkeeping publish, which is exactly what async does — so `MCORE_INFER_INCR_ATTN_STATE` is inert in this arm and the +3.68% is *net of losing that gate*. Making the fast path publish-deferral-aware is the obvious follow-up. New best = **80.94%** of vLLM 33,994.5 measured on this node; applying the same-node +3.68% to REBASE-S11 implies ~28,281 tok/s = **~83.2%** | +| QWEN-028 | 2026-07-30 | The other 48 block boundaries QWEN-027 could not reach. QWEN-027 fused `self_attn_bda` + the standalone `pre_mlp_layernorm`; the matching `mlp_bda` + next-layer *input* norm was ruled unreachable because that norm lives inside the TE `LayerNormLinear` QKV GEMM. That ruling was wrong for this config: at **TP=1** `InferenceLayerNormColumnParallelLinear.forward` runs the norm as its own `_te_rms_norm_kernel` launch immediately before the GEMM, so the preceding layer can compute it. The fused kernel already returns `(normed, new_residual)` — hand the normed half to the next layer's `linear_qkv` and skip its norm launch | `megatron/core/inference/fused_add_rmsnorm.py` (`can_use_fused_add_rmsnorm_qkv`, shared `_tensors_compatible`, env `MCORE_FUSED_ADD_NORM_QKV`), `megatron/core/tensor_parallel/inference_layers.py` (`prenormed_input` hand-off consumed in the `tp_size == 1` path), `megatron/core/transformer/transformer_layer.py` (guarded fused path at the `mlp_bda` site; `_next_layer_qkv` holder), `megatron/core/transformer/transformer_block.py` (`_wire_fused_add_norm_qkv_chain`) | **27,179.2 tok/s** (tpot 9.419 ms, avg_latency 9,361.3 ms, p99 9,661.1 ms) | **+2.42% vs same-node CLEANBASE-S12 26,538.2**; arms **fully separated** (min ON 27,027.8 > max OFF 26,713.9) | **residual bit-exact, norm bf16-ulp** (same kernel as QWEN-027). Coherence: **2 of 3** prompts byte-identical to the legacy arm; the France prompt diverges at a low-confidence branch (" So, what is the capital of Italy?" vs " So, the capital of Italy is Rome."), both fluent and factually correct — the same signature QWEN-026 showed | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/qwen028-on-1785380...` | **Accepted, and it over-delivered: +2.42% against a ~1.3% expectation** (47 boundaries × (add 1.5 + norm 2.9 − fused 4.1 µs) ≈ 14 µs/step of device time plus 47 × ~0.78 µs of node cost ≈ 50 µs of a 9,646 µs step ≈ 0.5%). Same pattern as QWEN-026/027: kernels on the serial dependency chain convert at ~3× their device-time arithmetic, because removing a launch also removes a graph node and a host dispatch gap. **The reusable lesson is about the earlier ruling, not the kernel:** "the norm is fused inside `LayerNormLinear`" was true of the *class name* and false of the *TP=1 code path*, and one look at the actual forward would have caught it two sessions earlier. 47 of 48 boundaries fuse — the last layer feeds `final_layernorm` and is left alone. Wiring stores the next layer's `linear_qkv` in a **plain list**, since assigning a Module to an attribute would register it as a second child and duplicate those parameters in `state_dict`. Heavily gated: TP=1, decode token count, RMSNorm, no bias, no offload, and `not layernorm_zero_centered_gamma` (the TE call it replaces hardcodes `zero_centered_gamma=False`). Default OFF; `MCORE_FUSED_ADD_NORM_QKV=1` | +| QWEN-030 | 2026-07-30 | Do QWEN-029 (async) and QWEN-028 (fused add+QKV-norm) compose, or do they cannibalize each other? They could overlap: async hides host bookkeeping behind GPU work, and part of QWEN-028's win is fewer host dispatches, which async would already be hiding | Both gates on together: `ASYNC_SCHED=async` + `MCORE_FUSED_ADD_NORM_QKV=1` + the 12 prior gates. No new code | **27,889.0 tok/s** (tpot 9.179 ms, iters 27,723 / 27,467 / 27,668 / 28,227 / 28,382) | **+5.09% vs same-node CLEANBASE-S12 26,538.2**; beats async alone (27,459 mean, +3.47%) by +1.56% and QWEN-028 alone (27,179, +2.42%) by +2.61% | same 2-of-3 byte-identical / France-branch divergence as QWEN-028 — async contributes no additional divergence (QWEN-029 alone was byte-identical) | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/async-qwen028-1785381339` | **Accepted — new best config.** They compose but **sub-additively**: +5.09% against +5.89% if the two were independent, so ~14% of the combined benefit is double-counted. That is the expected signature of two levers that partly attack the same resource (host dispatch), and it is the reason to always measure the *combination* rather than sum accepted deltas — the ledger's individual percentages cannot be added. Note the run shows a clear warm-up ramp (27,723 → 28,382 across 5 iters, +2.4%), steeper than the legacy arms, so the 5-iter mean **understates** the steady state under async | + +| QWEN-031 | 2026-07-30 | Async scheduling drives a dedicated copy stream for the bookkeeping H2D transfer, but the harness pins `CUDA_DEVICE_MAX_CONNECTIONS=1`, which collapses all streams onto one hardware work queue and can serialize that copy behind compute — defeating the overlap async exists to create | `CUDA_DEVICE_MAX_CONNECTIONS=8` with `ASYNC_SCHED=async` + the 12 gates. Env only, no code | 27,522.5 tok/s (tpot 9.301 ms) | **+0.23% vs the async mean 27,459** (two samples 27,515.0 / 27,403.5, which themselves differ by 0.41%) — the delta is **smaller than the noise between two runs of the identical config**, so this is a null result, not a small win | n/a (no numerical change) | session `qwen-gap1` job 5699934 on `nvl72067-T13` | **Rejected — no effect.** The serialization hypothesis was reasonable but wrong here: the bookkeeping H2D is small and evidently already fits alongside compute on one queue, so more queues buy nothing. Worth recording because `CUDA_DEVICE_MAX_CONNECTIONS=1` is a load-bearing setting in *training* configs and the instinct to raise it for overlap is strong — for this decode workload it is inert. **Methodological note: the only reason this is callable as a null rather than a +0.23% win is that two same-config samples existed to size the noise floor at ~0.4%.** Single-sample deltas below ~0.5% on this workload are unresolvable | + +| QWEN-032 | 2026-07-30 | `initialize_attention_state` is the largest actionable host phase (204 µs/step) and the code says why: the incremental fast path declines any caller that defers the bookkeeping publish, which is exactly what async scheduling does. Teaching the fast path to honor the deferred publish should recover most of that phase | `dynamic_context.py`: `_incremental_attention_state_update` takes the publish flags and returns `Tuple[bool, Optional[Event]]`, mirroring the full path's conditional publish at its tail; `initialize_attention_state` drops the two gate conditions and forwards the event | 28,178.8 tok/s (tpot 9.085 ms) | +1.04% vs QWEN-030 27,889.0 | coherence identical to QWEN-030 on all 3 prompts | session `qwen-gap1` job 5699934 | **The +1.04% is not real, and the diagnostic is why this entry matters.** A decline-reason tally (`MCORE_INFER_INCR_DIAG`) showed the fast path advancing on **0.1% of calls** (3 in 3,000), with **97.5% declining because `_request_layout_version` differed**. So the mechanism this change unlocks essentially never fires, and the throughput delta has no cause behind it — comparing iteration-by-iteration, the two runs' final iterations are 28,382 vs 28,405, i.e. the means differ only through the warm-up ramp. **Kept because it is a genuine prerequisite** (the gate really did exclude async) and it is what made the diagnostic possible, but it must not be credited with a gain. Lesson: a plausible mechanism plus a delta above the noise floor is still not evidence — instrument the mechanism itself, because here the two agreed and were both wrong | +| QWEN-033 | 2026-07-30 | The 97.5% decline points at one line: `resolve_requests()` bumps `_request_layout_version` **unconditionally**, invalidating the incremental cache every step. In steady-state decode that bump is provably spurious — with an all-ones active mask no rows move (`survivor_idxs == dst_idxs`), no KV blocks are released, the stale slice is empty and `total_request_count` is unchanged. Every real layout change in that method requires at least one finished request | `dynamic_context.py`: the bump moves from the top of `resolve_requests` into the `finished_idxs.numel() > 0` branch | 27,938.0 tok/s (tpot 9.163 ms) | **+0.18% vs QWEN-030 27,889.0 — no gain.** Best-config samples now read 27,889 / 28,179 / 27,769 / 27,938 (mean 27,944, spread 1.5%), so this sits inside the noise band | **bit-exact, verified under async**: 132 confirmation lines across ranks, up to 90 verified incremental steps each, **zero mismatches** against full recomputation. The specific risk — async's prepare-before-resolve reordering letting the cache key match while the true state differs — did not materialize | session `qwen-gap1` job 5699934 | **Mechanism fixed, no speedup — and that is the finding.** Fast-path engagement went **0.1% → 81%** of calls, so ~204 µs/step of host work really was removed, and throughput did not move. **Therefore that host work was already fully hidden behind GPU execution by async scheduling.** This independently confirms the host-side breakdown's implication and **retires host idle as a lever** — the line of attack that drove several sessions and three deadlocked nsys attempts. Corollary: the 1,598 µs/step of in-step GPU idle measured in the fresh trace should be treated as substantially a profiling artifact (`--cuda-graph-trace=node` plus nsys's own host overhead inflate exactly these host-bound gaps), not as 1.6 ms of recoverable time. **Kept anyway**: it makes a gate that ships in the tuned config actually functional instead of silently inert, at zero measured cost and with bit-exactness proven | + +### VLLM-SAMENODE: the first apples-to-apples reference (2026-07-30) + +Every "% of vLLM" figure in this ledger before now compared an mcore run on one node against a vLLM run on another, and this session measured node-to-node variance at ~2.7% — large enough to swamp most individual results. So the mcore session was ended and vLLM was run on **the same node** (`nvl72067-T13`, job 5701338), through the identical harness: same client (`static_benchmark.py`), same gsm8k dataset, BS256, OSL1024, 2 warmup + 5 timed iterations. + +**vLLM DP4+EP4: 34,415.7 tok/s, tpot 7.438 ms** (iters 34,095 / 35,141 / 34,495 / 33,765 / 34,615). + +| | tok/s | ms/step | % of vLLM | +|---|---|---|---| +| vLLM DP4+EP4 (same node) | 34,415.7 | 7.438 | 100% | +| mcore best config (mean of 4 samples) | 27,944 | 9.16 | **81.2%** | +| mcore best single run | 28,178.8 | 9.085 | 81.9% | +| mcore at session start (CLEANBASE-S12) | 26,538.2 | 9.646 | 77.1% | + +The session moved mcore from **77.1% → 81.2%** of same-node vLLM. Remaining gap: **1.72 ms/step (18.8%)**. Note this is *worse* than the 83.2% previously inferred from cross-node arithmetic — the honest number was slightly flattering before, which is exactly why the same-node reference was worth a run. + +### Host-side breakdown under async (the phase timing finally ran) + +The `perf_counter` shim (zero-nsys, built after three nsys host-capture deadlocks) produced its first usable report — 4,096 steps under the best config. The two largest entries are **not** actionable: `_forward` (13,996 µs/step) and `sampling` (7,233 µs/step) each exceed the 9,300 µs step, which is only possible because they are blocking waits on GPU work summed across two threads. Reading them as host cost would be the classic profiler mistake. + +Actionable pure-host cost, per step: `initialize_attention_state` **204 µs** (largest, and this is *with* `MCORE_INFER_INCR_ATTN_STATE` on), `prepare_requests` 80, `resolve_requests` 62, `_ep_establish_consensus` 58 (538 µs/call but only on ~10.6% of steps), `active_request_mask` 56, `async_sched_transfer_bookkeeping_to_gpu` 50, `drain_zmq_socket` 37. Summed with the smaller ranges, host work totals **~900 µs/step against a 9,300 µs step (~10%)**. + +**That total is the useful conclusion, and it redirects the campaign:** host work is far smaller than GPU-busy time, so a scheduler with good overlap should hide essentially all of it. Async captured only ~3.5%, so some overlap is imperfect — but the ceiling on *all* remaining host-side work is under 10%, and most of it is already hidden. **The remaining gap to vLLM is therefore predominantly GPU-side**, which retires "host idle is the top lever" — the ranking that had driven the last several sessions and that three failed nsys attempts were meant to refine. + +### Trace staleness: why the ranking had to be re-derived + +The gap-ranking used to pick targets came from `mcore-final-jul27` (QWEN-025, 10 gates). QWEN-026/027/028 then removed **precisely the kernels that ranking pointed at**, so it can no longer be used to choose the next target. Re-measuring the old trace made this concrete: + +| kernel (in the QWEN-025 trace) | launches | total device time | per call | status now | +|---|---|---|---|---| +| `direct_copy_kernel_cuda` (bf16) | 151,680 | 416.6 ms | 2.75 µs | **gone** (QWEN-026) | +| `triton_poi_fused_add_copy__0` | 151,680 | 200.1 ms | 1.32 µs | **~halved** (QWEN-027/028) | +| `rmsnorm_fwd_general` (q/k norm) | 2 per layer | — | ~2.3 µs | **gone** (QWEN-026, fused to 1) | + +The `direct_copy` finding is the useful one and explains QWEN-026's outsized +2.9%: timeline-neighbour analysis places two of these copies per layer between the QKV GEMM and the two QK norms, i.e. **the strided Q/K split views being materialized contiguous because TE's norm demands contiguous input**. At 2 × ~3 µs × 48 layers that is ~288 µs/step (~3% of an 8.85 ms step) of pure data movement. The fused QK-norm kernel reshapes the split views with no copy (`[-1, head_dim]` is a view when the last dim is contiguous), so it deleted the two copies *and* one norm launch — the copies, not the norm launches, were most of the win. **Generalizable: when a fusion replaces a vendor kernel that has a layout precondition, count the layout-conversion copies as part of the prize; they are often larger than the kernels being fused.** Next target selection requires a fresh trace of the current best config. + +## Session 2 (2026-07-23) — conclusion & recommendation + +**Best shippable config stays: legacy async / nvls dispatcher / vLLM grouped-GEMM +backend / FC1+SwiGLU fusion (QWEN-002) = 22,398.9 tok/s (65.9% of vLLM 33,994.5).** + +Every accessible knob was swept and rejected (QWEN-003…009). Key learnings: +- Decode at OSL1024 is **compute/comm-bound, not idle-bound** (async serial +0.85%). +- The gap is **structural**: mcore runs the MoE decode as a chain of discrete + kernels — router gemm → softmax → topk → count/scatter/metadata → **exposed + NVLS AllGather-V dispatch** → grouped GEMM (FC1+SwiGLU, FC2) → moe_sum → + **exposed NVLS ReduceScatter-V combine** — all on the per-layer critical path. + vLLM/TRT-LLM does the equivalent as one fused MoE (fused routing, cutlass + grouped bmm, fused finalize) with ~~0 exposed comm~~ — **the comm half of this claim + is false, see VLLM-COMM-S16**: vLLM runs the same AllGather+ReduceScatter pair per + layer per step. The fused-routing/finalize half stands. +- Built-in fusions (`--moe-router-fusion`, `--moe-permute-fusion`) and + `sampling-backend=flashinfer` are wired to training / non-full-graph paths and + are **incompatible** with the `inference_optimized` top-8 + full-iteration-graph + contract. + +**To actually close the 1.5× gap (large, multi-session work), in priority order:** +1. **Eliminate exposed NVLS comm** (~16% + it gates the critical path): overlap + dispatch/combine with expert GEMM via chunked/pipelined experts, or fuse the + combine ReduceScatter-V with the FC2 epilogue + moe_sum (finalize fusion). +2. **Fuse the routing chain** honoring the inference top-8 contract: one kernel + for softmax+top8+count+scatter-metadata (cuts ~18% of small serial kernels). +3. **A TRT-LLM-style single fused MoE decode kernel** (dispatch+groupedGEMM+ + SwiGLU+finalize) — the real vLLM-parity path; QWEN-001's full fusion was + slower, so this needs a cutlass/CUTE grouped-bmm core, not Triton. + +## Detailed records + +### VLLM-BASELINE — DP4 with expert parallelism + +| Field | Value | +|---|---| +| Date | 2026-07-22 | +| Hypothesis | Establish the fresh vLLM BS256 Nsight target | +| Code revision | `808c475352de6c3693b182048f174736af82356e`; skill files untracked, Megatron source clean | +| Changed files | none | +| Runtime flags | TP1, DP4, `--enable-expert-parallel`, max model length 4096, max sequences 512 | +| Image | `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/images/87e4947c6ce36433.sqsh` | +| Checkpoint / tokenizer | `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/checkpoints/qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/DP4/EP enabled | +| Workload | gsm8k, BS256, OSL128, one BS8/OSL32 warmup request, one timed request | +| Job / run | `5547673`; `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/runs/vllm-qwen30b-nsys-20260722-093437` | +| Throughput | 23,606.708 tok/s | +| Latency / TPOT | 1,368.153 ms / 10.844 ms-token | +| Correctness | Benchmark completed 256/256 requests | +| Nsight artifacts | `vllm_profile.nsys-rep`; `vllm_profile.sqlite` in the run directory | +| Result | vLLM target established | +| Next action | Record the matching mcore EP4 profile | + +### MCORE-BASELINE — EP4/TP1 + +| Field | Value | +|---|---| +| Date | 2026-07-22 | +| Hypothesis | Establish the fresh mcore EP4 BS256 Nsight starting point | +| Code revision | `808c475352de6c3693b182048f174736af82356e`; skill files untracked, Megatron source clean | +| Changed files | none in Megatron source; profile harness sets `max_requests=256` | +| Runtime flags | TP1, PP1, EP4, ETP1, NVLS dispatcher, vLLM grouped GEMM, inference-optimized transformer, full-iteration inference CUDA graphs | +| Image | Cog dev image `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/images/ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/EP4 | +| Workload | gsm8k, BS256, OSL128, one BS8/OSL32 warmup request, one timed request | +| Job / run | `5553135`; `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/runs/qwen-30b-nsys-20260722-161020` | +| Throughput | 12,346.092 tok/s | +| Latency / TPOT | 2,590.135 ms / 20.735 ms-token | +| Correctness | Checkpoint loaded and benchmark completed 256/256 requests | +| Nsight artifacts | `mcore_profile.nsys-rep`; `mcore_profile.sqlite` in the run directory | +| Result | mcore reaches 52.30% of vLLM; 47.70% below target | +| Next action | Analyze the timed-request A/B windows before changing code | + +### QWEN-001 — MoE decode mega-fusion (FC1+SwiGLU+FC2+topk-reduce) + +| Field | Value | +|---|---| +| Date | 2026-07-22 | +| Hypothesis | Fusing the whole MoE expert path into one Triton kernel (removing the `bounded_silu_mul` + `_moe_sum` kernels and two intermediate HBM round-trips) beats the 4-kernel `vllm_fused_moe` path, which is 40.5% of decode GPU-busy time | +| Code revision | branch `perf/moe-fused-decode-gemm` off `808c475352de6c3693b182048f174736af82356e`, dirty | +| Changed files | `megatron/core/inference/moe/fused_moe_decode.py` (new kernel), `dev/moe_fused/harness.py` (new standalone correctness+timing harness) | +| Runtime flags | standalone microbench; Qwen3-30B decode shapes H=2048, moe_ffn=768, 32 local experts, top-8, 256 valid tokens | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | n/a (synthetic weights, reference = production `vllm_fused_moe`) | +| Hardware / layout | OCI `oci-hsg`, 1×GB200 (session `qwen-moe-kernel`) | +| Workload | microbench, 10 warmup + 100 timed CUDA-event iters | +| Job / run | session `qwen-moe-kernel`, exec runs `manual1784770098/200/321` | +| Throughput | not measured end-to-end (rejected at microbench) | +| Latency / TPOT | kernel: reference 175 µs vs fused 258 µs (best sweep 0.80×; larger tiles OOM shared memory) | +| Correctness | Pass — max_abs_diff 2.6e-5, max_rel 0.15 on a 1e-4 floor, `allclose(rtol=2e-2,atol=2e-2)` True | +| Nsight artifacts | none (microbench) | +| Result | Rejected — fused kernel is consistently 20–50% slower than the reference in every same-run comparison | +| Next action | Root cause: one CTA per token-block serializes the H=2048 FC2 output loop (vs the reference's N-parallel multi-CTA GEMMs), 3× atomic traffic to `out`, and shared-memory pressure caps tile sizes; the HBM/launch savings are negligible under CUDA graphs. Pivot to either (a) partial FC1+SwiGLU epilogue fusion only, or (b) the 11.5% exposed NVLS all-gatherv/reduce-scatter-v communication | + +### QWEN-002 — FC1+SwiGLU epilogue fusion + +| Field | Value | +|---|---| +| Date | 2026-07-22 | +| Hypothesis | The 4-kernel MoE path (FC1→2N intermediate→`bounded_silu_mul`→FC2→reduce) wastes a full HBM round-trip of the `[num_valid, 2N]` intermediate and a whole kernel launch. Computing gate & up in the same FC1 program and applying `SiLU(gate)*up` in the fp32 epilogue — writing the `[num_valid, N]` activated intermediate directly — removes both while keeping FC2's N-parallel tiling | +| Code revision | branch `perf/moe-fused-decode-gemm`, dirty | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (add `FUSE_SWIGLU` constexpr to `_fused_moe_kernel`, `fuse_swiglu` to `_invoke_fused_moe_kernel`, `fuse_fc1_activation` path in `vllm_fused_moe`), `megatron/core/transformer/moe/experts.py` (`_vllm_forward` passes `fuse_fc1_activation=True`), `dev/moe_fused/harness_fc1.py` (new A/B harness) | +| Runtime flags | vLLM grouped-GEMM backend, SwiGLU; microbench at Qwen3-30B decode shapes H=2048, moe_ffn=768, 32 local experts, top-8, 256 valid tokens | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | n/a for microbench (synthetic weights; reference = unfused `vllm_fused_moe`) | +| Hardware / layout | OCI `oci-hsg`, 1×GB200 (session `qwen-moe-kernel`) | +| Workload | microbench, 10 warmup + 200 timed CUDA-event iters, 3 repeats | +| Job / run | session `qwen-moe-kernel`, exec runs `fc1c…`/`fc1t…` | +| Throughput | **22,741.9 tok/s** (fused) vs **22,617.7 tok/s** (same-env, fusion off) → **+0.55%**. gsm8k BS256 OSL1024, 2 warmup + 5 timed iters | +| Latency / TPOT | e2e TPOT 11.257 (fused) vs 11.319 ms/tok (off) → −0.55%. MoE-path kernel microbench: reference 174–182 µs vs fused 141–143 µs → **1.24–1.27×** (3 repeats) | +| Correctness | Pass — fused vs unfused `vllm_fused_moe`: max_abs_diff 3.9e-5, allclose; e2e coherence prompts coherent (2+2=4, capital of France = Paris) | +| Nsight artifacts | none (in-session A/B benchmark, not profiled) | +| Result | Accepted — correct, MoE kernel 1.25×, e2e +0.55% throughput / −0.55% TPOT, no regression. Unfused path byte-identical (all edits guarded by the `FUSE_SWIGLU` compile-time constexpr) | +| Next action | The 1.25× MoE-kernel win only nets +0.55% e2e at OSL1024, so MoE FC1 activation is not the wall-time bottleneck at the throughput regime — re-profile the OSL1024 decode (not the OSL128 profile trace) to find the true dominant cost before the next change. NOTE: at OSL1024 mcore already reaches ~22.6k tok/s; the ledger's 52% "gap" is an artifact of the OSL128 single-request profile workload | + +### QWEN-011 — fused MoE indirection-table build (lever #1) + +| Field | Value | +|---|---| +| Date | 2026-07-24 | +| Hypothesis | The decode "routing/permute kernel storm" is dominated by the 5 tiny serial kernels of `_moe_align_block_size_cuda_graphable` (init + count + prefix + fill + scatter) run per MoE layer × 48. Merging init + prefix + fill into one `_prefix_fill_init_kernel` cuts it to 3 kernels/layer (−96 launches/step) and narrows the vLLM gap without touching numerics. | +| Code revision | branch `perf/moe-fused-align` off `perf/moe-routing-fusion` (which carries QWEN-002), dirty | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_prefix_fill_init_kernel`, `_moe_align_block_size_fused`, env `MCORE_MOE_FUSED_ALIGN`, wired into `vllm_fused_moe`), `megatron/core/transformer/moe/experts.py` (`MCORE_FUSE_FC1_ACT` env gate), `dev/moe_fused/harness_align.py` (new A/B harness) | +| Runtime flags | `MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; nvls dispatcher, vLLM grouped-GEMM backend, full-iteration CUDA graphs | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 10 warmup + 200 timed CUDA-event iters (correctness) | +| Job / run | session `qwen-fuse` runs `fusedalign` (e2e) + `acb5ea1b` (microbench) | +| Throughput | **22,696.4 tok/s** vs 22,269.5 (fusion-on) / 22,241.5 (clean) → **+1.92% / +2.04%** | +| Latency / TPOT | avg_latency 11,257 ms vs 11,494 ms (clean) | +| Correctness | **Bit-exact** vs the 5-kernel path: max_abs_diff 0.0, allclose True; e2e coherence prompts coherent (2+2=4, Paris) | +| Nsight artifacts | in-session profile `sessions/qwen-moe-kernel/prof/1784930600/mcore_profile.{nsys-rep,sqlite}` (BS256/OSL128, fused-align ON). Kernel-name counts **confirm the mechanism**: `_init_sorted_ids_kernel`=0, `_prefix_sum_kernel`=0, `_fill_expert_block_ids_kernel`=0 (all removed), `_prefix_fill_init_kernel`=`_count_local_tokens`=`_scatter_token_indices_kernel`=76032 (align now 3 kernels/layer vs 5). microbench MoE-call 142.65→138.57 µs (1.029×) | +| Result | **Accepted** — correct (bit-exact), e2e +1.9% with tight iter variance (22.65–22.71k), no regression. Profile confirms −2 indirection kernels/MoE-layer/step. Default OFF (env-gated) so the untouched path stays byte-identical; enable with `MCORE_MOE_FUSED_ALIGN=1`. | +| Next action | Reduces align 5→3; the remaining serial dep (count → prefix_fill_init → scatter) can only merge further via a cooperative-grid launch (risk: CUDA-graph capture). Profile to confirm per-step kernel-count drop, then target the FC1/FC2 grouped-GEMM efficiency gap (lever #2, needs a cutlass/CUTE grouped bmm — Triton launch count is already minimal at 2/layer). | + +### QWEN-012 — grouped-GEMM roofline / padding decision gate + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | The ~40% of decode GPU time in `_fused_moe_kernel` is recoverable GEMM inefficiency, so a cutlass/CuTe BF16 grouped GEMM is worth writing. Measure achieved TFLOP/s, achieved weight bandwidth against a *measured* ceiling, and the indirection-table padding waste before writing any kernel. | +| Code revision | branch `perf/moe-fused-align` off `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, dirty (carries QWEN-002 + QWEN-011) | +| Changed files | `dev/moe_fused/harness_roofline.py` (new, analysis only — no Megatron source change) | +| Runtime flags | microbench at one EP4 rank of decode: hidden 2048, moe_ffn 768, 128 global / 32 local experts, top-8, `num_tokens_hint`=256 (= `local_tokens*ep_size` from `InferenceAllGatherDispatcherBase._get_host_valid_tokens_estimate()`), max_tokens 1024 | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`) | +| Checkpoint / tokenizer | n/a (synthetic weights at production shapes) | +| Hardware / layout | OCI `oci-hsg`, 1×GB200 of the 4-GPU node, session `qwen-cutlass` (job 5598012) | +| Workload | microbench, 20 warmup + 300 timed CUDA-event iters, 3 repeats (spread <1%) | +| Job / run | session `qwen-cutlass` exec `c05fd5f771b44bf7a002886a1242c1c4` | +| Throughput | not applicable (analysis) | +| Latency / TPOT | FC1 (fused SwiGLU) 52.25 µs, FC2 19.88 µs, **total 72.13 µs** per MoE layer per rank | +| Correctness | n/a (no functional change) | +| Nsight artifacts | none (CUDA-event microbench + torch profiler) | +| Result | **Memory-bound, not FLOP-bound — Stage 4 rejected before implementation.** (1) Measured bandwidth on this GB200: d2d copy 6.851 TB/s aggregate, Triton vectorized streaming read **6.081 TB/s** (torch `.sum()` gives 4.05 TB/s and is a poor proxy — do not use it as the ceiling). (2) Every decode step must read all 32 local experts' weights: 302.0 MB per layer per rank ⇒ **49.66 µs bandwidth floor**, vs 72.13 µs measured = **1.45× off roofline**. (3) Achieved 63.4 (FC1) / 83.2 (FC2) TFLOP/s on valid FLOPs — ~3% of GB200 BF16 dense peak; 246/323 TFLOP/s even counting padded rows. (4) Padding waste is large but almost free: `num_tokens_post_padded` vs 527 real local token-expert pairs is 784/1024/2048/4096 at BLOCK_SIZE_M 16/32/64/128 (32.8%/48.5%/**74.3%**/87.1% dead rows), yet dropping BLOCK_SIZE_M from the production 64 to 16 buys only ~1.2×, because each expert's weights are read exactly once per M-tile row-block regardless. (5) A 63-config tile sweep found 58.93 µs (1.225×) with one shared config and 57.24 µs (1.26×) with per-GEMM configs — i.e. **most of the 1.45× ceiling is reachable in Triton**, leaving ≤1.15× for any hand-written kernel. | +| Next action | Do not write a CUTLASS/CuTe grouped GEMM. Land the tile retune (QWEN-013), then redirect to the next-ranked levers from PROFILE-DECODE: exposed EP comm (~16%) and the routing chain (~18%). | + +### QWEN-013 — per-GEMM decode tile retune + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | QWEN-012's 1.45× roofline gap is tile-quantized SM occupancy, not the GEMM implementation. `_get_default_config` picks `BLOCK_SIZE_N=128` for both passes, leaving FC1 at ceil(768/128)=6 N-tiles × 32 M-tiles = 192 CTAs on 148 SMs (1.3 waves, ~30% tail idle). FC1 (N=768) needs a *small* `BLOCK_SIZE_N` to manufacture enough CTAs; FC2 (N=2048) has plenty and prefers a *large* one. vLLM shares one config across both passes and cannot express this. | +| Code revision | branch `perf/moe-fused-align` off `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, dirty | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_get_decode_tuned_configs`, `_TUNE_DECODE_GEMM` env gate, `config_fc1`/`config_fc2` threaded into the two `_invoke_fused_moe_kernel` calls and their grid sizing), `dev/moe_fused/harness_gemmtune.py` (new A/B + profiler + CUDA-graph harness), `dev/moe_fused/harness_align.py` (fix a module-vs-function import bug, see below) | +| Runtime flags | `MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; nvls dispatcher, vLLM grouped-GEMM backend, `transformer_impl=inference_optimized`, full-iteration CUDA graphs. Tuned configs: shared `BLOCK_SIZE_M=16`, `BLOCK_SIZE_K=64`, `GROUP_SIZE_M=1`; FC1 `BLOCK_SIZE_N=64`, 4 warps, 3 stages; FC2 `BLOCK_SIZE_N=256`, 8 warps, 4 stages. Applies only for `num_tokens_hint <= 512`; prefill keeps upstream's heuristic. | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 20 warmup + 300 timed iters × 3 repeats and a 50-iter torch-profiler window (attribution) | +| Job / run | session `qwen-cutlass2` (job 5600047) run `gemmtune-1785010791`; reference `ref0-1785010433` (CLEANBASE-S4) | +| Throughput | **23,636.0 tok/s** vs **22,657.7** same-session reference → **+4.32%**. Per-iter 23,553–23,668 tok/s | +| Latency / TPOT | avg_latency 10,802 ms vs 11,282 ms; TPOT 10.831 vs 11.299 ms/tok (−4.14%). Per-kernel GPU time: `_fused_moe_kernel` (both GEMMs) 77.65 → **61.89 µs** (1.255×); whole `vllm_fused_moe` call 100.64 → **84.31 µs** (1.194×) | +| Correctness | **Bit-exact** — max_abs_diff 0.0 and max_rel 0.0 at 128/256/384/512 valid tokens. Expected: `BLOCK_SIZE_K` is unchanged, so the fp32 reduction order along K is identical and only output-tile partitioning changes. e2e coherence prompts coherent (2+2=4, Paris, 3+2 cows = 5 cows) | +| Nsight artifacts | none (torch-profiler per-kernel attribution instead; timing composition change is fully explained by the two GEMM launches) | +| Result | **Accepted** — bit-exact, +4.32% e2e with tight variance, no regression. Default OFF (`MCORE_MOE_GEMM_TUNE`) so the untouched path stays byte-identical. New best mcore = **69.53%** of the vLLM 33,994.5 baseline (was 66.65%). | +| Next action | Two follow-ups, in order: (1) the remaining GEMM headroom is only 1.15× (57–62 µs vs the 49.66 µs floor), so stop here on the GEMM and attack **exposed EP comm (~16%)**; (2) `_get_decode_tuned_configs` is hand-tuned at one shape — if other MoE geometries adopt it, replace the constants with a small autotune keyed on (N, E, M). **Harness caveat worth knowing:** `import megatron.core.inference.moe.vllm_fused_moe as vfm` yields the *function*, not the module, because `moe/__init__.py` does `from .vllm_fused_moe import vllm_fused_moe`. Patching module flags on it silently no-ops, which makes an A/B compare one path against itself. This invalidated my first retune measurement (and affects QWEN-011's *microbench* numbers, though not its e2e result, which used the real env var). Both harnesses now use `importlib.import_module` and assert the object is the module. With the fix, the align A/B re-measures at 148.67 → 124.13 µs eager wall (1.198×). | + +### QWEN-014 — flashinfer `cutlass_fused_moe` backend A/B + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | `--inference-grouped-gemm-backend flashinfer` is already wired (`InferenceGroupedMLP._flashinfer_forward`, `experts.py:1057`) and accepts BF16, so QWEN-010's "blocked for SwiGLU" note is wrong and this backend has never actually been benchmarked. Enabling it should beat the vLLM Triton grouped GEMM, since it fuses permute + both GEMMs + activation into one launch. | +| Code revision | branch `perf/moe-fused-align`, dirty. No Megatron source change — the two integration bugs found were reproduced and fixed *in the harness only*, because the measurement then showed the backend is not worth adopting. | +| Changed files | `dev/moe_fused/inspect_flashinfer.py` (new), `dev/moe_fused/harness_flashinfer.py` (new) | +| Runtime flags | microbench at one EP4 rank: hidden 2048, moe_ffn 768, 128 global / 32 local experts, top-8, valid 256, max_tokens 1024, ep_size 4, ep_rank 0 | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh`; flashinfer **0.6.14**, torch 2.13.0a0, CUDA 13.3, GB200 sm100. `flashinfer_jit_cache` / `flashinfer_cubin` are **not** installed, so the first call downloads cubins from `edge.urm.nvidia.com` (~2 min from this cluster; it does have egress). Budget that for CUDA-graph warmup if this path is ever enabled. | +| Checkpoint / tokenizer | n/a (synthetic weights at production shapes; reference = production `vllm_fused_moe`) | +| Hardware / layout | OCI `oci-hsg`, 1×GB200 of the 4-GPU node, session `qwen-cutlass2` (job 5600047) | +| Workload | microbench: numerics single-shot, then 300 CUDA-graph replays × 3 repeats | +| Job / run | session `qwen-cutlass2` execs `7590adcb6e754ed8a7c32a58decb2f9b`, `659d263321e8df73...` | +| Throughput | not measured e2e — rejected at microbench | +| Latency / TPOT | CUDA-graph device time, valid=256: `vllm_fused_moe` default tiles **97.97 µs**, `vllm_fused_moe` QWEN-013 tiles **83.63 µs**, `cutlass_fused_moe` **90.81 µs** → **1.079× vs default, 0.921× vs QWEN-013**. (Eager wall-clock flatters cutlass to 1.17× because it is 1 launch vs 6 at ~12 µs/launch; that number is meaningless under full-iteration CUDA graphs.) | +| Correctness | Passes only after fixing the weight order. `[gate\|up]` (mcore's buffer as-is): max_abs 7.14e-4, max_rel **4.92**. `[up\|gate]`: max_abs 3.87e-5, max_rel 0.171 — bf16 rounding noise. So the kernel wants **up\|gate** (w3\|w1). Both pass a loose `allclose(2e-2)`, which is exactly why this would have shipped silently wrong. | +| Nsight artifacts | none (CUDA-graph event timing + torch profiler) | +| Result | **Rejected on measurement — but QWEN-010's stated reason was wrong, and two real mcore bugs were found.** (1) `_resolve_flashinfer_activation_type` (`experts.py:937-951`) maps `F.silu` → `ActivationType.Silu` **without consulting `config.gated_linear_unit`**, unlike `_resolve_mcore_activation_type` right below it which does. With `Silu` the kernel hard-fails: `fc1_expert_weights.size(1) == fc2_expert_weights.size(2) * mInnerDimMultiplier (1536 vs. 768)` — a non-gated activation expects fc1 out == ffn, but a SwiGLU fc1 emits 2×ffn. The correct enum is `ActivationType.Swiglu`, which works. **This — not any kernel limitation — is why the flashinfer backend appears "blocked for SwiGLU".** (2) The kernel expects fc1 as `[up\|gate]`; `_build_concatenated_weights` produces TE's `[gate\|up]`, so even with the enum fixed the results are numerically wrong. Fixing both is contained (an enum branch + a one-time reordered fc1 buffer, ~200 MB/layer extra since it cannot share storage with the TE `param.data` views), but at 0.921× it would still lose to QWEN-013. | +| Next action | Do not adopt the backend for this workload. Do fix the two bugs as a correctness matter on their own merit, since `--inference-grouped-gemm-backend flashinfer` is currently either a hard crash or silently wrong for every gated-activation MoE: guard `_resolve_flashinfer_activation_type` on `config.gated_linear_unit` (→ `Swiglu`), and reorder fc1 to `[up\|gate]` for that path. File as a separate bug-fix PR, not a perf change. | + +### QWEN-015 — TRT-LLM-Gen BF16 fused MoE (`trtllm_bf16_routed_moe`) probe + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | vLLM's BF16 SM100 oracle selects `FLASHINFER_TRTLLM` → `trtllm_bf16_moe` / `trtllm_bf16_routed_moe`, so this is the literal kernel vLLM wins with. The `_routed_` variant takes already-computed routing results, which matches mcore's contract, and it fuses SwiGLU + the topk-weighted finalize. Dropping it in behind an env gate should be the highest-value integration. | +| Code revision | branch `perf/moe-fused-align`, dirty. **No Megatron change made** — blocked at the probe stage. | +| Changed files | `dev/moe_fused/harness_flashinfer.py` (probe only) | +| Runtime flags | `num_experts=128, top_k=8, intermediate_size=768, local_expert_offset=0, local_num_experts=32, routing_method_type=RoutingMethodType.TopK, do_finalize=True`, both `use_shuffled_weight`/`weight_layout` combinations | +| Image | Cog dev image `ceecf5c304a5d8bd.sqsh`; flashinfer 0.6.14 (exposes `trtllm_bf16_moe`, `trtllm_bf16_routed_moe`, `convert_to_block_layout`, `reorder_rows_for_gated_act_gemm`) | +| Checkpoint / tokenizer | n/a (synthetic weights at production shapes) | +| Hardware / layout | OCI `oci-hsg`, 1×GB200, session `qwen-cutlass2` | +| Workload | microbench probe (numerics/latency never reached) | +| Job / run | session `qwen-cutlass2` exec `7590adcb6e754ed8a7c32a58decb2f9b` | +| Throughput | not reached | +| Latency / TPOT | not reached | +| Correctness | not reached | +| Nsight artifacts | none | +| Result | **Blocked on weight layout — and it is not the contract mismatch the plan anticipated.** `weight_layout=MajorK, use_shuffled_weight=False` is rejected by the launcher itself (`trtllm_fused_moe_kernel_launcher.cu:770`: `BF16 Moe: weight_layout must be BlockMajorK`). `weight_layout=BlockMajorK, use_shuffled_weight=True` then indexes `weights.size(3)` and throws `IndexError: Index 3 out of bounds for tensor with 3 dimensions`. So the kernel requires **4-D pre-shuffled block-major** weights; mcore's `[E, 2*ffn, H]` 3-D concatenated buffer cannot be passed in any configuration. Notably the dense top-8 `routing_map`, the `ep_size`/`ep_rank` semantics, and the NVLS symmetric-memory output tensor were *not* the blockers — the API accepts all of those. | +| Next action | Deprioritized, with a concrete recipe if revisited: build a one-time load-time weight pass using flashinfer's own helpers — `reorder_rows_for_gated_act_gemm` on fc1 (which also resolves the gate/up interleave that QWEN-014 found) then `convert_to_block_layout` on both — producing separate 4-D buffers. That breaks `_build_concatenated_weights`' storage sharing with TE's `param.data` views, so it roughly doubles expert-weight residency for the reordered copies (~14.5 GB/rank today) and needs its own `--inference-grouped-gemm-backend` value. **Priority is low now**: QWEN-012 measured the grouped GEMM at 1.15× off the weight-bandwidth floor after QWEN-013, and QWEN-014 measured the sibling cutlass kernel at 0.921× vs the retuned Triton path, so the fused-MoE ceiling here is roughly the ~22 µs of routing + finalize this kernel would absorb, not a GEMM win. Attack exposed EP comm (~16%) first. | + +## Session 4 (2026-07-25) — conclusion & recommendation + +**Best shippable config is now: nvls dispatcher / vLLM grouped-GEMM backend / +`MCORE_FUSE_FC1_ACT=1` (QWEN-002) / `MCORE_MOE_FUSED_ALIGN=1` (QWEN-011) / +`MCORE_MOE_GEMM_TUNE=1` (QWEN-013 + QWEN-013b) = 23,646.0 tok/s += 69.56% of vLLM 33,994.5** (was 66.65% at the start of this session). + +The session was chartered to pursue "the grouped-GEMM / cutlass lever". The +measurement says that lever is nearly exhausted, and says so three +independent ways: + +1. **Roofline (QWEN-012).** The decode grouped GEMM must read 302 MB of expert + weights per layer per rank against a *measured* 6.081 TB/s streaming-read + ceiling — a 49.66 µs floor. It ran at 72.13 µs (1.45× off) and now runs at + ~57 µs (1.15× off). Achieved FLOPs are ~3% of BF16 peak, so this is a + bandwidth problem and no GEMM implementation can fix it. The dramatic-looking + 74.3% indirection-table padding is nearly free for the same reason. +2. **A hand-written kernel is not the lever (QWEN-013).** The entire 1.45× gap + was tile-quantized SM occupancy, recovered by giving FC1 and FC2 their own + `BLOCK_SIZE_N` — 12 lines, bit-exact, +4.36% e2e. +3. **The vendor kernels do not beat it (QWEN-014, QWEN-015).** flashinfer's + `cutlass_fused_moe` is 0.921× the retuned Triton path in CUDA-graph device + time, and `trtllm_bf16_routed_moe` cannot accept mcore's weight layout at all. + +**Remaining gap to vLLM is 1.44×, and it is no longer in the GEMM.** +PROFILE-TUNED confirms it from the trace side: the MoE GEMM is now 32.9% of +decode GPU time and sits only 229 µs/step (2.3%) above its bandwidth floor, +while exposed EP comm is 11.7%, routing/permute 15.8%, and GPU idle 20.6%. +Per PROFILE-DECODE the decode critical path is +attn → router → **exposed NVLS AllGather-V dispatch** → grouped GEMM → +**exposed NVLS ReduceScatter-V combine**. In priority order: + +1. **Exposed NVLS comm (~16%, and it gates the critical path).** ~~vLLM's + equivalent has *zero* exposed comm.~~ **False — see VLLM-COMM-S16**: vLLM runs the + same AllGather+ReduceScatter at the same per-layer cadence, so this is not pure + deficit and was over-prioritized here. Overlap dispatch/combine with the expert + GEMM via chunked/pipelined experts, or fuse the combine ReduceScatter-V into + the FC2 epilogue + `_moe_sum`. This is now the single highest-value target. +2. **The routing chain (~18%).** After QWEN-011 the align is 3 kernels; the + residual `_count_local_tokens` (8.2 µs) + `_moe_sum` (8.6 µs) + scatter + (3.1 µs) is ~20 µs of the 84 µs MoE call. Fusing `_moe_sum` into the FC2 + epilogue is the contained piece. +3. **Only then** revisit a fully fused MoE (QWEN-015's recipe), whose remaining + upside is the routing/finalize overhead, not the GEMM. + +## Session 5 (2026-07-25) — exposed NVLS EP comm + +### QWEN-016 — EP comm decision gate: exposed, latency-bound, small prize + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | The 929 µs/step of exposed NVLS EP comm is a worthwhile lever. Gate it first: measure how much is genuinely exposed, whether it is latency- or bandwidth-bound, and what the floor is under perfect overlap. | +| Code revision | branch `perf/moe-fused-align`, dirty. **No Megatron change retained** — analysis only. | +| Changed files | `dev/moe_fused/probe_comm.py`, `dev/moe_fused/analyze_comm.py`, `dev/moe_fused/analyze_comm_skew.py`, `dev/moe_fused/harness_comm.py` (all new, analysis harnesses; nothing under `megatron/`) | +| Runtime flags | trace side: none (re-analysis of PROFILE-TUNED). Microbench: `hidden=2048 topk=8 local_tokens=64 per_rank_max=2048 rsv_dtype=float32`, CTA sweep `4,8,16,32,64,128` | +| Image | cog dev image for `oci-hsg`, venv `envs/megatron_lm/dd356431262b5db4` (sqsh tag not re-verified this session) | +| Checkpoint / tokenizer | n/a (trace re-analysis + synthetic-payload microbench at production shapes) | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), EP4/TP1 | +| Workload | trace: PROFILE-TUNED steady-state decode, stream 257, steps 32–96, 48 layers, 4 ranks. Microbench: NVLS AGV/RSV at the production decode shape under CUDA-graph replay, median of repeats | +| Job / run | session `qwen-comm`, Slurm job 5601961 (`batch_long`, 8 h); exec `08e6a692b5164fa1929086fcc0e315d1` (CTA sweep), `04bdfdd605aa49348ff3d031aba70c89` (decomposition) | +| Throughput | not applicable (no Megatron change measured) | +| Latency / TPOT | not applicable | +| Correctness | microbench gate: AGV max\|diff\| 0.000e+00 and RSV max rel err 0.000e+00 over 64 trials vs `all_gather_into_tensor` / analytic rank-scaled reduction | +| Nsight artifacts | source trace `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite` | +| Result | **Gate answered on all three questions; the prize is smaller than its 11.7% share suggests.** (1) *Exposed*: the per-step union of comm intervals equals their sum — there is no concurrent compute, so it is 100% exposed. Per-step comm totals 928–1219 µs across the four ranks. (2) *Latency-bound, decisively*. Decomposing the production-shape collectives at 128 CTAs: AGV 6.57 µs = launch 0.72 + **barrier 5.08** + transfer 0.77; RSV 7.88 µs = launch 0.72 + **barrier 5.03** + transfer 2.13. Payloads are only 268 KB/rank (AGV) and 512 KB/rank (RSV); against the 900 GB/s/dir NVLink floor (0.894 / 1.748 µs) the RSV transfer is already at 82% of peak and the AGV transfer is *below* the unicast floor because multimem multicast pays one egress for all peers. Byte movement is therefore 127 µs/step out of ~693 µs — **fusing or batching bytes cannot win; only removing or hiding barriers can.** (3) *Floor*. Splitting the trace into intrinsic cost (median release − last arrival) and inter-rank arrival skew: AGV 6.78 µs/kernel = 5.97 intrinsic + 0.76 skew; RSV 15.30 = 7.19 + 8.11. Per step: 1060 µs total = **632 µs intrinsic + 428 µs skew**. The skew is not comm work — it is ranks waiting for the slowest rank's expert GEMM (ranks differ in how many experts receive tokens), so it is a routing-balance problem, not a collective problem. The skew-free microbench independently lands at 693 µs/step, corroborating the 632 µs figure. So the recoverable critical path is **632–693 µs of a 9,933 µs step = 6.4–7.0%**; perfect elimination gives 1.068–1.075× → 25,250–25,420 tok/s = 74.3–74.8% of vLLM 33,994.5. Of that, ~485 µs is 96 symmetric-memory barriers × 5.05 µs. Sub-verdict on the two proposed approaches: **(b) fusing RSV into the FC2 epilogue** removes a launch and the transfer but *keeps the barrier*, ceiling ≈ 2.9 µs/layer = 137 µs/step = **1.4%**. **(a) chunked/pipelined experts** hides a collective behind GEMM but adds one barrier per extra chunk (5.05 µs) against a 14.4 µs/layer collective cost, so a 2-chunk split is worth at best ≈ 9 µs/layer = 430 µs/step = **4.3%**, and only if CUDA-graph capture can express concurrent streams across the MoE dependency chain — which QWEN-008 already identified as the binding constraint. Separately, the CTA count is **already optimal**: the sweep gives AGV 13.54/9.68/8.45/7.89/8.07/6.78 µs and RSV 37.37/21.12/13.43/9.31/8.25/8.05 µs at 4/8/16/32/64/128 CTAs, so the shipped `MAX_NUM_BLOCKS=128` is the best point and reducing CTAs to shrink the per-CTA barrier is strictly worse. A prepared `MCORE_NVLS_{AGV,RSV}_CTAS` override was therefore reverted unused. | +| Next action | **Redirect off this lever.** A 6.4–7.0% ceiling of which 1.4–4.3% is realistically reachable, inside graph-captured dispatcher code, is worse value than routing/permute (1251 µs/step, 15.8%, spread over 242 kernels at 5.2 µs each — a launch-count problem where fusion converts directly into wall time). Take the ledger's item 2: fuse `_moe_sum` into the FC2 epilogue and collapse the residual `_count_local_tokens` / scatter chain. | + +### QWEN-017 — load-poll symmetric-memory barrier (rejected) + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | `symm_mem_sync`'s wait loop spins on a system-scope `atom.cas`, one full uncached read-modify-write per attempt, which sets how quickly a rank notices its peers arrived. Replacing it with an `ld.acquire.sys` poll plus a single clearing store should cut the 5.05 µs barrier and thus ~485 µs/step. | +| Code revision | branch `perf/moe-fused-align`, dirty. Change **reverted** after measurement. | +| Changed files | `megatron/core/inference/communication/torch_symm_triton/barrier.py` (added `_wait_signal_ldpoll` + `MCORE_SYMM_BARRIER_LDPOLL` gate, default off) — reverted; `dev/moe_fused/harness_comm.py` (correctness gate) — retained | +| Runtime flags | `MCORE_SYMM_BARRIER_LDPOLL=0` vs `=1`, production shape, 128 CTAs, 3 alternating reps each | +| Image | cog dev image for `oci-hsg`, venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | n/a (collective microbench at production shapes) | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), EP4/TP1 | +| Workload | NVLS AGV + RSV under CUDA-graph replay, median of repeats, 64-trial correctness gate before each timing block | +| Job / run | session `qwen-comm`, job 5601961; exec `04bdfdd605aa49348ff3d031aba70c89` | +| Throughput | not run e2e — rejected at the microbench stage | +| Latency / TPOT | per-step comm (48 layers): baseline 695.7 / 692.2 / 693.3 µs (mean 693.7); ldpoll 710.0 / 710.2 / 709.1 µs (mean 709.8). **+2.3% regression**, no overlap between the two triplets | +| Correctness | bit-exact in both variants: AGV max\|diff\| 0.000e+00, RSV max rel err 0.000e+00 over 64 trials | +| Nsight artifacts | none (microbench uses CUDA events) | +| Result | **Rejected — measured regression.** The barrier-only kernel is unchanged between variants (5.75–5.83 µs at both 256 and 512 threads/CTA, against a 0.72 µs empty kernel), so the ~5.05 µs is the 4-way system-scope flag round trip itself, not polling granularity. The cheaper poll actually costs slightly more, consistently on RSV (7.88 → 8.19 µs). Combined with the CTA-count result in QWEN-016, the barrier looks like a hardware/driver latency floor that a Triton-level rewrite does not move. | +| Next action | Do not pursue further barrier micro-optimisation. Any future attempt on this lever must remove barriers (fewer collectives) or hide them behind compute, not make each one cheaper. | + +## Session 6 (2026-07-25) — the routing/permute chain + +### QWEN-018 — routing/permute decision gate: per-kernel breakdown and fusion ceilings + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | Routing/permute (1251 µs/step, 15.8%, 242 kernels at 5.2 µs) is a launch-count problem, so removing kernels converts directly into wall time. Gate it first: attribute the 1251 µs to individual kernel names, split each into fixed launch cost vs real work, measure the inter-kernel dispatch gap, and compute a wall-time ceiling per candidate fusion. Reject any candidate below ~1%. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty. **No Megatron change** — analysis only. | +| Changed files | `dev/moe_fused/analyze_routing.py` (new, analysis only; nothing under `megatron/`) | +| Runtime flags | n/a (re-analysis of the PROFILE-TUNED trace, which was captured with `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`) | +| Image | n/a (login-node `python3` 3.12.13 + stdlib `sqlite3`) | +| Checkpoint / tokenizer | n/a | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, EP4/TP1 (trace); analysis on the login node | +| Workload | PROFILE-TUNED decode window 192.7–195.5 s, device 3, stream 257; 40 steady-state steps averaged, one MoE layer dumped in launch order | +| Job / run | session `qwen-comm`, Slurm job 5601961 (allocation held, analysis ran on the login node); source trace `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite` | +| Throughput | not applicable (analysis) | +| Latency / TPOT | step wall 9927.2 µs, GPU-busy 7900.0 µs, idle 2027.2 µs (20.4%), 1362 kernels — reproduces PROFILE-TUNED (9.933 / 7.884 / 2.049 ms, 1362) from an independent script | +| Correctness | n/a (no functional change) | +| Nsight artifacts | `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite`; `dev/moe_fused/analyze_routing.py` output | +| Result | **The routing category is not one problem, it is two pathological kernels plus four cheap ones — and it is not launch-bound.** Per-step, per-kernel (48 launches each unless noted; `WALL` = device time + the dispatch gap that follows): `_moe_sum_kernel` 7.79 µs/k → **400.3 µs**; `_count_local_tokens_kernel_persistent` 7.52 → **387.4**; `gatherTopK` (router top-8) 6.05 → **316.9**; `_scatter_token_indices_kernel` 2.66 → **153.8**; `triton_per_fused__softmax_prep` (router softmax) 1.94 → **119.9**; `_prefix_fill_init_kernel` 1.33 → **90.9**; plus the `torch.zeros` fill of the count buffer (`vectorized_elementwise_kernel`, grid=1) 0.74 → **≈62**. Total **≈1531 µs/step**. Three measurements set the ceilings. (a) *The dispatch gap is 0.55 µs*, uniform across kernels in the graph, so a removed launch is worth `duration + 0.55 µs` — more than the kernel time alone, but far less than the 5.2 µs/kernel average would suggest if you assumed the whole thing were overhead. (b) *The fixed floor per launch is 1.27 µs* (0.72 µs empty kernel from QWEN-016 + 0.55 µs gap), so of the 1531 µs only **366 µs is fixed cost** and **1165 µs is in-kernel work** — pure launch-count reduction can win at most 3.7% and only by removing *every* routing launch. (c) *The 20.4% idle is now explained*: 862 µs of it is two once-per-step host gaps (537 µs after `index_elementwise_kernel`, 325 µs after `CatArrayBatchedCopy_vectorized`, i.e. sampling/detokenize between graph replays) and ~750 µs is 1362 × 0.55 µs of intra-graph node dispatch; there is no large unexplained residue. **Candidate ceilings.** *(1) Fuse `_moe_sum` into the FC2 epilogue* (the ledger's first-ranked): removes 48 × (7.79 + 0.55) = **400 µs = 4.03% gross**. But the topk slots of one token land in different experts' blocks and therefore different M-tiles, so accumulating into `out` needs cross-CTA fp32 atomics (~4 MB/layer of RMW: 512 local pairs × 2048 × 4 B) *and* a zeroing pass over `out[0:valid_tokens]` that today does not exist because `_moe_sum` writes rather than accumulates. That gives back ~70 µs, leaving ~3.3%, it is the exact shape QWEN-001 measured at 0.68–0.80×, and it stops being bit-exact. Held as the fallback, not built first. *(2) Cooperative-grid merge of count → prefix_fill_init → scatter*: the three are serially dependent, so a merge needs a grid-wide sync and keeps all the work; it recovers only the fixed cost, 2 × 48 × 1.27 = **122 µs = 1.23%**, before paying two `grid.sync()`s per layer against a kernel boundary that costs just 0.55 µs — and it needs cooperative launch under CUDA-graph capture. **Gated out on arithmetic.** *(3) New, and the one the breakdown actually points at: fold the token count and its `torch.zeros` fill into `_prefix_fill_init_kernel`.* `_prefix_fill_init_kernel` already has every CTA redundantly recompute the whole 32-wide cumsum in registers, and the count vector is its only input — so recomputing the histogram per CTA needs **no grid sync at all**. Removes 48 × (7.52 + 0.55) + 48 × (0.74 + 0.55) = **449 µs = 4.52% gross**, integer-exact. **Why `_count_local_tokens` costs 7.52 µs to bucket 2048 int32s**: it is neither launch-bound (0.72 µs floor) nor bandwidth-bound (8 KB), it is starved — with `BLOCK_SIZE=1024` and 2048 valid pairs `total_blocks=2`, so exactly **2 of its 152 CTAs receive work**, and each issues 1024 global atomics contending on 32 counters. **This corrects QWEN-003's conclusion.** QWEN-003 replaced those atomics with `tl.histogram` but left `BLOCK_SIZE=1024` untouched, so both variants ran on 2 CTAs; it measured a wash and inferred "the cost is per-launch fixed overhead, not atomic contention, so an in-kernel rewrite cannot help". The device-time measurement says the opposite: 6.8 of the 7.52 µs is in-kernel. | +| Next action | Build candidate (3) behind `MCORE_MOE_FUSED_COUNT` (default off, requires `MCORE_MOE_FUSED_ALIGN`) → QWEN-019. Two levers the gate surfaced but did not pursue, in order: the router pair `gatherTopK` + softmax is **436.8 µs/step = 4.4%** in two kernels (a hand-written fused top-8 router honouring the `inference_optimized` dense top-8 contract; QWEN-009 only ruled out TE's *built-in* fusion, not a hand-written one), and `_moe_sum` at 7.79 µs moves ~6.3 MB, i.e. ~1.05 µs at the measured 6.08 TB/s — it is **7× off its own bandwidth floor**, so restructuring it in place (hoist the 16 per-token scalar `routing_map` loads out of the K loop, parallelise over K as well as tokens) is worth up to 48 × 5.8 = 278 µs = 2.8% *without* the atomics that fusing into FC2 would require. | + +### CLEANBASE-S6 — fresh same-session reference + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | Establish a same-session OSL1024 reference at the session-4/5 best config before any session-6 A/B. Session 5 skipped this and session-to-session drift is ~1.5%. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b) | +| Changed files | none | +| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`; nvls dispatcher, vLLM grouped-GEMM backend, `transformer_impl=inference_optimized`, full-iteration CUDA graphs, `CUDA_DEVICE_MAX_CONNECTIONS=1` | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters | +| Job / run | session `qwen-comm`, Slurm job 5601961; exec `023b75121915449badd4b39ca610f470`; run dir `sessions/qwen-comm/e2e/ref-s6-1785024868` | +| Throughput | **23,264.4 tok/s** (per-iter 23,019.4 / 23,402.0 / 23,189.7 / 23,457.4 / 23,258.9) | +| Latency / TPOT | avg_latency 10,974.7 ms, p50 10,972.5, p99 11,346.9; TPOT 11.004 ms/tok | +| Correctness | Benchmark completed 5/5 iters at 256 requests | +| Nsight artifacts | none (un-profiled throughput run) | +| Result | Reference for session 6 = 68.44% of vLLM 33,994.5. −1.61% vs QWEN-013b's 23,646.0, consistent with the documented ~1.5% session drift; all A/Bs this session are against this number, not against 23,646.0. | +| Next action | Verified the synced snapshot (`workspaces/megatron_lm/6b9355b187072223`) contains no session-6 kernel, so this is exactly the accepted session-5 code. Proceed to QWEN-019. | + +### QWEN-019 — fold the token count into the indirection-table build + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | QWEN-018's chosen candidate. The decode indirection-table build spends 4 launches producing a 32-element count vector and then consuming it: a `torch.zeros` fill (0.74 µs), `_count_local_tokens_kernel_persistent` (7.52 µs), `_prefix_fill_init_kernel` (1.33 µs), `_scatter_token_indices_kernel` (2.66 µs). The count's only consumer already has **every CTA redundantly recompute the whole cumsum in registers**, so recomputing the histogram per CTA as well removes the first two launches with no grid sync and no change to the counts. Ceiling 48 × (7.52 + 0.55 + 0.74 + 0.55) = 449 µs = 4.52% of the step. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_count_prefix_fill_init_kernel`, `_moe_align_block_size_count_fused`, `_USE_FUSED_COUNT` / `MCORE_MOE_FUSED_COUNT` gate, `_FUSED_COUNT_MAX_TOKENS`, three-way `align_fn` selection), `dev/moe_fused/harness_countfuse.py` (new A/B harness) | +| Runtime flags | `MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. The fused-count path is restricted to `num_tokens_hint <= 512`; above that the redundant per-CTA read stops being free and the atomic count kernel is used. | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 300 CUDA-graph replays × 3 repeats at 128/256/384/512 tokens (timing) and 16 table-equality cases (correctness) | +| Job / run | session `qwen-comm`, job 5601961; microbench exec `d9f7f74a868241d69eb4cc1f4062d40b`; e2e exec `3a807f77358644c89e62866749519802`, run dir `sessions/qwen-comm/e2e/countfuse-1785025539` | +| Throughput | **23,964.4 tok/s** vs **23,264.4** same-session reference → **+3.01%**. Per-iter 23,566.6–24,304.6 | +| Latency / TPOT | avg_latency 10,642.7 ms vs 10,974.7 (−3.03%); TPOT 10.682 vs 11.004 ms/tok (−2.93%). CUDA-graph replay device time at 256 tokens: align call **16.38 → 10.24 µs (1.600×)**, whole `vllm_fused_moe` call **86.02 → 79.96 µs (1.076×)**, i.e. 6.06 µs/layer × 48 = 291 µs/step = 2.93% of the 9.93 ms step — the e2e result matches the microbench prediction to within 0.1 pp. Per-kernel (eager profiler, 100 iters): count 4.76 + zeros 1.20 + prefix_fill 1.76 = 7.72 µs replaced by a single 3.73 µs kernel; align total device time 10.17 → 6.11 µs | +| Correctness | **Bit-exact, and verified two ways.** (1) Table equality over 16 cases (valid ∈ {128,256,384,512} × BLOCK_M ∈ {16,64} × local_expert_start ∈ {0,32}): `num_tokens_post_padded` identical, `expert_ids` identical, and the sorted multiset of `sorted_token_ids[0:npp]` identical (sorted because the scatter's atomics permute rows within an expert block on both paths). (2) Whole-MoE output at all four token counts: max_abs 0.0 **and max_rel 0.0** — per QWEN-014's lesson the relative error is checked, not a loose `allclose`. e2e coherence prompts coherent (2+2=4, Paris, 3+2 cows = 5 cows) | +| Nsight artifacts | none (CUDA-graph event timing + torch profiler; the change is a launch-count/kernel-time change fully explained by the two removed launches) | +| Result | **Accepted** — bit-exact, +3.01% e2e with tight variance, no regression. Default OFF so untouched paths stay byte-identical. New best mcore = **23,964.4 tok/s = 70.49%** of the vLLM 33,994.5 baseline. Realised 291 of the 449 µs ceiling; the shortfall is the new kernel's own cost, 1.33 → 3.73 µs (eager) as each of the 32 CTAs now histograms all 2048 valid pairs instead of reading a 32-element vector. **The mechanism confirms QWEN-018's correction of QWEN-003**: the count kernel's cost was in-kernel, not per-launch, and the fix was to stop running it on 2 of 152 CTAs — not to change how it reduces. | +| Next action | The same breakdown ranks `_moe_sum` next: 7.79 µs/layer to move ~6.3 MB is 7× its own bandwidth floor, and it can be attacked without the cross-CTA atomics that fusing it into FC2 would need → QWEN-020. After that, the router pair (`gatherTopK` + softmax, 436.8 µs/step = 4.4% in two kernels) is the largest remaining routing item. | + +### QWEN-020 — predicate the locality test in the topk reduction + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | QWEN-018 measured `_moe_sum_kernel` at 7.79 µs/layer to move ~6.3 MB — 7× the ~1.05 µs its own traffic implies at the 6.08 TB/s of QWEN-012 — so it is neither launch- nor bandwidth-bound. The suspected cause is the `if lid >= 0 and lid < num_local_experts` guard: a uniform scalar branch gated on a dependent global load of `routing_map`, which prevents slot `t`'s data load from overlapping slot `t+1`'s index load and makes the CTA walk the 8 topk slots serially. Predicating the guard into the load mask, and widening `BLOCK_K` from 1024 to the full hidden size so the per-token index loads are issued once instead of `NUM_K_BLOCKS` times, should recover most of that gap **while staying bit-exact** — the reduction order and the fp32 arithmetic are unchanged and masked-off slots contribute an exact 0.0. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019) | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_moe_sum_kernel_fast`, `_USE_FAST_MOE_SUM` / `MCORE_MOE_SUM_FAST` gate, `_FAST_MOE_SUM_MAX_BLOCK_K = 2048`, dispatch in `_moe_sum`), `dev/moe_fused/harness_moesum.py` (new A/B harness) | +| Runtime flags | `MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. One new variable vs QWEN-019. | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 300 CUDA-graph replays × 3 repeats at 128/256/384/512 valid tokens, plus whole-MoE output equality at the same four shapes | +| Job / run | session `qwen-comm`, job 5601961; microbench exec `af464a419c284937a106aae25196359a`; e2e exec `af49c9f54f4a497d8808f50614fe6bf6`, run dir `sessions/qwen-comm/e2e/moesum-1785026515` | +| Throughput | **24,403.7 tok/s** vs **23,964.4** (QWEN-019) → **+1.83%**; vs **23,264.4** (CLEANBASE-S6) → **+4.90%**. Per-iter 24,442.9 / 24,137.8 / 24,580.4 / 24,607.1 / 24,257.1 | +| Latency / TPOT | avg_latency 10,445.6 ms (p50 10,447.9, p99 10,806.4); TPOT **10.490** vs 10.682 ms/tok. Microbench, CUDA-graph replay of the whole `vllm_fused_moe` call: 128 tok 73.10 → 71.59 µs (1.021×), 256 tok 80.69 → 78.77 (1.024×), 384 tok 86.81 → 83.94 (1.034×), 512 tok 102.68 → 98.56 (1.042×). Eager per-kernel at 256 tokens (100 iters): `_moe_sum_kernel` **8.14 → 5.97 µs (1.36×)**, every other kernel in the call unchanged (`_fused_moe_kernel` 62.18 → 61.88, `_count_prefix_fill_init` 4.18 → 4.18, `_scatter_token_indices` 3.07 → 3.03) | +| Correctness | **Bit-exact.** Whole-MoE output at valid ∈ {128, 256, 384, 512}: max_abs 0.0 **and max_rel 0.0** (QWEN-014's lesson — relative error, not a loose `allclose`). e2e coherence prompts coherent (2+2 = 4, Paris, 3+2 cows = 5 cows). Benchmark completed 5/5 iters | +| Nsight artifacts | none (CUDA-graph event timing + torch profiler; the change alters one kernel's internals and no launch counts, which the per-kernel table isolates directly) | +| Result | **Accepted** — bit-exact and positive on every shape measured. New best mcore = **24,403.7 tok/s = 71.79%** of the vLLM 33,994.5 baseline. Default OFF; enable with `MCORE_MOE_SUM_FAST=1`. **Honest accounting of the size of the win:** the microbench says 2.17 µs/layer × 48 = **104 µs/step ≈ 1.0%** of the step, and the graph-replay whole-call delta at 256 tokens is smaller still (1.92 µs × 48 = 92 µs), whereas the e2e TPOT moved 0.192 ms/step (192 µs). The e2e gain is therefore about **2× the microbench prediction**, and QWEN-019's per-iteration range (23,566–24,305) overlaps this run's (24,138–24,607), so part of the measured +1.83% is run-to-run variance rather than kernel time. The change is kept because it is bit-exact and every isolated measurement of it is a strict improvement, but the defensible attribution is ~1%, not 1.83%. QWEN-018's estimated ceiling for restructuring this kernel in place was 278 µs (2.8%); the predication recovered 104 µs of it, so `_moe_sum` at 5.97 µs is still ~6× its bandwidth floor and the remaining gap is the strided `[token, slot, K]` gather itself, not the branch | +| Next action | The router pair is now the largest routing item the QWEN-018 breakdown left standing: `gatherTopK` 6.05 µs + softmax 1.94 µs = 436.8 µs/step = 4.4% in two kernels, for what is only a `[256, 128]` fp32 softmax and a top-8 select → QWEN-021. | + +### QWEN-021 — fused softmax + top-8 router selection + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | QWEN-018 left the router pair as the largest untouched routing item: `gatherTopK` 6.05 µs + the compiled softmax 1.94 µs = 436.8 µs/step (4.4%) — for nothing more than a `[256, 128]` fp32 softmax and a top-8 select over 128 experts. `torch.topk` runs a multi-pass radix select sized for large `n`; at 128 candidates per row, one CTA per token can hold the whole row in registers, softmax it, and pick the top 8 by 8 max-then-mask passes. Ceiling: replace ~8.0 µs + 0.55 µs of dispatch gap per layer with a single ~2 µs kernel ⇒ ~312 µs/step ≈ 3.1%. **This is not QWEN-009**, which ruled out TE's *built-in* fused router because it emits a dense 128-expert map while `inference_optimized` needs the dense top-8 contract; this kernel is written to the top-8 contract directly. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020) | +| Changed files | `megatron/core/inference/moe/router_topk.py` (new: `_softmax_topk_kernel`, `fused_softmax_topk`, `can_use_fused_softmax_topk`, `MCORE_ROUTER_FUSED_TOPK` gate, `FUSED_ROUTER_TOPK_MAX_TOKENS = 1024`), `megatron/core/transformer/moe/router.py` (`InferenceTopKRouter._forward` takes the fused path when the contract matches), `dev/moe_fused/harness_routertopk.py` (new A/B harness) | +| Runtime flags | `MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. One new variable vs QWEN-020. The fused path is taken only for softmax + pre-softmax + no groups + no scaling factor + no expert bias + no router replay + ≤1024 tokens; anything else (including prefill chunks above 1024 tokens) falls back to the compiled torch path unchanged. | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 500 CUDA-graph replays × 3 repeats at 128/256/384/512 tokens × 128 experts, correctness at those four shapes × 2 seeds | +| Job / run | session `qwen-comm`, job 5601961; microbench exec `56ae1d844a86427ea60ddce2c73d903c`; e2e exec `2041011b410b45709d455514b32c85f3`, run dir `sessions/qwen-comm/e2e/routertopk-1785027029` | +| Throughput | **25,352.9 tok/s** vs **24,403.7** (QWEN-020) → **+3.89%**; vs **23,264.4** (CLEANBASE-S6) → **+8.98%**. Per-iter 25,012.5 / 25,092.4 / 25,405.9 / 25,625.0 / 25,642.4 — the whole range sits above QWEN-020's whole range (24,137.8–24,607.1), so unlike QWEN-020 this win is outside run-to-run variance | +| Latency / TPOT | avg_latency 10,041.4 ms (p50 10,039.1, p99 10,439.5); TPOT **10.097** vs 10.490 ms/tok → **−393 µs/step** against a predicted ~312 µs. CUDA-graph replay of the router selection alone: 128 tok 16.40 → 4.10 µs, 256 tok **16.41 → 4.10 µs (4.00×)**, 384 tok 18.43 → 4.10 (4.49×), 512 tok 20.47 → 4.11 (4.99×) — the fused kernel is flat in token count over this range while the torch pair is not. Eager per-kernel at 256 tokens: `gatherTopK` 9.12 + `unrolled_elementwise` 2.76 + `softmax_warp_forward` 2.09 + `vectorized_elementwise` 1.57 = **15.54 µs in 4 kernels → 2.04 µs in 1** | +| Correctness | **Bit-exact probabilities and identical expert sets** at 128/256/384/512 tokens × 2 seeds each: the selected expert id multiset per token is equal, max_abs 0.0 and max_rel 0.0 on the probabilities after aligning by expert id, and no duplicate expert ids. (`torch.topk` is called with `sorted=False` during inference, so the *order* of the k results is unspecified on both sides; the fused kernel returns descending score order with ties broken toward the lower expert id.) e2e coherence prompts coherent (2+2 = 4, Paris, 3+2 cows = 5 cows); benchmark 5/5 iters | +| Nsight artifacts | none (CUDA-graph event timing + torch profiler; a 4→1 launch change with per-kernel attribution on both sides) | +| Result | **Accepted — the largest single win of session 6.** New best mcore = **25,352.9 tok/s = 74.58%** of the vLLM 33,994.5 baseline. Default OFF; enable with `MCORE_ROUTER_FUSED_TOPK=1`. The measured e2e gain (393 µs/step) slightly exceeds the 312 µs predicted from QWEN-018's two-kernel accounting because the profile's "router" attribution missed the two elementwise kernels (the `.type_as` cast and one more) that the fused kernel also absorbs — 4 launches removed, not 2. Session total so far: 23,264.4 → 25,352.9 = **+8.98%**, all of it bit-exact. | +| Next action | The last item in the QWEN-018 breakdown that clears the gate is `_scatter_token_indices_kernel` at 2.66 µs + 0.55 µs gap × 48 = 154 µs = 1.6%. Now that QWEN-019 made every CTA stream the pairs anyway, the scatter can be a second streaming pass inside the same kernel — no atomics, no grid sync — taking the indirection-table build to a single launch → QWEN-022. | + +### QWEN-022 — one-launch indirection-table build (scatter folded in) + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | The last routing item above the gate is `_scatter_token_indices_kernel`: 2.66 µs + 0.55 µs dispatch gap × 48 = 154 µs/step ≈ 1.6%. QWEN-019 already has CTA `e` stream every valid pair to build a private histogram, so a second streaming pass over the same (L2-resident, 16 KB) pairs lets that CTA place its own rows itself at `excl_e + written + exclusive_cumsum(is_mine)` — no global atomics, no grid sync, and a deterministic table instead of an atomically permuted one. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020 + QWEN-021) | +| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_align_single_kernel`, `_moe_align_block_size_single`, `_USE_FUSED_SCATTER` / `MCORE_MOE_FUSED_SCATTER` gate, align dispatch), `dev/moe_fused/harness_scatterfuse.py` (new A/B harness) | +| Runtime flags | `MCORE_MOE_FUSED_SCATTER=1` on top of the QWEN-021 set (`MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`). One new variable. | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters, **run twice**; microbench 300 replays × 3 repeats at 128/256/384/512 tokens, 16 table-equality cases | +| Job / run | session `qwen-comm`, job 5601961; microbench execs `c7cfd653ca6f4bc1849649746fb1cb48` (first, with a faulty check) and `46060ef1233f407fb929ab82755517d4` (corrected); e2e execs `a31dd240c22d47588ffa58e264180f32` (`e2e/scatterfuse-1785027598`) and `a5557671044f45709e0a933e7b255741` (`e2e/scatterfuse-rep2-1785027955`) | +| Throughput | run 1 **25,441.4 tok/s** (+0.35% vs QWEN-021's 25,352.9), run 2 **25,495.9 tok/s** (+0.56%); mean **25,468.7 = +0.46%**. Per-iter run 1 24,920.5–25,791.6, run 2 24,870.8–25,827.4 — both ranges overlap QWEN-021's (25,012.5–25,642.4) | +| Latency / TPOT | TPOT 10.062 and 10.041 ms/tok vs 10.097 (−35 and −56 µs/step) against a microbench prediction of 125 µs/step. CUDA-graph replay at 256 tokens: align call **10.08 → 8.23 µs (1.224×)**, whole MoE call **82.17 → 79.56 µs (1.033×)**. At other shapes: 128 tok 1.045×, 384 tok 1.007×, **512 tok 0.997× (a slight loss)** — the extra streaming pass scales with pair count while the removed launch does not. Eager per-kernel: `_count_prefix_fill_init` 3.74 + `_scatter_token_indices` 2.38 = 6.12 µs in 2 kernels → `_align_single_kernel` 4.91 µs in 1 | +| Correctness | **Bit-exact whole-MoE output** at 128/256/384/512 (max_abs 0.0 and max_rel 0.0), plus table equality over 16 cases (valid × BLOCK_M ∈ {16,64} × local_expert_start ∈ {0,32}): identical `num_tokens_post_padded`, identical `expert_ids`, identical row multiset, and — the stronger check — every non-sentinel row sits in a block whose expert id actually owns that pair. e2e coherence prompts coherent; both benchmark runs 5/5 iters. **Note on the first harness run:** it reported a table MISMATCH at BLOCK_M=16, which was a bug in the *check*, not the kernel — an expert's rows span several BLOCK_M blocks and the two paths order them differently across that range, so comparing per 16-row block flags a difference that does not exist. The whole-MoE output was bit-exact in that same run, which is what exposed the faulty check. | +| Nsight artifacts | none (CUDA-graph event timing + torch profiler) | +| Result | **Accepted, marginally, and recorded as a weak win rather than a 1.2% one.** Both e2e runs beat QWEN-021, so the sign is reliable; the size (+0.46%) is a third of the microbench prediction and sits inside the per-iteration spread, so the defensible claim is "small but positive". Kept because it is bit-exact, removes a launch, and makes the table deterministic. Default OFF. Best mcore = **25,495.9 tok/s = 75.00%** of the vLLM 33,994.5 baseline (mean of the two runs 25,468.7 = 74.92%). **This is the point where the routing lever runs out**: every routing kernel QWEN-018 listed has now been fused or rewritten except `_moe_sum`, whose remaining gap needs the FC2-epilogue fusion QWEN-018 gated as ~3.3% net and QWEN-001-shaped. | +| Next action | Re-profile at the new configuration before choosing the next lever — the step has lost ~10% of its wall time and roughly 700 µs of routing, so the QWEN-018 ranking is stale → PROFILE-S6. | + +### PROFILE-S6 — re-profile at the session-6 configuration + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | The QWEN-018 ranking is stale: the step has lost ~10% of its wall time and roughly 820 µs of routing since it was taken. Re-profile with all six gates on and re-rank before choosing the next lever. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020 + QWEN-021 + QWEN-022) | +| Changed files | none (`dev/moe_fused/profile_insession.sh`) | +| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, **OSL128** under nsys (`--trace=cuda,nvtx --cuda-graph-trace=node`), 1 iter after a BS8 warmup; analysis window device 1, 184.9–186.9 s (~220 steady-state decode steps) | +| Job / run | session `qwen-comm`, job 5601961; exec `1a4e8dc90c744a1c9ae0d977ec5ad16e`; run dir `sessions/qwen-comm/prof/s6-all-1785028240` | +| Throughput | 10,454.6 tok/s at OSL128 under nsys (profiling overhead + short OSL; not comparable to the OSL1024 ledger numbers) | +| Latency / TPOT | one steady-state decode step: **wall 9.097 ms** (PROFILE-TUNED: 9.933), **GPU-busy 7.183 ms** (7.884), **idle 1.914 ms = 21.0%** (2.049 ms = 20.6%), **1170 kernels/step** (1362) | +| Correctness | n/a (profile capture) | +| Nsight artifacts | `qwen-comm:prof/s6-all-1785028240/mcore_profile.nsys-rep` and `.sqlite`; `forward_pass.py` and `dev/moe_fused/analyze_routing.py --device 1 --window 184.9,186.9` output | +| Result | **The routing lever is spent.** Per forward pass: MoE expert GEMM 2469.1 µs / 96 kernels (27.1%), dense GEMM 1083.6 / 241 (11.9%), comm 1062.9 / 96 (11.7%), attention 923.8 / 96 (10.2%), elementwise 487.6 / 253 (5.4%), norm 449.5 / 193 (4.9%), **MoE routing/permute 430.4 µs / 98 kernels (4.7%)** — down from 1251 µs / 242 kernels, a 66% cut, and it is now the smallest GPU category except sampling. Counting dispatch gaps, the routing family (including `_align_single_kernel`, which the categorizer files under "other") is ~665 µs/step. Its two survivors are `_moe_sum_kernel_fast` (5.57 µs × 48 = 294 µs with gaps) and `_align_single_kernel` (5.13 × 48 = 272 µs); `_softmax_topk_kernel` costs 1.53 µs × 48 = 99 µs, versus 436.8 µs for the pair it replaced. **The idle is now the ranked next lever and it is host-side, not GPU-side**: of 1914 µs, ~1237 µs sits in three once-per-step gaps — **548 µs** after `index_elementwise_kernel`, **386 µs** after `vectorized_elementwise_kernel`, **303 µs** after `CatArrayBatchedCopy_vectorized` — i.e. sampling, detokenize and scheduling on the host between graph replays. The remainder (~644 µs) is 1170 × 0.55 µs of intra-graph node dispatch, which only fewer kernels can shrink. | +| Next action | Attack the ~1237 µs/step of host gaps, but **profile the host first** — this is not the same experiment as QWEN-005a/QWEN-007 (`async-sched-mode=serial`, which crashed and then measured only +0.85%) or QWEN-004 (`sampling-backend=flashinfer`, incompatible with the full-iteration graph). Capture with `--trace=cuda,nvtx,osrt` plus Python sampling over the decode window and attribute the three gaps to concrete host functions before proposing a change; at 13.6% of the step this is now worth more than any remaining GPU-side category except the (closed) expert GEMM. | + +### PROFILE-HOST-S6 — host-visible capture for the inter-replay gaps + +| Field | Value | +|---|---| +| Date | 2026-07-25 | +| Hypothesis | PROFILE-S6's ~1237 µs/step of host gaps between CUDA-graph replays (548 µs after `index_elementwise_kernel`, 386 µs after `vectorized_elementwise_kernel`, 303 µs after `CatArrayBatchedCopy_vectorized`) are unattributed because every existing trace is GPU-only. A trace with `osrt` + CPU sampling + Python sampling will name the host functions occupying them and classify each gap as launch overhead, host synchronization, sampling/detokenization, scheduler bookkeeping, or ZMQ round-trip. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (session-6 set unchanged: `vllm_fused_moe.py`, `experts.py`, `router.py`, untracked `router_topk.py`). **No `megatron/` edits made by this task** — verified with `git diff --stat -- megatron`. | +| Changed files | `dev/moe_fused/profile_host_insession.sh` (new — host-visibility variant of `profile_insession.sh`), `dev/moe_fused/dispatch_host_profile.sh` (new — waits for the session to reach `running`, then dispatches the capture detached) | +| Runtime flags | All seven gates: `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | +| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | +| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1, nvls dispatcher, vllm grouped-GEMM, `transformer_impl=inference_optimized`, `full_iteration_inference` CUDA graphs | +| Workload | gsm8k, BS256, OSL128 under nsys, 1 iter after a BS8/OSL32 warmup (identical to PROFILE-S6 except for the trace flags) | +| nsys version | 2026.3.1.117-263137992252v0 — verified on-node that `--python-sampling`, `--python-sampling-frequency`, `--backtrace`, `--samples-per-backtrace`, `--cpuctxsw`, `--osrt-threshold` are all supported before committing to the command line | +| nsys command line | `nsys profile --trace=cuda,nvtx,osrt --sample=process-tree --backtrace=fp --samples-per-backtrace=1 --cpuctxsw=process-tree --python-sampling=true --python-sampling-frequency=1000 --osrt-threshold=1000 --cuda-graph-trace=node --force-overwrite=true -o /mcore_host_profile $PYBIN -m torch.distributed.run --nproc-per-node 4 ... -m examples.inference.launch_inference_server ...` (server args identical to `profile_insession.sh`). The script preflights this flag set against `python -c print(...)` and degrades to `--trace=cuda,nvtx,osrt --sample=process-tree --backtrace=fp --python-sampling=true` and then to `--trace=cuda,nvtx,osrt --sample=process-tree` if nsys rejects it, so an unsupported flag cannot abort the run. | +| Job / run | Intended session `qwen-comm` job 5601961 — **PREEMPTED** (`raw_state=PREEMPTED`, controller dead) before any capture ran; note `cog jobs get` still reported `RUNNING` while `cog session status` correctly reported `preempted`. Replacement: session `qwen-host`, job **5607600**, node `nvl72166-T15`, partition `batch`, `--time 02:00:00` (started 2026-07-26T06:53 UTC, expires ~08:53 UTC), 4 GPUs. Capture dispatched detached as exec **`d920241ea00a41618c63a75fefcb9ea2`** at 06:53:51 UTC; run dir **`sessions/qwen-host/prof/hosts6-1785048883`**. Exec log: `sessions/qwen-host/exec/runs/d920241ea00a41618c63a75fefcb9ea2/stdout.log`. | +| Throughput | n/a (profile capture) | +| Latency / TPOT | n/a (profile capture) | +| Correctness | n/a (profile capture) | +| Nsight artifacts | **Expected at `sessions/qwen-host/prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}` — dispatched and confirmed running, but not confirmed complete.** At hand-off the exec had reached the nsys flag preflight; the server load + BS8 warmup + BS256/OSL128 benchmark + sqlite export had not finished. The script ends with a sanity check that prints row counts for `CUPTI_ACTIVITY_KIND_KERNEL`, `CUPTI_ACTIVITY_KIND_RUNTIME`, `OSRT_API`, `COMPOSITE_EVENTS`, `SAMPLING_CALLCHAINS`, `PYTHON_SAMPLING_CALLCHAINS`, `PYTHON_SAMPLING_STRING`, `SCHED_EVENTS`, `GENERIC_EVENTS` and a `HOST_TABLES_PRESENT=` line — a capture with only `CUPTI_ACTIVITY_KIND_KERNEL` is a failed capture and must be re-run. | +| Result | **Inconclusive — capture not confirmed, no attribution performed.** The task lost its node: session `qwen-comm` job 5601961 was preempted roughly six hours into an eight-hour `batch_long` allocation, before the first capture could be dispatched, and the whole capture had to be re-queued from scratch on a new allocation that did not schedule inside the remaining budget. The deliverable that *was* produced is the reusable capture path: an nsys flag set validated against the installed 2026.3.1 on-node, a profile script that preflights and degrades rather than aborting, and a built-in host-table sanity check. Two operational findings worth carrying: (1) `cog jobs get --job-id` reported `RUNNING` for a preempted session — use `cog session status` to decide whether a session is usable; (2) an unrelated background `rsync` writing into the worktree caused two `WorktreeChangedError` dispatch failures, so no local process may write inside the repo while a `cog session exec` snapshot is in flight. | +| Next action | Check `sessions/qwen-host/exec/runs/d920241ea00a41618c63a75fefcb9ea2/stdout.log` for `PROFILE DONE` and `HOST_TABLES_PRESENT=True`. If present, run Workflow C on the new `.sqlite`, then window the `OSRT_API` / `CUPTI_ACTIVITY_KIND_RUNTIME` / `PYTHON_SAMPLING_CALLCHAINS` tables to the three inter-replay gaps and attribute them. If the exec never ran, re-dispatch `dev/moe_fused/profile_host_insession.sh` on a fresh session — it is unchanged and ready. | + +### HOSTGAP-S6 — host attribution of the three inter-replay gaps + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | PROFILE-S6's three largest inter-kernel gaps (after `index_elementwise_kernel`, `vectorized_elementwise_kernel`, `CatArrayBatchedCopy_vectorized`, ~1237 µs/step together) can be attributed to named host functions from the PROFILE-HOST-S6 capture, and each can be classified as CPU-bound host execution, blocking synchronization, driver launch latency, allocator activity, a genuine GPU dispatch gap, or profiler overhead. | +| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (session-6 set). **No `megatron/` edits** — analysis only. | +| Changed files | `dev/moe_fused/analyze_hostgaps.py` (new — `gaps` brackets every inter-kernel gap in the kernel-union complement by the kernel ending at its start and the kernel starting at its end; `attrib` aggregates CUPTI runtime, OSRT, NVTX and native sampling backtraces over all instances of one bracketed gap class) | +| Runtime flags | n/a (analysis of existing traces) | +| Image | n/a | +| Checkpoint / tokenizer | n/a | +| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 (as captured) | +| Workload | Analysis of the BS256/OSL128 decode loop in two traces of the same seven-gate configuration | +| Job / run | session `qwen-host` job 5607600, exec `d920241ea00a41618c63a75fefcb9ea2`, `sessions/qwen-host/prof/hosts6-1785048883` | +| Throughput | n/a. For the record, the host-traced benchmark itself reports 6,631.7 tok/s / TPOT 38.60 ms (vs PROFILE-S6's 10,454.6 tok/s) — but that overhead lands almost entirely **outside** the decode loop (model load, chunked prefill, graph capture); the steady-state decode step is only 2.1% slower under host tracing. | +| Latency / TPOT | Steady-state decode step: **9.134 ms** (GPU-only) vs **9.330 ms** (host trace) = 1.021×. Idle 1965.7 vs 2128.0 µs = 1.083×. | +| Correctness | n/a (analysis) | +| Nsight artifacts | Host: `qwen-host:prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}` (615 MB / 5.9 GB). GPU-only reference: `nsys_trace/mcore_s6_tuned_osl128.sqlite`. Analysis: `dev/moe_fused/analyze_hostgaps.py` | + +**Trace provenance and window — two corrections that had to be made first.** +The sqlite is intact (`pragma quick_check` ok) and complete despite the offline +`QdstrmImporter` recovery: all four ranks are present with ~479 k kernels each +spanning 75.0–293.0 s, four engine PIDs with ~665 k CUDA-runtime events each, +and OSRT/sampling covering 0–300 s. Two things were *not* as expected: + +1. **The densest kernel region is not the decode loop.** Seconds 196–241 look + like a steady loop (a `_fused_metadata_kernel` every ~146 ms, 48 layers per + period) but it is **CUDA-graph capture**: 25 distinct capture streams, no + BS256 LM-head GEMM, `index_elementwise_kernel` count **zero**, and + `_multimem_all_gatherv_3tensor_kernel` averaging **1807 µs** (vs 15.8 µs in + the GPU-only trace) because the symmetric-memory barrier absorbs inter-rank + skew while the other ranks are capturing. The real BS256/OSL128 decode loop + is a 1.3 s burst at **291.67–292.98 s** (146 anchor intervals, period 9.064 + ms), preceded by the BS8/OSL32 warmup at 287.51–287.84 s (44 steps, 7.53 ms). + Analysing 217–241 s would have produced a 16×-inflated fiction. +2. **`PYTHON_SAMPLING_*` and the graph tables are absent** despite + `--python-sampling=true`. Attribution therefore uses NVTX ranges + + `CUPTI_ACTIVITY_KIND_RUNTIME` + `OSRT_API` + native `SAMPLING_CALLCHAINS`, + which turned out to be sufficient because the engine emits per-phase NVTX. + +**Method.** Kernel-union complement per anchor-to-anchor step, anchored on the +once-per-step LM-head GEMM `nvjet_sm100_tst_512x64_64x3_2x1_2cta_v_bz_TNT` on +all four devices in both traces (forcing the same anchor matters: auto-detection +picked `vectorized_elementwise_kernel` on devices 0/2 of the host trace, which +straddles G1 and silently dropped it). Each gap is bracketed by both neighbours. +Windows: GPU-only 184.9–186.9 s (126 steps), host 291.75–292.95 s (117 steps). +The GPU-only trace is the source of truth for magnitude, the host trace for +attribution. + +**Per-step gap budget (GPU-only trace, µs/step, per rank).** + +| dev | step ms | busy ms | idle µs | G1 | G2A | G2B | G3 | G4 | +|---|---:|---:|---:|---:|---:|---:|---:|---:| +| 0 | 9.135 | 7.195 | 1940.2 | 559.0 | 63.7 | 259.1 | 303.8 | 48.8 | +| 1 | 9.135 | 7.181 | 1953.7 | 566.8 | 62.4 | 264.0 | 305.7 | 49.9 | +| 2 | 9.134 | 7.141 | 1993.8 | 582.7 | 63.1 | 275.8 | 308.6 | 54.1 | +| 3 | 9.133 | 7.158 | 1975.0 | 576.9 | 63.1 | 273.0 | 303.6 | 52.7 | +| **mean** | **9.134** | **7.169** | **1965.7** | **571.4** | **63.1** | **267.9** | **305.4** | **51.4** | + +Every rank is within 4% of every other on every row — **no rank asymmetry**, as +expected when the four EP ranks are barrier-coupled by the multimem collectives. +The five bracketed gaps are **1259.3 µs/step = 13.79% of the step and 64.1% of +the idle**; the residual 706.4 µs is 1158 × ~0.61 µs of intra-graph node +dispatch. Host-trace equivalents (mean of ranks) are G1 641.2, G2 378.5, G3 +308.2 µs — inflation 1.12× / 1.14× / 1.01×, so the host trace neither invents +nor hides these gaps. + +**Gap 1 — `index_elementwise_kernel` → `vectorized_elementwise_kernel`.** +571.4 µs/step GPU-only (host: median 635.4, mean 640.9, p10 585.0, p90 676.1, +116 instances, 0.99/step). Verdict **(i) CPU-bound Python**. NVTX covers 72.5%: +`update_requests` 183.6 µs (28.6%), `initialize_attention_state` 151.7 (23.7%), +`active_request_mask` 78.0 (12.2%), `transfer_samples_to_cpu` 41.2 (6.4%), +`sampling` tail 10.5 (1.6%); the remaining 27.5% is an unlabelled engine window +between `update_requests` and `initialize_attention_state` (median 152.2 µs) +that is itself 39.6%-leaf `_PyEval_EvalFrameDefault` with `at::_ops::_to_copy` +in 34% of stacks. Only **55.1 µs (8.6%)** of the gap is inside any CUDA API +(`cudaMemcpyAsync` ×2 = 41.7 µs, `cudaLaunchKernel` 8.6, `cudaStreamSynchronize` +**3.4 µs = 0.5%**), and only 32.9 µs (5.1%) inside any syscall on the engine +thread — so ~86% is plain userspace execution. All 161 samples landing in these +windows have `threadState = Running`; of the 142 on the engine thread, 100% have +`_PyEval_EvalFrameDefault` and `_PyObject_MakeTpCall` somewhere in the stack, +with `at::native::copy_` in 12.7%. Source: +`text_generation_controller.py:1750` (`transfer_samples_to_cpu`, D2H at `:1723`), +`:1756`–`1805` (`active_request_mask`), `:1807`–`1811` (`update_requests` → +`dynamic_context.py:3515`), `:644`–`649` (`initialize_attention_state` → +`dynamic_context.py:2128`, whose own comment at `:643` reads "100% CPU +computation"). + +**Gap 2 — `vectorized_elementwise_kernel` → `vectorized_gather_kernel`, twice +per step, bimodal.** Instance **2B** (267.9 µs GPU-only; host median 301.1, mean +310.4) is NVTX-covered 97.7% by `initialize_attention_state` 159.9 µs (51.5%), +`forward_pass` head 99.0 (31.9%) and `transfer_bookkeeping_to_gpu` 44.3 (14.3%); +CUDA API is 35.9 µs (11.5%), engine-thread syscalls 7.0 µs (2.3%), and all 76 +samples are Running with 100% `_PyEval_EvalFrameDefault` — verdict **(i) +CPU-bound Python**, the same bookkeeping chain as G1 continued past the first +GPU op it emits. Instance **2A** (63.1 µs GPU-only; host median 70.8) sits +entirely inside the `sampling` NVTX range and is 34.4% CUDA API +(`cudaLaunchKernel` 15.6 µs = 21.8%, `cudaStreamSynchronize` 8.7 µs = 12.1%), +with `at::_ops::index_Tensor::call` / `at::native::index_kernel` in 55.6%/44.4% +of its (only 9) samples — verdict **(i)+(iii)**, advanced-indexing dispatch plus +a short blocking stream sync. Source: `text_generation_controller.py:652`–`654` +(`transfer_bookkeeping_to_gpu` → `dynamic_context.py:2411`), `:790`/`:1864` +(`forward_pass`), `:1902`/`:2018`–`2023` (`sampling`; `torch.argmax` at `:2019`, +`.cpu()` at `:2022`). + +**Gap 3 — `CatArrayBatchedCopy_vectorized` → `rmsnorm_fwd_tuned_kernel` (the +next graph replay).** 305.4 µs/step GPU-only (host median 311.5, mean 314.3, 117 +instances). Verdict **(iii) driver graph-launch latency**. It is 100% inside +`forward_pass`, and **63.5% of it is a single `cudaGraphLaunch` call**: 132 calls +in the window, median **199.1 µs**, p10 188.6, p90 211.9, unimodal (so it is not +a graph-switch/upload artifact — one graph id, 140, carries 135,528 of the +window's kernels). Plus `cudaMemcpyAsync` 18.4 µs (5.8%) and ~96 µs of Python. +Engine-thread syscalls are 0.7%; all 78 samples are Running, 56.6% with +`cuGraphLaunch` as the deepest resolved frame and 57.9% with +`at::cuda::CUDAGraph::replay()` in the stack. At 1158 nodes per replay that is +**0.172 µs of host submit time per graph node**, on the critical path before the +first node executes. + +**Aggregate, and the profiling-overhead correction.** GPU-busy is 7.169 ms = +78.5% of the 9.134 ms step; host gap is 1965.7 µs = **21.5%**. Splitting that +idle by confidence: + +| Component | µs/step | % of step | Confidence | +|---|---:|---:|---| +| G1 + G2A + G2B + G4 (host Python + sampling dispatch) | 953.8 | 10.4% | Real — samples Running in the interpreter | +| G3 minus its `cudaGraphLaunch` (Python + H2D in `forward_pass`) | 111.4 | 1.2% | Real | +| `cudaGraphLaunch` inside G3 | 194.0 | 2.1% | **Artifact-suspect** | +| Intra-graph inter-node dispatch (1158 × ~0.61 µs) | 706.4 | 7.7% | **Artifact-suspect** | +| **Total idle** | **1965.7** | **21.5%** | | + +Overhead was corrected by comparing the two traces step-for-step rather than by +assuming a factor: the host capture adds only **2.1%** to the decode step +(9.134 → 9.330 ms) and 8.3% to the idle, even though it costs 37% of +whole-benchmark throughput, because its cost falls on load/prefill/capture. The +one caveat that does **not** cancel is that *both* traces used +`--cuda-graph-trace=node`, which forces CUPTI to instrument every graph node; +the 194 µs `cudaGraphLaunch` and the 706 µs of inter-node dispatch — 900 µs/step, +9.9% of the step, 45.8% of the idle — are exactly the quantities that flag +distorts, and this analysis cannot separate them from the real cost. A +`--cuda-graph-trace=graph` control capture settles it for the price of one +profile run. + +**What the recoverable upside is worth.** Taking the current best mcore of +25,495.9 tok/s (QWEN-022, 75.00% of vLLM's 33,994.5) and scaling by +`9134 / (9134 − Δ)` — an upper bound that assumes perfect removal and that the +OSL128 composition transfers to the OSL1024 throughput regime: + +| Removed | Δ µs | ceiling tok/s | % of vLLM | Share of the remaining gap | +|---|---:|---:|---:|---:| +| G1 | 571.4 | 27,197 | 80.0% | 20.0% | +| G2 (both) | 331.0 | 26,454 | 77.8% | 11.3% | +| G3 | 305.4 | 26,378 | 77.6% | 10.4% | +| G1 + G2B (the serial bookkeeping chain) | 839.3 | 28,067 | 82.6% | 30.3% | +| All five bracketed gaps | 1259.3 | 29,573 | 87.0% | 48.0% | +| All idle | 1965.7 | 32,488 | 95.6% | 82.3% | + +**The finding that reframes the lever.** G1 and G2B are not a scheduling +artifact that overlap could hide. The chain is +`graph(N) → logits(N) → argmax(N) → D2H tokens(N) → active_request_mask(N) → +update_requests(N) → initialize_attention_state(N+1) → H2D → graph(N+1)`, and +every link needs the previous one's result: the host cannot run step N+1's +bookkeeping during step N's graph replay because that bookkeeping consumes step +N's sampled tokens. This is a genuine serial data dependency, and it explains +QWEN-007 cleanly — `async-sched-mode=serial` measured +0.85% not because the +idle was illusory but because with a single request stream there is nothing to +overlap it *with*. The lever is therefore to make the CPU work cheaper or move +it onto the GPU, not to hide it. + +**Ranked candidate next levers** (upside is a measured ceiling, not a prediction; +none of these has been measured): + +1. **Make `initialize_attention_state` incremental** — ~311 µs/step (152 µs in + G1 + 160 µs in G2B), pure CPU by construction. In steady-state decode the + request set is fixed for 128 steps and only the sequence lengths increment, + so most of the per-step attention metadata is recomputed unchanged. Ceiling + ~3.4% (→ ~77.6% of vLLM); realistic maybe half. **Risk low-medium**: CPU-only, + no numerics exposure, but correctness depends on invalidating the cache on + every request add/pause/finish. Falsified if a fresh profile shows the G1 and + G2B windows unchanged. +2. **Cut `update_requests` + `active_request_mask` + the unlabelled engine + window** — 183.6 + 78.0 + ~175 = ~437 µs/step of Python over 256 requests. + Ceiling 4.8% (→ ~78.8%). **Risk low**: no numerics, no guards, purely + additive with everything else. +3. **`--cuda-graph-trace=graph` control capture** — no throughput at all, but it + decides whether the 900 µs/step (9.9% of the step, 45.8% of the idle) of + graph machinery is real or instrumentation. **Risk none, cost one profile + run.** This gates lever 4 and should be run before it. +4. **Reduce graph node count** — each removed node is worth ~0.172 µs of host + submit plus ~0.61 µs of GPU inter-node dispatch ≈ 0.78 µs, so 100 nodes ≈ + 78 µs ≈ 0.85%. **Risk medium** and sharply diminishing; this is the fusion + campaign that QWEN-019…022 already mined. +5. **Move per-step request bookkeeping onto the GPU** (persistent device-side + batch state, no D2H→Python→H2D round trip per step), which is what vLLM does. + Ceiling is most of the 839 µs chain, ~9.2%. **Risk high** — a large engine + rewrite, multi-session. + +Explicitly *not* recommended: re-litigating `async-sched-mode=serial`. The +dependency analysis above says the idle it targets is not overlappable with a +single request stream, which is consistent with its measured +0.85%. + +| Field | Value | +|---|---| +| Result | **Supported — all three gaps attributed with numbers, and one taxonomy verdict each.** G1 (571 µs) and G2B (268 µs) are CPU-bound Python in the engine's per-step bookkeeping chain; G2A (63 µs) is sampling-path launch plus a short stream sync; G3 (305 µs) is 63.5% a single 199 µs `cudaGraphLaunch`. No rank asymmetry. The idle budget is 1065 µs of real host work and 900 µs of graph machinery whose reality is gated on a `--cuda-graph-trace=graph` control. The campaign-level conclusion is that the largest remaining lever is a serial data dependency, so it must be made cheaper rather than hidden. | +| Next action | Run the `--cuda-graph-trace=graph` control capture (lever 3, one profile run) to fix the denominator, then implement lever 1 (incremental `initialize_attention_state`) and measure it e2e at OSL1024 against CLEANBASE-S6 with the full seven-gate set. | + +### Session 6 — conclusion & recommendation + +The routing/permute lever is closed, and it delivered. Against the same-session +reference CLEANBASE-S6 (23,264.4 tok/s), the four accepted changes take mcore to +**25,495.9 tok/s = 75.00%** of the fixed vLLM OSL1024 baseline of 33,994.5, a +**+9.6%** session gain — and every one of them is bit-exact, so none of it was +bought with numerics. + +| Change | Mechanism | e2e | +|---|---|---:| +| QWEN-019 `MCORE_MOE_FUSED_COUNT` | count + zero-fill folded into the table build (4 launches → 2) | +3.01% | +| QWEN-020 `MCORE_MOE_SUM_FAST` | topk-reduction locality test predicated instead of branched | ~+1% (measured +1.83%, partly variance) | +| QWEN-021 `MCORE_ROUTER_FUSED_TOPK` | softmax + top-8 select in one CTA-per-token kernel (4 launches → 1) | **+3.89%** | +| QWEN-022 `MCORE_MOE_FUSED_SCATTER` | scatter folded in as a second streaming pass (2 launches → 1) | +0.46% | + +Reproduce the best configuration with all seven gates: +`MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 +MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 +MCORE_MOE_FUSED_SCATTER=1`. + +Two methodological points worth carrying forward. First, QWEN-018's gate was +right to reject the two candidates the ledger had ranked first and second on +arithmetic alone — the win came from the third candidate the *measurement* +surfaced, and from the router pair the gate flagged but did not rank. Second, +microbench-to-e2e transfer was reliable for launch removals (QWEN-019 predicted +2.93% and got 3.01%; QWEN-021 predicted ~3.1% and got 3.89%) and unreliable for +in-kernel rewrites (QWEN-020 and QWEN-022 both landed at roughly a third to +double their predictions), so keep repeating e2e runs whose predicted effect is +under ~1%. + +PROFILE-S6 re-ranks the step: routing is now 4.7% and the largest addressable +item is **1237 µs/step of host gaps between graph replays** (21.0% idle overall). +That is a host-side problem and needs a host-side profile before any code +change. + +Append records using this exact structure: + +```markdown +### QWEN-NNN — short name + +| Field | Value | +|---|---| +| Date | YYYY-MM-DD | +| Hypothesis | One measurable claim | +| Code revision | Commit SHA and clean/dirty state | +| Changed files | Exact paths, or `none` for baseline | +| Runtime flags | Exact non-default flags | +| Image | Immutable image path/tag | +| Checkpoint / tokenizer | Exact paths | +| Hardware / layout | Cluster, GPUs, TP/PP/EP/ETP/DP | +| Workload | Dataset, batch, OSL, warmups, timed iterations | +| Job / run | Slurm job ID and run directory | +| Throughput | tokens/s | +| Latency / TPOT | ms / ms-token | +| Correctness | Prompt outputs and benchmark status | +| Nsight artifacts | `.nsys-rep`, `.sqlite`, analysis output | +| Result | Supported / rejected / inconclusive | +| Next action | One prioritized follow-up | +``` + +## Session 7 (2026-07-26) — host-side bookkeeping + +### CLEANBASE-S7 — same-session seven-gate reference + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | Establish a same-session OSL1024 reference before measuring a sub-1% host-side lever | +| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty with the session-6 change set | +| Changed files | none | +| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | +| Image | `agents-space/images/ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | `agents-space/checkpoints/qwen3-30b-a3b-mcore` / `-hf` | +| Hardware / layout | OCI `oci-hsg`, 1 node 4×GB200, TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed | +| Job / run | job 5613090, `sessions/qwen-attnstate/e2e/ref-s7-1785084185` | +| Throughput | **25,805.9 tok/s** — per-iter 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 | +| Latency / TPOT | 9,867.5 ms / 9.920 ms/tok | +| Correctness | Coherent on all three temperature-0 prompts; benchmark 5/5 | +| Nsight artifacts | none (throughput run) | +| Result | Clean reference at **75.91%** of vLLM 33,994.5. Spread across the five iterations is 0.23%, so a 0.5% effect is resolvable in this session | +| Next action | A/B the incremental `initialize_attention_state` lever against this reference | + +### QWEN-023 — incremental `initialize_attention_state` + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | The ~311 µs/step HOSTGAP-S6 attributed to `initialize_attention_state` is recomputation of step-invariant metadata; caching it behind a request-layout version counter removes most of it without changing a single generated token | +| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty (session-6 set plus this change) | +| Changed files | `megatron/core/inference/contexts/dynamic_context.py`, `megatron/core/inference/contexts/attention_context/mha_metadata.py`, `dev/moe_fused/harness_attnstate.py` (new) | +| Runtime flags | seven gates as CLEANBASE-S7, plus `MCORE_INFER_INCR_ATTN_STATE=1` | +| Image / checkpoint / hardware / workload | identical to CLEANBASE-S7 | +| Job / run | job 5613090, `e2e/incr-attn-on-1785085393` and `e2e/incr-attn-on-rep2-*`; host microbenchmark exec `a5d20392e75848b9855366cd7844c850` | +| Throughput | **26,032.1** (per-iter 26,052.6 / 26,018.7 / 26,075.7 / 25,992.1 / 26,021.4) and **26,092.5** on an independent repeat pair | +| Latency / TPOT | 9,780.4 ms / 9.834 ms/tok; repeat 9.811 ms/tok (−92 µs/step vs reference on two-run means) | +| Correctness | 130 consecutive decode steps bit-identical under `MCORE_INFER_INCR_ATTN_STATE_VERIFY=1`; temperature-0 coherence output byte-identical to gate-OFF | +| Nsight artifacts | none needed — the effect is host wall time, measured directly | +| Result | **Accepted, +0.94%** (two OFF/ON pairs, non-overlapping per-iteration distributions) | +| Next action | Cut `update_requests` + `active_request_mask` + the unlabelled engine window (~437 µs/step, HOSTGAP-S6 lever 2) with the same version-counter machinery | + +**Pre-optimization CPU breakdown.** Measured with `MCORE_INFER_ATTN_PROF=1` +(monotonic-ns marks around each phase) and with a standalone host harness +(`dev/moe_fused/harness_attnstate.py`) that drives a real +`DynamicInferenceContext` at BS256 through 400 steady-state decode steps with no +model loaded. The two agree to within 12%. Per call, 256 active requests: + +| Phase | µs/call | What it is | +|---|---:|---| +| `slices` | 73.9 | build + pad the active-request slices, sampling-metadata copies, logit indices | +| `mhameta` | 32.2 | KV/query lengths, both cumsums, block table, `set_state_data` | +| `xfer_inner` | 26.9 | H2D copy of the bookkeeping buffer | +| `graphmatch` | 12.8 | batch dimensions + `match_graph_config` | +| `tokenpad` | 12.0 | stamp the token padding slots | +| `maxlen` | 12.0 | max-seqlen scalars | +| `pre` + `paddims` + `tail` | 1.8 | pending Mamba ops, padded dimensions, epilogue | +| **TOTAL** | **171.5** | | + +The cost is **concentrated, not thin**: `slices` alone is 43%, and inside it two +statements — `build_active_slices` 29.3 µs and `pad_active_slices` 26.7 µs — +are a third of the whole call. Every one of those top statements recomputes a +value that is provably identical to the previous decode step. That is what made +the lever viable; had the profile come back flat across fifty 3 µs statements +the design below would not have been worth building. + +**Design.** A single monotonic `_request_layout_version` counter is bumped by +every path that changes *which* request occupies a slot, a request's sampling +metadata, or its KV block table. `initialize_attention_state` takes the fast +path only when the current cache key equals the stored one, where the key is +`(layout_version, total_request_count, paused_request_count, active_token_count, +kv_block_allocator.total_avail, chunked_prefill_request_id, +num_speculative_tokens)`. The version is the guard; the other six are +independent structural sentinels, so a missed bump would additionally have to +coincide with an unchanged request count, token count and block-allocator +occupancy to escape detection. The fast path then recomputes only the KV +sequence lengths and their cumsum — statements copied character-for-character +from the full path so the buffers stay bit-identical — re-stamps the token +padding sentinels, rebinds `state_data` through a new +`MHAMetadata.restore_state_data` (needed because `reset_attention_state` clears +the max-seqlen scalars every step), and issues the H2D transfer. Everything +else is reused. + +**Invalidation audit.** Thirteen mutation paths, each bumping the version: + +| Path | Event covered | +|---|---| +| `add_request` | request add, including a chunked-prefill chunk | +| `resume_paused_requests` (inside `if resume_request_count > 0`) | resume | +| `_swap_book_keeping_tensors` | pause/resume slot swap | +| `_move_book_keeping_tensors` | slot compaction | +| `resolve_requests` | end-of-step finish/retire | +| `release_memory_blocks_from_request_indexes` | evict, block release | +| `prepare_requests` (inside `if num_new_blocks > 0`) | new KV block allocated mid-decode | +| `add_dummy_requests_parallel` (inside `if requests`) | dummy padding requests | +| `add_dummy_requests_for_cudagraph_capture` | graph capture padding | +| `add_dummy_requests_for_expert_parallel_step` | EP dummy step | +| `initialize_all_tensors` | tensor state (re)allocation | +| `reset_tensors`, `reset_metadata` | context reset | + +Advancing a request's KV length is deliberately *not* a bump — that is the one +thing the fast path recomputes. Four conditions bail out to the full path +outright rather than relying on the key: `construct_graph_dimensions is not +None` (any graph-dimension change), `is_expert_parallel_dummy_cuda_graph_step`, +`is_hybrid_model` (Mamba state), and `num_prefill_requests != 0`. The cache is +additionally never *stored* while `is_creating_cuda_graphs` is set or when the +step did not use a CUDA graph. The class-scope declaration of every cache +attribute means any construction path starts cold and invalid. + +**Correctness evidence.** `MCORE_INFER_INCR_ATTN_STATE_VERIFY=1` runs the fast +path, snapshots everything it produced, then recomputes the whole thing from +scratch and asserts equality of: both MHA query-length buffers and cumsums, both +KV-length buffers and cumsums, the block table, `active_request_last_token_idxs`, +`active_logit_idxs`, all seven `active_request_metadata` tensors, the entire GPU +bookkeeping buffer, the padded token/request counts, the padded batch +dimensions, the graph-selection flag, both max-seqlen scalars, and the +`(address, shape, dtype)` binding of every `state_data` view. **130 consecutive +decode steps passed with zero mismatches.** End to end, the three temperature-0 +coherence completions are byte-identical between gate ON and gate OFF. + +**Performance.** Host CPU per call **174.1 → 51.7 µs, a 3.37× reduction**; the +coarse breakdown collapses to `graphmatch=0.2 slices=1.1 tokenpad=0.1 +mhameta=0.4 maxlen=0.1`, with the 48.2 µs residual being the KV-length +recompute plus the H2D bookkeeping copy, which cannot be cached. E2e: + +| Run | Gate | Throughput | Per-iter | +|---|---|---:|---| +| `ref-s7` | OFF | 25,805.9 | 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 | +| `incr-attn-on` | ON | **26,032.1** | 26,052.6 / 26,018.7 / 26,075.7 / 25,992.1 / 26,021.4 | +| `ref-s7-rep2` | OFF | 25,831.7 | 25,710.8 / 25,893.9 / 25,863.9 / 25,875.4 / 25,815.4 | +| `incr-attn-on-rep2` | ON | **26,092.5** | 26,157.8 / 26,085.1 / 26,096.3 / 26,106.8 / 26,017.0 | + +Pairwise the lever is +0.88% and +1.01%; on the two-run means (OFF 25,818.8, ON +26,062.3) it is **+0.94%**. Across all twenty timed iterations the slowest ON +iteration (25,992.1) is still faster than the fastest OFF iteration (25,893.9), +so the two distributions do not overlap at all and the sign is certain despite +the effect being under 1%. TPOT 9.915 → 9.823 ms/tok, i.e. **−92 µs/step**. +New best = **76.67%** of vLLM 33,994.5. + +**Why it delivered 0.94% and not the 3.4% ceiling.** The harness measures one +call at 174.1 µs and the fast path removes 122.4 µs of it, but the per-step +saving observed in TPOT is 92 µs — about three quarters of a single call's +worth, against HOSTGAP-S6's 311 µs/step spread over two windows. So the ceiling +was overstated for two compounding reasons: the H2D bookkeeping transfer is +~27 µs of every call and is not cacheable (the GPU buffer must be rewritten +every step), and the second of the two per-step windows is evidently not all +`initialize_attention_state` doing cacheable work. The honest read is that the +remaining host-side upside lives in `update_requests` and the engine window, +not here. + +### CGTRACE-CONTROL — is the graph machinery real? + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | The ~900 µs/step HOSTGAP-S6 attributed to CUDA-graph machinery is an artifact of `--cuda-graph-trace=node`; capturing with `=graph` will shrink it | +| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty | +| Changed files | `dev/moe_fused/profile_insession.sh` (`CUDA_GRAPH_TRACE` env, bounded nsys stop + qdstrm recovery), `dev/moe_fused/analyze_cgtrace.py` (new) | +| Runtime flags | seven gates plus `MCORE_INFER_INCR_ATTN_STATE=1`, identical in both captures | +| Workload | gsm8k, BS256, OSL128, 1 warmup request + 1 timed request | +| Job / run | job 5613090, `prof/cgt-node-1785086713`, `prof/cgt-graph-1785086964` | +| Nsight artifacts | `mcore_profile.{nsys-rep,sqlite}` in both run dirs; analysis exec `d6f572f2613645cdbb05226dbc8e5e98` | +| Result | **Hypothesis rejected — the machinery is real** | +| Next action | Treat HOSTGAP-S6's ranking as sound; the next lever is `update_requests` (lever 2), not graph-node reduction (lever 4), on measured size | + +| Metric (same rank, steady decode window) | `=node` | `=graph` | Δ | +|---|---:|---:|---:| +| Step period | 8882.5 µs | 8945.1 µs | −0.71% | +| Host `cudaGraphLaunch` (median) | 190.0 µs | 184.7 µs | +2.8% | +| Traced kernels per step | 1169 | 54 (graph interior hidden) | — | +| Steady steps analyzed | 91 | 88 | — | + +The step period is the number that matters, and it does not move: if CUPTI's +per-node instrumentation were adding ~900 µs to a 9 ms step, removing it would +have shortened the step by 10%, and instead the graph-mode capture came out +0.71% *slower*, which is ordinary run-to-run variance. The 190 µs +`cudaGraphLaunch` survives the control almost unchanged, so it is genuine +driver-side submit cost for a 1158-node graph. Every wall-time attribution in +PROFILE-S6 and HOSTGAP-S6 therefore stands as measured. + +One limit worth recording: `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` in the graph-mode +capture holds only 417 rows for the whole run (6 inside the analyzed window), +so the GPU-side "inter-node dispatch" half of the 900 µs cannot be +independently confirmed from this control. The step period and the host submit +cost are settled; that component is not. + +### CLEANBASE-S8 — same-session eight-gate reference + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | Establish a same-session OSL1024 reference before measuring the second host-side lever | +| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty with the session-7 change set | +| Changed files | none | +| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1 MCORE_INFER_INCR_ATTN_STATE=1` | +| Image | `agents-space/images/ceecf5c304a5d8bd.sqsh` | +| Checkpoint / tokenizer | `agents-space/checkpoints/qwen3-30b-a3b-mcore` / `-hf` | +| Hardware / layout | OCI `oci-hsg`, 1 node 4×GB200 (`nvl72039-T01`), TP1/PP1/EP4/ETP1 | +| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed | +| Job / run | job 5616264, `sessions/qwen-updreq/e2e/ref-s8-1785099726` (and `ref-s8-rep2-1785101268`) | +| Throughput | **25,944.9 tok/s** — per-iter 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7. Repeat: **25,983.9** — per-iter 26,021.5 / 25,976.0 / 26,037.9 / 25,983.6 / 25,901.2 | +| Latency / TPOT | 9,820.6 ms / 9.867 ms/tok; repeat 9.852 ms/tok | +| Correctness | Coherent on all three temperature-0 prompts; benchmark 5/5 | +| Nsight artifacts | none (throughput run) | +| Result | Clean reference at **76.32%** of vLLM 33,994.5 (two-run mean 25,964.4 = 76.38%) | +| Next action | A/B the reduced-op `update_requests` lever against this reference | + +The first run's iteration 1 is a 1.3% cold outlier; iterations 2–5 span 0.23%. +The repeat run does not show it and spans 0.53% across all five. Two-run mean +25,964.4 tok/s is the number the lever below is judged against. + +### QWEN-024 — reduced-op post-sampling bookkeeping + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | HOSTGAP-S6 lever 2. The ~437 µs/step of host work after the sampled tokens return is three different things, not one: measure `update_requests`, `active_request_mask` and the unlabelled engine window separately, name the window, and only then decide whether any of it is attackable | +| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty (session-7 set plus this change) | +| Changed files | `megatron/core/inference/contexts/dynamic_context.py` (`_write_decode_token_bookkeeping_fast`, `_write_token_bookkeeping_reference`, `_verify_decode_token_bookkeeping`, `_decode_req_idx_arange` cache, two count reductions, env `MCORE_INFER_VEC_UPDATE_REQS`, `MCORE_INFER_VEC_UPDATE_REQS_VERIFY`), `megatron/core/inference/text_generation_controllers/text_generation_controller.py` (`_empty_finished_idxs`, `_empty_finished_ids`, no-finisher short-circuit), `dev/moe_fused/harness_updreq.py` (new) | +| Runtime flags | eight gates as CLEANBASE-S8, plus `MCORE_INFER_VEC_UPDATE_REQS=1` | +| Image / checkpoint / hardware / workload | identical to CLEANBASE-S8 | +| Job / run | job 5616264, `e2e/vecupd-on-1785100576` and `e2e/vecupd-on-rep2-*`; harness execs `eaca22426f654bfb938bd68c13c7bff8` (gate OFF breakdown), `94604a2f79824182bef77456c9a2b7ee` (equivalence, verify, gate ON) | +| Throughput | **26,128.2** and **26,270.3 tok/s** on two independent OFF/ON pairs | +| Latency / TPOT | 9,745.4 ms / 9.798 ms/tok; repeat 9,691.7 ms / 9.745 ms/tok (−88 µs/step vs the reference on two-run means) | +| Correctness | 400-step two-context equivalence with 30 mid-batch terminations and a KV block-boundary crossing; 300 consecutive steps under `MCORE_INFER_VEC_UPDATE_REQS_VERIFY=1`; temperature-0 coherence output byte-identical to gate-OFF | +| Nsight artifacts | OSL128 A/B pair, `prof/g1-off-1785102888` and `prof/g1-on-1785103157`, `mcore_profile.{nsys-rep,sqlite}`; analysis exec `20bc3a28e2344c9b8f51aae8c90f6982`; window finder `dev/moe_fused/find_decode_window.py` (new) | +| Result | **Accepted, +0.90%** (two OFF/ON pairs: +0.71% and +1.10%), and the G1 window shrank by the predicted amount | +| Next action | The remaining host-side upside is `post_process_requests`, which this record shows is not of the same kind; see the ranked next levers below | + +**Pre-optimization CPU breakdown — and what the unlabelled window actually is.** +Measured with `dev/moe_fused/harness_updreq.py`, which drives a real +`DynamicInferenceContext` at BS256 through a synthetic steady-state decode loop +and calls the real `DynamicInferenceEngine.post_process_requests` against a stub +engine holding 256 real `DynamicInferenceRequest` objects. No model is loaded, +so a full breakdown takes about 90 seconds. Per step, 256 active requests: + +| Region | µs/step | HOSTGAP-S6 | Nature | +|---|---:|---:|---| +| `active_request_mask` | 46.9 | 78.0 | ~11 whole-tensor CPU ops on 256 elements | +| `update_requests` | 137.6 | 183.6 | ~22 whole-tensor CPU ops on 256 elements | +| `post_process_requests` | 187.9 | ~152.2 (unlabelled) | per-request Python object churn | +| **TOTAL** | **372.4** | **~413.8** | | + +The harness runs 10–20% cheaper than the in-server measurement across all three +regions, which is the expected direction (no server threads, no allocator +pressure, no profiler), and the three ratios agree to within 15% of each other. + +**The unlabelled engine window is `DynamicInferenceEngine.post_process_requests`, +reached through `async_bookkeep`.** It appeared unlabelled in HOSTGAP-S6 because +the controller emits NVTX through `torch.cuda.nvtx.range_push` directly, which is +always live, while the engine emits through `megatron.core.utils.nvtx_range_push`, +which is inert unless `_nvtx_enabled` is set — so the engine's own `bookkeeping` +and `detokenization` ranges were simply never pushed in that capture. The +measured 187.9 µs/step also matches the window's 152.2 µs median once the +harness discount is applied. Its `cProfile` composition, over 100 steps: + +| Item | calls/step | What it is | +|---|---:|---| +| `post_process_requests` body | 1 | 61% of total time is straight-line interpreter work in the per-request loop | +| `builtins.len` | 1,718.8 | ~6.7 length probes per request per step | +| `get_request` | 256 | `self.requests[id].record[-1]` dict + list indexing | +| `list.append` | 256 | `active_request_ids.append` | +| `builtins.isinstance` | 256 | the `if not isinstance(tokens, list)` scalar wrap | +| `_check_stop_words_for_request_post_append` | 256 | returns immediately (no stop words configured) | +| `DynamicInferenceRequestRecord.__getitem__` | 256 | record `[-1]` | +| tensor ops | 3 | three `.tolist()` calls for the whole step | + +**The falsification criterion, honored.** The pre-agreed rule was: vectorize if +the cost is a uniform per-request loop expressible as tensor ops, and stop if it +is per-request Python object churn. The measurement splits the chain cleanly in +half and the two halves fall on opposite sides of that line. + +`post_process_requests` is **churn, and it was left alone**. It is 187.9 µs/step +— the single largest item in the chain — and it does essentially no tensor work: +three `.tolist()` calls for the entire step, against 256 dict lookups, 256 list +appends, 256 method calls and ~1,719 `len()` probes. Nothing here vectorizes. +Making it cheaper means adding a fast path through the request state machine — +token appending, termination, stop words, log-prob routing — which is exactly the +EOS/length handling the task flagged as dangerous, for a benefit that would have +to be re-argued from scratch. It is not touched by this change. + +The other half is **neither** of the two anticipated cases. `update_requests` and +the `active_request_mask` block contain no per-request Python loop at all; they +are already vectorized. Their cost is ~33 whole-tensor operations on 256-element +CPU tensors at 2–16 µs each, and roughly half of that work is *provably dead* in +the decode regime. Three statements in `update_requests` cost 30.0 µs/step +between them and, at `num_speculative_tokens == 0`, compute nothing: + +| Statement | µs/step | Why it is dead at one generated token per request | +|---|---:|---| +| `token_to_pos_ids` write | 16.0 | `repeat_interleave(1)` is identity, and it adds `torch.arange(1).repeat(256)` — a 256-element zero vector | +| `raw_positions` + `crosses_boundary` + `.any()` | 10.1 | builds a `[256, 1]` tensor whose only consumer is an `else` branch that `num_speculative_tokens == 0` makes unreachable | +| `token_to_request_idx` write | 7.3 | `repeat_interleave(1)` is identity over a `torch.arange` that is a pure function of the request-slot bounds | + +So the lever here is **op elimination**, not vectorization. That is a third +category the dichotomy did not anticipate, and it is the low-risk one: it changes +no per-request semantics, touches no termination logic, and every removed +operation is removed because its result is bit-identical to a cheaper one. + +**Design.** One env gate, `MCORE_INFER_VEC_UPDATE_REQS`, default OFF. Note that +the mechanism is deliberately unlike QWEN-023: nothing is cached across steps +except a `torch.arange` that depends only on the request-slot bounds, because +`update_requests` consumes the current step's freshly sampled tokens and there is +no cross-step reuse to exploit. Five changes: + +1. The per-token bookkeeping tail of `update_requests` is factored into + `_write_token_bookkeeping_reference`, kept verbatim, and a new + `_write_decode_token_bookkeeping_fast` used only when + `num_speculative_tokens == 0`. The fast form drops the three dead statements + above, updates `request_last_kv_block_offset` without the defensive clone + (whose only other consumer was `raw_positions`), and reuses a cached + `torch.arange(paused_request_count, total_request_count)` invalidated by a + bounds compare. +2. `finished_request_count` is derived as `numel() - active_request_count` + instead of a second full comparison and reduction. The mask is 0/1 by + construction — it is an `&` of two `.byte()` comparisons in the controller — + so the complement is exact. +3. `active_requests_requiring_new_block` stays a bool tensor instead of being + cast with `.byte()`. Every downstream use (`torch.nonzero`, `== 0`, scalar + assignment, `sum`) is dtype-agnostic. +4. Its population count uses `int(tensor.sum())` rather than `(tensor == 1).sum().item()`. +5. In the controller, when the mask says nothing finished this step, the + `torch.nonzero` scan and the advanced-index gather it feeds are skipped in + favour of cached empty tensors of the same dtype and device. + +**Correctness evidence.** Three independent checks, all on the eight-gate config. + +*Two-context equivalence, 400 steps.* Two `DynamicInferenceContext` objects are +built identically and driven through the same scripted mask sequence, one with +the gate off and one on, comparing after every step: `total_request_count`, +`paused_request_count`, `active_token_count`, nine request bookkeeping tensors +(`request_ids`, `request_kv_length_offsets`, `request_query_lengths`, +`request_output_lengths`, `request_last_kv_block_offset`, +`request_last_kv_block_id`, `request_kv_block_counts`, +`request_to_kv_block_ids`, `request_in_prefill_status_tensor`), all six +`token_to_*` tensors, and the KV block allocator's `total_avail`, +`active_count` and `paused_count`. The script **terminates three requests +mid-batch every 37 steps** — 30 terminations in total, leaving 226 of 256 active +— and runs past the 256-token block boundary so the pause / new-block / resume +branches all execute under both gates. Zero mismatches. This is the check that +directly answers the off-by-one-in-termination risk: the finished-request count, +the finished index set, and the resulting slot compaction are all compared +element-wise on the steps where requests actually finish. + +*Verify mode, 300 steps.* `MCORE_INFER_VEC_UPDATE_REQS_VERIFY=1` runs the fast +tail, snapshots the six token buffers plus `request_last_kv_block_offset` and +`active_token_count`, restores the inputs, runs the reference tail, and asserts +equality. 300 consecutive steps, zero mismatches. + +*End to end.* The three temperature-0 coherence completions are byte-identical +between gate ON and gate OFF. + +**Performance.** Host CPU per step, from the harness at BS256: + +| Region | Gate OFF | Gate ON | Δ | +|---|---:|---:|---:| +| `active_request_mask` | 46.9 | 34.9 | −12.0 | +| `update_requests` | 137.6 | 74.0 | −63.6 (1.86×) | +| `post_process_requests` | 187.9 | 184.7 | unchanged, as intended | +| **chain total** | **372.4** | **293.7** | **−78.7** | + +E2e: + +| Run | Gate | Throughput | Per-iter | +|---|---|---:|---| +| `ref-s8` | OFF | 25,944.9 | 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7 | +| `vecupd-on` | ON | **26,128.2** | 25,856.1 / 26,141.3 / 26,233.9 / 26,141.6 / 26,272.3 | +| `ref-s8-rep2` | OFF | 25,983.9 | 26,021.5 / 25,976.0 / 26,037.9 / 25,983.6 / 25,901.2 | +| `vecupd-on-rep2` | ON | **26,270.3** | 26,291.7 / 26,281.1 / 26,279.5 / 26,302.1 / 26,197.4 | + +Pairwise the lever is **+0.71%** and **+1.10%**; on the two-run means (OFF +25,964.4, ON 26,199.3) it is **+0.90%**. Nine of the ten ON iterations are +faster than every one of the ten OFF iterations; the single exception is the +cold first iteration of the first ON run (25,856.1), and both OFF runs show the +same cold-first-iteration pattern. Restricted to iterations 2–5 the two +distributions do not overlap at all (slowest ON 26,141.3 > fastest OFF +26,041.7). TPOT 9.860 → 9.771 ms/tok on two-run means, i.e. **−88 µs/step**. +New best = **77.28%** of vLLM 33,994.5 on the best run, **77.07%** on the +two-run mean. + +**Why it delivered less than the 4.8% ceiling.** The ceiling assumed the whole +437 µs/step chain could vanish. Half of it is `post_process_requests`, which this +change deliberately does not touch, so the reachable half was ~185 µs/step of +which the harness says 78.7 µs was removed — 42% of the reachable half, 18% of +the whole chain. TPOT moved 9.860 → 9.771 ms/tok on two-run means, i.e. +−88 µs/step, against a harness-measured saving of 78.7 µs. The two agree to +within 12%, in the direction the harness discount predicts, which is the +strongest evidence that the host saving is what produced the throughput change +rather than session drift. + +**Profile confirmation — HOSTGAP-S6's falsification test for this lever class.** +HOSTGAP-S6 said a lever of this kind is falsified if a fresh profile shows the G1 +window unchanged. Two BS256/OSL128 captures were taken back to back in the same +session at the same eight-gate code, differing only in the gate, and analysed +with `analyze_hostgaps.py gaps` on the forced LM-head anchor +`nvjet_sm100_tst_512x64_64x3_2x1_2cta_v_bz_TNT`. The G1 bracket is now +`index_elementwise_kernel -> vectorized_gather_kernel` rather than +`-> vectorized_elementwise_kernel`, because QWEN-023 changed which kernel the +bookkeeping chain emits first; it is the same gap class (largest gap, ~1.0 +instances per step). + +| Metric (mean of devices 0 and 3) | Gate OFF | Gate ON | Δ | +|---|---:|---:|---:| +| G1 window | 640.6 µs | 562.3 µs | **−78.3 µs** | +| Step period | 8.9425 ms | 8.8530 ms | −89.5 µs (−1.00%) | +| Total idle | 1757.0 µs | 1679.2 µs | −77.8 µs | +| GPU-busy | 7.2015 ms | 7.2045 ms | +3.0 µs (noise) | +| Kernels per step | 1169 | 1169 | 0 | +| G3 (`CatArrayBatchedCopy` → `rmsnorm`) | 295.6 / 293.0 | 296.8 / 293.2 | unchanged | +| G2A (`vectorized_elementwise` → `vectorized_gather`) | 67.1 / 60.4 | 65.9 / 63.6 | unchanged | + +Four independent measurements agree: the harness says 78.7 µs of host work was +removed, the G1 window shrank 78.3 µs, the step period shortened 89.5 µs and +TPOT shortened 88 µs. GPU-busy time and the per-step kernel count are identical, +and the two gaps this change does not touch did not move, so the saving is +host-side and localized exactly where it was designed to be. Not falsified. + +**Ranked next levers, updated.** (Lever 1 below was then implemented in the same +session and accepted as QWEN-025; the decision-gate measurement that justified +starting it is recorded here.) + +1. **`post_process_requests` fast path** — 187.9 µs/step harness, ~152 µs/step + in-server, now the largest single host item in the decode loop. A fast path + for the common decode case (no speculative tokens, no stop words, no log + probs, not chunked prefill, request not finishing) could collapse the loop + body to a few operations. Ceiling ~1.7% e2e. **Risk high** — it is the request + termination state machine. Falsified if a `cProfile` of the fast path does not + cut the per-step call count below ~500. + + **Decision-gate measurement, taken with the time left in this session** + (harness exec `4c0a584d22704d4ab1d8c5a7223bc8c2`, gate ON, BS256, 200 + timed calls against the real engine method and against an in-harness + prototype of the decode fast path). The prototype implements the common + decode case only and removes the per-request dict lookup, the record `[-1]` + indexing, the `isinstance` scalar wrap, the `_check_stop_words_...` call and + the repeated `len()` probes, resolving `(request, token_limit)` pairs once + per request-id set instead of once per step: + + | Path | µs/step | + |---|---:| + | `post_process_requests`, as shipped | 213.8 | + | decode fast-path prototype | 33.8 | + | **reducible** | **180.0 (84%)** | + + So the lever clears its own gate decisively: 180 µs/step against an 8.85 ms + step period is **~2.0% e2e**, roughly twice what QWEN-024 delivered, and it + is the largest remaining host item. The mechanism is neither caching + (QWEN-023) nor op elimination (QWEN-024) but a third thing — collapsing a + per-request Python loop body — so it needs its own correctness argument, and + the termination path is exactly where that argument is hard. Note the + prototype is a *measurement*, not an implementation: it does not handle stop + words, log probs, chunked prefill, speculative tokens or eviction, and those + must fall back to the existing loop. +2. **Reduce graph node count** (HOSTGAP-S6 lever 4, now unblocked by + CGTRACE-CONTROL) — ~0.78 µs per removed node, so 100 nodes ≈ 0.85%. + **Risk medium**, sharply diminishing. +3. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — + ceiling ~9.2%, **risk high**, multi-session. + +### QWEN-025 — decode fast path for `post_process_requests` + +| Field | Value | +|---|---| +| Date | 2026-07-26 | +| Hypothesis | The decision-gate measurement in QWEN-024 says 84% of `post_process_requests` is reducible without any tensor work, by collapsing the per-request loop body. If that translates, it is worth ~2.0% e2e — twice QWEN-024 | +| Code revision | branch `perf/moe-fused-align`, dirty (session-7 set, QWEN-024, and this change) | +| Changed files | `megatron/core/inference/engines/dynamic_engine.py` (`_post_process_requests_decode_fast`, `_ppr_cache_epoch` invalidation at the three request-record mutation sites, env `MCORE_INFER_FAST_POST_PROCESS`, `..._VERIFY`), `dev/moe_fused/harness_updreq.py` (`run_ppr_equivalence`, in-process gate A/B) | +| Runtime flags | eight gates plus `MCORE_INFER_VEC_UPDATE_REQS=1`, plus `MCORE_INFER_FAST_POST_PROCESS=1` | +| Image / checkpoint / hardware / workload | identical to CLEANBASE-S8 | +| Job / run | job 5616264, e2e `ppr-off-1785104484` / `ppr-on-1785104751` / `ppr-off-r2-1785105441` / `ppr-on-r2-1785105667` / `ppr-off-r3-1785105923` / `ppr-on-r3-1785106159`; harness execs `e620b4429e3945e6a4b0ab61d731c372` (equivalence), `81ab8e8ee09840eebf825420633b2efc` (gate A/B + verify) | +| Throughput | **26,361.7 tok/s** on the three-pair ON mean; best single run **26,430.7** | +| Latency / TPOT | 9,658.7 ms / 9.7111 ms/tok on the ON mean, against 9.8060 OFF — **−94.9 µs/step** | +| Correctness | 400-step two-engine equivalence including 10 finisher steps and 200 steps past the token limit; 300 steps under `MCORE_INFER_FAST_POST_PROCESS_VERIFY=1`; temperature-0 coherence byte-identical | +| Nsight artifacts | OSL128 A/B pair `prof/g2-off-1785106875` and `prof/g2-on-1785107105`; analysis exec `0ec0392b8c1d47c5a099cee585d777a4`. Two follow-on host-visibility captures (`prof/hosts9-*`, `prof/hosts9b-*`) failed to import; see the note under the next-lever list | +| Result | **Accepted, +0.98%** on three-pair means (+0.59% on a best-4-of-5 trim), ON wins all three pairs | +| Next action | Host-side decode work is now 148 µs/step against 372 µs before QWEN-024; the next lever is graph node count, not the host chain | + +**Why this was started despite QWEN-024's stop verdict.** QWEN-024 declined to +touch `post_process_requests` because it is per-request Python object churn and +therefore not *vectorizable* — that was the pre-agreed falsification criterion +and it still holds. What the QWEN-024 decision-gate measurement then showed is +that being unvectorizable is not the same as being irreducible: an in-harness +prototype of the plain decode case ran at 33.8 µs/step against the real +method's 213.8 µs, i.e. 84% of the cost is dict lookups, record indexing, an +`isinstance` scalar wrap, a stop-word call that returns immediately, and ~1,719 +`len()` probes, none of which the plain decode case needs. That is a third +mechanism, distinct from QWEN-023's caching and QWEN-024's op elimination: +collapsing a loop body. + +**Design — the safety property comes first.** One env gate, +`MCORE_INFER_FAST_POST_PROCESS`, default OFF, dispatching to +`_post_process_requests_decode_fast` at the top of `post_process_requests`. The +fast path returns `None` and the full reference loop runs whenever anything is +not the plain decode case. The declining conditions are: speculative decoding +or accepted tokens, log probs, top-n log probs, finished routing block ids, +token event tracking, any pending stop-word state, chunked prefill in flight, a +TPOT sample due this step (`step_time > 0`), any eviction, **or any request +finishing this step**. + +That last condition is the design's whole safety argument. The task flagged an +off-by-one in EOS/length handling as the dangerous failure mode, and the honest +way to remove that risk is not to test it harder but to never enter it: the +fast path is structurally unreachable on any step where a request finishes, so +the termination state machine — pop, future resolution, routing +reconstruction, status, finish event — is only ever executed by the original, +well-tested code. At BS256 steady-state decode essentially every step has no +finisher, so nearly all the work is still avoided; the OSL1024 benchmark takes +the fast path on all but a handful of its ~1,024 steps. + +What remains in the fast path is one bounded append per request: + +```python +if num_generated < limits[i]: + generated.append(token) +``` + +which is the exact one-token specialization of the reference trim — the +reference computes `keep = num_tokens_to_generate - len(generated_tokens)` and +slices `tokens[:keep]`, which for a single token keeps it iff the request is +below its limit and drops it entirely otherwise. The first-token TTFT sample is +preserved inline on the `num_generated == 0` branch. + +The `(request, token_limit)` pairs are resolved once per active request-id set +rather than once per step, keyed on `(epoch, tuple(request_ids))`. The epoch is +bumped at all three sites that can change which object `record[-1]` resolves +to — the two `record.checkpoint()` calls (recompute-suspend and eviction) and +the `RequestEntry` insertion in `_add_request` — so a stale resolution is not +representable rather than merely unlikely. A configured stop word on any +request also declines, since the post-append scan is the one piece of +per-request work here that is not a bounded append. + +**Correctness evidence.** Three independent checks. + +*Two-engine equivalence, 400 steps.* Two stub engines are built with identical +request sets and driven through the same token stream, one gate OFF and one ON, +compared after every step on the full observable result: returned active id +list, finished record count and contents, `finished_request_count`, the live +request-id set, and per request `generated_tokens`, `generated_length`, `status` +and whether TTFT is set. Three requests are finished mid-batch every 37 steps +(10 finisher steps, 30 requests finished, all 10 correctly declined by the fast +path), and `num_tokens_to_generate` is set to 200 against 400 steps so every +surviving request spends 200 consecutive steps *at* its token limit — the +append-suppression branch — and both gates agree that it holds exactly 200 +tokens at the end. Zero mismatches. This is the check that answers the +off-by-one risk directly, on both sides of the limit and on the finish steps. + +*Verify mode, 300 steps.* `MCORE_INFER_FAST_POST_PROCESS_VERIFY=1` independently +recomputes the expected post-state of every request from the pre-state and the +sampled token before running the fast loop, then asserts equality. 300 +consecutive steps, zero mismatches. + +*End to end.* The three temperature-0 coherence completions are byte-identical +between gate ON and gate OFF, character for character across all three prompts. + +**Performance.** Host CPU per step at BS256, both arms measured in the same +process against the same engine so the comparison carries no process-to-process +noise: + +| Region | Gate OFF | Gate ON | Δ | +|---|---:|---:|---:| +| `post_process_requests` | 211.0 | 27.2 | **−183.9 (7.77×)** | + +And the whole post-sampling chain, against the pre-QWEN-024 starting point: + +| Region | Before QWEN-024 | After QWEN-024 | After QWEN-025 | +|---|---:|---:|---:| +| `active_request_mask` | 46.9 | 34.9 | 34.1 | +| `update_requests` | 137.6 | 74.0 | 73.7 | +| `post_process_requests` | 187.9 | 184.7 | **40.3** | +| **chain total** | **372.4** | **293.7** | **148.0** | + +The pre-registered falsification test for this lever was "falsified if a +`cProfile` of the fast path does not cut the per-step call count below ~500." +Measured: **261 calls/step** (7,831 over 30 steps), of which 256 are the single +`len()` per request, against ~1,719 `len()` probes alone before. Not falsified. + +E2e, three OFF/ON pairs run back to back in one session: + +| Pair | Gate | Throughput | Per-iter | +|---|---|---:|---| +| 1 | OFF | 26,054.0 | 26,251.6 / 26,249.3 / 26,241.1 / 26,274.0 / 25,283.4 | +| 1 | ON | **26,313.1** | 26,104.8 / 26,306.3 / 26,431.3 / 26,418.0 / 26,307.9 | +| 2 | OFF | 26,101.0 | 26,183.2 / 26,218.8 / 26,253.2 / 26,195.0 / 25,664.3 | +| 2 | ON | **26,430.7** | 26,442.2 / 26,416.0 / 26,428.8 / 26,426.4 / 26,440.3 | +| 3 | OFF | 26,164.4 | 25,883.8 / 26,270.7 / 26,270.5 / 26,228.0 / 26,173.0 | +| 3 | ON | **26,341.2** | 26,404.3 / 26,333.4 / 26,243.2 / 26,376.0 / 26,349.9 | + +Pairwise **+0.99%, +1.26%, +0.68%**; on three-pair means (OFF 26,106.5, ON +26,361.7) **+0.98%**. The task set +0.5% as the noise floor, so the headline +number clears it, but each OFF run contains one slow outlier iteration (25,283.4 +and 25,664.3 in the fifth position, 25,883.8 in the first) which flatters the +mean comparison. Dropping the slowest iteration from every run gives OFF +26,234.0 against ON 26,388.7, **+0.59%** — still above the floor, and the +honest lower bound. + +The distribution-level evidence is stronger than either mean. Pooling all 15 OFF +and 15 ON iterations, **13 of the 15 ON iterations are faster than every one of +the 15 OFF iterations** (the two exceptions are 26,104.8, the cold first +iteration of pair 1, and 26,243.2). ON also has visibly lower spread — pair 2's +five ON iterations span 0.10% — because the removed work included the variance, +not just the mean. ON wins all three pairs. +New best = **77.75%** of vLLM 33,994.5 on the best run, **77.55%** on the +three-pair mean. + +**Profile confirmation, and why only part of the host saving converts.** Two +BS256/OSL128 captures back to back, same code, differing only in the gate, +analysed with `analyze_hostgaps.py gaps` on the forced LM-head anchor, mean of +devices 0 and 3: + +| Metric | Gate OFF | Gate ON | Δ | +|---|---:|---:|---:| +| G1 (`index_elementwise` → `vectorized_gather`) | 555.5 µs | 499.4 µs | **−56.1 µs** | +| Total idle | 1717.5 µs | 1639.6 µs | −77.9 µs | +| Step period (median) | 8.8645 ms | 8.8050 ms | −59.5 µs | +| GPU-busy | 7.202 ms | 7.219 ms | +17 µs (noise) | +| Kernels per step | 1169 | 1169 | 0 | +| G3 (`CatArrayBatchedCopy` → `rmsnorm`) | 302.1 | 294.9 | −7.2 | +| G2A (`vectorized_elementwise` → `vectorized_gather`) | 65.8 | 61.5 | −4.3 | + +The saving is host-side and localized: the kernel count is bit-identical, GPU-busy +does not move, and the idle reduction is concentrated in G1. + +But note the conversion ratio, which is the interesting result and is *worse* +than QWEN-024's. The harness says 183.9 µs of host work was removed; G1 shrank +56.1 µs, total idle 77.9 µs, the profiled step period 59.5 µs, and e2e TPOT +94.9 µs. QWEN-024 removed 78.7 µs of host work and got 78.3 µs out of G1 — very +nearly 1:1 — because `update_requests` sits directly in the serial dependency +between the sampled tokens returning and the next step's launches. Only about a +third to a half of this change converts, which says most of +`post_process_requests` was already partly overlapped with GPU execution: it +runs from `async_bookkeep`, so the interpreter was working through it while the +device still had queued work, and only the exposed tail was ever on the critical +path. **The corollary for future levers: harness-measured host CPU savings are +an upper bound, and how much converts depends on where in the step the work +sits, not on how much of it there is.** The 2.0% projection from the +decision-gate measurement assumed a 1:1 conversion and was therefore too +optimistic by roughly a factor of two; the measured +0.98% is what it is worth. + +**Ranked next levers, after QWEN-025.** + +1. **Reduce graph node count** (HOSTGAP-S6 lever 4, unblocked by + CGTRACE-CONTROL) — ~0.78 µs per removed node, so 100 nodes ≈ 0.85%. Now the + top-ranked lever by expected value. **Risk medium**, sharply diminishing. + Falsified if removing nodes does not move the profiled step period. +2. **Re-attribute the residual G1** — 499 µs/step remains in G1 but the whole + measured host chain is now only 148 µs/step, so the majority of G1 is + something HOSTGAP-S6's attribution folded in but this campaign has not + isolated. Pure measurement, no code risk, and it is what decides whether any + further host-side lever exists at all. Blocked on the nsys host-capture + failure recorded above; unblock it by bisecting the host flag set against a + short capture, starting with `--cpuctxsw=process-tree`. +3. **G3, the `CatArrayBatchedCopy` → `rmsnorm` gap** — 294.9 µs/step and + *untouched by the last three levers*, which all moved G1. HOSTGAP-S6 + attributed it to a single 199 µs `cudaGraphLaunch` for a 1158-node graph, + which makes it the same lever as graph node count rather than an independent + one. **Risk unknown**; confirm the attribution still holds before treating it + as separate. +4. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — + ceiling now lower than the original ~9.2%, since QWEN-024 and QWEN-025 + together removed 224 µs/step of the host chain. **Risk high**, multi-session. + +Note the residual G1 is still 499 µs/step against a measured host chain of only +148 µs/step, so most of what is left in G1 is *not* the post-sampling +bookkeeping chain any more. Attributing the rest of G1 is a prerequisite for +claiming any further ceiling on host-side work. + +**Attempted and failed in the leftover time: re-attributing G1 at the new +config. Host-visibility capture is currently broken, and two plausible causes +are now ruled out.** Three captures with `profile_host_insession.sh` at the +QWEN-025 config, all unrecoverable in the same way — nsys finalization +deadlocked after the target exited, the bounded stop expired, and the +intermediate qdstrm failed to import with +`QuadDCommon::IncompleteFileException` from `verifyHeader`: + +| Run | OSL | `--python-sampling` | qdstrm | Outcome | +|---|---:|---|---:|---| +| `prof/hosts9-1785107698` | 128 | on | 366 MB | import failed | +| `prof/hosts9b-1785109992` | 96 | on | 347 MB | import failed, stream byte-stable 30 s first | +| `prof/hosts9c-1785110873` | 96 | **off** | 219 MB | import failed | + +*Hypothesis 1, copy-while-growing race — falsified.* The recovery path was +changed to wait for the source qdstrm size to hold steady before copying. The +second capture copied a stream that had been byte-stable for 30 s and still +failed `verifyHeader`, so the deadlocked nsys never writes the stream's +terminating section at all. The file is structurally incomplete, not truncated +in transit, and no amount of waiting recovers it. The settle loop is retained +only because it is harmless and documents the ruled-out cause. + +*Hypothesis 2, event volume from the Python sampler — falsified.* Dropping +`--python-sampling` (new `PYTHON_SAMPLING=0` knob) cut the capture from 347 MB +to 219 MB and it failed identically. Neither the sampler nor sheer volume is the +trigger. + +This does *not* affect any QWEN-025 result: the GPU-only script +(`profile_insession.sh`) finalized cleanly twice in this same session and +produced the `g2-off`/`g2-on` pair the profile table above is built from. The +deadlock is specific to the host-visibility flag set, and it is not universal — +HOSTGAP-S6 obtained a usable host trace from this same script on an earlier +allocation, so something environmental differs. Remaining untested suspects, in +order: `--cpuctxsw=process-tree`, then `--sample=process-tree` itself, then +`osrt` tracing. **Do not spend a third session's leftover time on blind retries +of the full flag set** — bisect the flags against a short capture first, since +each full attempt costs ~14 minutes. + +**QWEN-026 landed: fused QK-RMSNorm, +2.9%, new best 79.87% of vLLM.** This is +the first attention-side lever to pay off and the largest win since QWEN-021, +and it over-delivered ~3× against its microbench ceiling — the two per-head +RMSNorm launches sit directly in the serial attention chain, so collapsing them +to one kernel removes a launch *and* a graph node per layer (×48) and cuts the +host dispatch gap between them, not just the ~2 µs/layer of device time the +microbench saw in isolation. **The lesson for ranking: a small-device-time +kernel that sits on the serial critical path and removes a graph node is worth +materially more than its isolated microbench, because graph-node count and +serial-chain launches both convert.** This makes "collapse two adjacent +small launches into one" a higher-value pattern than the raw device-time +numbers suggest — the opposite of QWEN-025's host-side levers, which converted +at only ~1/3. + +The one caveat is that the fused norm is **bf16-ulp, not bit-exact** (TE's +internal rsqrt/reduction differs), so one of the three coherence prompts flips a +single near-tie greedy token. Both continuations are fluent and factually +correct, so it was accepted, but a bit-exact variant (matching TE's rsqrt +instruction so the two-call path is reproduced exactly) would remove the last +correctness caveat and is a cheap follow-up if byte-identity is ever required. + +**Ranked next levers, after QWEN-026.** + +1. **Fuse the other adjacent small-launch pairs on the serial path** — QWEN-026 + proved the pattern converts at ~3× microbench. The remaining per-layer + small-kernel pairs that sit back-to-back in the decode graph are the + input-layernorm/pre-MLP-layernorm `rmsnorm_fwd_tuned` pair and the + `triton_poi_fused_add_copy` residual-add kernels flanking each block. Target + the pair with the shortest serial gap between them first. **Risk medium**; + the same decode-only gating and coherence bar apply. Falsified if collapsing + the pair does not move the profiled step period the way QWEN-026 did. +2. **Reduce graph node count** (HOSTGAP-S6 lever 4) — still live at ~0.78 µs per + removed node; QWEN-026 removed 48 nodes as a side effect and that clearly + helped, so a direct campaign to prune redundant nodes is now better-evidenced + than before. **Risk medium**, sharply diminishing. +3. **Re-attribute the residual G1** — unchanged from the QWEN-025 ranking; still + blocked on the nsys host-capture failure. Bisect the host flag set against a + short capture, starting with `--cpuctxsw=process-tree`. +4. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — + **Risk high**, multi-session; unchanged. + +The most productive next step is to re-profile at the new eleven-gate config +(all ten gates + QWEN-026) so the next adjacent-launch pair is chosen from the +current serial gaps rather than the pre-QWEN-026 trace. + +**QWEN-027 landed: fused add+RMSNorm (self_attn_bda + pre_mlp_layernorm), ++1.37%, new best 80.98% of vLLM.** PROFILE-S9 chose the target exactly as +predicted — the residual-add and the next norm were the two dominant small +kernels and adjacent on the serial path — and the fusion converted cleanly. Two +things worth carrying forward. First, the **input RMSNorm is not a separate +module**: for the non-MLA path `linear_qkv = column_parallel_layer_norm_linear`, +so that norm lives inside the TE `LayerNormLinear` QKV GEMM and cannot be +intercepted at the layer level. Only the standalone `pre_mlp_layernorm` was +fusible, so QWEN-027 removed 48 of the 96 candidate boundaries; the other 48 +(the `mlp_bda` + next-layer input norm) would require either fusing across the +layer boundary or reaching inside TE's `LayerNormLinear`. Second, the add+norm +kernel does **not** regress at high token counts (flat ~4.1 µs vs QWEN-026's +1-row-per-CTA design that lost above 256 tokens), so the same pattern is safe to +apply more widely if the prefill path is ever targeted. + +**Ranked next levers, after QWEN-027.** + +1. **Re-profile at the twelve-gate config and re-rank.** Two norm+add fusions + have now landed; the serial chain has shifted again. Cheapest high-value + step before picking the next lever. Specifically check whether + `_fused_qk_rmsnorm` (QWEN-026) and the RoPE/append-kv kernels around it are + now the tallest adjacent small-kernel run, since attention is the remaining + untouched small-kernel cluster. +2. **Fuse the `mlp_bda` + next-layer input norm boundary** — the other 48 + boundaries QWEN-027 could not reach, because the input norm is inside + `LayerNormLinear`. Options: (a) fuse the `mlp_bda` add into the *start* of the + TE QKV GEMM's norm (needs a TE hook or a mcore-side prenorm that bypasses + `LayerNormLinear`'s internal norm), or (b) accept a small numeric change by + moving the input norm out of `LayerNormLinear` into a standalone fused + add+norm like QWEN-027. **Risk medium–high**; (b) changes the QKV GEMM path. +3. **Reduce graph node count** (HOSTGAP-S6 lever 4) — QWEN-026+027 removed ~96 + nodes/step as a side effect and both converted at ~3×; a direct node-pruning + pass is now well-evidenced. **Risk medium**, diminishing. +4. **Re-attribute the residual G1** — unchanged; blocked on the nsys + host-capture failure. Bisect the host flag set against a short capture. + +A bit-exact variant of both fused norms (matching TE's rsqrt instruction) would +remove the last coherence caveat and is a cheap follow-up if byte-identity is +ever required — QWEN-027's residual is already bit-exact, only the norm output +is one bf16 ulp off. + +## Session 13 (2026-07-30) — measurement session: why the block is latency-bound + +No optimization was accepted this session. It replaced the campaign's measurement +basis, which had been quietly wrong, and produced the first quantified path to +parity. Everything below is same-node (`nvl72151-T13`) unless stated. + +### CLEANBASE-S13 — same-node reference, and the tightest drift control yet + +| Arm | Throughput | tpot | +|-----|-----------|------| +| `ep01_ref` (13 gates + async) | **28,529.1 tok/s** | 8.973 ms | +| `ep06_ref2` (repeat, end of allocation) | **28,606.9 tok/s** | 8.949 ms | + +0.27% apart, so on this node anything above ~0.3% is a real effect. Note this node +is ~2.1% faster than `nvl72067-T13`, where the best config measured 27,944 and the +vLLM reference 34,415.7 — **the 81.2% figure remains the only same-node +comparison; 28,529 must not be divided by 34,415.7.** + +### NSYS-UNUSABLE — nsys cannot measure this workload, in any mode + +The hypothesis was that `--cuda-graph-trace=node` was inflating the trace by timing +~1,800 kernel nodes per step, and that `=graph` would collapse each replay to one +range and restore fidelity. It failed on every count: + +| Arm | Throughput | vs unprofiled | +|-----|-----------|---------------| +| unprofiled | 28,469.5 | — | +| `--cuda-graph-trace=graph` | 20,557.2 | **−27.8%** | +| `--cuda-graph-trace=node` | 20,174.0 | **−29.1%** | + +1. The mode is worth only 1.3 points. **nsys costs ~28% here regardless**, and the + overhead is host-side. +2. On nsys 2026.3.1 `=graph` does **not** collapse replays: all 224k kernel rows + still appear, with `graphId` NULL. The 98 rows in + `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` are graph *instantiations* from warmup. +3. Step reconstruction from the trace again produced an impossible 71.8 ms "step" + against a known 12.45 ms — the same divisor trap as earlier sessions. +4. Worst: the host-side overhead desynchronizes the EP ranks, and the NVLS + collectives are spin-waits, so they inflate **~120×** + (`_multimem_all_gatherv_3tensor` reads 849 µs/call; 48 of those would be 40 ms + against a 9 ms step). An earlier trace measured the same kernel at ~7 µs. + +**Consequence: every absolute idle, gap, and barrier-wait number this campaign read +off an nsys trace is unsafe**, including HOSTGAP-S6's inter-replay gaps and the +"1.6 ms/step idle" and "4.4× barrier spread" that motivated sessions 7–12. Relative +per-call durations of *compute* kernels remain usable. Do not profile this workload +with nsys to obtain a busy/idle split; use STEPGPU below. + +### STEPGPU-S13 — the block graph is 92.8% of the step (`step_gpu_timing.py`) + +New env-gated instrumentation (`MCORE_INFER_STEP_GPU_TIMING`) wraps +`TransformerBlock.__call__` on the CUDA-graph path with CUDA events, held in a ring +of 64 pairs and read back only when the ring wraps, so no step synchronizes on its +own work. It costs 2.3% and, unlike nsys, **validates itself**: the step period it +reports (8.97 ms) matches the benchmark's tpot (8.973 ms). + +| Quantity | Value | Share | +|----------|-------|-------| +| step period | 8.97 ms | 100% | +| transformer-block CUDA graph | ~8.33 ms | **92.8%** | +| everything outside the graph | **0.59 ms** | 6.6% | + +Held within 0.2 points across all 4 ranks over 4,032 steps, and reproduced at 92.8% +in a second sample. **This retires host-side optimization quantitatively.** +Embedding, logits GEMM, sampling, bookkeeping and every host gap together are +0.59 ms; the gap to vLLM is ~1.53 ms/step. Third confirmation after QWEN-032/033, +and the first that is a measurement rather than an inference. + +### EPSKEW-S13 — the EP skew hypothesis is falsified + +All 4 ranks are processes on one node, so `perf_counter_ns` is directly comparable +between them; the spread of step-entry timestamps *is* the arrival skew the NVLS +spin-waits absorb. Median spread **24 µs**, p90 46 µs — 0.27% of the step. Rank 1 is +mildly structurally late (last on 42% of steps, first on 7%) but by only ~21 µs +median. The "4.4× barrier wait spread" was an nsys artifact per NSYS-UNUSABLE. +**Host arrival skew is not a lever.** Do not revisit without new evidence. + +### BSSCALE-S13 — 32× the tokens costs 1.70× the block time + +Batch-size sweep with the block timer on. Valid points are exact CUDA-graph buckets; +BS192 is padded into the 256 bucket, so the curve is a staircase, not smooth. + +| BS | step | block | µs/layer | +|----|------|-------|----------| +| 8 | 5.563 | 5.030 | 105 | +| 16 | 5.896 | 5.361 | 112 | +| 32 | 6.290 | 5.770 | 120 | +| 64 | 7.081 | 6.542 | 136 | +| 128 | 7.140 | 6.585 | 137 | +| 256 | 9.151 | 8.553 | 178 | + +**The block is latency-bound, not compute-bound.** The batch-independent floor is +~5.03 ms/step = **105 µs/layer = 55% of the whole step**; the marginal cost of the +other 248 tokens is only 73 µs/layer. For scale, vLLM's *entire* step is 7.44 ms +while mcore's floor plus non-graph overhead is already ~5.6 ms. + +### CHAINCOST-S13 — fusion pays, but *not* per kernel removed + +All 13 gates off vs on, at two batch sizes: + +| | block ON | block OFF | delta | +|---|---------|-----------|-------| +| BS8 | 5.030 | 6.295 | 1.265 ms/step = 26.4 µs/layer | +| BS256 | 8.553 | 10.317 | 1.764 ms/step = 36.7 µs/layer | + +At BS8 the block is nearly pure latency, so 26.4 µs/layer over the ~10 kernels the +gates remove suggests ~2.6 µs of serial chain latency per kernel — well above those +kernels' own 1–2 µs device time. The 13 gates are jointly worth **16.1%** of +throughput. + +**The per-kernel decomposition did not survive calibration — it is refuted.** +Six single-gate arms at BS8, against the all-gates-ON block of 5.030 ms: + +| arm | block | delta | µs/layer | kernels/layer | µs/kernel | +|-----|-------|-------|----------|---------------|-----------| +| `MCORE_FUSED_QK_NORM=0` | 5.291 | +0.260 | 5.4 | 1 | 5.4 | +| `MCORE_ROUTER_FUSED_TOPK=0` | 5.292 | +0.261 | 5.4 | 3 | 1.8 | +| `MCORE_FUSED_ADD_NORM=0` | 5.698 | +0.668 | 13.9 | 2 | 7.0 | +| `MCORE_FUSED_ADD_NORM_QKV=0` | 5.375 | +0.345 | 7.2 | 1 | 7.2 | +| `MCORE_MOE_GEMM_TUNE=0` | 4.762 | **−0.268** | −5.6 | count-neutral | — | +| `MCORE_MOE_SUM_FAST=0` | 5.073 | +0.043 | 0.9 | count-neutral | — | + +Least squares over the four chain-reducing arms gives **slope 0.30 µs/layer per +kernel removed with a 7.46 µs/layer intercept** — i.e. kernel count explains almost +nothing and the per-gate constant dominates. Per-kernel cost ranges 1.8–7.2 µs with +no consistency. **There is no usable per-kernel coefficient, and any projection of +the form "remove N kernels, gain N × c" is invalid.** The earlier ~2.6 µs/kernel +figure and the "25 → 15 kernels buys 1.25 ms and most of the path to parity" +projection are both withdrawn. + +Two further cautions from the same data: + +- **Single-gate deltas are strongly sub-additive and must not be summed.** Just four + gates sum to 31.9 µs/layer, already *more* than all thirteen together (26.4), so + these arms are contending for a shared resource — the same signature QWEN-030 saw. +- The measurement itself is sound: six arms produced six clearly distinct outcomes + (−5.6 to +13.9 µs/layer), so there is no "any gate off costs the same" artifact, + and the earlier suspicion of BS8 graph-node quantization was unfounded. The + identical QK-norm / router-topk pair was coincidence. + +**GEMM-TUNE-SMALLBATCH (new, actionable):** `MCORE_MOE_GEMM_TUNE=1` is a **~5% +regression at BS8** — turning it off raised throughput 1438.0 → 1509.6 tok/s. The +QWEN-013 tiles were tuned for 256-token decode and QWEN-013b added an *upper* bound +(fall back above M=384) but never a lower one, so at 8 tokens the tuned tiles are +simply the wrong shape. A lower-bound guard is a small, low-risk fix. Irrelevant to +the BS256 parity goal; relevant to anyone serving small batches. + +**This finally explains the campaign's recurring puzzle.** QWEN-026, 027 and 028 +each over-delivered ~3× their device-time arithmetic and each was recorded as a +surprise. Removing a launch removes a *chain link*, not just a duration — in a +latency-bound block that is the expected multiplier, not luck. + +**The 105 µs/layer floor remains unattributed.** A split into ~66 µs of serial chain +and ~39 µs of NVLS collectives was arithmetically tempting and consistent with +QWEN-016's "exposed, latency-bound" finding, but it was built on the per-kernel +figure the calibration then refuted, so it is withdrawn. Attributing the floor needs +a direct measurement, not a coefficient — the obvious candidate is an env-gated +ablation that stubs one component at a time inside the graph and reads the block +timer, since `step_gpu_timing.py` now makes that a reliable readout. + +### NCCL-HANG — `inference_moe_token_dispatcher_type=nccl` is non-functional here + +With async scheduling and the 13 gates, the server starts and then every request +times out. It hung for 2 h and cost half the allocation, which also lost the +nvls-vs-nccl comparison. Two process lessons: job scripts now carry +`timeout --signal=INT 1200`, and a watcher that reports only *state changes* cannot +distinguish a hung job from a healthy one — it needs a log-mtime staleness check. + +### Session 13 — conclusion & recommendation + +The target is no longer expensive kernels; it is **per-layer latency**. Ranked by +evidence: + +What is certain: the prize is inside the transformer block (92.8% of the step), the +block is latency-bound rather than throughput-bound (32× tokens for 1.70× time, +~55% batch-independent), and neither host work nor EP arrival skew is worth another +hour. What is *not* known is what the 105 µs/layer floor is made of — and this +session showed that guessing at it from coefficients produces retractions. + +1. **Attribute the 105 µs/layer floor by direct ablation.** Env-gated stubs for one + component at a time (NVLS dispatch, NVLS combine, expert GEMM, attention) read + through the now-trustworthy block timer. Output is garbage under ablation, which + is fine — only the timing is wanted. This is the prerequisite for sizing + anything else, and it is the cheapest remaining experiment. +2. **The 2 NVLS collectives** — the leading suspect for the floor, on the strength + of BSSCALE (cost is barrier latency, not bytes) and QWEN-016. Levers: fewer + barriers (fuse dispatch/combine, or hoist across layers) or overlap with + compute. QWEN-017 already failed at replacing the barrier itself; overlap is + unexplored. **Size it via item 1 before building.** +3. **More fusion** — still plausible, and the aggregate (13 gates = 16.1%) says the + direction has paid. But the per-gate value is idiosyncratic, not proportional to + kernels removed, so each candidate must be sized on its own rather than from a + chain-length model. +4. **Do not** pursue host-side work, EP arrival skew, or nsys-derived idle. +5. Small, separate: add a lower-bound guard to `MCORE_MOE_GEMM_TUNE` + (GEMM-TUNE-SMALLBATCH). + +## Session 14 (2026-07-30/31) — the floor, attributed at last + +Session 13 ended unable to say what the ~105 µs/layer floor was made of, having +twice retracted an inference built on per-kernel coefficients. This session +measured it directly and the answer reorders the whole campaign. Node +`nvl72115-T08`; all comparisons within it. + +### ABLATE-HARNESS — how the floor was measured (`floor_ablation.py`) + +New env-gated module plus six call-site gates. Each deletes one component from +the *captured* block graph and the change is read from `step_gpu_timing`. Three +design points are what make it trustworthy: + +* **Ablate only while capturing** (`torch.cuda.is_current_stream_capturing()`). + Warmup runs eagerly first, so every real kernel executes at least once and + every preallocated buffer holds plausible values; capture then takes the + ablated branch, so the choice is resolved once and no replay gains + data-dependent control flow. +* **Substitute cached zeros, never uninitialized memory.** Zeros keep the + residual finite, so the next layer's router still spreads tokens across + experts. Garbage would collapse routing and silently change the grouped-GEMM + tile balance of the components *not* being ablated. Caching by shape matters + too: a fresh `torch.zeros` per replay would add a multi-MB fill to the + measurement it is meant to isolate. +* **`hit()` logs the first time each site takes its ablated branch.** This + immediately caught a dead gate (see ATTN-BIG) that would otherwise have read + as "attention is free". + +Deltas are **chain** deltas: they include the serialization a component imposed +on its neighbours, which in a latency-bound block is the quantity that matters, +and they are deliberately not comparable to a profiler's kernel duration. + +Drift control: two baselines 0.4% apart on block time (7.639 / 7.607 ms) but +**2.5% apart on throughput** (27,626 / 28,321 tok/s). Block time is the primary +metric this session; throughput differences under ~2.5% mean nothing here. + +### FLOOR-S14 — the composition, and which parts are addressable + +Baseline block **7.623 ms = 158.8 µs/layer** at BS256/OSL1024. + +| component removed | block | delta | µs/layer | % block | tput | +|---|---|---|---|---|---| +| expert GEMMs (FC1+SwiGLU+FC2) | 4.804 | −2.819 | −58.7 | −37.0% | +48.3% | +| attention (`flash_decode_and_prefill`) | 5.756 | −1.867 | −38.9 | −24.5% | +27.4% | +| topk reduce (`_moe_sum`) | 7.079 | −0.544 | −11.3 | −7.1% | +7.9% | + +**GEMM-ROOFLINE — the expert GEMM is done, and this is the session's most +consequential finding.** Per layer per rank it streams 32 local experts' FC1+FC2 +weights = **302 MB**, so 14.5 GB per step. Against QWEN-012's measured 6.08 TB/s +that is a 49.7 µs/layer floor, and it measures 58.7 — **84% of achievable +bandwidth**. Only 19.3 GFLOP/layer, so it is nowhere near compute-bound; it is +weight-streaming-bound. Consequences: + +1. The largest single component, 37% of the block, is **essentially optimal in + bf16 and cannot be tuned further.** Stop looking at grouped-GEMM tiling for + throughput at BS256. +2. **vLLM pays the same 2.83 ms/step**, because it streams the same weights in + the same precision. So the gap to vLLM does *not* live here, and every prior + session that ranked "MoE grouped-GEMM 40.5%" as the top target was ranking by + size rather than by addressability. +3. The only lever on it is fewer weight bytes: FP8/NVFP4 experts, which is a + precision decision, not a kernel one. + +Subtracting it: our addressable non-GEMM cost is ~4.80 ms of block + 0.59 ms +outside ≈ **5.39 ms/step**, against vLLM's ~4.6 ms on the same arithmetic. **The +whole remaining gap is ~0.8 ms/step of non-GEMM work**, and attention plus the +topk reduce are 2.41 ms of that region — enough room to close it without +touching the GEMM. + +**ATTN-BIG — attention is 38.9 µs/layer and roughly 3× off its own roofline.** +This is the new top target and it was invisible until now. At 64 tokens/rank and +~600 average context it reads ~79 MB of paged KV per layer per rank, a 13 µs +bandwidth floor against 38.9 µs measured = **33% of 6.08 TB/s**. Unlike the +GEMM there is real headroom, and it is 24.5% of the block with a +27.4% +throughput ceiling. + +Process note: the first attention arm reported **zero** delta, and only the +`hit()` log revealed why — the gate was on `_run_core_attention`, which the +*static* engine uses. Dynamic batching calls `flash_decode_and_prefill`. Without +the hit counter this would have been recorded as "attention is free", which is +exactly the class of false negative that cost sessions 7–12. + +**Two arms are invalid; recording them so they are not re-read as results.** +`MCORE_ABLATE_NVLS_DISPATCH` and `..._COMBINE` both hit the 1500 s timeout +(rc=124). Prefill runs the real collective eagerly while replays skip it, which +almost certainly desynchronizes the symmetric-memory barrier sequence between +ranks. Their block numbers (−0.66 and −3.26 ms) were measured in a degraded +state and **must not be quoted**; the NVLS collectives remain unmeasured. +`MCORE_ABLATE_MOE_ALIGN` is invalid for a different reason: replaying warmup's +align tables also shrank `num_tokens_post_padded`, so the GEMMs did less work and +the arm measured that instead of the align cost. + +### QWEN-034 — FC2 epilogue topk reduce (atomics): rejected, −12.5% + +Hypothesis, sized off the ablation rather than guessed: the topk reduce is worth +11.3 µs/layer of chain, so folding it into FC2's epilogue should recover most of +it and drop the `[num_valid, K]` intermediate as well. FC2 already loads the +routing weight (`MUL_ROUTED_WEIGHT`) and row *r* belongs to token `r // topk`, so +the change is small: scale the fp32 accumulator and `tl.atomic_add` into the +`[max_tokens, K]` output instead of storing. + +| | block | throughput | +|---|---|---| +| baseline | 7.623 | ~27,974 | +| `MCORE_FUSE_FC2_REDUCE=1` | 8.995 | 24,468 | + +**+1.372 ms/step (+28.6 µs/layer), −12.5% throughput.** Coherence passes and the +fused form is strictly *more* accurate (no bf16 truncation of each slot before +the sum), but it is far slower: the atomics cost ~40 µs/layer against the 11.3 it +was meant to save, and the destination also needs an 8.4 MB zero-fill per layer. +Kept default-off as a documented negative. + +Why it failed, and the general lesson: a topk slot's contributions land on one +output row from *different* tiles, so the reduction cannot be kept in registers +and every element becomes a global read-modify-write. This is the same wall +QWEN-001's mega-fusion hit. **Epilogue fusion pays when the consumer is +elementwise on the tile (QWEN-002's SwiGLU, QWEN-013's squared-relu) and loses +when it reduces *across* tiles.** That rule now has two independent +confirmations and should be applied before building, not after. + +### Session 14 — conclusion & recommendation + +The campaign's target list is now ordered by *addressability* rather than size: + +1. **Attention, 38.9 µs/layer at 33% of its bandwidth roofline** — the largest + component with real headroom, +27.4% ceiling. Start by identifying which + FlashAttention generation decode actually resolves to (`flash_attention_version` + has no CLI flag; a benchmarking env override and a one-shot log were added) and + A/B the alternatives, then look at paged-KV over-read and the decode split-K. +2. **The topk reduce, 11.3 µs/layer** — real but *not* via epilogue atomics + (QWEN-034). A separate kernel that reduces with better locality, or folding it + into the combine, are the remaining options. +3. **The NVLS collectives are still unmeasured** — the ablation needs to skip the + collective in warmup too, so the barrier sequence stays consistent across + ranks, before any number from it is trusted. +4. **Do not** pursue grouped-GEMM tuning for BS256 throughput (GEMM-ROOFLINE), + host-side work, EP arrival skew, or nsys-derived idle. +5. Precision (FP8/NVFP4 experts) is the *only* lever on the 37% the GEMM + occupies, and it is a correctness decision that needs its own evaluation. + +## Session 15 (2026-08-01) — FlashAttention-2 beats FlashAttention-4 on decode + +All arms: BS256, OSL1024 unless noted, 4×GB200 node `nvl72151-T14`, job `5765929`, +session `qwen-attn`, thirteen accepted gates on, `run_rebased.sh`. Block-graph GPU +time from `step_gpu_timing.py` (CUDA events, no profiler), median over the whole +decode range. + +### QWEN-035 — pin FlashAttention 2 on the decode path: **+3.5%** + +| arm | block-graph GPU | throughput | +| --- | --- | --- | +| `ga00_ref` FA4 (mcore's auto choice) | 7.695 / 7.707 ms | 28,149.6 | +| `ga05_ref2` FA4, repeat | 7.751 / 7.698 ms | 28,098.1 | +| `ga02_fa2` FA2 pinned | 7.412 / 7.400 ms | 28,999.2 | +| `ga04_fa2b` FA2 pinned, repeat | 7.398 / 7.384 ms | 29,225.9 | +| `ga11_fa2_five` FA2, 5 timed iters | — | **29,132.8** | + +Repeat spread is 0.3% on block time and 0.2–0.8% on throughput, so the 4.1% block +separation is roughly ten times the noise floor. `ga11_fa2_five` is the headline +number: 5 timed iterations at 28,705.6 / 29,287.2 / 29,199.4 / 29,230.4 / +29,249.3 tok/s — 0.30% spread across iterations 2–5, with iteration 1 the usual +ramp outlier. tpot 8.787 ms/tok, avg_latency 8,733.2 ms. It agrees with both +3-iteration runs (mean of all three: 29,119). Coherence is identical on two of +three probes and equally correct on the third — expected, since the two backends +accumulate in a different order. **85.7% of vLLM's 33,994.5** on the 5-iteration +headline run, up from 82.7%. + +The mechanism is in `flash_decode_and_prefill`, not in the kernels' relative +quality: mcore's auto-preference is FA4 → FA3 → FA2, but **FA4 has no dedicated +decode kernel on this path**. FA4 decode is routed into `flash_attn4_varlen_func` +— the varlen *prefill* interface, called with `max_seqlen_q = 1` — while FA2 goes +to `flash_attn_with_kvcache`, the purpose-built flash-decoding kernel with +split-KV and combine. So the auto-resolution order silently costs 3.5% on decode. +This is arguably an upstream default bug rather than a tuning knob. FA3 is not +installed in this container (`HAVE_FA3=False`), so `ga01_fa3` asserted out; the +FA3 leg of the comparison is untested. + +Enable with `MCORE_FLASH_ATTN_VERSION=2` (the env override added in Session 14 +because `flash_attention_version` has no CLI flag). Config still wins over env. + +The gain is **entirely inside attention**, and the ablation proves it: with +attention stubbed out, the block floor is the same under both backends (5.858 ms +FA4 vs 5.884 ms FA2, inside noise), so attention alone goes 1.843 → 1.510 ms/step, +−18%. + +### ATTN-ROOFLINE-S14-RETRACTED — attention is *not* 33% of roofline + +Session 14 recommended attention as "the largest component with real headroom, +33% of its bandwidth roofline, 3× available". **That roofline was wrong** and the +recommendation it produced should not be reused. + +KV-cache reads alone, per layer per rank, are +`B(256) × S × 2(K,V) × H_kv(4) × D(128) × 2 B`. At OSL1024 with ~100-token gsm8k +prompts the average KV length across the decode is ~600, giving **315 MB/layer** +→ **39.3 µs/layer at 8 TB/s HBM**. The measured attention chain delta under FA2 is +31.5 µs/layer, and a chain delta *understates* a kernel's duration whenever any of +it overlaps neighbouring work. Attention is therefore at or above its bf16 +bandwidth bound, not a third of the way to it. There is no 3× there, and the only +lever on it is reading fewer bytes (FP8 KV cache — no knob exists in mcore today). + +### FLOOR-CONST-S15 — the non-attention floor is KV-length independent + +| arm | OSL | attention | block-graph GPU | +| --- | --- | --- | --- | +| `ga06_fa2_na` | 1024 | ablated | 5.880 / 5.888 ms | +| `ga09_o512_na` | 512 | ablated | 5.833 / 5.871 ms | + +Halving the output length halves the average KV length and changes the ablated +floor by 0.6%, i.e. not at all. **The 5.85 ms floor contains no KV-dependent +work**, so it is pure weight-streaming, collectives, and chain latency, and it can +be reasoned about independently of the decode position. Its OSL512 attention-free +throughput is 34,432 tok/s. + +The matching OSL512 *reference* arm (`ga08_o512`) is unusable: rc=124, the server +stalled after ~200 decode steps and hit the 1500 s timeout, so its single +4.93 ms report is from a degraded run. The intended KV-length scaling test of +attention is therefore still untested; only the floor half of the pair survived. + +### Rejected: KV-cache block size 512 + +`ga07_fa2_b512` (`--inference-dynamic-batching-block-size 512`, default 256): +block 7.413 / 7.390 ms, 29,122.9 tok/s — indistinguishable from FA2 at the default +block size. Paged-KV indirection granularity is not a cost here. + +### NVLS-UNABLATABLE — the collectives cannot be sized by subtraction, ever + +The Session 14 recommendation was to fix the NVLS ablation by suppressing the +collective in warmup too, so the symmetric-memory barrier sequence stays +consistent across ranks. That fix was made (`nvls_off()`, an unconditional gate +replacing the capture-only one) and it worked as intended: `ga10_nonvls` ran to +completion, rc=0, no hang, and both gates fired on all four ranks +(`nvls_dispatch` 4 hits, `nvls_combine` 4 hits). + +**The result is that removing both collectives makes the block 13.5% slower:** +8.391 / 8.407 ms against the 7.39 ms FA2 baseline, 25,950.2 tok/s against 29,113. + +Removing two collectives cannot make their own absence expensive, so the arm is +confounded, and the reason is structural rather than a gating bug: +**`multimem_all_gatherv_3tensor` produces the work descriptor for everything +downstream.** `dispatch_preprocess` reads `routing_map` back out of the gathered +symmetric buffer (`agv_r["tensor"].view(global_max, topk)`), and that routing map +is what determines how many token-expert assignments — hence how many grouped-GEMM +tiles — the experts execute. Skip the gather and the buffer still holds whatever +warmup's *prefill* gather left there, which has a different and denser assignment +distribution than steady-state decode. So the experts do a different amount of +work, and the measured block time says nothing about the collective. + +This is the same work-descriptor confound that invalidated the Session 14 MoE-align +ablation, and it is not fixable by better gating. **Subtractive ablation is the +wrong instrument for the dispatch collective**, and by extension for anything that +computes a work descriptor. Four attempts have now failed on this (two hangs, one +align confound, this one); stop trying. + +The right instrument is non-invasive in-situ timing: a second CUDA event pair +recorded around the collective inside the captured graph, read back through the +same wrapped ring `step_gpu_timing.py` already uses for the whole block. Graph- +captured event records replay, so this measures the real duration per replay while +leaving the work descriptor and the barrier sequence completely untouched. That is +the next experiment, and it needs no ablation at all. + +### Session 15 — conclusion & recommendation + +Where the 7.39 ms block actually goes, with roofline where it is known: + +| component | measured | bf16 bandwidth floor | verdict | +| --- | --- | --- | --- | +| attention | 1.51 ms | ~1.89 ms (315 MB/layer) | **at the bound** | +| expert GEMMs | ~2.16 ms | 1.81 ms (302 MB/layer) | 84% of bound | +| dense QKV/O proj | — | 0.23 ms (37.8 MB/layer) | negligible | +| **unaccounted** | **~3.4 ms** | n/a | **the target** | + +Two thirds of what is left is neither attention nor weight streaming. That ~3.4 ms +— 46% of the block — is NVLS dispatch/combine, routing/permute, norms, and chain +latency, and it is the only place a 14% end-to-end gap can still come from. The +early OSL1024 profile already noted mcore carrying 11.5% exposed EP comm against +vLLM's zero, which points the same way. + +1. **Attribute the ~3.4 ms remainder with in-situ event timing, not ablation.** + NVLS-UNABLATABLE settles the method question: the dispatch collective produces + the routing map that sizes all downstream GEMM work, so subtracting it changes + the work and the timing is meaningless. Add a per-site CUDA event pair inside + the captured graph (same ring-buffer readback as `step_gpu_timing.py`) around + the dispatch, the combine, and the expert GEMM. Nothing else should be built + before those three numbers exist. +2. **Overlap, not replacement, for the NVLS barriers.** QWEN-017 already failed at + replacing the barrier itself; overlapping dispatch/combine with expert compute + is unexplored and is the natural lever if item 1 sizes them large. +3. **Precision is the only lever on the bandwidth-bound 3.7 ms** (attention + + expert GEMMs together). FP8 KV cache would need implementing — no mcore knob + exists — and would roughly halve attention. FP8/NVFP4 experts likewise. Both + are correctness decisions needing their own evaluation. +4. **Raise the FA4 decode path upstream.** Independent of this campaign: FA4 + resolving to a prefill kernel for decode costs every mcore inference user 3.5%. +5. **Do not** pursue attention kernel tuning (at its bound), grouped-GEMM tuning, + KV block size, host-side work, EP arrival skew, or nsys-derived idle. + +## Session 16 (2026-08-02) — the block, attributed + +Job `5777945` on `nvl72117-T17`, session `qwen-insitu`, thirteen gates + FA2, +BS256/OSL1024. New module `megatron/core/inference/insitu_timing.py`. + +### INSITU-METHOD — external event records, and the two probes that were needed + +Subtraction is dead (NVLS-UNABLATABLE) and nsys is unusable here, so components are +timed *in place* with an event pair inside the captured graph. Getting there took two +standalone probes, and both earned their keep by failing: + +1. **Plain `torch.cuda.Event.record()` during capture does not work.** It raises + nothing and does become a graph node, so it looks fine — and then + `elapsed_time` fails with `cudaErrorInvalidValue`, because an ordinary + event-record node exists only for intra-graph ordering and carries no + host-readable timestamp. The first instrumented run duly reported every site as + zero. Had the probe not existed, the natural reading of that run would have been + "the sites cost nothing". +2. **`cudaEventRecordWithFlags(..., cudaEventRecordExternal)` does work**, via + ctypes since PyTorch exposes no flags argument. It must target the *capturing* + stream; `torch.cuda.graph()` picks its own side stream unless passed one, and + getting that wrong returns `cudaErrorIllegalState` (401) rather than anything + descriptive. + +Two properties make the output trustworthy. Summed sites reconstruct the whole graph +to **1.3%** in the probe, so the decomposition is additive. And the instrumentation +costs **2.0%** of block time (7.543 ms instrumented vs 7.396 ms control, same +config, same node), which is small and, more importantly, measured rather than +assumed. Each event pair carries ~2.2 µs of its own overhead, read directly from a +`_calib` site wrapping an empty region in the same graph and subtracted from every +other site — without that, a 8 µs component reads ~30% high. + +### FLOOR-ATTRIB-S16 — where the 7.4 ms block goes + +Rank 0, per-layer medians net of pair overhead, scaled ×48. Two sampled layers (8 +and 24) with their spread, so an atypical layer is visible rather than silent. + +| component | µs/layer | ms/step | % of block | layer spread | +| --- | --- | --- | --- | --- | +| `expert_gemm_fc1` | 37.6 | 1.806 | 24.0% | 6% | +| `expert_gemm_fc2` | 26.5 | 1.272 | 16.9% | 18% | +| `attn_core` | 25.2 | 1.208 | 16.0% | 3% | +| `nvls_combine` | 14.1 | 0.677 | 9.0% | 81% | +| `nvls_dispatch` | 13.9 | 0.665 | 8.8% | 5% | +| `moe_sum` | 7.9 | 0.380 | 5.0% | 9% | +| `moe_align` | 7.8 | 0.373 | 5.0% | 3% | +| **sum of sites** | | **6.381** | **84.6%** | | +| unattributed (norms, router, QKV/O proj, RoPE) | | 1.180 | 15.6% | | + +**The NVLS collectives are finally sized: 1.34 ms/step, 17.8% of the block.** After +two hangs and two work-descriptor confounds, this is the first real number for them, +and it lands at the low end of the 1.5–2 ms the barrier-count argument predicted. +`nvls_combine`'s 81% layer spread (18.2 µs at layer 8 versus 10.0 at layer 24) and +its cross-rank variation are the expected signature of a spin-wait absorbing +arrival skew, not measurement noise. + +**Expert GEMMs are 3.08 ms, 40.8% — larger than the ~2.16 ms previously estimated**, +and the split is the interesting part. FC2 costs 70% of FC1 while reading half the +weight bytes (101 vs 201 MB/layer/rank) and doing half the FLOPs. Either FC2's +tiling is leaving bandwidth on the table or it is not weight-bound at all. + +**Attention is 25.2 µs/layer, not the 31.5 the chain delta implied.** Chain deltas +overstate, as suspected — one more reason not to quote them as component costs. + +### HBM-BW-S16 — the roofline denominator was wrong + +Measured achievable HBM bandwidth on this GB200 (184 GiB, 152 SMs): **6.85 TB/s** +sustained on a large bf16 copy, at both 2 GiB and 8 GiB working sets — 86% of the +8 TB/s spec figure, which is normal for a copy kernel and is the right number for a +roofline denominator. + +**Every roofline percentage in Sessions 14–15 used 8 TB/s and is therefore too +generous**, including "expert GEMMs at 84% of bound". Worse, redoing attention with +the correct denominator makes it *impossible*: 315 MB/layer of KV at 6.85 TB/s needs +45.9 µs/layer, and attention measures 25.2. A measured cost below its own computed +bound means the byte estimate is wrong, not that the kernel is superhuman — the +likely suspects are the assumed ~600-token average KV length and how many requests +are actually active per replay. **No attention roofline claim should be quoted until +that is resolved by counting bytes from the code rather than from the config.** + +### FULL-ATTRIB-S16 — the accounting closes, and the dense projections are the surprise + +Three more sites (`qkv_proj`, `out_proj`, `router_topk`) take the block from 84.6% +attributed to **101.8%** — sum of sites 7.747 ms against a 7.608 ms block. The 1.8% +over-count is the instrumentation's own inflation plus the constant-bias +approximation, so the block is now effectively fully accounted for. Ranks 0 and 2 +agree within 6% on every site. + +| component | µs/layer | ms/step | % of block | bf16 roofline @6.85 TB/s | % of bound | +| --- | --- | --- | --- | --- | --- | +| `expert_gemm_fc1` | 36.7 | 1.763 | 23.2% | 29.4 µs (201 MB) | 80% | +| `expert_gemm_fc2` | 26.1 | 1.253 | 16.5% | 14.7 µs (101 MB) | **56%** | +| `attn_core` | 24.6 | 1.180 | 15.5% | see HBM-BW-S16 | unresolved | +| `nvls_dispatch` | 17.1 | 0.819 | 10.8% | n/a (barrier) | n/a | +| `nvls_combine` | 13.5 | 0.649 | 8.5% | n/a (barrier) | n/a | +| `out_proj` | 12.1 | 0.581 | 7.6% | 2.4 µs (16.8 MB) | **20%** | +| `qkv_proj` | 10.6 | 0.508 | 6.7% | 3.1 µs (21 MB) | **29%** | +| `moe_align` | 8.7 | 0.418 | 5.5% | n/a | n/a | +| `moe_sum` | 7.9 | 0.377 | 5.0% | n/a | n/a | +| `router_topk` | 4.1 | 0.197 | 2.6% | n/a | n/a | +| **sum** | | **7.747** | **101.8%** | | | + +The norm fusions are not instrumented and the sum already exceeds the block, which +is its own small result: after QWEN-026/027/028 the norms no longer occupy +measurable standalone time. + +**The dense attention projections are the finding: 1.089 ms/step (14.3%) at 20-29% +of roofline.** Session 15 dismissed them at "0.23 ms, negligible" — that estimate +was 4.7× low because it was a *bandwidth* estimate for GEMMs that turn out to be +latency-bound. At M=256 tokens, `qkv_proj` (N=5120, K=2048) needs 3.1 µs on +bandwidth and ~2.4 µs on FLOPs, and takes 10.6; `out_proj` takes 12.1 against +2.4. Two skinny GEMMs, each ~3-5× off both of their own bounds. **~0.8 ms of +headroom, against a total gap of 1.14 ms** — and unlike the collectives, nothing +about this is cross-rank or barrier-shaped. + +Corrected headroom, all measured against 6.85 TB/s: + +| target | measured | bound | headroom | +| --- | --- | --- | --- | +| dense projections (qkv + out) | 1.089 ms | ~0.27 ms | **~0.8 ms** | +| NVLS collectives (overlappable, not removable) | 1.468 ms | n/a | up to 1.47 ms | +| `expert_gemm_fc2` | 1.253 ms | 0.787 ms | ~0.47 ms | +| `expert_gemm_fc1` | 1.763 ms | 1.44 ms | ~0.32 ms | + +### VLLM-COMM-S16 — vLLM runs the same AllGather+ReduceScatter, so "0 exposed comm" was wrong + +`PROFILE-OSL1024`, `PROFILE-DECODE` and the Session 8 recommendation all assert that +vLLM's MoE path has **0 exposed comm** and that mcore's dispatch/combine is therefore +pure deficit. Direct query of `nsys_trace/vllm_baseline_osl1024.sqlite` falsifies this. +vLLM's steady state contains exactly two collectives, and they are the same two +mcore uses: + +| | kernel | launches/rank | µs/launch | ms/step | +| --- | --- | --- | --- | --- | +| dispatch | `ncclDevKernel_AllGather_RING_LL` | 48,752 | 20.9 | 1.005 | +| combine | `ncclDevKernel_ReduceScatter_Sum_bf16_RING_LL` | 48,740 | 20.6 | 0.987 | + +48,752 = 48 layers × 1016 decode steps exactly — one pair per MoE layer per step, the +same cadence as mcore's `NVLSAllGatherVDispatcher`. There is **no all-to-all anywhere in +the trace**: no DeepEP, no pplx-kernels, no `nccl*AllToAll`. Note also that neither +framework uses an all-to-all dispatcher, so the recurring "all-to-all" shorthand for +this path (including in `QWEN-005b`) is a misnomer — it is AllGatherV in, ReduceScatterV +out, every token broadcast to every rank. + +The only difference is transport: NCCL ring/LL for vLLM, NVLink multimem on symmetric +memory for mcore. On the one clean measurement each way, mcore is *ahead*: 30.6 µs/layer +(`FULL-ATTRIB-S16`, in-graph events, unprofiled) against vLLM's 41.5 µs/layer. + +Two caveats bound how hard that comparison can be leaned on. mcore's collectives read +433 µs and 91 µs per launch **in the profile** — 25× nsys spin-wait inflation, per +`NSYS-OVERHEAD-S12`; only the in-situ number is usable. And vLLM's 41.5 µs is itself +measured under nsys, where RING_LL also spins, so its true cost lies somewhere between +the ~1.2 µs wire floor for ~1 MB over NVLink5 and 41.5 µs, and this trace cannot pin it. + +What survives is the structural claim, which is the useful one: **mcore has no +architectural disadvantage in the MoE collective** — same algorithm, same launch +cadence, better transport. And since both sides sit 15-30× above the wire time for the +bytes they move, **both are barrier-latency bound, not bandwidth bound**, which is why +QWEN-017's transport swap failed and why overlap, not a faster wire, is the only lever +left here. Correspondingly, the 1.47 ms is *not* free money relative to vLLM: the +differential is unmeasured and is certainly far smaller than 1.47 ms. + +### GAP-DECOMP-S17 — compute is at parity; the gap is kernel count. (Idle magnitude partly retracted — see GRAPH-LAUNCH-S17) + +> **Read the caveat first.** The per-component *kernel* comparison below is sound and is +> what identified the real lever. The **idle magnitude and its attribution are not**: +> `GRAPH-LAUNCH-S17` shows that node-level graph tracing inflates the host side of a +> graph launch ~20×, by an amount comparable to the entire idle budget, and mcore pays +> more of it than vLLM by having more graph nodes. Treat "77% of the gap is GPU idle" as +> unproven. The surviving claim — mcore runs ~42% more kernels per step, and cutting them +> pays — is confirmed by QWEN-036. + +First like-for-like decomposition of both sides: same window methodology, both traces +`--cuda-graph-trace=node`, steady-state decode only, one rank, step count derived from +the one-per-layer collective. Scripts: `nsys_trace/busy_union.py`, `gap_profile.py`, +`stall_detail.py`. GPU busy is the **union** of kernel intervals, not the sum — the sum +overcounts whenever streams overlap, which is exactly where vLLM wins. + +| per step, rank 0 | mcore | vLLM | delta | +| --- | --- | --- | --- | +| wall | 8.571 ms | 7.456 ms | **+1.115** | +| GPU busy (interval union) | 6.926 ms | 6.674 ms | +0.252 | +| GPU idle (wall − union) | **1.645 ms** | **0.783 ms** | **+0.862** | +| sum-of-kernel-durations | 6.970 ms | 7.030 ms | −0.060 | +| kernel launches | 1163 | 817 | +346 | + +**77% of the gap is GPU idle. mcore's compute is at parity or better.** Both profiled +wall times track the unprofiled throughput (mcore 8.571 vs 8.787, vLLM 7.456 vs 7.53), +so these windows are representative and nsys is not distorting the totals. + +Per-component, profiled against profiled (ms/step): + +| component | mcore | vLLM | delta | +| --- | --- | --- | --- | +| expert GEMM | 2.502 (`_fused_moe_kernel`, 96) | 2.411 (2× cutlass `bmm`, 94) | +0.09 | +| attention | 0.796 (FA2 sm100, 48) | 0.820 (`fmhaSm100f`, 48) | −0.02 | +| collectives | 0.847 | 1.629 | **−0.78** | +| dense GEMM + splitK reduce | 1.036 | 0.910 | +0.13 | +| routing / align / sum / topk | 0.559 | 0.702 | −0.14 | +| norms | 0.515 (193 launches) | 0.426 (189) | +0.09 | +| misc elementwise / rotary / kv | 0.643 (241+96) | 0.125 (48) | **+0.52** | + +**This retracts the Session 16 target list.** `FULL-ATTRIB-S16` put the dense projections +at 20-29% of roofline with ~0.8 ms of headroom; in the clean trace they are 1.036 ms +against vLLM's 0.910 — **0.13 ms apart, not 0.8**. The in-situ event pairs inflated every +site they measured (~2.2 µs/pair × 2 pairs × 48 layers), which is why the instrumented +sites summed to 7.747 ms against a clean 6.97 ms. **Do not size a target from in-situ +numbers; size it from a clean trace and only use in-situ for apportionment.** Likewise +`expert_gemm_fc2` at "56% of roofline" is not a vLLM deficit: vLLM's FC2 equivalent is +only 0.09 ms cheaper across *both* expert GEMMs. + +Where the 1.645 ms of idle sits (`gap_profile.py`): + +| gap band | mcore | vLLM | +| --- | --- | --- | +| >100 µs | 0.840 ms (2.1 gaps) | 0.399 ms (1.0 gap) | +| 20-100 µs | 0.163 ms (3.7) | ~0 | +| 5-20 µs | 0.070 ms (5.8) | 0.002 ms | +| 2-5 µs | ~0 | 0.343 ms (95.6) | +| <1 µs | **0.535 ms (1033 gaps)** | 0.038 ms (240) | + +Two co-equal levers, and one non-lever: + +1. **Two once-per-step host stalls, 0.776 ms** — a ~492 µs gap after + `at::native::index_elementwise_kernel` and a ~289 µs gap after + `CatArrayBatchedCopy_vectorized` (`torch.cat`), once each per step. vLLM's single + equivalent stall is 0.30 ms, so **~0.48 ms is available**. On the rare steps where + this blows out to 4 ms, the API trace shows the mechanism: **278 `cudaMemcpyAsync` + + 270 `cudaStreamSynchronize` in one gap** — a per-request D2H sync storm that should be + one batched copy. This is the between-step Python phase that HOST-IDLE-S13/14/15 + hunted and missed; GPU-side gap analysis found it where host sampling could not. +2. **Serialization, ~0.5 ms** — mcore has 1033 sub-µs gaps against vLLM's 240, and gets + essentially *no* stream concurrency (union 6.926 ≈ sum 6.970) while vLLM recovers + 0.36 ms through overlap (union 6.674 vs sum 7.030). Driven by kernel count: the +346 + excess launches are **~241 standalone elementwise kernels** (`elementwise_kernel` 96, + `triton_poi_fused_add_copy__0` 96, `vectorized_elementwise` 49 — the residual adds, + which vLLM fuses into its RMSNorm, cf. the `add` in + `triton_red_fused__to_copy_add_mean_mul_pow_rsqrt`), **90 `splitKreduce` kernels** that + vLLM has none of, and 96 rotary/append-kv against vLLM's 48. Caveat: node-level graph + tracing costs per node, so an unknown part of the 0.535 ms is CUPTI overhead that mcore + pays more of by having more nodes. The kernel-count difference underneath is real. +3. **Not a lever: the collectives.** mcore is 0.78 ms/step *ahead* on comm, consistent + with `VLLM-COMM-S16`. It is the single largest thing mcore is winning, and the reason + the compute totals come out even despite the elementwise sprawl. + +### Session 16 — conclusion & recommendation + +The gap to vLLM is ~1.14 ms/step (29,133 vs 33,994 tok/s at a ~7.97 ms step). The +block is now fully attributed, and there is more identified headroom (~1.6 ms +excluding the collectives) than gap. Ordered by headroom-per-unit-risk: + +0. **The dense attention projections, ~0.8 ms of headroom** (FULL-ATTRIB-S16) — the + cheapest target and the one this session found by accident. Two skinny + M=256 GEMMs at 20-29% of their own bounds, latency-bound rather than + bandwidth-bound, with no cross-rank behaviour to reason about. Start here. +1. **The NVLS collectives, 1.47 ms — but demoted, see VLLM-COMM-S16.** vLLM runs the + same AllGather+ReduceScatter at the same cadence and pays for it too, so this is + not 1.47 ms of deficit against vLLM; the differential is unmeasured and much + smaller. Removal is impossible (they carry the tokens) and replacing the transport + already failed once (QWEN-017), and both sides are barrier-latency bound rather + than bandwidth bound, so the only lever is **overlap**: chunk the expert GEMM and + overlap combine of chunk *i* with the GEMM of chunk *i+1*, or prefetch the next + layer's dispatch. The 81% layer spread says part of this cost is skew absorption, + so some of it may be recoverable by better arrival balance rather than by + overlap. Worth doing eventually, but it is no longer the item that "alone would + close the gap." +2. **FC2, ~0.3–0.6 ms of apparent inefficiency** — 26.5 µs/layer against 14.8 µs at + the corrected roofline, and against FC1 doing twice the work for 1.4× the time. + Cheaper to investigate than item 1 and independently useful. Note + `MCORE_MOE_GEMM_TUNE` already tunes FC1/FC2 tiles separately, so this is a + question about that tuning, not about its absence. +3. **Recount the attention KV bytes** before any attention work, per HBM-BW-S16. +4. Precision (FP8 KV cache, FP8/NVFP4 experts) remains the only lever on whatever + genuinely is bandwidth-bound, and still needs its own correctness evaluation. + +### GRAPH-LAUNCH-S17 — the 191 µs `cudaGraphLaunch` was a profiler artifact; real cost is 9 µs + +`GAP-DECOMP-S17` decomposed mcore's 1.645 ms/step of GPU idle and found +`cudaGraphLaunch` inside those gaps at **4 calls/step × 191.6 µs = 761 µs/step**, which +would have made it the single largest item in the idle budget. Measuring it without a +profiler (`megatron/core/inference/graph_launch_timing.py`, `MCORE_GRAPH_LAUNCH_TIMING=1`, +`perf_counter` around `torch.cuda.CUDAGraph.replay`) shows the hot graph replaying in +**median 9.0-10.0 µs, p90 ~12 µs** over 1404 replays. The profiled figure was inflated +**~20×** by `--cuda-graph-trace=node`, which makes the driver do per-node bookkeeping +inside the launch. + +Graph launch is therefore ~40 µs/step, not 761 — **not a lever**, and the two 460-1000 µs +replays that do exist are `n=1` first-touch of each freshly captured bucket, not steady +state. + +**This partially retracts `GAP-DECOMP-S17`.** Its per-component *kernel* numbers stand — +kernel durations are barely perturbed, and that half of the analysis is what identified +the kernel-count gap. But the **magnitude and attribution of the 1.645 ms idle are not +trustworthy**, because the host side of a node-traced graph launch is inflated by an +amount that happens to be the same order as the whole idle budget, and mcore pays more of +it than vLLM by having more nodes. Do not quote "77% of the gap is GPU idle" without this +caveat. What survives, and is confirmed by QWEN-036 below, is the weaker and still +actionable claim: **mcore runs ~42% more kernels per step than vLLM and that costs real +throughput.** + +Method note for the next session: a profiler cannot be used to size an overhead the +profiler itself creates. Any host-side cost read out of a node-traced profile must be +re-measured with `perf_counter` before it becomes a target. + +### FUSION-INERT-S17 — the residual-add fusion was written, gated on, and never firing + +`QWEN-026/027/028` built `fused_add_rmsnorm` for both transformer-block boundaries. +`MCORE_FUSED_ADD_NORM=1` was in the standing gate list and the config that produced +`nsys_trace/mcore-final-jul27` is labelled "all ten gates on" — yet that trace contains +**zero** fused-add-norm kernels and exactly the two-kernel pattern the module's own +docstring describes as the thing it replaces (`triton_poi_fused_add_copy__0` at 96/step = +2/layer, plus 193 standalone TE RMSNorms). + +`megatron/core/inference/fusion_diag.py` (`MCORE_FUSION_DIAG=1`) reports each site's +guard conjunction once per distinct verdict. Two findings: + +1. Both sites **do engage** at 256 tokens. The blocked verdicts are all at + `(512, 1, 2048)` — prefill, correctly excluded by `MCORE_FUSED_ADD_NORM_MAX_TOKENS=256`. +2. The second site, `mlp_bda+next_layer_qkv_norm`, needs `MCORE_FUSED_ADD_NORM_QKV=1`, + which **was never in any gate list**. Worse, the two sites are *coupled*: + `_forward_pre_mlp_layernorm` only sets `self.mlp_norm_manager` when the attention-side + fusion did *not* fire, and the MLP-side guard requires `mlp_norm_manager is None`. So + the MLP-side fusion is eligible only on layers where the attention-side one already + fired — a dependency neither gate's name nor docstring mentions. + +Diagnostic lesson: report guards **once per distinct verdict, not once per site**. The +first version reported only the first call, which is prefill, and therefore said "BLOCKED" +about a fusion that engages fine in decode — the exact wrong conclusion. + +### QWEN-036 — enable the second fusion boundary (`MCORE_FUSED_ADD_NORM_QKV`): **+2.8%** + +Same node, same session, same allocation, back-to-back arms; only the one gate differs. +Node `nvl72104-T17`, Slurm `5792909`, BS256 / OSL1024 / 2 warmup + 5 timed. + +| arm | gates | per-iter throughput (tok/s) | mean | TPOT | +| --- | --- | --- | --- | --- | +| `a1base` | standing 12 + FA2 | 26803, 26871, 27111, 27176, 27947 | **27,182** | 9.420 ms | +| `a2qkv` | + `MCORE_FUSED_ADD_NORM_QKV=1` | 27589, 27718, 28092, 28172, 28200 | **27,954** | 9.159 ms | + +**+2.84%, −0.261 ms/step.** 3 of 5 `a2qkv` iterations beat *every* `a1base` iteration, and +the arms' means are separated by more than either arm's spread, so this clears the ~4% +per-iteration noise. Coherence checked in both arms. + +Note the absolute numbers: `a1base` is 27,182 on this node against the 29,133 recorded for +the same configuration in Session 16. **Node-to-node variation is larger than most of the +wins in this ledger** — roughly 7%. Never compare an arm against a historical number; +always re-measure the baseline in the same allocation. The gap to vLLM cannot be restated +from this session because no same-node vLLM run was taken. + +This is the first confirmation of the `GAP-DECOMP-S17` strategy: the change removes +kernels rather than making any kernel faster, and it converted directly into throughput. + +### QWEN-037 / QWEN-038 — two env-only follow-ups, both neutral + +Same allocation as QWEN-036, so each is directly comparable to the `a2qkv` arm +(27,954 tok/s) without re-running a baseline. + +| arm | change | mean tok/s | vs `a2qkv` | +| --- | --- | --- | --- | +| `b1maxtok` | `MCORE_FUSED_ADD_NORM_MAX_TOKENS=512` | 27,862 | −0.33% | +| `b2nosplitk` | `CUBLASLT_WORKSPACE_SIZE=1` | 27,811 | −0.51% | + +**QWEN-037** aimed at the 9.2 boundaries/step that still fall back to the unfused path +(96 possible − 86.8 fused). Raising the token ceiling did not recover them and cost +slightly, so those fallbacks are not token-count rejections. One is layer 47, whose +`mlp_bda` has no next layer to donate a norm to; the rest are most likely the coupling in +`FUSION-INERT-S17` (a layer is eligible on the MLP side only if the attention side fired). + +**QWEN-038** tried to delete the ~76 `splitKreduce` launches/step (0.185 ms of kernel +time) by starving the cuBLASLt workspace that split-K algorithms require. Neutral. Two +readings, and the second is the important one: either the heuristic ignored the limit, or +it switched to non-split-K kernels that gave back exactly what the reduce cost. **The run +did not verify which**, so "split-K removal doesn't pay" is not yet established — only +"this way of attempting it doesn't." A profile of the arm would settle it. + +### LAUNCH-VS-WORK-S17 — the QWEN-036 win came from removed GPU work, not removed launches + +Worth pinning down, because the two readings of QWEN-036 point at completely different +next targets, and the launch-count reading is the wrong one. + +The fusion removed **111 launches/step** *and* **~0.386 ms/step of kernel time** in the +four affected kernel families: + +| family | before (n × µs) | after (n × µs) | +| --- | --- | --- | +| `triton_poi_fused_add_copy__0` | 96 × 1.27 | 9.2 × 2.35 | +| `at::native::elementwise_kernel<128,4>` | 96 × 2.56 | 52.1 × 3.11 | +| `rmsnorm_fwd_tuned` | 85.1 × 2.73 | 9.8 × 3.36 | +| `rmsnorm_fwd_general` | 96 × 1.95 | 8.3 × 3.50 | +| `_fused_add_rmsnorm_kernel` | — | 86.8 × 1.79 | +| **total** | **0.787 ms** | **0.401 ms** | + +The measured end-to-end saving was **0.261 ms/step — 68% of the 0.386 ms of kernel time +removed**, which is what partial overlap predicts. It is *not* explained by launch count: +QWEN-038 attacked 76 launches carrying 0.185 ms and returned nothing, and +`GRAPH-LAUNCH-S17` already showed launch submission costs ~10 µs, not the ~180 µs the +launch-bound reading needs. + +**So size every remaining candidate by the GPU work it deletes, not by how many kernels it +deletes.** Under that rule the routing-mask kernel — 48 launches/step but only +48 × 0.97 µs = 0.047 ms — is worth at most ~0.5% and is *not* the next lever, despite +being the most attractive item on a launch-count ranking. This also retires the last +operational use of `GAP-DECOMP-S17`'s "+346 excess launches" framing: the launch delta is +a symptom of unfused work, and the work is what costs. + +Ranked by kernel time, the remaining decode candidates are `_moe_sum_kernel_fast` +(48 × 5.81 = 0.279 ms), `_align_single_kernel` (39.5 × 4.67 = 0.185 ms), the residual +`elementwise_kernel<128,4>` (52.1 × 3.11 = 0.162 ms, identity not yet established), and +`_fused_qk_rmsnorm` (0.105 ms). Caveat on all four: these averages come from a profile +whose window was not restricted to steady-state decode, so they are inflated by prefill +instances and are ranking hints, not budgets. + +### QWEN-039 — raise `CUDA_DEVICE_MAX_CONNECTIONS` from 1 to 8: **+0.9%, provisional** + +`GAP-DECOMP-S17` observed that mcore gets essentially no stream concurrency (interval +union 6.926 ms ≈ sum-of-durations 6.970 ms) while vLLM recovers 0.36 ms through overlap. +The harness pins `CUDA_DEVICE_MAX_CONNECTIONS=1` (`dev/moe_fused/run_e2e_cfg.sh`), which +puts every stream on one hardware channel and serializes work that has no data dependency +— notably the shared-expert stream against the main stream. Env-only test: + +| arm | `CUDA_DEVICE_MAX_CONNECTIONS` | mean tok/s | vs `a2qkv` | +| --- | --- | --- | --- | +| `a2qkv` | 1 | 27,954 | — | +| `c2conn4` | 4 | 27,999 | +0.16% | +| `c1conn8` | 8 | **28,214** | **+0.93%** | + +**Provisional, not confirmed.** All 5 of `c1conn8`'s iterations sit above `a2qkv`'s mean, +which is suggestive, but the two arms' per-iteration spreads overlap (553 and 611) and the +allocation expired before a repeat could be run. **Re-run `a2qkv` and `c1conn8` back to +back before believing +0.9%**, and note that `=1` is often set deliberately to keep +collectives from being starved of channels, so a confirmation should check the collective +timings too, not just throughput. + +Session best configuration: standing 12 gates + FA2 + `MCORE_FUSED_ADD_NORM_QKV=1` + +`CUDA_DEVICE_MAX_CONNECTIONS=8` = **28,214 tok/s, +3.80%** over the same node's baseline +of 27,182. + +### BUDGET-S17 — first trustworthy per-kernel decode budget, and how to window a trace + +Every per-kernel table before this one was taken over a window that included model load, +graph capture, warmup and the gaps *between* benchmark iterations. That inflates +`wall/step` without bound and pollutes every average with prefill instances — it is why an +earlier pass of this same trace reported `wall/step=199 ms` and a 6.5 ms +`_fused_metadata_kernel`. `nsys_trace/steady_window.py` locates the window from the cadence +of the one-per-layer-per-step collective, takes the last *N* whole steps, and **warns when +the span still contains an inter-iteration gap**. Narrowing 150 → 60 steps moved +`wall/step` from 34.6 ms to 7.217 ms and `_fused_metadata_kernel` from 13,182 µs to 116 µs. + +> **Always run the window guard.** Two sessions' worth of component sizing was distorted by +> unwindowed spans. If `wall/step` does not land near the throughput-derived step time, +> the window is wrong and nothing else in the table means anything. + +Config: both fusion boundaries on, FA2, BS256, **OSL128** (`wall/step` 7.217 ms; the +OSL1024 regime is 9.16 ms, the difference being KV-length-dependent attention). 60 steps, +rank 0, `sum-of-durations` 5.627 ms/step over 1025 launches. + +| ms/step | n/step | avg µs | group | +| --- | --- | --- | --- | +| 1.949 | 96 | 20.31 | `_fused_moe_kernel` — expert GEMMs, **35% of all kernel time** | +| 0.769 | 34-14 | 3.4-6.7 | dense projection GEMMs (`nvjet_*`) | +| 0.759 | 48+48 | 9.33/6.49 | NVLS collectives (reduce-scatter 0.448 + all-gather 0.311) | +| 0.558 | 34-14 | 4.3-10.4 | FlashAttention decode + splitkv combine | +| 0.218 | 48 | 4.54 | `_moe_sum_kernel_fast` | +| 0.217 | 48 | 4.52 | `_align_single_kernel` | +| 0.179 | 48+34 | 2.05/2.33 | `cublasLt::splitKreduce` | +| 0.168 | 95 | 1.77 | `_fused_add_rmsnorm_kernel` (the QWEN-036 fusion) | +| 0.127 | 48 | 2.66 | `at::native::elementwise_kernel<128,4>` — 1/layer, **identity unknown** | +| 0.085 / 0.084 / 0.070 / 0.062 / 0.036 | 48 each | ~0.8-1.8 | qk_rmsnorm, rotary, softmax_topk, append_kv, mask_routing_padding | + +Sized by `LAUNCH-VS-WORK-S17`'s rule (removed GPU work × ~68% conversion, against the +9.16 ms OSL1024 step), the ranked remaining candidates are: + +1. **`_fused_moe_kernel`, 1.949 ms — the only large target left.** No fusion helps here; + it is bandwidth-bound (FC1 at 80% of achievable HBM, FC2 at 56%). The lever is fewer + bytes, i.e. **FP8 expert weights**, worth up to ~0.9 ms. This is the one item whose + payoff justifies a large implementation. +2. `_moe_sum_kernel_fast` + `_align_single_kernel`, 0.435 ms combined → ~1.5-3% if folded + into neighbours. Both are already single-kernel-per-layer, so this means changing what + they are fused *with*, not fusing them together. +3. `elementwise_kernel<128,4>`, 0.127 ms → ~0.9%. **Identify it first**: one per layer at + 2.66 µs, and if it is a stray cast or copy it may be removable outright rather than + fused. Cheapest next step in the list and the only one that starts with a question + rather than a kernel. +4. `_mask_routing_padding_kernel`, 0.036 ms → ~0.3%. Explicitly **not worth it**, recorded + here only because a launch-count ranking puts it near the top (48/step). + +### FP8-WEIGHTS-S18 — candidate 1 above is dead: fewer bytes does not mean less time + +Weight-only fp8 (e4m3, one fp32 scale per output channel, activations and accumulation +unchanged) was implemented end to end — `moe/fp8_experts.py`, an `FP8_WEIGHTS` path in +`_fused_moe_kernel`, quantized buffers built in `_build_concatenated_weights` — and it +**loses at every tile configuration**, so the 0.9 ms above is not available this way. + +| GEMM | best bf16 | best weight-only fp8 | fp8 penalty | +| --- | --- | --- | --- | +| FC1 (+SwiGLU epilogue) | **36.24 µs** @ M16 N64 K64 w4 s3 | 40.08 µs @ M32 N64 K128 w8 s4 | +10.6% | +| FC2 | **18.86 µs** @ M32 N256 K64 w8 s4 | 22.24 µs @ M32 N256 K128 w4 s3 | +17.9% | + +Numerics were never the problem: cosine similarity 0.9989, mean relative error 6.6e-2 +concentrated in near-zero outputs, weight bytes 302.0 → 151.5 MB/layer exactly as intended. + +**Why halving the bytes cannot help here.** With bf16 activations every fp8 weight tile must +be widened back to bf16 in registers before the MMA, one convert per weight element. The +best bf16 configs run at 5.55 TB/s (FC1) and 5.34 TB/s (FC2) — ~90% of this part's +achievable HBM — so the byte saving is worth at most ~45% of each kernel, and the converts +cost more than that. The fp8 arms land at 40/22 µs against an fp8 *bandwidth* floor of +~18/9 µs, i.e. once the bytes are halved the kernel is no longer bandwidth-bound at all. +The only version of this lever that can pay is **w8a8**, where activations are also fp8 and +the tensor cores consume both operands directly with nothing converted. + +> **Two measurement traps, both of which produced a wrong answer first.** +> 1. An earlier harness reported fp8 at 0.995× and was read as "no effect". It was timing +> **Python**: four Triton launches per call, ~153 µs/call, host-bound with the GEMM +> invisible underneath. Any A/B of a <50 µs kernel must time **graph replay**, not a +> launch loop — `harness_gemmgrid.py` captures one launch and replays it. +> 2. `_get_decode_tuned_configs`' shipped tiles were re-derived independently by this sweep +> and are **already optimal**, including the shared `BLOCK_SIZE_M=16`: M32 makes FC2 13% +> faster (21.33 → 18.86 µs) but FC1 10% slower, and the two share one indirection table, +> so M16 wins on the sum (57.6 vs 58.8 µs). Retuning tiles is closed. + +Combined with QWEN-014 (flashinfer cutlass, 0.921×) and QWEN-010 (torch grouped_mm, 0.82×), +**every dtype-preserving and vendor-kernel route into the expert GEMM has now been measured +and none beats the shipping Triton path.** The candidate list above should be read as: item 1 +requires w8a8 or nothing; items 2-4 are the only cheap ones left. + +### STATUS-S18 — mcore at 87.9% of vLLM; the remaining gap is packing, not work + +| | tok/s | ms/step | vs vLLM | +| --- | --- | --- | --- | +| vLLM DP4+EP (`VLLM-BASELINE` / `PROFILE-OSL1024`) | 33,994.5 | 7.53 | — | +| **mcore, current best** (`y2pdl`) | **29,898.3** | **8.56** | **87.9%** (vLLM 1.137x) | +| mcore at session-18 start | 28,594.3 | 8.95 | 84.1% | +| mcore at session-2 start (`SESSION2-BASE`) | 22,398.9 | 11.71 | 65.9% | + +Session 18 closed 0.35 ms of the 1.38 ms/step gap (**+4.56%**), from three changes that +stack: `QWEN-040` bf16 MoE combine (+2.5%), `QWEN-041` flashinfer trtllm-gen decode +(+1.7%), `QWEN-042` PDL on that kernel (+0.3%). + +**Attention is now at parity, and that closes the largest work-bucket gap.** In the final +trace it runs 22.9 us/launch against vLLM's 25.3 us -- both engines now execute the same +trtllm-gen kernel family, and what is left of the difference is window seqlen, not +kernel quality. The 0.409 ms/step deficit `GAP-S18` measured is gone. + +What remains, in order: + +1. **Serialization, ~0.3 ms.** mcore's sum-of-durations (5.922 ms) and interval union + (5.877 ms) differ by 0.8%, so **almost nothing overlaps**; vLLM hides 4.8% of its work. + mcore's 0.905 ms/step of collectives is fully exposed against compute. This is now the + single largest structural difference and it is a scheduling change, not a kernel. +2. **Dispatch overhead, 0.416 ms.** 806 sub-microsecond gaps per step, at 963 launches + against vLLM's 810. Priced by `BALLAST-S18` at ~0.5 us of gap per node, so each fusion + still pays twice. +3. **Work vLLM does not do at all**: splitK reduce 0.167 ms/step (82 launches, no vLLM + equivalent) and the 0.164 ms of copies whose origin is still open after four attempts + (`COPY-ID-S18` -- `SplitAlongDim` eliminated). + +> **Reading the trace tables above.** Per-bucket ms/step may only be compared *within* one +> trace. This window's attention is 22.9 us/launch against the harness's 45.1 us at KV +> length 512, which puts its average KV length near 260 -- so its absolute times are +> cheaper than the run average and not comparable to the earlier traces'. The throughput +> column is the authority; the buckets rank levers, they do not size them. + +### QWEN-042 — PDL on the flashinfer decode kernel: **+0.30%** + +| arm | gates | tok/s | vs control | +| --- | --- | --- | --- | +| x2fi | control (flashinfer decode) | 29,802.9 | — | +| y1pdl | + `MCORE_FLASHINFER_PDL=1` | 29,888.2 | +0.29% | +| y2pdl | repeat | 29,898.3 | +0.32% | + +Programmatic Dependent Launch lets the kernel's prologue start while its predecessor +drains, aimed at the 0.416 ms/step of sub-microsecond gaps. Small but reproducible across +two arms, and above the 0.5% iteration spread only because both arms agree; kept on the +strength of the repeat rather than the single measurement. + +### QWEN-041 — Blackwell-native decode attention via flashinfer trtllm-gen: **+2.61%** + +| arm | gates on top of standing set (incl. `MCORE_NVLS_RS_BF16=1`) | tok/s | vs control | +| --- | --- | --- | --- | +| x1base | control (FA2 decode) | 29,043.8 | — | +| **x2fi** | **`MCORE_FLASHINFER_DECODE=1`** | **29,802.9** | **+2.61%** | + +`GAP-S18` put attention 0.409 ms/step behind vLLM at an *identical* launch count, which +rules out anything Megatron does around the call and points at the kernel generation: +Megatron's decode runs FA2's `flash_attn_with_kvcache`, whose kernels predate Blackwell, +while vLLM reaches `fmhaSm100f` through flashinfer. flashinfer 0.6.14 turned out to be +**already installed** in the same venv as flash-attn, so this needed no new dependency. + +Measured in isolation first (`dev/moe_fused/harness_attn.py`), at B=256, 32 q heads, +4 kv heads, D=128, under graph replay so the number is device time and not Python: + +| KV len | FA2 | trtllm-gen | delta | +| --- | --- | --- | --- | +| 512 | 59.6 us | 45.1 us | −24.3% | +| 1024 | 105.8 us | 80.7 us | −23.7% | +| 2048 | 199.3 us | 151.9 us | −23.8% | + +A flat ~24% across the range, with outputs matching FA2 to bf16 tolerance (rel ~3e-3). +Two properties made the integration small: with `kv_layout="NHD"` the kernel takes +**exactly the paged layout Megatron already has**, and unlike flashinfer's wrapper APIs +the trtllm-gen entry point needs no host-side `plan()`, so it captures in the decode +graph. Page size is irrelevant to it (256/128/64 all within 0.1 us), so Megatron's +256-token pages stay -- and FA2 is the stricter of the two, *requiring* pages be a +multiple of 256. + +The e2e gain (+2.61%) is smaller than the kernel gain (24% of a 1.36 ms bucket would be +0.33 ms of an 8.8 ms step, or 3.8%) because the removed device time partly un-hides host +work, the same ~68% conversion `LAUNCH-VS-WORK-S17` measured for the add-norm fusion. + +> **Lesson.** The gap said "attention, same launch count, more time", which reads as *our +> kernel is the wrong generation for this GPU* -- and the right kernel was already sitting +> in the venv. Before treating a kernel-quality gap as a porting project, check what the +> reference implementation calls and whether it is installed: two sessions of FA2-vs-FA4 +> flag flipping never left the flash-attn package, where the answer was not. + +### QWEN-040 — reduce the MoE combine in bf16, not fp32: **+2.37%** + +| arm | gates on top of standing set | tok/s | vs control | +| --- | --- | --- | --- | +| s1base | control (FA2) | 28,594.3 | — | +| **t1rsbf16** | **`MCORE_NVLS_RS_BF16=1`** | **29,271.8** | **+2.37%** | +| t2rsbf16 | repeat of t1 | 29,354.2 | +2.66% | +| t3rsconn8 | + `CUDA_DEVICE_MAX_CONNECTIONS=8` | 29,346.3 | +0.0% vs t2 | +| s2fa4 | FA4 (version pin removed) | 27,716.1 | −3.07% | +| s3conn8fa4 | FA4 + `CUDA_DEVICE_MAX_CONNECTIONS=8` | 27,812.6 | +0.35% vs s2fa4 | + +Found by following `COPY-ID-S18`'s post-reduce-scatter copy back to its cause rather +than trying to fuse it away. The copy is `output.to(torch.bfloat16)` in +`NVLSAllGatherVDispatcher.combine`, and it exists because the `ep_rsv` symmetric buffer +is allocated fp32 — so the MoE writes fp32 through `_moe_sum`'s `out=`, the combine +reduce-scatter moves **twice the NVLink bytes it needs to**, and every layer pays a cast +on the way back to a bf16 residual stream. Allocating the buffer bf16 fixes all three at +once: `_moe_sum`'s `tl.store` casts on the way in for free, and `output.to(bfloat16)` +becomes a no-op that returns its argument, so the cast kernel disappears. + +The precision cost is far smaller than "reduce in bf16" suggests, and this is the part +worth remembering: `multimem.ld_reduce` **accumulates in f32 regardless** — `REDUCE_F32` +only selects whether the operands it loads are `f32` or `bf16x2`. So the change halves +the bytes on the step's largest collective while keeping f32 accumulation in hardware, +and it lands where vLLM already is (`ncclDevKernel_ReduceScatter_Sum_bf16`). Greedy +decode stays coherent on all three probes and the five iterations span 0.4%. + +Two arms, 29,271.8 and 29,354.2, put the mean at 29,313 (**+2.52%**), and TPOT drops from +~8.95 to 8.75 ms — 0.23 ms/step against the 0.29 ms predicted from the trace, so the +accounting holds. Stacking `CUDA_DEVICE_MAX_CONNECTIONS=8` on top adds nothing (29,346), +which retires `QWEN-039`: that flag's earlier +0.4-0.9% was relieving contention on the +same fp32 reduce-scatter this change shrinks, so the two are one win, not two. + +Left env-gated and **off by default**: it is a numerics change, and enabling it by +default needs an accuracy run (lm-eval or equivalent), not three greedy-decode probes. + +**Mechanism confirmed by launch count, not by time.** A trace captured with the gate on +shows `bfloat16_copy_kernel` instances falling from 258,564 to 4,572 whole-trace (−98%) +and the copy bucket falling from 107 to 59 launches/step — exactly 48, one per layer, +the cast that is now a no-op. Counts are the right evidence here because the two traces' +*times* are not comparable: the second window's expert GEMM is 27% cheaper (2.436 → +1.777 ms/step) and its attention launches went 50 → 63, neither of which a +reduce-scatter dtype can cause. The two 60-step windows simply sit at different average +sequence lengths, so per-bucket ms/step may only be compared *within* a trace. The +throughput number in the table above, measured without a profiler attached, is what +sizes the win. + +> **Lesson.** A "removable copy" is usually a symptom. Three sessions treated the +> per-layer copies as fusion targets worth ~0.06 ms each; the copy was actually pointing +> at a **buffer dtype** decision upstream worth 2.5%. When a copy shows up next to a +> collective, check what dtype the collective is moving before trying to fuse the copy. + +### FA4-RECHECK-S18 — the FA2 pin survives upstream's split-KV fix + +`QWEN-0xx` pinned `MCORE_FLASH_ATTN_VERSION=2` because FA4 hardcoded `num_splits=1` and +so could not split KV across SMs at decode. The rebase brought `num_splits=0` (auto), so +the rejection was re-tested per the rebase rule — and FA4 is **still 3.07% slower** +(27,716 vs 28,594). `num_splits` was therefore not the dominant term; the FA4 path's +routing through the varlen interface is. Attention stays the largest work-bucket gap +(+0.409 ms/step) with no env-level lever left, so closing it needs a different kernel +(flashinfer/TRT-LLM-gen decode), not a flag. + +### GAP-S18 — the gap is two-thirds packing, one-third work, and the expert GEMM is done + +First decomposition built from a **fresh** mcore trace (current gate set, OSL1024, 60 +steady steps, no window warnings) against the vLLM baseline trace, both bucketed by +`nsys_trace/compare_budget.py` and both measured for interval-union busy by +`nsys_trace/union_window.py`. Every prior version of this comparison used the Jul 27 +mcore trace, which predates the FA2 pin and both add+norm fusions and therefore +overstated attention, norm and copies. + +| | vLLM | mcore | delta | +| --- | --- | --- | --- | +| GPU busy, interval union | 6.989 ms | 7.478 ms | +0.489 | +| exposed gap (true step − union) | ~0.54 ms | ~1.43 ms | **+0.89** | +| step time, unprofiled | 7.53 ms | 8.91 ms | +1.38 | + +Per-bucket device time (ms/step, launches/step), mcore minus vLLM: + +| bucket | vLLM | mcore | delta | +| --- | --- | --- | --- | +| attention | 1.214 (48) | 1.623 (50) | **+0.409** | +| elementwise / copy | 0.004 (1) | 0.232 (107) | +0.228 | +| splitK reduce | 0 (0) | 0.204 (95) | +0.204 | +| MoE finalize | 0.215 (47) | 0.261 (48) | +0.046 | +| expert GEMM | 2.376 (94) | 2.436 (96) | **+0.060** | +| norm | 0.251 (95) | 0.272 (145) | +0.021 | +| MoE routing | 0.475 (142) | 0.367 (144) | −0.108 | +| collective | 1.607 (94) | 1.035 (96) | **−0.572** | +| **sum of durations** | **7.341 (810)** | **7.521 (1027)** | **+0.180 (+217)** | + +Three conclusions, in order of how much they should change what gets worked on: + +1. **Only 0.18 ms of the 1.38 ms gap is extra kernel work.** The expert GEMM is at + parity (+0.06 ms) and mcore's collectives are *half a millisecond faster* than + vLLM's. Four sessions of MoE GEMM work, `FP8-WEIGHTS-S18` included, were aimed at a + bucket that had already converged. +2. **mcore overlaps nothing.** Sum-of-durations 7.521 vs union 7.478 means 0.6% of its + work runs concurrently with other work; vLLM hides 0.352 ms (4.8%). mcore's 1.035 ms + of collectives is fully serialized against compute. Matching vLLM's overlap ratio is + worth ~0.3 ms and is a scheduling change, not a kernel. +3. **The gaps are launch-shaped.** 892 of mcore's 934 per-step gaps are sub-microsecond + and sum to 0.456 ms — consistent with `BALLAST-S18`'s 1.34 µs/node (~0.85 µs kernel + + ~0.5 µs gap) across 1027 launches vs vLLM's 810. Every removed launch pays twice. + +### COPY-ID-S18 — both copy families named from the trace, after three failed hooks + +The 107 copy launches/step were finally identified **positionally in the trace**, by +aggregating each copy's immediate predecessor and successor, after three attempts to +name them with a `TorchDispatchMode` hook produced nothing usable: + +| n/step | kernel | sits between | what it is | +| --- | --- | --- | --- | +| 47 | `vectorized_elementwise_kernel<8>` (bf16 copy) | `_multimem_reduce_scatter_v` → `_fused_add_rmsnorm` | materializes the MoE output that the next kernel immediately reads | +| 48 | `elementwise_kernel<128,4>` | QKV GEMM (`nvjet_*`) → `_fused_qk_rmsnorm` | contiguous copy of the strided Q/K slices | + +The first one turned out to be a dtype symptom, not a fusion target, and is fixed in +`QWEN-040` (+2.5%). The second is still open, and the guess in the first version of this +entry -- "make `_fused_qk_rmsnorm` stride-aware" -- was wrong: that kernel is *already* +stride-aware and takes no-copy `[-1, head_dim]` views. So the copy is **upstream of the +norm**, between the QKV GEMM and it, and at 2.66 us it is about the size of a full +round-trip of the QKV tensor (256 tok x 5120 x 2 B, read+write, at ~3 TB/s). + +The prime suspect was `SplitAlongDim` in `Attention.get_query_key_value_tensors` +materializing q/k/v instead of returning views. **Tested and eliminated**: taking the +`torch.split` branch instead (`MCORE_QKV_SPLIT_VIEWS=1`, arm `u1qkvsplit`) gives +29,307 tok/s against the 29,313 control -- exactly neutral. Either TE's split already +returns views, or `torch.split` + the `[sq, b, np, hn]` reshape materializes the same +bytes. The copy's source remains open; what is now known is that it is QKV-sized, sits +between the QKV GEMM and the norm, and is not the split. + +A fourth attempt at naming it with `copy_trace` landed in prefill again -- it reported +`rmsnorm.py:182 op_forward` doing `.contiguous()` on a strided `(512, 1, 4, 128)` K +tensor, which is TE's norm op on a 512-token prefill, not the decode copy. That is a +real finding for prefill but not the one being chased, and it is the fourth time this +hook has answered a different question than the one asked. + +> **Lesson on the failed hook.** `copy_trace` was armed three ways: on the first layer +> forward (caught the one-time lazy expert-weight consolidation and reported 32 weight +> copies per layer as if per-step), on `inference_context.is_decode_only()` (never fired +> — under graph capture the context does not reach the layer as a keyword), and finally +> on a call counter. Meanwhile the trace already contained the answer. **When a kernel +> needs identifying and a profile exists, read the profile's neighbours first**; +> instrumenting the framework to re-derive what the trace already recorded cost three +> allocation slots here. + +## Session 19 (2026-08-03) — the QKV copy, named by arithmetic + +`COPY-ID-S18` left one kernel unattributed after four attempts: a 2.8 us +`elementwise_kernel<128,4>` between the QKV GEMM and the q/k norm, once per layer per +step, QKV-sized. It is the `reshape`, and it is provable on a laptop in ten seconds with +no GPU, no cluster, and no profiler: + +``` +q slice [sq, b, ng, (np/ng) * hn] strides (.., .., 1280, 1) +q reshape [sq, b, np, hn] strides (.., .., 128, 1) COPY MATERIALIZED +``` + +Megatron's QKV projection writes `[sq, b, ng, (np/ng + 2) * hn]` -- each group's k and v +head sit *between* consecutive groups' q heads. Merging the group axis with the head axis +therefore requires the group stride to equal `(np/ng) * hn`; it equals +`(np/ng + 2) * hn`. For Qwen3-30B that is **1024 required against 1280 actual**, so +`Tensor.reshape` cannot return a view and silently copies the whole query instead. + +### Why four attempts missed it + +All four tested *the split*: `SplitAlongDim` against `torch.split` (`u1qkvsplit`, exactly +neutral), then three `copy_trace` arms. The neutral split result was read as "the copy is +elsewhere", when it was itself the answer -- both split branches feed the same reshape, so +of course they measured the same. The one line never questioned was the one carrying a +comment that explained it as a shape change. + +> **Lesson.** A layout claim is decidable from strides alone, and `reshape` reports a copy +> as silently as it reports a view. When a profile shows a copy the size of a known +> tensor, walk the *strides* of every reshape/view/permute between that tensor's producer +> and its consumer before instrumenting anything. Four GPU sessions guessed at this; one +> CPU-only `python3 -c` settled it. + +### The fix is free because the norm already pays for the write + +The fused q/k norm allocates its output anyway, so it can repack on the way through: read +the query with the two strides the grouped layout needs (one per group, one per head +inside a group) and write the contiguous `[sq, b, np, hn]` result the rest of attention +wants. Same values in the same order, so the result is **bit-identical** -- the right +acceptance bar here, because a head-permutation bug would produce individually plausible +values in every head and sail past a tolerance check. + +| | reshape + norm | grouped norm | +|---|---:|---:| +| kernel, 256 tokens (graph replay) | 12.30 us | **8.22 us** | +| launches per layer | 2 | **1** | +| per step, 48 layers | — | **0.196 ms saved** | + +### QWEN-043 — grouped-read q/k norm: **+1.5%** + +Four alternating arms on one node, so node drift shows up as base-to-base spread: + +| Arm | Gates | Steady-state tok/s | +|---|---|---:| +| `h2base` | standing set | 29,828 | +| `h3grp` | `+MCORE_GROUPED_QK_NORM=1` | 30,312 | +| `h4base` | standing set | 29,858 | +| `h5grp` | `+MCORE_GROUPED_QK_NORM=1` | 30,284 | + +Base-to-base spread **0.10%**, so the effect is ~15x the noise. Pooled +**29,845 -> 30,298 tok/s (+1.52%)**, i.e. 8.58 -> 8.45 ms/step -- 0.13 ms of the +microbenchmark's 0.196 ms lands end-to-end. Generated text is identical character for +character, as bit-exactness requires. Shipped behind `MCORE_GROUPED_QK_NORM=1`; requires +the fused q/k norm, and declines when the query needs the tensor-parallel head slice. + +**Status: 30,298 against vLLM's 33,994.5 = 89.1% of vLLM** (from 87.9% at session start, +65.9% at the campaign's start). Remaining gap 0.90 ms/step. + + +### The trace confirms it, and nothing else moved + +Diffing the fresh trace against the previous one, same anchor and window, is the cleanest +attribution in this campaign: + +| bucket | before | after | delta | +|---|---:|---:|---:| +| elementwise / copy | 0.164 ms, 59 launches | **0.042 ms, 11 launches** | **-0.122 ms, -48** | +| every other bucket | — | — | within +/-0.012 ms, 0 launches | +| TOTAL | 5.923 ms, 963 | 5.804 ms, 916 | -0.120 ms, -48 | + +Exactly one launch per layer left the copy bucket and nothing else changed, which also +re-validates the window: the buckets the change cannot touch held still, the check +`union_window.py` demands after the 27%-seqlen-skew incident. The 0.122 ms of device time +matches the 0.13 ms measured end-to-end. + +### The last mcore-only bucket, attributed + +`splitKreduce` is the one category vLLM has none of (82 launches, 0.173 ms/step). Its trace +neighbours name both parents: + +| n/step | parent GEMM | successor | +|---:|---|---| +| 48 | `nvjet_..._2cta_h_bz_splitK_TNT` / `..._4x1_v_bz_splitK_TNN` | `_softmax_topk_kernel` | +| 34 | `nvjet_..._2cta_h_bz_splitK_TNT` | `_fused_add_rmsnorm_kernel` | + +So it is the **router GEMM** (one per layer, exactly 48) and the **attention output +projection** (34 of 48 layers; the rest pick a non-splitK algorithm). The router case is +not a mis-tuned heuristic: at M=256, K=2048, N=128 the output is a couple of tiles, so +splitting K is the only way cuBLASLt gets parallelism, and the reduce is the price of it. +That reframes the target -- not "stop cuBLASLt from splitting K", but "fuse the router", +which would collapse GEMM + splitK reduce + softmax-topk + padding mask (4 launches/layer, +~192 total) into one. A fused router must beat GEMM+reduce on its own terms, so it needs a +microbenchmark before any integration: 256 CTAs each reading the whole 512 KB weight is +128 MB of L2 traffic, and blocking tokens instead leaves 4 CTAs on 148 SMs. + +### QWEN-044 — eight rows per CTA in the q/k norm: **+0.35%** + +The grouped norm moves ~4.5 MB in 8.22 us -- about 570 GB/s on a GB200, because one CTA +per 128-wide row means 9,216 single-warp CTAs. A rows x warps sweep, keeping only +bit-exact candidates: + +| rows/CTA | warps | us | GB/s | vs 1x1 | +|---:|---:|---:|---:|---:| +| 1 | 1 | 8.21 | 574 | 1.00x | +| 2 | 2 | 6.17 | 765 | 1.33x | +| 4 | 4 | 6.15 | 767 | 1.34x | +| **8** | **8** | **4.12** | **1146** | **2.00x** | + +In-tree the same shape measures **6.15 us**, not 4.12 -- a 2 us gap between the harness +kernel and the shipped one at identical rows and warps that is **still unexplained** after +three hypotheses, each tested and each measuring 6.15 us unchanged: + +| Hypothesis | Why it was plausible | Result | +|---|---|---| +| Runtime output row stride blocks store vectorization | the harness hardcoded `HN`, the kernel took a stride argument | rejected, 6.15 us | +| `where(is_q, rows, 0)` in the store address hides contiguity | the harness addressed stores off the raw row index | rejected, 6.15 us | +| Two separate load address arrays instead of one shared | the harness built a single `where`-selected offset array | rejected, 6.15 us | + +**Closed in session 20: there is no gap.** The PTX comparison this pointed to was run +(`dev/moe_fused/harness_qknorm_ptx.py`) and measured the shipped kernel and the harness +kernel at **4.12 us each, ratio 1.00x** -- same answer for q and k, 21 registers, no +spills, `.v2` vectorized loads and stores. The shipped kernel had already reached the +harness number; the 6.15 us readings were stale, taken before the store-stride fix landed +and then carried forward as if current. The three hypotheses above were each rejected for +the right reason: every one of them was measuring a kernel that was already fast. + +The lesson is measurement hygiene, not Triton. A number that survives three +disconfirmations without moving is more likely stale than robust, and re-measuring the +baseline is cheaper than the next hypothesis. Nothing left to recover -- item closed. + +End-to-end, both arms from the same file so the launch shape is the only variable: +1x1 controls **30,341 / 30,399** against 8x8 **30,482** tok/s, i.e. **+0.35%** on a +control-to-control spread of 0.19%. Note the microbenchmark predicted 1.17% +(2.06 us x 48 layers on an 8.44 ms step) and a third of that landed -- a reminder that a +kernel-level win on a non-critical-path kernel is an upper bound, not a forecast. + +### QWEN-045 — fusing the router GEMM: **-2.37%, rejected** + +The router chain is four launches per layer: gating GEMM, a cuBLASLt splitK reduce (the +logits have to reach memory before top-k can read them), fused softmax+topk, and the +padding mask. `kernel_neighbors.py` attributed 48 of the 82 splitK reduces per step to +this GEMM, so collapsing all four into one Triton kernel that keeps the logit tile in +registers looked like the largest single launch-count win available. + +The microbenchmark endorsed it. Sweeping BLOCK_M x BLOCK_K x warps at the decode shape, +keeping only configs whose expert sets matched exactly: + +| BLOCK_M | BLOCK_K | warps | us | vs reference | +|---:|---:|---:|---:|---:| +| 32 | 128 | 8 | 16.41 | 0.94x | +| 16 | 128 | 4 | 12.68 | 1.21x | +| **16** | **256** | **4** | **12.31** | **1.25x** | + +Expert sets bit-exact, probabilities within 7.45e-09, and the folded padding mask correct +(200 real rows, the other 56 getting -1). Against the reference's 15.34 us that is +0.146 ms/step, ~1.7% of an 8.35 ms step. The winning shape is a *narrow* BLOCK_M with the +widest K tile -- it maximises CTA count, and the whole-weight read each CTA does is then +what hides the K loop. Note the in-tree defaults were initially the 0.94x config, which +would have tested a kernel slower than baseline; they were corrected before the A/B. + +End to end it lost, consistently, in both replicate pairs: + +| arm | tok/s | +|---|---:| +| control | 30,689.4 | +| fused router | 29,981.5 | +| control | 30,697.5 | +| fused router | 29,951.2 | + +**-2.37%** on a control-to-control spread of 8 tok/s. The sign is opposite to the +prediction and the magnitude is larger, so this is not noise and not a tuning problem. + +The mechanism worth carrying forward: the four launches it replaced were not serial dead +time. cuBLASLt's GEMM and its splitK reduce overlap with neighbouring kernels inside the +CUDA graph, whereas a 16-CTA Triton kernel occupies few SMs *and* forces every CTA to +stream the full weight matrix, which evicts what its neighbours are reading. Isolated +device time credits the fused kernel for work it removed from the critical path only in +a benchmark where nothing else was running. + +**Rule this establishes:** an isolated-kernel win does not transfer when the kernel it +replaces is one that overlaps well. Before fusing across a cuBLAS call, measure the +*ablation* ceiling (delete the work and run e2e) rather than the replacement's device +time -- the ablation prices the critical path, the microbenchmark prices the kernel. +Retained behind `MCORE_FUSED_ROUTER=0`. + +### QWEN-046 — the padding mask: **+1.02% available, fusion not yet working** + +Applying the rule from QWEN-045 to the cheapest member of the router chain. +`mask_routing_padding` writes -1 into every topk slot of the CUDA-graph padding rows so +those tokens route to no expert; it is one launch per layer per step (48/step) over a +256x8 int64 tensor, which is almost entirely launch overhead. Ablating it outright +(`MCORE_ABLATE_ROUTE_MASK=1`, correctness-breaking by construction, valid only as a +ceiling) priced the critical path directly: + +| arm | tok/s | +|---|---:| +| control | 30,662.5 | +| mask ablated | 30,965.1 | +| control | 30,661.0 | +| mask ablated | 30,983.4 | + +**+1.02%** on a control-to-control spread of 1.5 tok/s -- among the tightest replicate +pairs measured in this project, and a lower bound besides: skipping the mask lets padding +rows route to real experts, which if anything *adds* expert GEMM work. + +The fusion target is the router's `_softmax_topk_kernel`, which is one CTA per token and +already stores that token's index row, so the sentinel is a scalar compare and a select +with no extra launch. It was implemented (module-level publish of the context's +`int32[1]` count, `MASK_PADDING` constexpr in the kernel, and a dispatcher skip keyed off +a tag the router leaves on the tensor it masked) and it **hung in warmup with the gate +off**, which localises the fault to the always-on parts rather than to the masking logic. +Reverting the three files reproduced the baseline to within 1 tok/s (30,661.9 against +30,662.5 and 30,661.0), so the hang is attributable to those edits alone. + +Design, suspects, and a cheapest-first diagnostic ladder are written up in +`dev/moe_fused/NOTES-route-mask-fusion.md` so the next attempt starts from the hang +rather than from the design. Leading suspect is the kernel signature change itself: it is +the only always-on edit that alters generated code, and if `tl.where(False, -1, best_idx)` +fails to fold, every row gets -1 and the NVLS all-gather-v barrier can deadlock on +rank-divergent counts. Unmeasured guess, listed first in the ladder. + +### Session 21 status + +| | tok/s | ms/step | vs vLLM | +|---|---:|---:|---:| +| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | +| **mcore, current best** | **~30,751** | **8.31** | **90.5%** | +| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | + +Session 21 added no throughput either, but it changed what the campaign believes about +where the remaining 9.5% is, which had been wrong since the collectives were first +measured in situ. + +- **The collective floor is not a deficit; it is a lead.** Reading vLLM's MoE source + (QWEN-051) showed it runs the same all-gather-v / reduce-scatter-v pair over plain + NCCL with symmetric memory off by default, and our own bucket table has had the + number all along: vLLM 1.607 ms/step against mcore 1.035. Every remaining plan that + started with "hide the collectives" — DBO included — was aimed at a bucket where we + are 0.57 ms/step ahead. Those are now closed rather than deferred. +- **The MoE kernel, not the MoE communication, is the last large lever.** flashinfer's + `trtllm_bf16_moe` measures 63.3 us/layer of device time against mcore's 78.0 at the + operating point, captures into a CUDA graph, and replays bit-identically. Estimated + **+0.5 to +0.7 ms/step (6-8%)**, which is the difference between 90.5% and ~96-98% of + vLLM. Unbuilt: it needs a gate/up half-swap at load and has to displace the router + top-k, the route mask and the dispatcher permutation. +- Two ledger corrections that had been blocking that lever for sessions: the BlockMajorK + weight requirement is a ~1 s load-time transform rather than an incompatibility, and + the flashinfer MoE previously rejected at 8% slower was a *different kernel* from the + one vLLM actually prefers. The run venv already ships the kernel, so shipping it needs + no new dependency. +- One self-inflicted false alarm, retracted the same session (QWEN-052): the flashinfer + decode gate looked inert because the container's system python lacks flashinfer, but + the e2e runner uses a venv that has it. The gate is live and the baseline is intact. + Attention nonetheless remains **+0.409 ms/step** behind vLLM *with* the trtllm-gen + kernel active, which makes it the second-largest known gap after the MoE kernel. + +### Session 20 status + +| | tok/s | ms/step | vs vLLM | +|---|---:|---:|---:| +| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | +| **mcore, current best** | **30,662–30,697** | **8.35** | **90.2%** | +| mcore, with the mask ablated (not shippable) | 30,974 | 8.27 | 91.1% | +| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | + +Session 20 added **no throughput**: both candidates were rejected, one on measurement +(-2.37%) and one on a hang. The code is unchanged from session 19; the higher number +against session 19's recorded 30,482 is node-and-day drift, not a code change, measured +across five controls in one job (30,689 / 30,697 / 30,662 / 30,661 / 30,662). + +What session 20 did produce, all of which outlives it: + +- A priced target: **+1.02%** sitting in the padding mask, with the ceiling measured + rather than estimated, and the implementation already designed and written up. +- A rule that would have saved this session's larger experiment (QWEN-045): price the + ablation ceiling before fusing across a well-overlapped cuBLAS call. +- One closed open item (QWEN-044's phantom 2 us) and one retired measurement habit. +- **14 hours of queue time recovered.** Two jobs sat `PENDING (Priority)` for 14 h and + 7 h with 315 nodes idle, purely because the sbatch template omitted `--qos`: the + default `normal` is priority 100 against a queue head near 350k. Resubmitted under + `--qos=interactive` (priority 700, <=4 nodes) and `--qos=short` (priority 200, <=2 h), + both started within seconds. Every job template in `skills/` now sets `--qos`, and + `skills/run-qwen-model/SKILL.md` carries the diagnostic: if nodes are idle and + `squeue -j -o '%Q'` is far below the head of `squeue -p batch -t PD -S -Q`, it is + QOS and not contention. + +### Session 19 status + +| | tok/s | ms/step | vs vLLM | +|---|---:|---:|---:| +| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | +| **mcore, current best** (`hftune`) | **30,482** | **8.40** | **89.7%** | +| mcore at session-19 start (`y2pdl`) | 29,898 | 8.56 | 87.9% | +| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | + +Session 19: **+1.95%** from two changes to one kernel, both bit-exact. Remaining gap +0.87 ms/step. + +**Where the remaining gap is not.** Per-launch, mcore is now at or better than vLLM in +every functional bucket, including attention and the expert GEMM. The residue is +launch count (916 against 810) and packing: the overlap ratio is 1.008x against vLLM's +1.050x. And the >100 us gap band is *not* the lever it looks like -- mcore shows +0.893 ms/step there against vLLM's 2.727 ms/step in its own trace, so that band is +dominated by profiler-inflated host work on both sides and cannot be compared across +traces captured with different `--cuda-graph-trace` settings. + +### QWEN-047 — the collectives re-priced in place: latency-bound, exposed, unhideable without microbatching + +Session 21 opened by re-pricing every per-layer component with the in-graph external-event +timer (`MCORE_INSITU_TIMING=1` plus `MCORE_INFER_STEP_GPU_TIMING=1`, which is what actually +drives the report — the timer alone prints nothing). Motivation: the trace analysis had put +~0.3 ms on matching vLLM's overlap ratio, but that measurement predated the bf16 +reduce-scatter, which halved the combine's bytes on the wire. + +Per-layer, scaled x48, net of the calibrated 2.2 us/pair event overhead: + +| component | ms/step | component | ms/step | +|---|---:|---|---:| +| expert_gemm_fc1 | 1.650 | nvls_dispatch | 0.491 | +| expert_gemm_fc2 | 1.183 | qkv_proj | 0.482 | +| attn_core | 1.152 | moe_align | 0.449 | +| nvls_combine | 0.548 | moe_sum | 0.321 | +| out_proj | 0.533 | router_topk | 0.253 | +| | | **SUM of sites** | **7.060** | + +Two results, both load-bearing. + +**The collectives are latency-bound, not bandwidth-bound.** Dispatch plus combine is +1.039 ms/step against 1.035 ms measured before bf16 halved the combine's bytes — a 0.4% +move for a 2x reduction in traffic. With `QWEN-016`'s CTA-count result and the barrier +finding, that is now three independent measurements saying the same thing: **these +collectives cannot be made faster, only hidden.** Do not spend another session on the +transport. + +**The block is a serial chain, confirmed a second way.** Sites sum to 7.060 ms against a +block total of ~7.09 — about 100%. Parts can only reconstruct the whole like that if +nothing overlaps, which independently corroborates the trace-derived 1.008x overlap ratio +using a completely different instrument. + +**Why the 1.04 ms is not collectable today.** Hiding it needs a second, independent stream +of work to hide it behind, and inside a layer everything is strictly dependent: attention +feeds the router, the router feeds dispatch, the all-gather *sizes* the expert GEMM (see +`floor_ablation.py`), and the GEMM feeds combine. The standard remedy is two interleaved +microbatches. The inference decode path has no such machinery — the only overlap primitive +there is the shared-expert stream, and Qwen3-30B-A3B has no shared expert, so that lever is +empty. **Dual-batch overlap is the single largest remaining prize (1.04 ms exposed against +a 0.82 ms gap) and it is an architecture project, not a tuning session.** + +One trap avoided: the same report shows 0.78 ms/step (9.8%) outside the block graph, which +looks like a fresh target and is not. It is the bucket already measured at 0.59 ms and +**quantitatively retired** — embedding, logits GEMM, sampling, bookkeeping and every host +gap — whose links form a true serial dependency (`graph(N) -> logits -> argmax -> D2H -> +update_requests -> initialize_attention_state(N+1) -> H2D -> graph(N+1)`). Flashinfer +sampling, the obvious lever there, is `QWEN-004`: it crashes under `full_iteration_inference` +capture. Check the ledger before re-opening anything outside the block. + +### QWEN-048 — route-mask fusion: the hang is fixed, the win is inside the noise + +`QWEN-046` left this implemented-but-hanging with a measured +1.02% ablation ceiling. The +retry skipped straight to rung 3 of that entry's diagnostic ladder rather than working up +it, because duplicating the kernel *eliminates* the leading suspect instead of testing it: +`_softmax_topk_mask_kernel` is now a separate `@triton.jit` function and +`_softmax_topk_kernel` is byte-identical to before, so the gate-off path provably cannot +have changed. **That fixed the hang** — the gate-off arm, which is what hung in +`QWEN-046`, now runs normally. The cause was therefore the constexpr branch in the shared +kernel, not the publish or the skip. + +Correctness (`dev/moe_fused/harness_routemask.py`, run with the venv interpreter and +`PYTHONPATH` set, or it cannot import megatron): indices bit-exact and probs delta 0.000e+00 +at real counts 256/200/137/1/0, padding rows all -1. + +Throughput did not follow. Two back-to-back pairs disagreed in sign: + +| pair | gate off | gate on | delta | +|---|---:|---:|---:| +| 1 | 30,412.1 | 30,816.5 | **+1.33%** | +| 2 | 30,564.2 | 30,347.1 | **-0.71%** | + +Four runs span 30,347-30,817 with the treatment at both extremes, which looked like a noise +floor as large as the effect. Six more alternating pairs (`s1a`..`s6b`, gate off then on, +one server start each) resolved it: + +| | n | delta | | t | +|---|---:|---:|---:|---:| +| 6 alternating pairs | 6 | **+224.5 +/- 44.7 tok/s** | **+0.74%** | 5.02 | +| all 8 pairs | 8 | +191.8 +/- 70.6 tok/s | +0.63% | 2.72 | + +**Accepted: +0.74%** (p ~ 0.004), landing just under the +1.02% ablation ceiling, which is +the right shape — the fused kernel still does the compare and the select that the standalone +launch was doing, it just stops paying for a launch to do it. + +The noise diagnosis was itself wrong, and correcting it changes how future A/Bs here should +be read. Gate-off across all 8 runs has sd **51.2 tok/s (0.17%)** — the baseline is extremely +reproducible. Gate-on has sd 183.2, but that is one anomalous run (30,347, the lowest of all +16, the `rmask2` that produced the -0.71%); the other seven average 30,750.9 with a tight +spread. So this node is *not* broadly noisy: it produces an occasional bad server start, +roughly 1 in 16. **A single outlier start, not a wide distribution, is what makes +single-pair A/Bs unsafe here.** Alternate arms and read the paired t, or one bad start will +invert the sign of any sub-1% result. + +Level after this change: **~30,751 tok/s, 90.5% of vLLM.** + +### QWEN-049 — dual-batch overlap is the wrong technique for a latency-bound decode + +`QWEN-047` left DBO as the largest remaining prize (1.04 ms exposed against a 0.82 ms +gap). A full architecture map of the decode path plus an economic check says **do not +build it.** + +**What it would cost.** The decode path is single-microbatch at three independent levels. +(1) One `DynamicInferenceContext` owns every fixed-address buffer the graph reads; nothing +is replicated per request, only indexed. (2) Graphs are keyed on the *identity* of that +context object (`ArgMetadata` stores non-tensors by value, `cuda_graphs.py:256`), so a +second context needs its own captured graph set. (3) The NVLS dispatcher keeps its +symmetric buffers and `_step_metadata` at **class** level, and the barrier +(`torch_symm_triton/barrier.py:65`) is a per-(block, sender) mailbox with no sequence +number, so two concurrent collectives sharing a handle either false-pass into silent +corruption or lose a token and hang. + +**Why it would not pay even if built.** Take `BSSCALE-S13`: step(128) = 7.140 ms, +step(256) = 9.151 ms, with a ~5.03 ms batch-independent floor. Splitting 256 into two 128s +doubles the count of every latency-bound operation, and the collectives are latency-bound, +so the comm to hide goes from 1.04 ms to 2.08 ms. Best case is +floor + 2 x marginal(128) = 5.03 + 4.22 = **9.25 ms against today's 9.15**. DBO hides +communication behind long GEMMs in a *compute-bound* regime; this decode is latency-bound, +so halving the batch costs about what the overlap returns. + +**And it was already gated out.** `QWEN-016` priced the whole lever: recoverable critical +path 6.4-7.0%, 2-chunk pipelining 4.3% best case, concurrent streams under graph capture +identified as the binding constraint, and a CTA sweep showing 128 is already optimal so +freeing SMs by shrinking the collective's grid is strictly worse. Its recorded next action +was "redirect off this lever". **Read QWEN-016 before touching comm overlap again.** + +### QWEN-050 — EP expert-load imbalance is real and worth nothing + +A promising-looking lever, killed by direct measurement. Recorded in full because the +causal story was strong and still wrong. + +**The evidence for it.** A 127M-assignment histogram (`expert_histogram.py`, new) shows +expert popularity is uneven and does not average out within a rank's contiguous block: +loads are **+6.9% / -2.9% / +2.1% / -6.1%** of the mean. In-situ timing appears to show +the barrier absorbing exactly that — per layer rank 0 spends 63.0 us in the expert GEMM +and 28.2 us in the collectives, rank 1 spends 56.8 and 34.5, and *both sum to 91.2*. The +rank ordering of GEMM time matches the histogram ordering exactly. An equal-cardinality +greedy partition takes the worst rank from +6.9% to +0.1%, so on the "critical path is the +busiest rank" model the prize is (max - mean) = 4.85 us/layer = 233 us/step = **2.8%**. + +**The measurement.** Permuting only the router's gating weight was the cheap probe, and it +**hangs** — it leaves expert weights in place, so numerics degrade, logits degenerate, +routing collapses onto a few experts and blows the per-rank token capacity. *A probe that +corrupts numerics cannot rely on data-dependent routing.* The fix is to make routing not +depend on numerics at all: `synthetic_routing.py` (new) substitutes a fixed index tensor +with a chosen per-rank split, so both arms have equally invalid numerics and differ only +in load distribution. + +| arm | per-rank split | tok/s | +|---|---|---:| +| `skew` | +5.9 / -2.2 / +2.6 / -6.2% | 29,656.4 | +| `balanced` | -0.6 / -0.3 / +0.4 / +0.5% | 29,583.7 | + +**Removing the entire imbalance is worth -0.25%, i.e. nothing.** A 2.8% effect would have +been ~830 tok/s, far outside any noise here, so one pair is enough to reject it. + +**Why the causal story failed.** The GEMM/collective anti-correlation is the barrier +equalizing the ranks, and that is exactly why rebalancing cannot help: it moves work +between ranks whose totals are *already* equalized, without shortening the thing that sets +the total. The binding constraint is the collective latency floor, not expert work. Note +also that several sites split cleanly along r0,r1 vs r2,r3 (`nvls_dispatch` 14.9/14.8 vs +10.2/10.1, `router_topk` 3.8/3.8 vs 5.3/5.2) — that is a two-superchip topology signature, +plus a per-rank calibration difference (2.2 vs 2.6-2.7 us/pair), and I initially read part +of it as load imbalance. **Per-rank in-situ deltas conflate load, topology and calibration; +do not attribute them to load without an independent load measurement.** + +Both diagnostics are kept, gated off: `MCORE_EXPERT_HISTOGRAM` and `MCORE_SYNTH_ROUTING`. + + +### QWEN-051 — vLLM's MoE communication is *worse* than ours; the one thing worth taking is a kernel vLLM itself does not run + +Prompted by "look at what vLLM's MoE communication does and reuse it." The answer to the +question as asked is that there is nothing to reuse, and the reason overturns a conclusion +this campaign has been steering by. + +**vLLM's default all2all backend is `allgather_reducescatter`** (`vllm/config/parallel.py:186`) +— the *same* all-gather-v + reduce-scatter-v pair mcore runs. It goes over plain NCCL: +`VLLM_USE_NCCL_SYMM_MEM` defaults to `False` (`vllm/envs.py:267`), so vLLM is not even using +the NVLS multimem path our dispatcher is built on. No stream overlap around dispatch or +combine; `supports_async()` defaults to `False` and only DeepEP/NIXL/MoRI override it; DBO is +hard-asserted to those same backends (`vllm/config/vllm.py:1536`) and so is unreachable from +`TP=1 DP=4 --enable-expert-parallel`. `naive` and `pplx` were removed and now silently fall +back to the default (`parallel.py:448`). + +The measured consequence was already in this ledger and had not been read as such: +**collective 1.607 ms/step (94 launches) in vLLM against 1.035 ms/step (96) in mcore.** +mcore's collectives are **0.572 ms/step faster than vLLM's**. Several attempts to attack +"the collective latency floor" were aimed at the one bucket where we are already ahead. +*When a bucket comparison shows you winning, stop optimising it — even when an in-situ +measurement says the bucket is large and fully serialized. Large and serialized is not the +same as improvable, and the competitor's number is the evidence that settles which.* + +The other half of the request is already done: `vllm_fused_moe.py` **is** vLLM's Triton fused +MoE, vendored, and is what `GEMM_BACKEND=vllm` selects. That is why expert GEMM sits at +parity (2.436 vs 2.376) and why the bucket launch counts nearly coincide — 288 in mcore +against 283 in vLLM across routing + expert GEMM + finalize. Both engines run the same +kernels. There is no unclaimed vLLM MoE code left to import. + +**What is left is `flashinfer.fused_moe.trtllm_bf16_moe`**: one kernel doing router, both +expert GEMMs, activation and finalize, and vLLM's *first-choice* unquantized backend on +Blackwell (`oracle/unquantized.py:68`), explicitly compatible with the AG/RS layout we +already have (`experts/trtllm_bf16_moe.py:207`). The trace proves vLLM fell back to Triton in +the run we benchmarked (283 launches, not ~50), so **vLLM's own numbers do not vouch for it** +— adopting it would be overtaking vLLM's configuration, not copying it. + +Two ledger notes had to be corrected before it could be priced: + +- The "4-D pre-shuffled BlockMajorK weights" blocker is **real for bf16 too** (the docstring + rejects any other `weight_layout`), not FP8-only as the vLLM source suggested. But it is + **not a blocker**: `reorder_rows_for_gated_act_gemm` → `shuffle_matrix_a` → + `convert_to_block_layout` round-trips mcore's 3-D `[E, 2*ffn, H]` in **0.02 s per layer, + ~1 s of load time for all 48**. A one-time startup transform had been recorded as a + structural incompatibility. +- The prior rejection at "90.81 vs 83.63 us, 8% slower" was `cutlass_fused_moe`, a + **different kernel**, which vLLM actively demotes when `dp_size > 1`. `trtllm_bf16_moe` had + never been measured. + +**Measurement.** Harnesses in `dev/moe_fused/`: `harness_trtllm_moe.py` (bake-off), +`harness_moe_breakdown.py` (per-kernel), `harness_moe_graph.py` (capture), +`harness_moe_sweep.py` (token sweep). The first bake-off was wrong twice over and both +errors are worth keeping: + +1. It ran with the production gate set unset, so the Triton path used its untuned default + config: 206.9 us/layer against a true 134.6. *A harness that imports the production kernel + does not inherit the production environment; copy the gate set from the last e2e script.* +2. Wall clock overstated the win (1.49x) because ~40% of both sides was host dispatch that + the decode graph removes. Device time is the only comparable number. + +Device time at T=256, the operating point (512 token-expert pairs per rank, matching +production): + +| | device us/layer | launches | +|---|---:|---:| +| mcore `vllm_fused_moe` | 78.0 | 5 | +| flashinfer `trtllm_bf16_moe` | 63.3 | 4 | + +**CUDA-graph capture succeeds and replay is bit-identical to eager** (max diff 0.000e+00) at +61.1 us/layer. That was the kill criterion: mcore's decode is one captured graph and the +kernel uses a cooperative launch reserving 144 of 152 SMs. + +A token sweep anchors the projection. mcore's per-launch grouped GEMM never reaches the +trace's 25.4 us — it asymptotes near 30 us at T=64 — so the harness carries a ~25-30% +constant overhead over in-graph execution and the raw delta must be discounted by it: + +| T | mcore total | GEMM/launch | flashinfer | speedup | raw ms/step | +|---:|---:|---:|---:|---:|---:| +| 64 | 69.5 | 29.9 | 59.3 | 1.17 | +0.49 | +| 128 | 71.5 | 31.2 | 63.2 | 1.13 | +0.40 | +| 256 | 78.0 | 33.0 | 63.3 | 1.23 | +0.71 | +| 512 | 99.3 | 41.3 | 67.6 | 1.47 | +1.53 | +| 1024 | 118.0 | 46.1 | 72.8 | 1.62 | +2.17 | + +**Estimate: +0.5 to +0.7 ms/step, 6-8%**, which would take mcore from 90.5% to roughly 96-98% +of vLLM. The win is *kernel quality, not launch reduction* — mcore's MoE chain is already +only 5 launches/layer thanks to the align/count/scatter fusions, so exactly one launch is +removed. An earlier "288 launches collapse to 48" framing was wrong. + +**Before integrating.** Numerics need a gate/up half-swap at load: flashinfer matches a torch +reference with `gate_first=False` where mcore matches `gate_first=True`, and the two disagree +by 1.6e-4 against a 1.5e-3 signal. This is the same silent-corruption trap the cutlass probe +hit. The real cost is integration, not the kernel: `trtllm_bf16_moe` routes from logits +itself, so it displaces the router top-k, the newly fused route mask and the dispatcher's +permutation, and has to interoperate with the NVLS AGV/RSV pair. + +**And it must be A/B'd e2e before being believed.** Two microbenchmark wins of this size have +already gone negative here (QWEN-045's fused router: 1.25x in isolation, -2.37% e2e). The +reason to expect a different outcome is mechanistic rather than hopeful: both prior failures +replaced *well-overlapped* work and lost the overlap, whereas mcore's MoE is measured at +essentially zero overlap (sum-of-durations / union = 1.003-1.008), so there is none to lose. + +Environment note: **no install is needed to ship this.** The e2e run venv +(`run_e2e_cfg.sh:36`) already carries flashinfer 0.6.14, which exports `trtllm_bf16_moe`, +`convert_to_block_layout`, `shuffle_matrix_a` and `reorder_rows_for_gated_act_gemm`. The +container's *system* python has no flashinfer, which is why the standalone harnesses needed +`persistent_cache/pyextra` (0.6.16.post1, `--no-deps`, since pruned to `flashinfer`, +`nccl4py` and `tvm_ffi` so it shadows nothing). Cubins are cached at +`persistent_cache/flashinfer_cubins`; the first call downloads for ~14 minutes. See QWEN-052 +for the interpreter confusion this distinction caused. + + +### QWEN-052 — RETRACTED: `MCORE_FLASHINFER_DECODE` is live; I probed the wrong interpreter + +Kept rather than deleted, because the false alarm is a better lesson than the correction. + +**The claim.** `flashinfer_decode.py` degrades on `ImportError` to `HAVE_FLASHINFER = False` +with no warning, and `pip show flashinfer` inside the mcore image `ceecf5c304a5d8bd.sqsh` +reports "Package(s) not found". Since 88 of 92 e2e scripts run in that image, the gate +looked inert, and the differential seemed to corroborate it: attention is still +**1.623 ms/step against vLLM's 1.214** on a current-gate-set trace, exactly the bucket the +gate was credited with closing (+2.6%). + +**Why it is wrong.** The e2e runner does not use the container's python. `run_e2e_cfg.sh:36` +sets `PYBIN=$VENV/bin/python` against a venv on lustre, and that venv **has flashinfer +0.6.14** — the very version this ledger recorded in QWEN-039. Checked inside the container +with the real interpreter: + +``` +torch 2.13.0a0+8145d630e8.nv26.06 +flashinfer 0.6.14 +HAVE_FLASHINFER = True +``` + +against `ModuleNotFoundError` from the container's `/usr/bin/python` two lines later. The +gate is live, the +2.6% is in the current baseline, and nothing regressed. + +**Rule.** *Probe the interpreter the workload actually launches, not the container default.* +`which python` inside the container answered a different question than the one being asked, +and the answer looked decisive enough to build a whole causal story on — one that a stale ++0.409 ms attention bucket then appeared to confirm. Read the runner for `PYBIN`/`VENV` +before concluding anything about what is installed. Corollary: coincidental corroboration +is the dangerous kind. The attention deficit is real but has some other cause, and reading +it as confirmation nearly cost an e2e A/B chasing a bug that does not exist. + +**The `ImportError` critique still stands on its own**, just with no victim here: a +capability gate that degrades silently makes a run script's gate set a statement of intent +rather than a record of what ran, and it should log once at init when enabled but +unavailable. Every other optional-import gate in `inference/` (`fused_add_rmsnorm`, +`fused_qk_norm`, `permute`, `vllm_fused_moe`, `router_topk`, `fused_router`) degrades on +**triton**, which is present at 3.7.0, so no other ledger entry is exposed either. + +**Useful by-product.** The run venv's flashinfer 0.6.14 already exports +`fused_moe.trtllm_bf16_moe`, `fused_moe.convert_to_block_layout`, `shuffle_matrix_a` and +`reorder_rows_for_gated_act_gemm` — everything QWEN-051's integration needs. **No install +is required to ship it.** The `persistent_cache/pyextra` install from QWEN-051 was only ever +needed because the standalone harnesses run under the container's system python; it has +since been pruned to `flashinfer`, `nccl4py` and `tvm_ffi` so it adds modules without +shadowing any the container provides. + +Also unchanged, and still open: attention is **+0.409 ms/step** against vLLM *with* the +trtllm-gen decode kernel active. That is a genuine remaining deficit rather than a +misconfiguration, and it is now the second-largest known gap after the MoE kernel. + +## Optimization rules + +1. Profile and classify before proposing a code change. +2. Change one performance variable at a time. +3. Preserve the fixed protocol. +4. Validate correctness before accepting throughput. +5. Revert regressions or correctness failures. +6. Record the result before beginning another experiment. +7. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines + once to confirm parity under identical conditions. diff --git a/skills/run-qwen-model/SKILL.md b/skills/run-qwen-model/SKILL.md new file mode 100644 index 00000000000..6910350d1bc --- /dev/null +++ b/skills/run-qwen-model/SKILL.md @@ -0,0 +1,153 @@ +--- +name: run-qwen-model +description: Run Qwen3-30B-A3B inference with Megatron-Core or vLLM on one OCI 4×GB200 node at batch size 256, and capture matching Nsight Systems profiles. Use for Qwen 30B mcore runs, vLLM runs, baseline benchmarks, or nsys profile capture. Performance optimization belongs to the qwen-model-optimizer agent. +--- + +# Run Qwen3-30B-A3B on OCI GB200 + +This skill only runs the fixed comparison workload. Use the +`qwen-model-optimizer` subagent for performance investigation or code changes. + +## Fixed comparison + +| Setting | Megatron-Core | vLLM | +|---|---|---| +| Cluster | OCI `oci-hsg`, one 4×GB200 node | same | +| Model | Qwen3-30B-A3B, BF16 | same HF weights | +| Batch | 256 gsm8k requests | same | +| Throughput OSL | 1024 | 1024 | +| Profile OSL | 128 (bounded trace) | 128 | +| Parallelism | **TP=1, PP=1, EP=4** | **TP=1, DP=4, expert parallel enabled** | +| Warmup / timed | 2 / 5 | 2 / 5 | + +Checkpoint paths: + +```bash +export QWEN30B_CKPT=/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/checkpoints/qwen3-30b-a3b-mcore +export QWEN30B_TOKENIZER=/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/checkpoints/qwen3-30b-a3b-hf +export QWEN30B_HF="$QWEN30B_TOKENIZER" +``` + +Before running: + +```bash +source ~/.cog/setup.env.oci-hsg +export COG_MEGATRON_REPO=/path/to/Megatron-LM +cog prepare-image --repo "$COG_MEGATRON_REPO" --cluster-name "$COG_CLUSTER_NAME" +cog ensure-env --repo "$COG_MEGATRON_REPO" \ + --cluster-name "$COG_CLUSTER_NAME" \ + --run-name qwen30b-env --gpus 4 --time 00:30:00 \ + --partition "$COG_BATCH_PARTITION" +``` + +Never download checkpoints. If either path is missing, ask the user. + +## If a cog command hangs, switch to `sbatch` after the second attempt + +These runs are long enough that a stuck launch is expensive. Every cog command +that touches your code syncs the workspace first; when that sync hangs there is +no job ID and no Slurm log, so it is indistinguishable from a slow start except +by the missing `job_id`. + +> **Two strikes, then hand-write the job.** After the second consecutive hang +> or broken pipe from the same client, stop retrying and submit `sbatch` +> directly against a tree staged on lustre. On the rebase measurement this +> converted a day of failed cog attempts into a completed benchmark in under 7 +> minutes. Full recipe — image path via `cog profile` (local-only, so it still +> works), staging to a fresh directory, container mounts, venv activation — is +> in `skills/cog-setup-and-help/SKILL.md`, section "Escape hatch: when +> workspace sync hangs". + +Three rules for any directly-submitted job here, all learned the hard way: + +- **Always pass an explicit `--qos`.** Omitting it takes the default `normal` + (priority 100) and buries you behind a queue whose head sits near 350k, even + when hundreds of nodes are idle: two jobs submitted without it sat `PENDING + (Priority)` for 14 hours and 7 hours, then both started *within seconds* of + resubmission under a correct QOS. **Pick by node count, not by walltime**, and + the name `interactive` is misleading — a QOS constrains resources, not how you + submitted, so it applies to `sbatch` exactly as it does to `salloc`: + `--qos=interactive` (priority 700, ≤4 nodes/user, **no walltime cap**) for + anything up to 4 nodes, which is every job in this skill; `--qos=short` + (priority 200, ≤2 h, ≤64 nodes) only when you need more than 4 nodes. For a + 1-node job `interactive` strictly dominates: 3.5× the priority and no time + limit. All three carry `DenyOnLimit`, so exceeding a limit is rejected at + submit rather than pending forever. Check your grants with + `sacctmgr -nP show assoc user=$USER format=Account,QOS` and the limits with + `sacctmgr -nP show qos format=Name,Priority,MaxWall,MaxTRESPU`. Diagnose a + stalled job by comparing `squeue -j -o '%Q'` against + `squeue -p batch -t PD -S -Q -o '%.8Q' -h | head`; if idle nodes exist and + your number is far below the head, it is QOS, not contention. +- **Gate the allocation on an `os.path.isfile` filesystem health check** over a + few of your changed files plus something deep in the shared venv, and `exit` + before the model load if any are unreadable. A node with a broken Lustre + client lists files it cannot open, and a 4×GB200 allocation is far too + expensive to discover that after checkpoint load. +- **A cascade of `ModuleNotFoundError` across unrelated packages is a bad node, + not a bad venv.** Exclude it (`--exclude=`) and resubmit; do not + pip-install overlay copies to work around it. + +`dev/moe_fused/rebased12.sbatch` is a working instance of this pattern if it is +still present in your tree — but it is untracked, so treat the cog skill's +inlined template as the source of truth. + +## 1. Megatron-Core inference + +Runs `examples.inference.launch_inference_server` with +`transformer_impl=inference_optimized`, full-iteration CUDA graphs, and the +fixed EP4/TP1 layout. + +```bash +EXPERIMENT_ID=MCORE-BASELINE \ +EXPERIMENT_HYPOTHESIS="Fresh EP4 mcore baseline" \ +QWEN30B_TP=1 QWEN30B_EP=4 QWEN30B_ETP=1 \ +BENCH_SIZES_OVERRIDE=256 BENCH_OUTPUT_TOKENS=1024 \ +NUM_WARMUP_ITERS=2 NUM_TIMED_ITERS=5 \ +bash skills/run-qwen-model/run_qwen_inference.sh \ + qwen3-30b-a3b --checkpoint "$QWEN30B_CKPT" +``` + +Do not add optimization flags to a baseline run. + +## 2. vLLM inference + +Runs `vllm serve` with TP1/DP4 and `--enable-expert-parallel`. + +```bash +EXPERIMENT_ID=VLLM-BASELINE \ +EXPERIMENT_HYPOTHESIS="Fresh vLLM DP4+EP baseline" \ +BENCH_BS=256 BENCH_OUTPUT_TOKENS=1024 \ +NUM_WARMUP_ITERS=2 NUM_TIMED_ITERS=5 \ +bash skills/run-qwen-model/run_qwen_vllm.sh +``` + +## Nsight Systems profiles + +Profiles use BS256 and OSL128 to bound trace size while preserving the same +parallel layouts. Both scripts export `.nsys-rep` and `.sqlite`. + +```bash +# mcore EP4/TP1 +PROFILE_BS=256 PROFILE_OSL=128 \ +QWEN30B_CKPT="$QWEN30B_CKPT" \ +bash skills/run-qwen-model/profile_qwen_mcore.sh + +# vLLM DP4+EP +PROFILE_BS=256 PROFILE_OSL=128 \ +bash skills/run-qwen-model/profile_qwen_vllm.sh +``` + +Run the profiles sequentially to avoid node contention. Record the run +directory, Slurm job, throughput, latency, TPOT, and trace paths in +`EXPERIMENTS.md`. + +For analysis, use `skills/nsight-system-analysis/SKILL.md` — Workflow C +(`scripts/forward_pass.py`) for the single-decode-step mcore/vLLM comparison, and +its Steps 1–6 for deeper interval-union attribution. Optimization decisions +belong to `skills/optimize-inference-siddharth/SKILL.md`. + +## Records + +`EXPERIMENTS.md` is the sole performance ledger. The first two records must be +the fresh vLLM and mcore nsys baselines. Append every later attempt, including +failures and regressions; never rewrite prior records. diff --git a/skills/run-qwen-model/profile_qwen_mcore.sh b/skills/run-qwen-model/profile_qwen_mcore.sh new file mode 100755 index 00000000000..c171648657f --- /dev/null +++ b/skills/run-qwen-model/profile_qwen_mcore.sh @@ -0,0 +1,149 @@ +#!/usr/bin/env bash +# Profile mcore Qwen3-30B-A3B inference under Nsight Systems (nsys) via cog and +# export a .sqlite for A/B analysis. Uses the fixed baseline: +# inference_optimized, TP=1, EP=4, vLLM grouped-GEMM, full-iteration CUDA +# graphs, BS=256. OSL stays short so the trace remains bounded. +# +# Usage: +# source ~/.cog/setup.env.oci-hsg +# export COG_MEGATRON_REPO=/path/to/Megatron-LM +# QWEN30B_CKPT=/lustre/.../qwen3-30b-a3b-mcore \ +# PROFILE_BS=256 PROFILE_OSL=128 \ +# bash skills/run-qwen-model/profile_qwen_mcore.sh +set -euo pipefail + +PROFILE_BS="${PROFILE_BS:-256}" +PROFILE_OSL="${PROFILE_OSL:-128}" + +_USER_REPO="${COG_MEGATRON_REPO:-}" +if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env.oci-hsg" +elif [[ -f "${HOME}/.cog/setup.env" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env" +else + echo "ERROR: no ~/.cog/setup.env — run cog-setup-and-help skill" >&2 + exit 1 +fi +if [[ -n "$_USER_REPO" ]]; then export COG_MEGATRON_REPO="$_USER_REPO"; fi + +: "${COG_MEGATRON_REPO:?COG_MEGATRON_REPO not set}" +: "${COG_SSH_HOST:?COG_SSH_HOST not set}" +export COG_ARTIFACTS_ROOT="${COG_ARTIFACTS_ROOT:-/lustre/fsw/portfolios/coreai/projects/coreai_dlalgo_mcore/mcore_ci}" + +CKPT_ABS="${QWEN30B_CKPT:-${COG_SCRATCH_ROOT}/checkpoints/qwen3-30b-a3b-mcore}" +TOKENIZER="${QWEN30B_TOKENIZER:-${COG_SCRATCH_ROOT}/checkpoints/qwen3-30b-a3b-hf}" + +NPROC=4 +ALLOC_GPUS="$NPROC" +if [[ -n "${COG_GPUS_PER_NODE_LIMIT:-}" ]] && [[ "$ALLOC_GPUS" -lt "$COG_GPUS_PER_NODE_LIMIT" ]]; then + ALLOC_GPUS="$COG_GPUS_PER_NODE_LIMIT" +fi + +QWEN_MODEL_ARGS="--model-provider gpt --num-layers 48 --hidden-size 2048 --ffn-hidden-size 6144 --num-attention-heads 32 --group-query-attention --num-query-groups 4 --kv-channels 128 --num-experts 128 --moe-router-topk 8 --moe-ffn-hidden-size 768 --moe-grouped-gemm --moe-router-dtype fp32 --moe-router-pre-softmax --moe-token-dispatcher-type alltoall --swiglu --normalization RMSNorm --norm-epsilon 1e-6 --position-embedding-type rope --rotary-base 1000000 --qk-layernorm --disable-bias-linear --untie-embeddings-and-output-weights --no-gradient-accumulation-fusion --make-vocab-size-divisible-by 1187 --tensor-model-parallel-size 1 --pipeline-model-parallel-size 1 --expert-model-parallel-size 4 --expert-tensor-parallel-size 1 --inference-moe-token-dispatcher-type nvls --inference-grouped-gemm-backend vllm" + +RUN_NAME="qwen-30b-nsys-$(date +%Y%m%d-%H%M%S)" +ARTIFACTS="$COG_ARTIFACTS_ROOT" +export COG_EXTRA_MOUNTS="$ARTIFACTS:$ARTIFACTS" + +echo "Submitting $RUN_NAME: mcore EP4/TP1 nsys profile BS=$PROFILE_BS OSL=$PROFILE_OSL ckpt=$CKPT_ABS" + +cog --pretty submit \ + --repo "$COG_MEGATRON_REPO" \ + --cluster-name "${COG_CLUSTER_NAME:-oci-hsg}" \ + --run-name "$RUN_NAME" \ + --gpus "$ALLOC_GPUS" --nodes 1 --ntasks-per-node 1 \ + --time 01:00:00 \ + --partition "${COG_BATCH_PARTITION:-batch}" \ + --command "$(cat </dev/null || true + +SERVER_LOG="\$RUN_DIR/server.log" +PROF_BASE="\$RUN_DIR/mcore_profile" + +if ! command -v nsys >/dev/null 2>&1; then + echo "ERROR: nsys not found in image" >&2; exit 1 +fi +nsys --version + +# Launch the server UNDER nsys. --cuda-graph-trace=node is required so kernels +# inside the full_iteration_inference CUDA graph are individually traced. +nsys profile \\ + --trace=cuda,nvtx,osrt \\ + --sample=none --cpuctxsw=none \\ + --cuda-graph-trace=node \\ + --force-overwrite=true \\ + -o "\$PROF_BASE" \\ + python -m torch.distributed.run --nproc-per-node $NPROC --log-dir "\$RUN_DIR/torchrun_logs" \\ + -m examples.inference.launch_inference_server \\ + --load "\$CKPT" \\ + --dist-ckpt-strictness log_unexpected \\ + --tokenizer-type HuggingFaceTokenizer \\ + --tokenizer-model "\$TOKENIZER" \\ + --no-use-tokenizer-model-from-checkpoint-args \\ + --micro-batch-size 1 --bf16 --te-rng-tracker --inference-rng-tracker \\ + --transformer-impl inference_optimized \\ + --inference-dynamic-batching \\ + --inference-dynamic-batching-unified-memory-level 0 \\ + --use-flashinfer-fused-rope \\ + --inference-dynamic-batching-max-tokens 4096 \\ + --enable-chunked-prefill \\ + --seq-length 4096 --max-position-embeddings 4096 --inference-max-seq-length 4096 \\ + --inference-dynamic-batching-buffer-size-gb 40 \\ + --inference-dynamic-batching-max-requests 256 \\ + --inference-dynamic-batching-num-cuda-graphs -1 \\ + --cuda-graph-impl local \\ + --cuda-graph-scope full_iteration_inference \\ + --inference-use-synchronous-zmq-collectives \\ + --inference-logging-step-interval 100 \\ + --port 5000 \\ + $QWEN_MODEL_ARGS \\ + > "\$SERVER_LOG" 2>&1 & +NSYS_PID=\$! + +READY=0 +for i in \$(seq 1 300); do + if grep -q "Running on http://0.0.0.0:5000" "\$SERVER_LOG" 2>/dev/null; then READY=1; break; fi + if ! kill -0 \$NSYS_PID 2>/dev/null; then echo "SERVER/NSYS DIED"; tail -100 "\$SERVER_LOG"; exit 1; fi + sleep 5 +done +if [[ "\$READY" != "1" ]]; then echo "SERVER TIMEOUT"; tail -100 "\$SERVER_LOG"; kill \$NSYS_PID 2>/dev/null; exit 1; fi +echo "===== SERVER READY (profiling) =====" + +# One short warmup call (BS=8) so CUDA graphs are captured and the allocator is +# warm; steady-state decode after this is what the analysis anchors on. +python -u tests/performance_tests/client/static_benchmark.py \\ + --server-url "http://localhost:5000/v1" --model qwen \\ + --batch-size 8 --dataset gsm8k --num-output-tokens 32 \\ + --num-iters 1 --num-warmup-iters 0 || true + +echo "===== PROFILED BENCHMARK BS=$PROFILE_BS OSL=$PROFILE_OSL =====" +python -u tests/performance_tests/client/static_benchmark.py \\ + --server-url "http://localhost:5000/v1" --model qwen \\ + --batch-size $PROFILE_BS --dataset gsm8k --num-output-tokens $PROFILE_OSL \\ + --num-iters 1 --num-warmup-iters 0 2>&1 | tee "\$RUN_DIR/profile_bench.log" + +# Stop nsys gracefully (SIGINT) so it finalizes and writes the .nsys-rep. +echo "===== stopping nsys =====" +kill -INT \$NSYS_PID 2>/dev/null || true +wait \$NSYS_PID 2>/dev/null || true + +ls -la "\$RUN_DIR"/mcore_profile.* || true +echo "===== exporting sqlite =====" +nsys export --type sqlite --force-overwrite=true \\ + --output "\$PROF_BASE.sqlite" "\$PROF_BASE.nsys-rep" +ls -la "\$PROF_BASE.sqlite" +echo "===== PROFILE DONE =====" +echo "REP=\$PROF_BASE.nsys-rep" +echo "SQLITE=\$PROF_BASE.sqlite" +EOF +)" diff --git a/skills/run-qwen-model/profile_qwen_vllm.sh b/skills/run-qwen-model/profile_qwen_vllm.sh new file mode 100755 index 00000000000..74229c98b81 --- /dev/null +++ b/skills/run-qwen-model/profile_qwen_vllm.sh @@ -0,0 +1,134 @@ +#!/usr/bin/env bash +# Profile vLLM Qwen3-30B-A3B inference under Nsight Systems (nsys) on oci-hsg and +# export a .sqlite for A/B analysis. Uses TP=1, DP=4, expert parallel and wraps +# `vllm serve` +# with nsys at fixed BS256, profiles a short OSL, and exports sqlite. +# +# Usage: +# source ~/.cog/setup.env.oci-hsg +# PROFILE_BS=256 PROFILE_OSL=128 bash skills/run-qwen-model/profile_qwen_vllm.sh +set -euo pipefail + +PROFILE_BS="${PROFILE_BS:-256}" +PROFILE_OSL="${PROFILE_OSL:-128}" + +if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env.oci-hsg" +else + echo "ERROR: ~/.cog/setup.env.oci-hsg not found" >&2; exit 1 +fi +: "${COG_SSH_HOST:?}"; : "${COG_SCRATCH_ROOT:?}" + +IMG="${VLLM_IMAGE:-$COG_SCRATCH_ROOT/images/87e4947c6ce36433.sqsh}" +HF_CKPT="${QWEN30B_HF:-$COG_SCRATCH_ROOT/checkpoints/qwen3-30b-a3b-hf}" +ACCOUNT="${COG_RUNTIME_ACCOUNT:-coreai_dlalgo_llm}" +PART="${COG_BATCH_PARTITION:-batch}" +EXP="$COG_SCRATCH_ROOT/runs/vllm-qwen30b-nsys-$(date +%Y%m%d-%H%M%S)" + +echo "Creating vLLM nsys sbatch on $COG_SSH_HOST: BS=$PROFILE_BS OSL=$PROFILE_OSL img=$IMG" + +# Push the sbatch to the cluster. Outer heredoc quoted so it is sent verbatim; +# $PROFILE_BS / $PROFILE_OSL / paths are substituted by expanding them into the +# stream via a small template below. +ssh -o BatchMode=yes "$COG_SSH_HOST" "mkdir -p '$EXP'; cat > '$EXP/vllm_nsys.sbatch'" </dev/null | head -1) +mkdir -p "\$EXP" +rm -f "\$EXP/vqdone_status" "\$EXP/vqserver.log" "\$EXP/vqprof.log" + +srun --container-image="\$IMG" \\ + --container-mounts=/lustre:/lustre \\ + --no-container-mount-home \\ + bash -c ' +set -x +export HF_HOME='"\$EXP"'/hf_home +export VLLM_WORKER_MULTIPROC_METHOD=spawn +export OUTLINES_CACHE_DIR=/tmp/outlines_\$\$ +export CUDA_DEVICE_MAX_CONNECTIONS=1 +EXP='"\$EXP"' +HF_CKPT='"\$HF_CKPT"' +MLM='"\$MLM_HOST"' +PROFILE_BS='"\$PROFILE_BS"' +PROFILE_OSL='"\$PROFILE_OSL"' +VENV=/opt/ray_venvs/nemo_rl.experience.sync_rollout_actor.SyncRolloutActor +PY=\$VENV/bin/python +VLLM=\$VENV/bin/vllm +BENCH=\$MLM/tests/performance_tests/client/static_benchmark.py +PROF_BASE=\$EXP/vllm_profile + +if ! command -v nsys >/dev/null 2>&1; then echo "ERROR: nsys not in vLLM image"; exit 3; fi +nsys --version + +nsys profile \\ + --trace=cuda,nvtx,osrt \\ + --sample=none --cpuctxsw=none \\ + --cuda-graph-trace=node \\ + --force-overwrite=true \\ + -o "\$PROF_BASE" \\ + \$VLLM serve "\$HF_CKPT" --served-model-name qwen \\ + --tensor-parallel-size 1 --data-parallel-size 4 --enable-expert-parallel \\ + --max-model-len 4096 --max-num-seqs 512 \\ + --gpu-memory-utilization 0.9 --trust-remote-code \\ + --port 5000 --host 0.0.0.0 \\ + > "\$EXP/vqserver.log" 2>&1 & +NSYS_PID=\$! + +READY=0 +for i in \$(seq 1 360); do + if grep -qE "Application startup complete|Uvicorn running on http://0.0.0.0:5000" "\$EXP/vqserver.log" 2>/dev/null; then READY=1; break; fi + if ! kill -0 \$NSYS_PID 2>/dev/null; then echo "SERVER/NSYS DIED"; tail -100 "\$EXP/vqserver.log"; exit 1; fi + sleep 5 +done +if [ "\$READY" != "1" ]; then echo "VLLM NOT READY"; tail -100 "\$EXP/vqserver.log"; kill \$NSYS_PID 2>/dev/null; exit 1; fi +echo "===== VLLM READY (profiling) =====" + +\$PY -u "\$BENCH" --server-url "http://localhost:5000/v1" --model qwen \\ + --batch-size 8 --dataset gsm8k --num-output-tokens 32 --num-iters 1 --num-warmup-iters 0 || true + +echo "===== PROFILED BENCHMARK BS=\$PROFILE_BS OSL=\$PROFILE_OSL =====" +\$PY -u "\$BENCH" --server-url "http://localhost:5000/v1" --model qwen \\ + --batch-size \$PROFILE_BS --dataset gsm8k --num-output-tokens \$PROFILE_OSL \\ + --num-iters 1 --num-warmup-iters 0 2>&1 | tee "\$EXP/vqprof.log" + +echo "===== stopping nsys =====" +kill -INT \$NSYS_PID 2>/dev/null || true +wait \$NSYS_PID 2>/dev/null || true +ls -la "\$EXP"/vllm_profile.* || true +echo "===== exporting sqlite =====" +nsys export --type sqlite --force-overwrite=true --output "\$PROF_BASE.sqlite" "\$PROF_BASE.nsys-rep" +ls -la "\$PROF_BASE.sqlite" +echo done > "\$EXP/vqdone_status" +echo "===== PROFILE DONE =====" +echo "SQLITE=\$PROF_BASE.sqlite" +' +echo "JOB DONE rc=\$?" +EOF + +if [[ "${DRY_RUN:-0}" == "1" ]]; then + echo "DRY_RUN=1 — wrote $EXP/vllm_nsys.sbatch on $COG_SSH_HOST, not submitting." + exit 0 +fi + +echo "Submitting vLLM nsys sbatch..." +JOBID=$(ssh -o BatchMode=yes "$COG_SSH_HOST" "cd '$EXP' && sbatch --parsable vllm_nsys.sbatch") +echo "vLLM nsys job: $JOBID" +echo "Run dir: $EXP" +echo "$JOBID" > /tmp/vllm_nsys_jobid.txt diff --git a/skills/run-qwen-model/run_qwen_inference.sh b/skills/run-qwen-model/run_qwen_inference.sh new file mode 100755 index 00000000000..24fdf170472 --- /dev/null +++ b/skills/run-qwen-model/run_qwen_inference.sh @@ -0,0 +1,334 @@ +#!/usr/bin/env bash +# Run Qwen inference (server + coherence + optional gsm8k benchmark) on a Slurm +# cluster via cog. See skills/run-qwen-model/SKILL.md. +set -euo pipefail + +MODEL="${1:-qwen3-30b-a3b}" +shift || true + +FULL_BENCH=0 +CHECKPOINT_OVERRIDE="" +NUM_WARMUP_ITERS="${NUM_WARMUP_ITERS:-2}" +NUM_TIMED_ITERS="${NUM_TIMED_ITERS:-5}" +EXPERIMENT_ID="${EXPERIMENT_ID:-unassigned}" +EXPERIMENT_HYPOTHESIS="${EXPERIMENT_HYPOTHESIS:-not-recorded}" +# Install the FA4 (flash-attn-4) beta into the run venv before launching. +# Requires an FA4-aware megatron (PR #5804: num_splits=0 for FA4 inference). +INSTALL_FA4="${INSTALL_FA4:-0}" +REQUIRE_FA4="${REQUIRE_FA4:-0}" +# Fixed Qwen3-30B baseline defaults. Override only for a recorded experiment. +INFERENCE_MOE_DISPATCHER="${INFERENCE_MOE_DISPATCHER:-nvls}" +SYNC_ZMQ_COLLECTIVES="${SYNC_ZMQ_COLLECTIVES:-1}" +QWEN30B_TP="${QWEN30B_TP:-1}" +QWEN30B_EP="${QWEN30B_EP:-4}" +QWEN30B_ETP="${QWEN30B_ETP:-1}" +MOE_ROUTER_FUSION="${MOE_ROUTER_FUSION:-0}" +MOE_PERMUTE_FUSION="${MOE_PERMUTE_FUSION:-0}" +BENCH_SIZES_OVERRIDE="${BENCH_SIZES_OVERRIDE:-256}" +BENCH_OUTPUT_TOKENS="${BENCH_OUTPUT_TOKENS:-1024}" +INFERENCE_GROUPED_GEMM_BACKEND="${INFERENCE_GROUPED_GEMM_BACKEND:-vllm}" +DYNAMIC_BATCHING_BUFFER_GB="${DYNAMIC_BATCHING_BUFFER_GB:-}" +DYNAMIC_BATCHING_MAX_TOKENS="${DYNAMIC_BATCHING_MAX_TOKENS:-4096}" +NUM_CUDA_GRAPHS="${NUM_CUDA_GRAPHS:--1}" +ENABLE_CHUNKED_PREFILL="${ENABLE_CHUNKED_PREFILL:-1}" +CUDA_GRAPH_SIZING_DISTRIBUTION="${CUDA_GRAPH_SIZING_DISTRIBUTION:-exponential}" +DYNAMIC_BATCHING_MAX_REQUESTS="${DYNAMIC_BATCHING_MAX_REQUESTS:-256}" +DYNAMIC_BATCHING_ASYNC_SCHED_MODE="${DYNAMIC_BATCHING_ASYNC_SCHED_MODE:-legacy}" +DYNAMIC_BATCHING_SAMPLING_BACKEND="${DYNAMIC_BATCHING_SAMPLING_BACKEND:-torch}" +EXTRA_SERVER_ARGS="${EXTRA_SERVER_ARGS:-}" +trap 'echo "Record ${EXPERIMENT_ID} in skills/run-qwen-model/EXPERIMENTS.md, including failures."' EXIT +while [[ $# -gt 0 ]]; do + case "$1" in + --full-benchmark) FULL_BENCH=1; shift ;; + --checkpoint) CHECKPOINT_OVERRIDE="$2"; shift 2 ;; + -h|--help) + sed -n '1,20p' "$0" + exit 0 + ;; + *) echo "Unknown arg: $1" >&2; exit 1 ;; + esac +done + +# Prefer oci-hsg env when present; fall back to default cog env. +_USER_REPO="${COG_MEGATRON_REPO:-}" +if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env.oci-hsg" +elif [[ -f "${HOME}/.cog/setup.env" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env" +else + echo "ERROR: no ~/.cog/setup.env — run cog-setup-and-help skill" >&2 + exit 1 +fi +# Caller-exported repo wins over the value baked into setup.env.*. +if [[ -n "$_USER_REPO" ]]; then + export COG_MEGATRON_REPO="$_USER_REPO" +fi + +: "${COG_MEGATRON_REPO:?COG_MEGATRON_REPO not set}" +: "${COG_SSH_HOST:?COG_SSH_HOST not set}" + +export COG_ARTIFACTS_ROOT="${COG_ARTIFACTS_ROOT:-/lustre/fsw/portfolios/coreai/projects/coreai_dlalgo_mcore/mcore_ci}" + +if ! ssh -o BatchMode=yes "$COG_SSH_HOST" "test -d '$COG_ARTIFACTS_ROOT/model'"; then + echo "ERROR: artifacts root '$COG_ARTIFACTS_ROOT' not reachable on $COG_SSH_HOST" >&2 + exit 1 +fi + +# --- Fixed Qwen3-30B-A3B model --- +case "$MODEL" in + qwen3-30b-a3b|30b|qwen-30b) + MODEL_TAG="qwen3-30b-a3b" + NPROC=4 + if [[ -n "$CHECKPOINT_OVERRIDE" ]]; then + CKPT_ABS="$CHECKPOINT_OVERRIDE" + elif [[ -n "${QWEN30B_CKPT:-}" ]]; then + CKPT_ABS="$QWEN30B_CKPT" + else + echo "ERROR: Qwen3-30B-A3B checkpoint required. Pass --checkpoint or set QWEN30B_CKPT" >&2 + exit 1 + fi + TOKENIZER="${QWEN30B_TOKENIZER:-/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/checkpoints/qwen3-30b-a3b-hf}" + QWEN_MODEL_ARGS="--model-provider gpt --num-layers 48 --hidden-size 2048 --ffn-hidden-size 6144 --num-attention-heads 32 --group-query-attention --num-query-groups 4 --kv-channels 128 --num-experts 128 --moe-router-topk 8 --moe-ffn-hidden-size 768 --moe-grouped-gemm --moe-router-dtype fp32 --moe-router-pre-softmax --moe-token-dispatcher-type alltoall --swiglu --normalization RMSNorm --norm-epsilon 1e-6 --position-embedding-type rope --rotary-base 1000000 --qk-layernorm --disable-bias-linear --untie-embeddings-and-output-weights --no-gradient-accumulation-fusion --make-vocab-size-divisible-by 1187 --tensor-model-parallel-size $QWEN30B_TP --pipeline-model-parallel-size 1 --expert-model-parallel-size $QWEN30B_EP --expert-tensor-parallel-size $QWEN30B_ETP --inference-moe-token-dispatcher-type $INFERENCE_MOE_DISPATCHER --inference-grouped-gemm-backend $INFERENCE_GROUPED_GEMM_BACKEND" + BENCH_DATASET="gsm8k" + BUFFER_GB=40 + if [[ "$QWEN30B_TP" -gt 1 ]]; then + SEQ_PARALLEL_FLAG="--sequence-parallel" + else + SEQ_PARALLEL_FLAG="" + fi + ;; + *) + echo "ERROR: unknown model '$MODEL'. Use qwen3-30b-a3b" >&2 + exit 1 + ;; +esac + +ALLOC_GPUS="$NPROC" +if [[ -n "${COG_GPUS_PER_NODE_LIMIT:-}" ]] && [[ "$ALLOC_GPUS" -lt "$COG_GPUS_PER_NODE_LIMIT" ]]; then + ALLOC_GPUS="$COG_GPUS_PER_NODE_LIMIT" +fi + +if [[ -n "$DYNAMIC_BATCHING_BUFFER_GB" ]]; then + BUFFER_GB="$DYNAMIC_BATCHING_BUFFER_GB" +fi + +if [[ -n "$BENCH_SIZES_OVERRIDE" ]]; then + BENCH_SIZES="$BENCH_SIZES_OVERRIDE" +elif [[ "$FULL_BENCH" == "1" ]]; then + BENCH_SIZES="16 64 256" +else + BENCH_SIZES="16" +fi +BENCH_ITERS="--num-iters $NUM_TIMED_ITERS --num-warmup-iters $NUM_WARMUP_ITERS" +BENCH_CLIENT='te''sts/performance_tests/client/static_benchmark.py' + +if [[ "$SYNC_ZMQ_COLLECTIVES" == "1" ]]; then + SYNC_ZMQ_FLAG="--inference-use-synchronous-zmq-collectives" +else + SYNC_ZMQ_FLAG="--no-inference-use-synchronous-zmq-collectives" +fi +FUSION_FLAGS="" +if [[ "$MOE_ROUTER_FUSION" == "1" ]]; then FUSION_FLAGS+=" --moe-router-fusion"; fi +if [[ "$MOE_PERMUTE_FUSION" == "1" ]]; then FUSION_FLAGS+=" --moe-permute-fusion"; fi +if [[ "$ENABLE_CHUNKED_PREFILL" == "1" ]]; then + CHUNKED_PREFILL_FLAG="--enable-chunked-prefill" +else + CHUNKED_PREFILL_FLAG="" +fi +MAX_REQUESTS_FLAG="" +if [[ -n "$DYNAMIC_BATCHING_MAX_REQUESTS" ]]; then + MAX_REQUESTS_FLAG="--inference-dynamic-batching-max-requests $DYNAMIC_BATCHING_MAX_REQUESTS" +fi + +RUN_SUFFIX=$(printf '%s' "$EXPERIMENT_ID" | tr '[:upper:]_' '[:lower:]-' | tr -cd '[:alnum:]-') +RUN_NAME="qwen-${MODEL_TAG}-${RUN_SUFFIX:-run}-$(date +%Y%m%d-%H%M%S)" +ARTIFACTS="$COG_ARTIFACTS_ROOT" +export COG_EXTRA_MOUNTS="$ARTIFACTS:$ARTIFACTS" +CODE_REVISION=$(git -C "$COG_MEGATRON_REPO" rev-parse HEAD) +if [[ -n "$(git -C "$COG_MEGATRON_REPO" status --porcelain)" ]]; then + CODE_STATE=dirty +else + CODE_STATE=clean +fi + +echo "Submitting $RUN_NAME: model=$MODEL_TAG nproc=$NPROC alloc_gpus=$ALLOC_GPUS ckpt=$CKPT_ABS" +echo "Experiment $EXPERIMENT_ID: $EXPERIMENT_HYPOTHESIS" +echo "Code: $CODE_REVISION ($CODE_STATE); warmups=$NUM_WARMUP_ITERS timed_iters=$NUM_TIMED_ITERS" + +cog --pretty submit \ + --repo "$COG_MEGATRON_REPO" \ + --cluster-name "${COG_CLUSTER_NAME:-oci-hsg}" \ + --run-name "$RUN_NAME" \ + --gpus "$ALLOC_GPUS" --nodes 1 --ntasks-per-node 1 \ + --time 01:00:00 \ + --partition "${COG_BATCH_PARTITION:-batch}" \ + --command "$(cat </dev/null || true + +SERVER_LOG="\$RUN_DIR/server.log" +BENCH_LOG="\$RUN_DIR/benchmark.log" + +cat > "\$BENCH_LOG" <<'META' +===== EXPERIMENT METADATA ===== +experiment_id=$EXPERIMENT_ID +hypothesis=$EXPERIMENT_HYPOTHESIS +code_revision=$CODE_REVISION +code_state=$CODE_STATE +model=$MODEL_TAG +checkpoint=$CKPT_ABS +tokenizer=$TOKENIZER +world_size=$NPROC +allocated_gpus=$ALLOC_GPUS +dataset=$BENCH_DATASET +batch_sizes=$BENCH_SIZES +num_output_tokens=$BENCH_OUTPUT_TOKENS +warmup_iters=$NUM_WARMUP_ITERS +timed_iters=$NUM_TIMED_ITERS +install_fa4=$INSTALL_FA4 +require_fa4=$REQUIRE_FA4 +inference_moe_dispatcher=$INFERENCE_MOE_DISPATCHER +sync_zmq_collectives=$SYNC_ZMQ_COLLECTIVES +tp=$QWEN30B_TP +ep=$QWEN30B_EP +etp=$QWEN30B_ETP +moe_router_fusion=$MOE_ROUTER_FUSION +moe_permute_fusion=$MOE_PERMUTE_FUSION +inference_grouped_gemm_backend=$INFERENCE_GROUPED_GEMM_BACKEND +dynamic_batching_buffer_gb=$BUFFER_GB +dynamic_batching_max_tokens=$DYNAMIC_BATCHING_MAX_TOKENS +num_cuda_graphs=$NUM_CUDA_GRAPHS +enable_chunked_prefill=$ENABLE_CHUNKED_PREFILL +cuda_graph_sizing_distribution=$CUDA_GRAPH_SIZING_DISTRIBUTION +dynamic_batching_max_requests=$DYNAMIC_BATCHING_MAX_REQUESTS +dynamic_batching_async_sched_mode=$DYNAMIC_BATCHING_ASYNC_SCHED_MODE +dynamic_batching_sampling_backend=$DYNAMIC_BATCHING_SAMPLING_BACKEND +extra_server_args=$EXTRA_SERVER_ARGS +META + +if [[ "$INSTALL_FA4" == "1" ]]; then + echo "===== Installing flash-attn-4 beta into run venv =====" | tee -a "\$BENCH_LOG" + PYBIN="\$(command -v python)" + echo "python=\$PYBIN" | tee -a "\$BENCH_LOG" + FA4_PKGS="flash-attn-4[cu13]==4.0.0b20 quack-kernels==0.5.3 nvidia-cutlass-dsl[cu13]==4.6.0.dev0 nvidia-cutlass-dsl-libs-base==4.6.0.dev0 nvidia-cutlass-dsl-libs-cu13==4.6.0.dev0 apache-tvm-ffi==0.1.12 torch-c-dlpack-ext==0.1.5" + set -f # keep word-splitting on \$FA4_PKGS but stop [cu13] being glob-expanded + if command -v uv >/dev/null 2>&1; then + # Install into the active run venv (same site-packages as flash_attn 2.x). + uv pip install --python "\$PYBIN" --prerelease=allow --no-deps \$FA4_PKGS 2>&1 | tee -a "\$BENCH_LOG" + else + "\$PYBIN" -m pip install --pre --no-deps --no-cache-dir \$FA4_PKGS 2>&1 | tee -a "\$BENCH_LOG" + fi + set +f + echo "----- FA4 verification -----" | tee -a "\$BENCH_LOG" + if ! python - <<'PYFA4' 2>&1 | tee -a "\$BENCH_LOG" +from importlib.metadata import version +import flash_attn +print("flash_attn", getattr(flash_attn, "__version__", "?")) +print("flash-attn-4", version("flash-attn-4")) +from flash_attn.cute import flash_attn_varlen_func # noqa: F401 +print("FA4_IMPORT_OK") +PYFA4 + then + echo "FA4 install/verify FAILED — aborting run so we do not silently benchmark without FA4." | tee -a "\$BENCH_LOG" + exit 1 + fi +fi + +if [[ "$REQUIRE_FA4" == "1" && "$INSTALL_FA4" != "1" ]]; then + echo "----- FA4 verification (no install) -----" | tee -a "\$BENCH_LOG" + if ! python - <<'PYFA4' 2>&1 | tee -a "\$BENCH_LOG" +from importlib.metadata import version +from packaging.version import Version +import flash_attn +from flash_attn.cute import flash_attn_varlen_func # noqa: F401 +installed = version("flash-attn-4") +print("flash_attn", getattr(flash_attn, "__version__", "?")) +print("flash-attn-4", installed) +assert Version(installed) >= Version("4.0.0b20") +print("FA4_IMPORT_OK") +PYFA4 + then + echo "Required FA4 beta is unavailable — aborting rather than benchmarking FA2." | tee -a "\$BENCH_LOG" + exit 1 + fi +fi + +python -m torch.distributed.run --nproc-per-node $NPROC --log-dir "\$RUN_DIR/torchrun_logs" \\ + -m examples.inference.launch_inference_server \\ + --load "\$CKPT" \\ + --dist-ckpt-strictness log_unexpected \\ + --tokenizer-type HuggingFaceTokenizer \\ + --tokenizer-model "\$TOKENIZER" \\ + --no-use-tokenizer-model-from-checkpoint-args \\ + --micro-batch-size 1 --bf16 --te-rng-tracker --inference-rng-tracker \\ + --transformer-impl inference_optimized \\ + $SEQ_PARALLEL_FLAG \\ + --inference-dynamic-batching \\ + --inference-dynamic-batching-unified-memory-level 0 \\ + --use-flashinfer-fused-rope \\ + --inference-dynamic-batching-max-tokens $DYNAMIC_BATCHING_MAX_TOKENS \\ + $MAX_REQUESTS_FLAG \\ + --inference-dynamic-batching-cuda-graph-sizing-distribution $CUDA_GRAPH_SIZING_DISTRIBUTION \\ + --inference-dynamic-batching-async-sched-mode $DYNAMIC_BATCHING_ASYNC_SCHED_MODE \\ + --inference-dynamic-batching-sampling-backend $DYNAMIC_BATCHING_SAMPLING_BACKEND \\ + $CHUNKED_PREFILL_FLAG \\ + --seq-length 4096 --max-position-embeddings 4096 --inference-max-seq-length 4096 \\ + --inference-dynamic-batching-buffer-size-gb $BUFFER_GB \\ + --inference-dynamic-batching-num-cuda-graphs $NUM_CUDA_GRAPHS \\ + --cuda-graph-impl local \\ + --cuda-graph-scope full_iteration_inference \\ + $SYNC_ZMQ_FLAG \\ + $FUSION_FLAGS \\ + --inference-logging-step-interval 100 \\ + --port 5000 \\ + $QWEN_MODEL_ARGS \\ + $EXTRA_SERVER_ARGS \\ + > "\$SERVER_LOG" 2>&1 & +SERVER_PID=\$! + +READY=0 +for i in \$(seq 1 240); do + if grep -q "Running on http://0.0.0.0:5000" "\$SERVER_LOG" 2>/dev/null; then READY=1; break; fi + if ! kill -0 \$SERVER_PID 2>/dev/null; then echo "SERVER DIED"; tail -80 "\$SERVER_LOG"; exit 1; fi + sleep 5 +done +if [[ "\$READY" != "1" ]]; then echo "SERVER TIMEOUT"; tail -80 "\$SERVER_LOG"; kill \$SERVER_PID 2>/dev/null; exit 1; fi + +echo "===== SERVER READY =====" | tee -a "\$BENCH_LOG" +python - <<'PYEOF' 2>&1 | tee -a "\$BENCH_LOG" +import json, urllib.request +URL="http://localhost:5000/v1/completions" +print("===== COHERENCE (temperature=0) =====") +for p in ["Question: What is 2+2? Answer:","The capital of France is","Q: 3 cows + 2 cows = ? A:","Once upon a time"]: + b=json.dumps({"model":"qwen","prompt":p,"max_tokens":48,"temperature":0.0}).encode() + try: + txt=json.loads(urllib.request.urlopen(urllib.request.Request(URL,data=b,headers={"Content-Type":"application/json"}),timeout=300).read())["choices"][0]["text"] + except Exception as e: + txt=f"" + print(f"PROMPT {p!r} -> {txt!r}") +PYEOF + +for BS in $BENCH_SIZES; do + echo "===== static_benchmark dataset=$BENCH_DATASET BS=\$BS OSL=1024 =====" | tee -a "\$BENCH_LOG" + python -u $BENCH_CLIENT \\ + --server-url "http://localhost:5000/v1" --model qwen \\ + --batch-size \$BS --dataset $BENCH_DATASET --num-output-tokens $BENCH_OUTPUT_TOKENS \\ + $BENCH_ITERS 2>&1 | tee -a "\$BENCH_LOG" +done + +kill \$SERVER_PID 2>/dev/null || true +wait \$SERVER_PID 2>/dev/null || true +echo "===== DONE =====" +tail -30 "\$BENCH_LOG" +EOF +)" diff --git a/skills/run-qwen-model/run_qwen_vllm.sh b/skills/run-qwen-model/run_qwen_vllm.sh new file mode 100755 index 00000000000..651c5d4f50c --- /dev/null +++ b/skills/run-qwen-model/run_qwen_vllm.sh @@ -0,0 +1,110 @@ +#!/usr/bin/env bash +# Run the vLLM Qwen3-30B-A3B comparison at one batch size on OCI. +set -euo pipefail + +if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then + # shellcheck disable=SC1091 + source "${HOME}/.cog/setup.env.oci-hsg" +else + echo "ERROR: ~/.cog/setup.env.oci-hsg not found" >&2 + exit 1 +fi + +: "${COG_SSH_HOST:?}" +: "${COG_SCRATCH_ROOT:?}" + +EXPERIMENT_ID="${EXPERIMENT_ID:-unassigned}" +EXPERIMENT_HYPOTHESIS="${EXPERIMENT_HYPOTHESIS:-not-recorded}" +BENCH_BS="${BENCH_BS:-256}" +BENCH_OUTPUT_TOKENS="${BENCH_OUTPUT_TOKENS:-1024}" +NUM_WARMUP_ITERS="${NUM_WARMUP_ITERS:-2}" +NUM_TIMED_ITERS="${NUM_TIMED_ITERS:-5}" +VLLM_MAX_NUM_SEQS="${VLLM_MAX_NUM_SEQS:-512}" +VLLM_GPU_MEMORY_UTILIZATION="${VLLM_GPU_MEMORY_UTILIZATION:-0.9}" + +IMG="${VLLM_IMAGE:-$COG_SCRATCH_ROOT/images/87e4947c6ce36433.sqsh}" +HF_CKPT="${QWEN30B_HF:-$COG_SCRATCH_ROOT/checkpoints/qwen3-30b-a3b-hf}" +RUN_SUFFIX=$(printf '%s' "$EXPERIMENT_ID" | tr '[:upper:]_' '[:lower:]-' | tr -cd '[:alnum:]-') +RUN_NAME="vllm-qwen30b-${RUN_SUFFIX:-run}-$(date +%Y%m%d-%H%M%S)" +RUN_DIR="$COG_SCRATCH_ROOT/runs/$RUN_NAME" +REMOTE_SCRIPT="$RUN_DIR/run.sbatch" + +ssh -o BatchMode=yes "$COG_SSH_HOST" "mkdir -p '$RUN_DIR'; cat > '$REMOTE_SCRIPT'" </dev/null | head -1) +srun --container-image="$IMG" --container-mounts=/lustre:/lustre --no-container-mount-home bash -c ' +set -euo pipefail +export HF_HOME="$RUN_DIR/hf_home" +export VLLM_WORKER_MULTIPROC_METHOD=spawn +export OUTLINES_CACHE_DIR=/tmp/outlines_\$\$ +VENV=/opt/ray_venvs/nemo_rl.experience.sync_rollout_actor.SyncRolloutActor +PY=\$VENV/bin/python +VLLM=\$VENV/bin/vllm +MLM='"\$MLM_HOST"' +BENCH=\$MLM/tests/performance_tests/client/static_benchmark.py +SERVER_LOG="$RUN_DIR/server.log" +BENCH_LOG="$RUN_DIR/benchmark.log" + +cat > "\$BENCH_LOG" < "\$SERVER_LOG" 2>&1 & +SERVER_PID=\$! +trap '"'"'kill "\$SERVER_PID" 2>/dev/null || true'"'"' EXIT + +READY=0 +for _ in \$(seq 1 360); do + if grep -qE "Application startup complete|Uvicorn running on http://0.0.0.0:5000" "\$SERVER_LOG"; then + READY=1 + break + fi + if ! kill -0 "\$SERVER_PID" 2>/dev/null; then + echo "VLLM SERVER DIED" + tail -100 "\$SERVER_LOG" + exit 1 + fi + sleep 5 +done +if [[ "\$READY" != 1 ]]; then + echo "VLLM SERVER TIMEOUT" + tail -100 "\$SERVER_LOG" + exit 1 +fi + +\$PY -u "\$BENCH" --server-url http://localhost:5000/v1 --model qwen \ + --batch-size "$BENCH_BS" --dataset gsm8k --num-output-tokens "$BENCH_OUTPUT_TOKENS" \ + --num-iters "$NUM_TIMED_ITERS" --num-warmup-iters "$NUM_WARMUP_ITERS" \ + 2>&1 | tee -a "\$BENCH_LOG" +' +EOF + +JOB_ID=$(ssh -o BatchMode=yes "$COG_SSH_HOST" "cd '$RUN_DIR' && sbatch --parsable '$REMOTE_SCRIPT'") +echo "experiment_id=$EXPERIMENT_ID" +echo "job_id=$JOB_ID" +echo "run_dir=$RUN_DIR" From af734982d79f5fd688bbe81baa7de2d77c19bd07 Mon Sep 17 00:00:00 2001 From: shanmugamr1992 Date: Thu, 6 Aug 2026 11:17:58 -0700 Subject: [PATCH 2/4] Qwen optimizer --- .cursor/agents/qwen-model-optimizer.md | 14 +- skills/cog-setup-and-help/SKILL.md | 101 +- skills/nightly-sync/SKILL.md | 2 +- skills/nsight-system-analysis/SKILL.md | 6 + .../references/decision-gates.md | 6 +- .../references/measuring.md | 8 + .../references/moe-inference.md | 31 + .../references/updating-this-skill.md | 4 +- skills/run-qwen-model/EXPERIMENTS.md | 3577 +---------------- skills/run-qwen-model/SKILL.md | 35 +- skills/run-qwen-model/profile_qwen_mcore.sh | 23 +- skills/run-qwen-model/profile_qwen_vllm.sh | 9 +- skills/run-qwen-model/run_qwen_inference.sh | 12 +- skills/run-qwen-model/run_qwen_vllm.sh | 14 +- 14 files changed, 243 insertions(+), 3599 deletions(-) diff --git a/.cursor/agents/qwen-model-optimizer.md b/.cursor/agents/qwen-model-optimizer.md index d23ce5bb5e3..bded8248a1b 100644 --- a/.cursor/agents/qwen-model-optimizer.md +++ b/.cursor/agents/qwen-model-optimizer.md @@ -59,6 +59,13 @@ Supporting skills, only when the task calls for them: ## Baseline gate +Before the first run, confirm you are benchmarking **this** checkout: +`~/.cog/setup.env*` is machine-wide and may name a different Megatron-LM tree, +in which case cog syncs that one and every number you record describes code you +did not change — with no error. Check `echo "$COG_MEGATRON_REPO"` against +`git rev-parse --show-toplevel`, and verify each run's recorded `CODE_REVISION` +matches local `HEAD`. + Do not modify Megatron-Core until both fresh baselines are recorded: 1. Run vLLM DP4+EP under Nsight Systems at BS256. @@ -84,9 +91,10 @@ Repeat until mcore reaches vLLM: .sqlite .sqlite --label-a mcore --label-b vllm ``` - (the current baseline trace paths, under `nsys_trace/`, are recorded in - `EXPERIMENTS.md`). It auto-isolates one decode step per engine and prints - wall time, GPU-busy vs idle, launch counts, + The current baseline trace paths are recorded in `EXPERIMENTS.md`; re-capture + rather than hunt for a trace you cannot open. `forward_pass.py` auto-isolates + one decode step per engine and prints wall time, GPU-busy vs idle, launch + counts, and a per-category Δ table. Same µs/kernel with more launches ⇒ the lever is fusion / fewer launches; higher µs/kernel on the same shape ⇒ a real kernel-selection finding. Fall through to the skill's Steps 1–6 (exposed diff --git a/skills/cog-setup-and-help/SKILL.md b/skills/cog-setup-and-help/SKILL.md index 467b98f2e25..a792cce09d9 100644 --- a/skills/cog-setup-and-help/SKILL.md +++ b/skills/cog-setup-and-help/SKILL.md @@ -33,11 +33,13 @@ cluster via cog" in four steps: > submit`, `cog session exec`, `cog ensure-env`, …) **must `source > ~/.cog/setup.env` first** and use the `$COG_*` variables it defines > (`$COG_SSH_HOST`, `$COG_RUNTIME_ACCOUNT`, `$COG_INTERACTIVE_PARTITION`, -> `$COG_BATCH_PARTITION`, `$COG_SCRATCH_ROOT`, `$COG_MEGATRON_REPO`, …). -> Never assume hardcoded hostnames, accounts, partitions, or repo -> paths — read the env file. If the file is missing, run Step 3a to -> populate it (asking the user for each value with the documented -> defaults) before proceeding. +> `$COG_BATCH_PARTITION`, `$COG_SCRATCH_ROOT`, …). +> Never assume hardcoded hostnames, accounts, or partitions — read the +> env file. If the file is missing, run Step 3a to populate it (asking +> the user for each value with the documented defaults) before +> proceeding. **The repo is the exception**: it is not in the env file, +> because cog resolves it from your current directory. See *Which +> checkout does cog deploy?* below. ## Prerequisites @@ -136,10 +138,10 @@ hardcode anything):** Step 3b. Confirm each variable is set (`echo $COG_SSH_HOST`, …) and only re-prompt for any that come up empty. 3. **If it does not exist**, ask the user for each variable in the - table below using `AskUserQuestion` (or equivalent). For every - variable except `COG_MEGATRON_REPO`, *offer the listed default* — - accept the default if the user doesn't override. `COG_MEGATRON_REPO` - has no default; require an answer. + table below using `AskUserQuestion` (or equivalent), *offering the + listed default* — accept the default if the user doesn't override. + Do not ask for a repo path: cog derives it from the working + directory, and persisting one here is a known footgun (below). | Variable | Prompt to the user | Default | |---|---|---| @@ -151,7 +153,7 @@ hardcode anything):** | `COG_BATCH_PARTITION` | Partition for long-running batch jobs | `batch` | | `COG_IMPORT_PARTITION` | Partition that runs the CPU `enroot import` job | `cpu` | | `COG_SCRATCH_ROOT` | Absolute path to your scratch root on the cluster (cog will store `.sqsh`, `.venv`, workspaces, run logs under this) | `/lustre/fsw/portfolios/coreai/users/${USER}/agents-space` | - | `COG_MEGATRON_REPO` | Absolute path to your local Megatron-LM checkout | **no default — must come from the user** | + | `COG_MEGATRON_REPO` | *Not stored.* Optional per-shell override; omit it and cog uses the checkout you are standing in | **omit** | 4. Write the answers to `~/.cog/setup.env`. **Each line must use `export KEY="value"`** — without the `export`, the variables exist @@ -172,7 +174,7 @@ hardcode anything):** export COG_BATCH_PARTITION="batch" export COG_IMPORT_PARTITION="cpu" export COG_SCRATCH_ROOT="/lustre/fsw/portfolios/coreai/users/${USER}/agents-space" - export COG_MEGATRON_REPO="/Users/${USER}/Megatron-LM" + # No COG_MEGATRON_REPO — cog resolves the repo from your cwd. EOF source ~/.cog/setup.env # Sanity check — both lines must show the value, not blank: @@ -183,6 +185,39 @@ hardcode anything):** 5. From this point on, every `cog` invocation in this skill (and in later turns) references `$COG_*` rather than literal values. +### Which checkout does cog deploy? + +**The one you are standing in.** `--repo` defaults to the current +directory, resolved to its git toplevel, which also makes it correct +inside a `git worktree add` worktree. So the normal usage is to `cd` into +the checkout you want and omit `--repo` entirely, as the examples below +do. + +Pass `--repo` explicitly only to deploy a checkout you are *not* in: + +```bash +cog submit --repo /path/to/other/Megatron-LM … +``` + +**Never persist a repo path in `~/.cog/setup.env`.** That file is +machine-wide, so a baked-in `COG_MEGATRON_REPO` overrides the cwd default +in every shell. The moment a second checkout exists, half your commands +deploy the wrong one — and nothing fails, because the other tree is a +valid repo. Cog will happily sync it, the job runs, and the numbers get +attributed to code you never sent. + +Two habits make this self-checking. Run cog from the tree you are editing +so cwd resolution does the work, and verify after any run that the +recorded revision matches your local `HEAD`: + +```bash +git rev-parse HEAD # must equal the CODE_REVISION in the run record +``` + +One trap to know: if your cwd is the **cog clone** rather than a +Megatron-LM checkout, cwd resolution finds the cog repo and `doctor` +fails the repo-profile check. That is a loud failure, not a silent one. + --- ## Step 3b — Register the cluster with cog @@ -219,7 +254,7 @@ Verify and run `doctor`: ```bash cog cluster ls -cog doctor --repo "$COG_MEGATRON_REPO" +cog doctor # run from inside your Megatron-LM checkout ``` `cog doctor` checks: local Python / cog version, SSH reachability, @@ -251,10 +286,9 @@ job. Skip if you'd rather see one combined job: ```bash source ~/.cog/setup.env -cog prepare-image --repo "$COG_MEGATRON_REPO" +cog prepare-image cog ensure-env \ - --repo "$COG_MEGATRON_REPO" \ --run-name env-warmup \ --gpus 1 \ --time 00:20:00 \ @@ -272,7 +306,6 @@ returns `cache_hit: true` instantly. ```bash source ~/.cog/setup.env cog submit \ - --repo "$COG_MEGATRON_REPO" \ --run-name verify-cog-basic \ --command 'python -m pytest tests/unit_tests/test_basic.py -v -o addopts=' \ --gpus 1 \ @@ -313,7 +346,6 @@ path: ```bash source ~/.cog/setup.env cog submit \ - --repo "$COG_MEGATRON_REPO" \ --run-name verify-cog-gpt \ --command 'python -m torch.distributed.run --nproc-per-node 8 --log-dir "$TORCHRUN_LOG_DIR" -m pytest -xvs tests/unit_tests/models/test_gpt_model.py::TestGPTModel::test_constructor' \ --gpus 8 \ @@ -353,8 +385,7 @@ torch.distributed.run` from the start avoids the trap. > > # Hold an 8-GPU node for 3 hours. > cog session start \ -> --repo "$COG_MEGATRON_REPO" \ -> --run-name iter-debug \ +> > --run-name iter-debug \ > --gpus 8 --nodes 1 --ntasks-per-node 1 \ > --time 03:00:00 \ > --partition "$COG_INTERACTIVE_PARTITION" @@ -404,7 +435,7 @@ it still works while sync is broken: ```bash source ~/.cog/setup.env -cog profile --repo "$COG_MEGATRON_REPO" --run-name fallback --pretty # sqsh_plan.sqsh_path +cog profile --run-name fallback --pretty # sqsh_plan.sqsh_path ``` Do **not** reach for `cog ensure-env` to discover the venv path here — it @@ -420,7 +451,7 @@ ssh "$COG_SSH_HOST" ls -d "$COG_SCRATCH_ROOT"/envs/megatron_lm/*/.venv ```bash STAGE="$COG_SCRATCH_ROOT/workspaces/megatron_lm/fallback_$(date +%s)/repo" ssh "$COG_SSH_HOST" mkdir -p "$STAGE" -rsync -a --exclude .git "$COG_MEGATRON_REPO"/ "$COG_SSH_HOST:$STAGE/" +rsync -a --exclude .git "$(git rev-parse --show-toplevel)"/ "$COG_SSH_HOST:$STAGE/" ``` > **Why a fresh path, not `rsync --delete` over the old one.** `--delete` @@ -578,6 +609,24 @@ cog logs app --lines 50 (e.g. user moved their Megatron-LM checkout), edit just that line in the file rather than wiping it. Sourcing the file is cheap — always do it at the start of any turn that runs a `cog` command. +- **Per-cluster overrides: `~/.cog/setup.env.`.** Some workflows + keep a variant per cluster (e.g. `~/.cog/setup.env.oci-hsg`) and source + *that* in preference to the base file. Two things to know. Scripts that + choose between them usually do it by mere **existence** (`if [[ -f + setup.env. ]]`), so once the variant exists it always wins and + the base file is effectively dead for those scripts, whichever cluster + you intended. And when you correct a variable, correct it in **every** + variant — `grep -n ~/.cog/setup.env*` shows them all at once. +- **A repo path in any `setup.env*` when you have more than one checkout.** + Do not put one there; see *Which checkout does cog deploy?* above. This is + the most dangerous value to persist, because nothing fails: cog syncs the + *other* checkout, the job runs, and you attribute its numbers to code you + never sent. Symptoms are a benchmark that ignores your edit and a + `CODE_REVISION` in the run record that does not match your `HEAD`. Note + that removing the variable does not merely fix the current mistake — it + converts the whole failure mode from silent to loud, since a command that + needed it now errors with `REPO_NOT_FOUND` instead of quietly deploying + the wrong tree. - **`cog` runs but `--repo` / `--partition` / `--ssh-host` flags receive empty strings.** The env file was written without `export`, so `source ~/.cog/setup.env` populated the variables in the current @@ -677,7 +726,7 @@ Output: `checks[]` (each with `name`, `status` ∈ {`ok`,`degraded`,`fail`}, `du Exit `0` if `overall` is `ok` or `degraded`; `1` if `fail`. ```bash -cog doctor --repo "$COG_MEGATRON_REPO" +cog doctor ``` ### `cog profile` — resolve the run plan without touching the cluster @@ -966,9 +1015,10 @@ See Steps 1-3 above (`cluster add --set-default`, populate ### 2. Warm image + env once per recipe ```bash +cd /path/to/your/Megatron-LM # cog deploys the checkout you are in source ~/.cog/setup.env -cog prepare-image --repo "$COG_MEGATRON_REPO" -cog ensure-env --repo "$COG_MEGATRON_REPO" \ +cog prepare-image +cog ensure-env \ --run-name env-warmup --gpus 1 --time 00:20:00 \ --partition "$COG_INTERACTIVE_PARTITION" ``` @@ -977,7 +1027,6 @@ cog ensure-env --repo "$COG_MEGATRON_REPO" \ ```bash cog submit \ - --repo "$COG_MEGATRON_REPO" \ --run-name test-basic \ --command 'python -m pytest tests/unit_tests/test_basic.py -v -o addopts=' \ --gpus 1 --time 00:10:00 \ @@ -991,11 +1040,11 @@ you'll run **more than 2 commands** against the same code, start a session instead of submitting again. ```bash -cog session start --repo "$COG_MEGATRON_REPO" \ +cog session start \ --session-handle iter-debug --gpus 8 --time 03:00:00 \ --partition "$COG_INTERACTIVE_PARTITION" -cog session exec --session-handle iter-debug --repo "$COG_MEGATRON_REPO" \ +cog session exec --session-handle iter-debug \ --command 'python -m pytest tests/unit_tests/test_basic.py -v' # … re-run as many times as needed … diff --git a/skills/nightly-sync/SKILL.md b/skills/nightly-sync/SKILL.md index cd3b85f2e0a..c86b75c396a 100644 --- a/skills/nightly-sync/SKILL.md +++ b/skills/nightly-sync/SKILL.md @@ -114,7 +114,7 @@ CODEOWNERS change after the PR lands is painful. **NEVER manually edit `uv.lock`.** It is a machine-generated lockfile. If it needs to change, it must be regenerated with `uv lock` inside a CUDA -container (see `.claude/skills/build-and-test/SKILL.md`). +container (see `skills/mcore-build-and-dependency/SKILL.md`). ### Git Source Reconciliation (pyproject.toml) diff --git a/skills/nsight-system-analysis/SKILL.md b/skills/nsight-system-analysis/SKILL.md index d4255354585..e555a56d7ff 100644 --- a/skills/nsight-system-analysis/SKILL.md +++ b/skills/nsight-system-analysis/SKILL.md @@ -94,6 +94,12 @@ python scripts/forward_pass.py mcore.sqlite vllm.sqlite --label-a mcore --label- **Caveats:** - `GPU-time (Σ durations)` per category **sums** kernel durations and so over-counts wall time when streams overlap — that's intentional for composition. Use the reported **GPU-busy (interval union)** for the true single-step wall figure, and the anchor **period** for the forward-pass time. - The built-in taxonomy is tuned for MoE decode; verify the category assignments once (as in Step 3's YAML verification) if a large "other/misc" bucket appears, and extend via `--yaml`. +- **Confirm the step boundary with a second anchor before quoting the period.** Re-run with `--anchor `; two independent anchors should agree to a few µs on both period and kernel count. If they disagree, auto-detection latched onto something that does not fire exactly once per step and every downstream number is wrong. +- **Reconcile the step-period ratio against the end-to-end ratio, and expect them to differ.** The step period measures only steady-state decode in the *densest* window; the benchmark's throughput and TPOT also carry prefill, scheduling, and client overhead. Report both and say which is which — a 2.31× step-period ratio next to a 1.52× throughput ratio is normal when a fixed prefill cost dilutes the shorter run, not evidence that one of them is wrong. + +**Sanity check the attribution arithmetic:** `Δ step wall` must equal `Δ GPU-busy + Δ idle` exactly. The per-category Δs sum to *less* than `Δ GPU-busy` (they are Σ-durations, so overlap differs). If the first identity fails, the two windows are not both one step. + +**Profile at the sequence length you intend to ship.** Per-step cost scales with context length, and the two engines need not scale alike — one MoE decode comparison measured 2.31× at OSL128 but 1.52× at OSL1024, because the faster engine's step grew 61% with context while the slower engine's grew 6%. A short OSL is cheaper to capture and can *invert* which levers matter, so treat a short-OSL profile as a lever-finder and re-measure the ceiling at the target OSL before committing to a change. For deeper attribution (source-level root cause, exposed comm, module-slicing) fall through to Steps 1–6 below, restricting the windows to the decode region this workflow identified. diff --git a/skills/optimize-inference-siddharth/references/decision-gates.md b/skills/optimize-inference-siddharth/references/decision-gates.md index c16f6a6052d..5db393ff6bd 100644 --- a/skills/optimize-inference-siddharth/references/decision-gates.md +++ b/skills/optimize-inference-siddharth/references/decision-gates.md @@ -10,9 +10,9 @@ optimization worked perfectly, how much would the step actually get faster?* If the answer is small, or if the mechanism you plan to attack is not the mechanism that is slow, you stop before writing production code. -Source: the Qwen3-30B-A3B EP4 campaign ledger -(`skills/run-qwen-model/EXPERIMENTS.md`), where three gates each killed a -multi-session effort, and a fourth picked the right candidate out of six. +Source: a Qwen3-30B-A3B EP4 optimization campaign on 4×GB200, where three gates +each killed a multi-session effort, and a fourth picked the right candidate out +of six. ## Why share is not headroom diff --git a/skills/optimize-inference-siddharth/references/measuring.md b/skills/optimize-inference-siddharth/references/measuring.md index 5270f206346..2537e5ff3e1 100644 --- a/skills/optimize-inference-siddharth/references/measuring.md +++ b/skills/optimize-inference-siddharth/references/measuring.md @@ -72,6 +72,14 @@ nsys profile --trace=cuda,nvtx --sample=none \ -o inference_profile ``` +**Grep the launch scripts for these flags before you submit, every time.** This +table being written down does not mean the harness obeys it: both +`profile_qwen_mcore.sh` and `profile_qwen_vllm.sh` shipped `--trace=cuda,nvtx,osrt` +long after the bisection was recorded here, and would each have burned an +allocation on an unusable capture. Removing `osrt` was the whole fix — both then +finalized first try. A banned flag costs a full profile round-trip to discover +and one `rg osrt` to prevent. + This is a real constraint on method, not a nuisance: **the documented way to attribute host time is exactly the way that breaks.** Two consequences, both with working substitutes below — attribute host-side idle from the CUDA-API rows in a diff --git a/skills/optimize-inference-siddharth/references/moe-inference.md b/skills/optimize-inference-siddharth/references/moe-inference.md index 167538e7404..826d3539872 100644 --- a/skills/optimize-inference-siddharth/references/moe-inference.md +++ b/skills/optimize-inference-siddharth/references/moe-inference.md @@ -37,6 +37,37 @@ Relevant files: [backends.py](megatron/core/models/backends.py) [moe_module_specs.py](megatron/core/models/gpt/moe_module_specs.py), [moe_layer.py](megatron/core/transformer/moe/moe_layer.py). +### `--inference-grouped-gemm-backend vllm` is not the kernel vLLM runs + +The flag name invites the wrong conclusion. The mapping is: + +| value | kernel | +|---|---| +| `vllm` | vLLM's **Triton** `_fused_moe_kernel` | +| `flashinfer` | FlashInfer **`cutlass_fused_moe`** | +| `torch` | `torch._grouped_mm` | + +vLLM itself, on a BS256 MoE decode step, does **not** dispatch to its Triton +kernel — it runs CUTLASS/TRT-LLM `bmm_Bfloat16_..._t128x8x128...`, whose N-tile +of **8** is shaped for the handful of tokens each expert sees during decode. So +selecting `vllm` reproduces vLLM's *slower* path while believing you matched it. +On one matched pair this cost **25.2 vs 10.7 µs/kernel at identical launch +counts** — 2.35×, the single largest row in the gap table. + +Two consequences worth internalizing: + +- **Identical launch counts with a large Δ means kernel selection, not fusion.** + Check the kernel *names* on both sides before reaching for a fusion lever; + `forward_pass.py`'s `µs/kernel` column versus `#` column is the tell. +- **The CUTLASS path is gated shut for gated activations.** + `transformer_config.py` rejects `flashinfer` whenever `gated_linear_unit` is + set, so every SwiGLU/GeGLU model — Qwen3 included — is excluded from it *by + validation, not by kernel capability*. That same restriction is why mcore pays + a standalone `_silu_mul_bounded_kernel` per layer while vLLM shows **zero** + activation kernels (its CUTLASS epilogue fuses SwiGLU). Treat the GEMM and the + activation as **one** lever: measured together they were 3137 vs 1026 µs/step, + 3.06×, 35% of the whole engine gap. + ## What each replaced component does differently ### `InferenceTopKRouter` diff --git a/skills/optimize-inference-siddharth/references/updating-this-skill.md b/skills/optimize-inference-siddharth/references/updating-this-skill.md index 29850d2d696..c6c8224d04a 100644 --- a/skills/optimize-inference-siddharth/references/updating-this-skill.md +++ b/skills/optimize-inference-siddharth/references/updating-this-skill.md @@ -189,6 +189,6 @@ discoverable rather than mysterious. | Date | Change | Source | |---|---|---| | 2026-07 | Initial skill: the five moves, hard rules 1-9, the CUDA-graph / MoE / Mamba-Triton / host-path / measuring references, commit log, review checklist | Siddharth Singh's 2026 inference work, 29 commits | -| 2026-07-28 | Added `decision-gates.md` (ceiling-before-building, per-launch fixed costs, three gates that each killed a multi-week effort) and `vllm-differential.md` (competitor-trace comparison). Hardened `measuring.md`: nsys flag combinations that deadlock finalization, the node-vs-graph trace control, `perf_counter` phase timing, union-busy idle accounting with gap-size decomposition, the same-session back-to-back A/B protocol, kernel-to-e2e conversion, ledger requirements. Added *What a wide capture costs you* to `cuda-graphs.md`. Added per-GEMM tile tuning, the measured backend comparison, the training-path-only fusion flags, and comm-vs-skew to `moe-inference.md`. Added hard rules 10 (A/B protocol) and 11 (non-bit-exact acceptance) plus the *flags that look like free wins* table | Qwen3-30B-A3B EP4 on 4×GB200, `skills/run-qwen-model/EXPERIMENTS.md` | +| 2026-07-28 | Added `decision-gates.md` (ceiling-before-building, per-launch fixed costs, three gates that each killed a multi-week effort) and `vllm-differential.md` (competitor-trace comparison). Hardened `measuring.md`: nsys flag combinations that deadlock finalization, the node-vs-graph trace control, `perf_counter` phase timing, union-busy idle accounting with gap-size decomposition, the same-session back-to-back A/B protocol, kernel-to-e2e conversion, ledger requirements. Added *What a wide capture costs you* to `cuda-graphs.md`. Added per-GEMM tile tuning, the measured backend comparison, the training-path-only fusion flags, and comm-vs-skew to `moe-inference.md`. Added hard rules 10 (A/B protocol) and 11 (non-bit-exact acceptance) plus the *flags that look like free wins* table | Qwen3-30B-A3B EP4 on 4×GB200 | | 2026-07-28 | Made the skill self-maintaining: this file, the editing authorization at the top of `SKILL.md`, the Step 6 feedback loop, the checklist's skill-maintenance block, and this log | — | -| 2026-08-03 | Added `scripts/` (`steady_window.py`, `union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) so the trace method is runnable rather than only described. **Corrected** the fp32 reduce-scatter buffer in `moe-inference.md`, which was documented as a reproducibility feature and costs 2.5% (the `atol=0` guarantee is about the fp32 accumulator, and `multimem.ld_reduce` accumulates in f32 regardless of operand width). Added to `measuring.md`: cross-trace window drift at differing KV length (27% bucket error, and the launch-count / untouched-bucket / seqlen-inference defenses), the sum÷union overlap ratio, and naming a kernel from its trace neighbours instead of framework instrumentation that cannot fire under graph replay. Added to `vllm-differential.md`: resolve the competitor's kernel to a *package* and check whether it is already installed (+2.6% from flashinfer `trtllm-gen`, after two sessions tuning inside flash-attn), and the packing split. Added thesis moves 5-width and 6, three routing-table rows, and the convention that a rejection records its mechanism and date so it can be re-tested | Qwen3-30B-A3B EP4 on 4×GB200, PR #6064, `skills/run-qwen-model/EXPERIMENTS.md` sessions 16-18 | +| 2026-08-03 | Added `scripts/` (`steady_window.py`, `union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) so the trace method is runnable rather than only described. **Corrected** the fp32 reduce-scatter buffer in `moe-inference.md`, which was documented as a reproducibility feature and costs 2.5% (the `atol=0` guarantee is about the fp32 accumulator, and `multimem.ld_reduce` accumulates in f32 regardless of operand width). Added to `measuring.md`: cross-trace window drift at differing KV length (27% bucket error, and the launch-count / untouched-bucket / seqlen-inference defenses), the sum÷union overlap ratio, and naming a kernel from its trace neighbours instead of framework instrumentation that cannot fire under graph replay. Added to `vllm-differential.md`: resolve the competitor's kernel to a *package* and check whether it is already installed (+2.6% from flashinfer `trtllm-gen`, after two sessions tuning inside flash-attn), and the packing split. Added thesis moves 5-width and 6, three routing-table rows, and the convention that a rejection records its mechanism and date so it can be re-tested | Qwen3-30B-A3B EP4 on 4×GB200, PR #6064 | diff --git a/skills/run-qwen-model/EXPERIMENTS.md b/skills/run-qwen-model/EXPERIMENTS.md index b81ed871c6d..97f1bd6a5aa 100644 --- a/skills/run-qwen-model/EXPERIMENTS.md +++ b/skills/run-qwen-model/EXPERIMENTS.md @@ -3,8 +3,9 @@ Goal: make Megatron-Core EP4 inference match or exceed the vLLM DP4+EP throughput on one OCI 4×GB200 node without correctness regressions. -This ledger starts from scratch. Append every experiment, including failures -and regressions. Never edit an earlier result after it is recorded. +This ledger starts empty. Append every experiment, including failures and +regressions. Never edit an earlier result after it is recorded — supersede it +with a new row. ## Fixed protocol @@ -27,3564 +28,40 @@ parallelism, or warmup/timed counts differ. ## Baselines -| ID | Engine | Throughput | Avg latency | TPOT | Job / run | Nsight trace | Status | -|---|---|---:|---:|---:|---|---|---| -| VLLM-BASELINE | vLLM DP4+EP | 23,606.7 tok/s | 1,368.2 ms | 10.844 ms/tok | 5547673 / `vllm-qwen30b-nsys-20260722-093437` | `vllm_profile.nsys-rep`, `.sqlite` | Pass | -| MCORE-BASELINE | mcore EP4/TP1 | 12,346.1 tok/s | 2,590.1 ms | 20.735 ms/tok | 5553135 / `qwen-30b-nsys-20260722-161020` | `mcore_profile.nsys-rep`, `.sqlite` | Pass | - -Fresh profile gap: mcore delivers 52.30% of vLLM throughput and is 47.70% -below the target. vLLM is 1.912× faster on this profile workload. - -Baseline order is mandatory: +Nothing recorded yet. Baseline order is mandatory: 1. Record `VLLM-BASELINE` with Nsight Systems. 2. Record `MCORE-BASELINE` with Nsight Systems. 3. Compute the absolute and percentage gap. 4. Only then modify Megatron-Core. -## Experiment index - -| ID | Date | Hypothesis | Changed files / flags | Throughput | Delta vs mcore baseline | Correctness | Job / run | Conclusion | -|---|---|---|---|---:|---:|---|---|---| -| VLLM-BASELINE | 2026-07-22 | Establish the fixed competitor target | none | 23,606.7 | n/a | Benchmark pass | 5547673 | Target established | -| MCORE-BASELINE | 2026-07-22 | Establish the fixed EP4 starting point | `max_requests=256` | 12,346.1 | baseline | Benchmark pass | 5553135 | Starting point established | -| QWEN-001 | 2026-07-22 | Single-kernel FC1+SwiGLU+FC2+topk-reduce mega-fusion beats the 4-kernel vLLM MoE path | `megatron/core/inference/moe/fused_moe_decode.py` (new), `dev/moe_fused/harness.py` (new) | microbench only | n/a (0.68–0.80× kernel) | Numerics pass (max_abs 2.6e-5, allclose) | session `qwen-moe-kernel` | Rejected — fused kernel 20–50% slower than reference; not integrated | -| QWEN-002 | 2026-07-22 | Fusing SiLU(gate)*up into the FC1 GEMM epilogue (removing bounded_silu_mul + the 2N round-trip) speeds up the decode MoE path without hurting FC2 tiling | `vllm_fused_moe.py` (FUSE_SWIGLU), `experts.py` (`fuse_fc1_activation=True`), `dev/moe_fused/harness_fc1.py`, `dev/moe_fused/run_e2e_insession.sh` (new) | **22,741.9 tok/s** (OSL1024) | **+0.55%** vs same-env fusion-off (22,617.7) | Coherent + numerics (max_abs 3.9e-5) | session `qwen-moe-kernel` in-session A/B | Accepted — MoE path 1.25×, e2e +0.55% throughput / −0.55% TPOT, no regression | -| PROFILE-OSL1024 | 2026-07-22 | Re-profile at the real throughput regime (BS256/OSL1024) to find the true vLLM→mcore gap and the dominant decode bottleneck | none (profiling only); `dev/moe_fused/profile_insession.sh` (new), `dev/moe_fused/vllm_osl1024_tput.sbatch` (new) | vLLM **33,994.5** vs mcore **~22,700** tok/s | mcore = **66.8%** of vLLM (**vLLM 1.50×**) | Both coherent | vLLM job `5555787` (tput) + `5555868` (nsys); mcore in-session `prof256a` | Gap is real & large at OSL1024. mcore decode: GPU 79% busy; MoE grouped-GEMM 41%, attn 22%, MoE routing 12% (49k tiny kernels), **exposed EP comm 11.5%**, norm/elt 10%, GPU idle 21%. vLLM uses TRT-LLM fused MoE (1-kernel routing + cutlass bmm + fused finalize, ~~0 exposed comm~~ — **comm claim falsified, see VLLM-COMM-S16**). Next target: routing-kernel storm + exposed EP comm (pending approval) | -| SESSION2-BASE | 2026-07-23 | Re-establish clean OSL1024 baseline in fresh session `qwen-opt` (fusion on, histogram off) before autonomous optimization campaign | none (config = current best) | **22,398.9 tok/s** | baseline for session 2 (−1.5% vs QWEN-002 run, within variance) | Coherent | session `qwen-opt` run `e2ebase` | Clean reference; 65.9% of vLLM 33,994.5 | -| QWEN-003 | 2026-07-23 | Replacing per-pair `atomic_add` in the MoE local-token count kernel with a `tl.histogram` variant (one atomic/bin/CTA) cuts the routing-kernel cost | `permute.py` (`_count_local_tokens_kernel_histogram`, env `MCORE_MOE_HISTOGRAM_COUNT`), `dev/moe_fused/harness_count.py` (new) | microbench only | **0.96× (wash)** | **EXACT integer match** vs reference | session `qwen-opt` `harness_count.py` | Rejected — count-kernel cost is per-launch fixed overhead, not atomic contention; in-kernel rewrite can't help. Default OFF. Real lever = fewer launches + less host-scheduling idle | -| QWEN-004 | 2026-07-23 | flashinfer sampling backend is faster than torch sampling | `run_e2e_cfg.sh` (`--inference-dynamic-batching-sampling-backend flashinfer`) | server crash | n/a | n/a | session `qwen-opt` `flashinfer` | Rejected — incompatible with `full_iteration_inference` CUDA graph capture: `RuntimeError: Generator not registered with the capturing graph` in `flashinfer.sampling`. torch sampling stays. | -| QWEN-005a | 2026-07-23 | `async-sched-mode=serial` overlaps host resolve with next forward, hiding the ~2.3ms/step (21%) GPU idle | `run_e2e_cfg.sh` (`--inference-dynamic-batching-async-sched-mode serial`) | server crash | n/a | n/a | session `qwen-opt` `serial` | Blocked by explicit guards: `ValueError: Async scheduling does not support expert parallelism` (+ separate MoE guard). Guards env-gated for experiment → QWEN-005b | -| QWEN-005b | 2026-07-23 | Guards were merely conservative; async serial works for MoE+EP if opened | env-gated EP+MoE guards in `dynamic_engine.py` + `text_generation_controller.py` (`MCORE_ALLOW_ASYNC_MOE`) | **hang** | n/a | server init OK but **first decode request hangs** (>2min) | session `qwen-opt` `asyncmoe` (cancelled) | Rejected — the guard encodes a real limitation: async serial + nvls alltoall + EP deadlocks on the first decode step. Patch reverted (tree clean). Would need real engine work to support. | -| QWEN-006 | 2026-07-23 | `nccl` AllGather/ReduceScatter inference dispatcher overlaps/costs less than `nvls` | `run_e2e_cfg.sh` (`--inference-moe-token-dispatcher-type nccl`) | **14,677 tok/s** | **0.66× (much worse)** | Coherent | session `qwen-opt` `nccl` | Rejected — nccl pads to worst-case per-rank token count (fixed-count AllGather), inflating comm volume ~2×. `nvls` variable-count stays the best dispatcher. | -| QWEN-007 | 2026-07-23 | `async-sched-mode=serial` (guards opened) hides the between-graph idle at the real OSL1024 regime | env-gated guards + `run_e2e_cfg.sh` | **22,589 tok/s** | **+0.85% (marginal)** | Runs at BS256 (hang was single-request-only) | session `qwen-opt` `async1024` | Marginal. Proves decode at OSL1024 is **NOT idle-bound** (the earlier "21% idle" was an OSL256 prefill artifact). Not worth shipping (unsupported path + tiny gain). Guards left env-gated default-off. | -| QWEN-008 | 2026-07-23 | `CUDA_DEVICE_MAX_CONNECTIONS=8` (was hardcoded 1) lets comm & compute overlap on separate HW queues, hiding exposed NVLS comm | `run_e2e_cfg.sh` (env override) | **~22,556 tok/s** | **~flat (noise)** | Coherent | session `qwen-opt` `maxconn8` | Reject — overlap is bounded by the full-iteration CUDA-graph structure / data deps, not connection count. No effect. | -| QWEN-010 | 2026-07-23 | `torch` grouped-GEMM backend (`torch.nn.functional.grouped_mm`, cuBLAS) beats the vLLM Triton fused-MoE backend on GB200 | `run_e2e_cfg.sh` (`--inference-grouped-gemm-backend torch`) | **18,434 tok/s** | **0.82× (worse)** | Coherent | session `qwen-opt` `gemmtorch` | Rejected — vLLM Triton backend stays best. (`flashinfer` cutlass backend is blocked for SwiGLU; only torch/vllm allowed.) All three backends now evaluated → vLLM is optimal. | -| QWEN-009 | 2026-07-23 | Built-in `--moe-router-fusion` (TE fused softmax+topk) + `--moe-permute-fusion` cut the routing critical path (~18%) | `run_e2e_cfg.sh` EXTRA_SERVER_ARGS | server crash | n/a | n/a | session `qwen-opt` `routperm`/`routfus` | Rejected — both crash: `AssertionError: hidden_size mismatch: 128 vs 8`. TE fused router emits a dense **128-expert** routing map, but `InferenceTopKRouter` (transformer_impl=inference_optimized) uses a dense **top-8** contract for the vLLM/nvls dispatcher. The built-in fusions are wired to the training MoE path only. A hand-written fused softmax+topk would need to honor the top-8 inference contract. | -| PROFILE-DECODE | 2026-07-23 | Get the TRUE per-step decode bottleneck (prior OSL256 totals were prefill-contaminated) | analysis of archived `mcore_osl256.sqlite`, pure-decode window (t0+220s, big-dispatch-free) | n/a | n/a | n/a | local sqlite | **Corrected model** (decode GPU-time share): MoE grouped-GEMM (`_fused_moe_kernel`) **~40% #1**, routing (count/moe_sum/topk/scatter/softmax/meta) ~18%, exposed comm (dispatch 122k + combine 320k) ~16%, attention ~13%, norm/elt ~10%. Kernels overlap across streams → **wall = per-layer critical path** (attn→router→dispatch→GEMM→combine). vLLM wins via TRT-LLM fused MoE (fused routing+finalize; the ~~0 exposed comm~~ half is **falsified, see VLLM-COMM-S16**). Explains QWEN-002 1.25× kernel → +0.55% e2e. | -| CLEANBASE-S3 | 2026-07-24 | Fresh clean un-profiled OSL1024 baseline (SwiGLU fusion OFF) in session `qwen-fuse` before session-3 fusion campaign | none (`MCORE_FUSE_FC1_ACT=0`) | **22,241.5 tok/s** | baseline (session 3) | Coherent | session `qwen-fuse` `cleanbase` | Clean reference; 65.4% of vLLM un-profiled 33,994.5. avg_latency 11,494 ms | -| QWEN-002-CONFIRM | 2026-07-24 | Re-measure QWEN-002 SwiGLU FC1-epilogue fusion (lever #3) cleanly at OSL1024 | `experts.py` (`MCORE_FUSE_FC1_ACT=1`), `vllm_fused_moe.py` FUSE_SWIGLU | **22,269.5 tok/s** | **+0.13% vs OFF** (wash) | Coherent | session `qwen-fuse` `qwen002on` | Confirms QWEN-002: MoE-kernel 1.25× but e2e wash at OSL1024 — FC1 activation is not the decode wall bottleneck. Kept on (free, exact). | -| CLEANBASE-S4 | 2026-07-25 | Fresh clean un-profiled OSL1024 reference in session `qwen-cutlass2` at the session-3 best config, before the grouped-GEMM/cutlass campaign | none (`MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`, vllm backend) | **22,657.7 tok/s** | baseline (session 4); −0.17% vs QWEN-011 (within drift) | Coherent | session `qwen-cutlass2` `ref0-1785010433` | Clean reference; 66.65% of vLLM un-profiled 33,994.5. avg_latency 11,282 ms, TPOT 11.299 ms/tok | -| QWEN-012 | 2026-07-25 | **Decision gate**: quantify grouped-GEMM headroom before writing any kernel — is the ~40% decode share recoverable inefficiency or irreducible expert-weight traffic? | none (analysis); `dev/moe_fused/harness_roofline.py` (new) | n/a (microbench) | n/a | n/a | session `qwen-cutlass` `c05fd5f7` | **The decode grouped GEMM is memory-bound, not FLOP-bound.** Weight traffic 302 MB/layer/rank vs a *measured* 6.081 TB/s streaming-read ceiling ⇒ 49.66 µs floor; production FC1+FC2 = 72.13 µs = **1.45× off roofline**. Achieved 63–83 TFLOP/s on valid FLOPs (~3% of GB200 BF16 peak). Padding waste is real (74.3% dead rows at BLOCK_M=64) but nearly **free**: cutting it 3.89×→1.49× buys only ~1.2×. **⇒ Stage 4 (hand-written CUTLASS/CuTe grouped GEMM) cannot win**: its entire ceiling is 1.45×, and 1.26× of that is reachable by Triton tile retuning alone (→ QWEN-013). Remaining levers are elsewhere (exposed EP comm ~16%, routing ~18%) | -| QWEN-013 | 2026-07-25 | Tiling, not the GEMM implementation, is the gap QWEN-012 found: FC1 (N=768) and FC2 (N=2048) want *opposite* BLOCK_SIZE_N, which vLLM's single shared config cannot express | `vllm_fused_moe.py` (`_get_decode_tuned_configs`, per-GEMM `config_fc1`/`config_fc2`, env `MCORE_MOE_GEMM_TUNE`), `dev/moe_fused/harness_gemmtune.py` (new) | **23,636.0 tok/s** | **+4.32% vs CLEANBASE-S4** | **Bit-exact** (max_abs 0.0 at 128/256/384/512 tokens) + coherent | session `qwen-cutlass2` `gemmtune-1785010791` | **Accepted** — GEMM GPU time 77.65→61.89 µs (1.255×), whole MoE call 100.64→84.31 µs (1.194×), e2e +4.32% with tight variance (23.55–23.67k). Bit-exact because `BLOCK_SIZE_K` is unchanged, so the fp32 K-reduction order is identical. Default OFF; enable with `MCORE_MOE_GEMM_TUNE=1`. Now **69.53%** of vLLM | -| QWEN-013b | 2026-07-25 | The QWEN-013 tuned tiles regress past the decode point, so the fallback threshold should be 384, not 512 | `vllm_fused_moe.py` (`_get_decode_tuned_configs` gate `M > 512` → `M > 384`) | **23,646.0 tok/s** | **+4.36% vs CLEANBASE-S4** (+0.04% vs QWEN-013, i.e. same) | Bit-exact (unchanged kernel) + coherent | session `qwen-cutlass2` `gemmtune2-1785012743` | **Accepted** — CUDA-graph device time vs default tiles is 1.289×/1.197×/1.120×/**0.952×** at 128/256/384/512 tokens, so 512 was a 5% regression inside the tuned range. Decode runs at 256 so e2e is unchanged (23,646.0 vs 23,636.0, within variance); the guard just removes a latent regression for other batch shapes. **Current best config.** | -| QWEN-014 | 2026-07-25 | The already-wired `--inference-grouped-gemm-backend flashinfer` (`flashinfer.fused_moe.cutlass_fused_moe`) accepts BF16 and beats the vLLM Triton path; the ledger's "blocked for SwiGLU" note (QWEN-010) is unsupported by the code | none shipped; `dev/moe_fused/inspect_flashinfer.py`, `dev/moe_fused/harness_flashinfer.py` (new) | microbench only | **0.921× vs QWEN-013** (1.079× vs default tiles) | Numerics pass only with a weight-layout fix (see record) | session `qwen-cutlass2` exec `7590adcb`/`659d2633` | **Rejected on measurement, but the ledger note was wrong.** The kernel *does* support BF16 gated SwiGLU; mcore's wiring is broken in two independent ways (an `ActivationType` mis-map that hard-fails, and a gate/up ordering mismatch that silently corrupts numerics — both root-caused below). After fixing both in the harness, CUDA-graph device time is **90.81 µs vs 83.63 µs** for the retuned Triton path — 8% slower — so fixing the backend would not win. Consistent with QWEN-012: everything is pinned near the weight-bandwidth floor. | -| QWEN-015 | 2026-07-25 | `flashinfer.fused_moe.trtllm_bf16_routed_moe` — the actual TRT-LLM-Gen kernel vLLM wins with — can be dropped in as a new grouped-GEMM backend | none (blocked before any Megatron change); probe in `dev/moe_fused/harness_flashinfer.py` | not reached | n/a | n/a | session `qwen-cutlass2` exec `7590adcb` | **Blocked on weight layout** — not a contract or EP problem. `use_shuffled_weight=False, weight_layout=MajorK` is rejected outright (`BF16 Moe: weight_layout must be BlockMajorK`), and `BlockMajorK` reads `size(3)` of the weights (`IndexError: Index 3 out of bounds for tensor with 3 dimensions`): it requires **4-D pre-shuffled block-major** weights, not mcore's 3-D `[E, 2*ffn, H]`. Path forward and why it is now low priority are in the record. | -| PROFILE-TUNED | 2026-07-25 | Confirm the decode timing composition after QWEN-013 and re-rank the remaining levers | none (profiling only) | n/a | n/a | n/a | session `qwen-cutlass2` `prof/tuned-1785013555` (`mcore_profile.nsys-rep`, `.sqlite`) | Workflow C on one steady-state decode step (BS256/OSL128, device 3, 302-step window): forward pass **9.933 ms**, GPU-busy 7.884 ms, **idle 2.049 ms (20.6%)**, 1362 kernels/step. Share of GPU time: **MoE expert GEMM 2611 µs (32.9%, 96 kernels, 27.2 µs/kern)** — down from ~40% pre-retune — routing/permute 1251 µs (15.8%, 242 kernels), dense GEMM 1092 µs (13.8%), **exposed EP comm 929 µs (11.7%, 96 kernels, 9.7 µs/kern)**, attention 913 µs (11.5%), elementwise 528 µs (6.7%), norm 462 µs (5.8%). **Corroborates QWEN-012/013 independently**: 2611 µs / 48 layers = 54.4 µs per layer for FC1+FC2, matching the ~57 µs microbench, and leaving only 2611−48×49.66 = **229 µs/step (2.3%)** above the weight-bandwidth floor. The GEMM lever is spent; comm + routing + idle are what remain. | -| QWEN-011 | 2026-07-24 | Fuse the MoE indirection-table build (lever #1): merge `_init_sorted_ids` + `_prefix_sum` + `_fill_expert_block_ids` into one `_prefix_fill_init_kernel`, cutting the 5-kernel routing storm to 3 kernels/layer (−96 launches/step) | `vllm_fused_moe.py` (`_prefix_fill_init_kernel`, `_moe_align_block_size_fused`, env `MCORE_MOE_FUSED_ALIGN`), `dev/moe_fused/harness_align.py` (new) | **22,696.4 tok/s** | **+1.92% vs QWEN-002-CONFIRM** (+2.04% vs clean) | **EXACT** (max_abs_diff 0.0, allclose) + coherent | session `qwen-fuse` `fusedalign`, branch `perf/moe-fused-align` | **Accepted** — bit-exact, MoE microbench 1.029×, e2e +1.9% with tight variance (22.65–22.71k). Reduces routing/permute kernel count (lever #1). Now 66.8% of vLLM. | -| QWEN-016 | 2026-07-25 | **Decision gate**: is the 929 µs/step of exposed NVLS EP comm a worthwhile lever — how much is genuinely exposed, is it latency- or bandwidth-bound, and what is the floor under perfect overlap? | none (analysis); `dev/moe_fused/probe_comm.py`, `analyze_comm.py`, `analyze_comm_skew.py`, `harness_comm.py` (new) | n/a (analysis) | n/a | Microbench gate bit-exact (AGV max abs 0.0, RSV max rel 0.0 over 64 trials) | session `qwen-comm` job 5601961, execs `08e6a692`/`04bdfdd6` | **Redirect off this lever.** 100% exposed (interval union = sum), but **latency-bound**: AGV 6.57 µs = 0.72 launch + **5.08 barrier** + 0.77 transfer; RSV 7.88 = 0.72 + **5.03** + 2.13. Bytes are only 127 µs/step of ~693 µs and RSV transfer is already 82% of the NVLink floor, so batching bytes cannot win. Of 1060 µs/step, 632 µs is intrinsic and 428 µs is inter-rank skew (ranks waiting on the slowest rank's expert GEMM — a routing-balance problem). Recoverable critical path is **6.4–7.0% of the step**; RSV-into-FC2 fusion keeps the barrier (≈1.4%) and 2-chunk pipelining adds a barrier per chunk (≈4.3%, and needs concurrent streams under graph capture — the QWEN-008 blocker). CTA count already optimal at the shipped 128. | -| QWEN-018 | 2026-07-25 | **Decision gate**: break the 1251 µs/step routing/permute category down by kernel name — launches, device µs, dispatch gap, and fixed-vs-work split — and compute a wall-time ceiling per candidate fusion before building anything | none (analysis); `dev/moe_fused/analyze_routing.py` (new) | n/a (analysis) | n/a | n/a | session `qwen-comm` job 5601961 (login-node re-analysis of `qwen-cutlass2:prof/tuned-1785013555`) | **Two of the six routing kernels are pathological, and neither is launch-bound.** Per steady-state decode step (40-step mean, device 3): `_moe_sum` 7.79 µs × 48, **`_count_local_tokens` 7.52 µs × 48**, `gatherTopK` 6.05, `_scatter_token_indices` 2.66, router softmax 1.94, `_prefix_fill_init` 1.33, plus a 0.74 µs `torch.zeros` fill = 1531 µs/step of wall including gaps. Inter-kernel dispatch gap is **0.55 µs**, so a removed launch is worth `duration + 0.55`, and the fixed floor per launch is only 0.72+0.55 = **1.27 µs**. Ceilings: **(1) `_moe_sum`→FC2 epilogue = 400 µs = 4.03% gross**, but needs cross-CTA fp32 atomics into `out` plus a zeroing launch — QWEN-001's exact failure shape — so ~3.3% net and possibly negative; **(2) cooperative-grid merge of count→prefix_fill→scatter = 122 µs = 1.23%** before paying two grid syncs ⇒ **gated out**; **(3) new: fold the count + its zeros-fill into `_prefix_fill_init_kernel` = 449 µs = 4.52% gross**, no grid sync needed and integer-exact ⇒ **chosen**. `_count_local_tokens` is slow because at `BLOCK_SIZE=1024` only **2 of its 152 CTAs get work**; QWEN-003 changed the reduction but not that, which is why it measured a wash | -| CLEANBASE-S6 | 2026-07-25 | Fresh same-session OSL1024 reference at the session-4/5 best config, in session `qwen-comm`, before the routing campaign | none (`MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`) | **23,264.4 tok/s** | baseline (session 6); −1.61% vs QWEN-013b's 23,646.0 (session drift) | Benchmark pass, 5/5 iters | session `qwen-comm` job 5601961, run `e2e/ref-s6-1785024868` | Clean reference; 68.44% of vLLM 33,994.5. avg_latency 10,975 ms, TPOT 11.004 ms/tok. Synced snapshot `workspaces/megatron_lm/6b9355b187072223` verified to **not** contain the session-6 kernel | -| QWEN-019 | 2026-07-25 | Fold the per-expert token count and the `torch.zeros` fill of its counter buffer into `_prefix_fill_init_kernel`, taking the decode indirection-table build from 4 launches to 2. No grid sync is needed because that kernel already has every CTA redundantly recompute the cumsum in registers | `vllm_fused_moe.py` (`_count_prefix_fill_init_kernel`, `_moe_align_block_size_count_fused`, env `MCORE_MOE_FUSED_COUNT`, `_FUSED_COUNT_MAX_TOKENS`), `dev/moe_fused/harness_countfuse.py` (new) | **23,964.4 tok/s** | **+3.01% vs CLEANBASE-S6** | **Bit-exact** (max_abs 0.0 *and* max_rel 0.0 at 128/256/384/512) + tables integer-identical + coherent | session `qwen-comm` job 5601961, run `e2e/countfuse-1785025539` | **Accepted** — align call 16.38→10.24 µs under CUDA-graph replay (1.600×), whole MoE call 86.02→79.96 µs (1.076×), e2e +3.01% with tight variance (23.57–24.30k). TPOT 11.004→10.682 ms/tok. Default OFF; enable with `MCORE_MOE_FUSED_COUNT=1` (requires `MCORE_MOE_FUSED_ALIGN=1`). New best = **70.49%** of vLLM 33,994.5 | -| QWEN-017 | 2026-07-25 | `symm_mem_sync` spins on a system-scope `atom.cas` (a full uncached RMW per attempt); an `ld.acquire.sys` poll + clearing store should cut the 5.05 µs barrier and thus ~485 µs/step | `communication/torch_symm_triton/barrier.py` (`_wait_signal_ldpoll`, env `MCORE_SYMM_BARRIER_LDPOLL`) — **reverted**; `dev/moe_fused/harness_comm.py` | microbench only | **−2.3% (regression)** | Bit-exact both variants (AGV max abs 0.0, RSV max rel 0.0 over 64 trials) | session `qwen-comm` job 5601961, exec `04bdfdd6` | **Rejected** — per-step comm 693.7 µs baseline vs 709.8 µs ldpoll (3 alternating reps each, no overlap). The barrier-only kernel is unchanged at 5.75–5.83 µs vs a 0.72 µs empty kernel, so the 5.05 µs is the 4-way system-scope flag round trip itself, not polling granularity — a hardware/driver floor a Triton rewrite does not move. Patch reverted; `megatron/` back to the session-4 state. | -| QWEN-020 | 2026-07-25 | `_moe_sum_kernel` spends 7.79 µs/layer moving ~6.3 MB — 7× its own bandwidth floor — because the per-topk-slot locality test is a uniform scalar branch on a dependent `routing_map` load, serialising the topk walk. Predicate it into the load mask and widen the K tile to the full hidden size | `vllm_fused_moe.py` (`_moe_sum_kernel_fast`, env `MCORE_MOE_SUM_FAST`, `_FAST_MOE_SUM_MAX_BLOCK_K`) | **24,403.7 tok/s** | **+1.83% vs QWEN-019**, **+4.90% vs CLEANBASE-S6** | **Bit-exact** (max_abs 0.0 *and* max_rel 0.0 at 128/256/384/512) + coherent | session `qwen-comm` job 5601961, run `e2e/moesum-1785026515` | **Accepted** — `_moe_sum` 8.14 → 5.97 µs/layer (1.36×), whole MoE call 80.69 → 78.77 µs at 256 tokens. TPOT 10.682 → 10.490 ms/tok. Same reduction order and same fp32 arithmetic (masked slots add an exact 0.0), so bit-exactness is structural, not luck. Default OFF; enable with `MCORE_MOE_SUM_FAST=1`. New best = **71.79%** of vLLM 33,994.5 | -| QWEN-021 | 2026-07-25 | The router's `torch.softmax` + `torch.topk(sorted=False)` pair costs 436.8 µs/step (4.4%) for a `[256, 128]` fp32 reduction; one CTA per token can do softmax in registers and select the top-8 by max-then-mask in a single kernel | `megatron/core/inference/moe/router_topk.py` (new), `megatron/core/transformer/moe/router.py` (`InferenceTopKRouter._forward` fused-path branch), env `MCORE_ROUTER_FUSED_TOPK`; `dev/moe_fused/harness_routertopk.py` (new) | **25,352.9 tok/s** | **+3.89% vs QWEN-020**, **+8.98% vs CLEANBASE-S6** | **Bit-exact probs and identical expert sets** at 128/256/384/512 × 2 seeds | session `qwen-comm` job 5601961, run `e2e/routertopk-1785027029` | **Accepted, largest single win of the session.** 4 kernels → 1: 16.41 → 4.10 µs under graph replay at 256 tokens (4.00×), eager per-kernel 15.54 → 2.04 µs. TPOT 10.490 → 10.097 ms/tok (−393 µs/step vs a predicted ~312 µs). Default OFF; `MCORE_ROUTER_FUSED_TOPK=1`. New best = **74.58%** of vLLM 33,994.5 | -| QWEN-022 | 2026-07-25 | Now that every CTA streams the routing pairs to build its own histogram (QWEN-019), the scatter can be a second streaming pass in the same kernel — no atomics, no grid sync — taking the decode indirection-table build from 2 launches to 1 | `vllm_fused_moe.py` (`_align_single_kernel`, `_moe_align_block_size_single`, env `MCORE_MOE_FUSED_SCATTER`), `dev/moe_fused/harness_scatterfuse.py` (new) | **25,495.9 / 25,441.4 tok/s** (two runs) | **+0.46% vs QWEN-021** (mean of two), inside the noise band | **Bit-exact** whole-MoE output at 128/256/384/512 + table equality (16 cases: `npp`, `expert_ids`, row multiset, and every row placed under an expert id that owns it) | session `qwen-comm` job 5601961, runs `e2e/scatterfuse-1785027598`, `e2e/scatterfuse-rep2-1785027955` | **Accepted, marginally.** Align call 10.08 → 8.23 µs and whole MoE call 82.17 → 79.56 µs at 256 tokens (1.033×), predicting 1.24%; e2e delivered **+0.35% and +0.55%** in two runs. Both runs beat QWEN-021, so the sign is reliable, but the size is not — the honest attribution is ~0.5%. The candidate is a **wash at 384 tokens (1.007×) and a slight loss at 512 (0.997×)**, since the second streaming pass scales with pairs while the removed launch does not. Default OFF; `MCORE_MOE_FUSED_SCATTER=1` (requires `MCORE_MOE_FUSED_COUNT=1`). New best = **75.00%** of vLLM 33,994.5 | -| PROFILE-S6 | 2026-07-25 | Re-profile at the session-6 configuration (all six gates on) and re-rank, since the step lost ~10% of its wall time and ~820 µs of routing | none (`profile_insession.sh`) | n/a (BS256/OSL128 profile run) | n/a | n/a | session `qwen-comm` job 5601961, `prof/s6-all-1785028240/mcore_profile.{nsys-rep,sqlite}` | **Routing is done as a lever.** One decode step: wall **9.097 ms** (was 9.933), GPU-busy 7.183 (was 7.884), idle 1.914 (21.0%), **1170 kernels** (was 1362). Routing/permute is now **430 µs in 98 kernels** (was 1251 µs in 242). New ranking: MoE expert GEMM 2469 µs (27.1%, closed), **host-side idle 1914 µs (21.0%)**, dense GEMM 1084, comm 1063 (closed), attention 924, elementwise 488, norm 450, routing 430. The idle is concentrated in **three host gaps totalling ~1237 µs/step**: 548 µs after `index_elementwise_kernel`, 386 µs after `vectorized_elementwise_kernel`, 303 µs after `CatArrayBatchedCopy_vectorized` — sampling/detokenize/scheduling between graph replays | -| CLEANBASE-S7 | 2026-07-26 | Fresh same-session OSL1024 reference at the seven-gate session-6 best config, in session `qwen-attnstate`, before the host-side `initialize_attention_state` lever | none (all seven gates on) | **25,805.9 tok/s** | baseline (session 7); +1.22% vs QWEN-022's 25,495.9 (session drift) | Coherent, 5/5 iters | session `qwen-attnstate` job 5613090, run `e2e/ref-s7-1785084185` | Clean reference; **75.91%** of vLLM 33,994.5. Per-iter 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 (spread 0.23%, the tightest reference this campaign has had). avg_latency 9,867 ms, TPOT 9.920 ms/tok | -| QWEN-023 | 2026-07-26 | HOSTGAP-S6's ~311 µs/step of `initialize_attention_state` is recomputation, not computation: at fixed BS256 the request set, sampling metadata, KV block table and CUDA-graph selection are identical for 128 consecutive decode steps, and only the KV sequence lengths advance. Recompute those and reuse the rest, invalidated by a request-layout version counter | `megatron/core/inference/contexts/dynamic_context.py` (`_bump_request_layout_version`, `_incr_attn_state_cache_key/_store/_snapshot/_verify`, `_incremental_attention_state_update`, env `MCORE_INFER_INCR_ATTN_STATE`, `MCORE_INFER_INCR_ATTN_STATE_VERIFY`, profiling gate `MCORE_INFER_ATTN_PROF`), `megatron/core/inference/contexts/attention_context/mha_metadata.py` (`restore_state_data`), `dev/moe_fused/harness_attnstate.py` (new) | **26,032.1 / 26,092.5 tok/s** (two OFF/ON pairs) | **+0.94%** (pairwise +0.88% and +1.01%) | **130 consecutive decode steps bit-identical** (cached vs freshly recomputed, every host and GPU bookkeeping buffer) + temperature-0 coherence output byte-identical to gate-OFF | session `qwen-attnstate` job 5613090, runs `e2e/incr-attn-on-1785085393`, `e2e/incr-attn-on-rep2-*` | **Accepted.** Host CPU per call **174.1 → 51.7 µs (3.37×)**; the residual is the H2D bookkeeping copy, which is irreducible. TPOT 9.915 → 9.823 ms/tok (−92 µs/step). Across all twenty timed iterations the two gates are **fully separated** (slowest ON 25,992.1 > fastest OFF 25,893.9), so the sign is certain even though the size is under 1%. Delivered ~28% of the 3.4% ceiling: the removable phases are ~72% of the call, and the H2D bookkeeping copy that remains cannot be cached. Default OFF; `MCORE_INFER_INCR_ATTN_STATE=1`. New best = **76.67%** of vLLM 33,994.5 | -| CGTRACE-CONTROL | 2026-07-26 | HOSTGAP-S6 could not tell whether the ~900 µs/step of CUDA-graph machinery (a 199 µs `cudaGraphLaunch` plus GPU-side inter-node dispatch) is real or an artifact of `--cuda-graph-trace=node` instrumenting all 1158 graph nodes. Capture the same workload both ways and compare | none (control); `dev/moe_fused/profile_insession.sh` (env `CUDA_GRAPH_TRACE`, plus the hardened nsys stop/qdstrm recovery ported from `profile_host_insession.sh`), `dev/moe_fused/analyze_cgtrace.py` (new) | n/a (two BS256/OSL128 profile runs) | n/a | n/a | session `qwen-attnstate` job 5613090, `prof/cgt-node-1785086713` and `prof/cgt-graph-1785086964` | **The machinery is real; node-mode instrumentation is not inflating the step.** Same rank, same steady-state decode window: step period **8882.5 µs (node) vs 8945.1 µs (graph)** — 0.71% apart and in the *wrong* direction for an instrumentation artifact — and host `cudaGraphLaunch` **190.0 vs 184.7 µs** (2.8%). Node mode traces 1169 kernels/step, matching PROFILE-S6's 1170. **HOSTGAP-S6's wall-time attributions therefore all stand**, and lever 4 (reduce graph node count, ~0.78 µs/node) is live rather than chasing a measurement artifact. Caveat: `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` is too sparsely populated (417 rows total, n=6 in-window on the analyzed device) to independently confirm the *GPU-side* inter-node dispatch component; only the step period and host submit cost are settled | -| HOSTGAP-S6 | 2026-07-26 | Attribute PROFILE-S6's three inter-kernel host gaps to concrete host functions using the recovered host-visibility trace | none (analysis); `dev/moe_fused/analyze_hostgaps.py` (new) | n/a (analysis) | n/a | n/a | session `qwen-host` job 5607600, `prof/hosts6-1785048883` | **All three gaps attributed; the largest is a true serial data dependency, not a scheduling artifact.** Per step (GPU-only trace, mean of 4 ranks, 9.134 ms step / 1965.7 µs idle): **G1 `index_elementwise`→`vectorized_elementwise` 571.4 µs** = `transfer_samples_to_cpu`+`active_request_mask`+`update_requests`+engine window+`initialize_attention_state`, **CPU-bound Python** (100% of samples Running in `_PyEval_EvalFrameDefault`, only 8.6% of the gap in CUDA API, 0.5% in `cudaStreamSynchronize`); **G2 `vectorized_elementwise`→`vectorized_gather` 331.0 µs in 2 instances** (267.9 µs inside `initialize_attention_state`+`transfer_bookkeeping_to_gpu`+`forward_pass` head, **CPU-bound Python**; 63.1 µs inside `sampling`, launch+`cudaStreamSynchronize`); **G3 `CatArrayBatchedCopy`→`rmsnorm_fwd` 305.4 µs** = **one `cudaGraphLaunch` costing 199.1 µs median** for a 1158-node graph (0.172 µs/node), driver-side, host Running in `cuGraphLaunch`. No rank asymmetry (≤4%). Idle splits **1065 µs real host Python (11.7% of step)** vs **900 µs graph machinery (9.9%)** that is `--cuda-graph-trace=node`-artifact-suspect. Async overlap cannot fix G1: the chain is data-dependent on the current step's sampled tokens, which explains QWEN-007's +0.85% | -| QWEN-025 | 2026-07-26 | QWEN-024's decision-gate measurement showed 84% of `post_process_requests` is reducible by collapsing the per-request loop body — unvectorizable is not the same as irreducible. Add a decode fast path that is structurally unable to touch the request termination state machine | `dynamic_engine.py` (`_post_process_requests_decode_fast`, `_ppr_cache_epoch` invalidation at all three record-mutation sites, env `MCORE_INFER_FAST_POST_PROCESS`, `..._VERIFY`), `dev/moe_fused/harness_updreq.py` (`run_ppr_equivalence`, in-process gate A/B) | **26,361.7 tok/s** (three-pair ON mean); best run 26,430.7 | **+0.98%** vs same-session gate-OFF (pairwise +0.99% / +1.26% / +0.68%; +0.59% on a best-4-of-5 trim) | 400-step two-engine equivalence with 10 finisher steps and 200 steps past the token limit; 300 steps under VERIFY; coherence byte-identical | session `qwen-updreq` job 5616264, runs `e2e/ppr-{off,on}{,-r2,-r3}-*`, profiles `prof/g2-off-1785106875` and `prof/g2-on-1785107105` | **Accepted.** `post_process_requests` **211.0 → 27.2 µs/step (7.77×)**; whole post-sampling host chain **372.4 → 148.0 µs/step** across QWEN-024+025. The fast path **declines any step on which a request finishes**, so the termination state machine is unreachable from it — that is the safety argument, not extra testing. cProfile 261 calls/step vs the pre-registered <500 threshold. 13 of 15 ON iterations beat every one of the 15 OFF iterations. G1 −56.1 µs, idle −77.9 µs, kernels/step unchanged at 1169. **Only ~1/3–1/2 of the 183.9 µs host saving converts** (vs ~1:1 for QWEN-024) because `async_bookkeep` was already partly overlapped with GPU work. Default OFF. New best = **77.75%** of vLLM 33,994.5 | -| PROFILE-S10 | 2026-07-27 | Re-profile at the twelve-gate config (ten gates + QWEN-026 + QWEN-027) and, for the first time this campaign, take a **clean 2 s decode window** (densest-window search, not a span fraction) to get an honest category breakdown and decide whether more small-kernel fusion is worth it | none (`profile_insession.sh` BS256/OSL256; `prof/s10-addnorm-1785188237/window_breakdown.py`, new) | n/a | n/a | n/a | session `qwen-nodecut`, `prof/s10-addnorm-1785188237/mcore_profile.{nsys-rep,sqlite}` | **Both fusions confirmed live and each removed ~48 kernels/step**: `rmsnorm_fwd_tuned` 104→56/step, `triton_poi_fused_add_copy` 103→55/step; `_fused_qk_rmsnorm` and `_fused_add_rmsnorm` each 55/step (1/layer). Clean-window (union-busy **81.6%**, idle **18.4%** = 360 µs-scale host gaps) device-time by category: **moe_expert_gemm 587 ms**, **comm/NVLS 247 ms**, **dense_gemm 236 ms**, **attention(flash) 211 ms**, routing 69, norm 67, elementwise 66, moe_sum 65, rope_kv 15. **Verdict: small-kernel fusion is now in diminishing returns.** The fusion-addressable buckets (norm/elementwise/routing) are each ~65-70 ms and partly overlapped on other streams, so each remaining fusion is ~0.5-1%; the gap to vLLM is now dominated by expert GEMM (compute-bound), NVLS comm, attention, and the 18.4% host idle (blocked on the nsys host-capture failure). Next levers should target those, not more small-kernel merges | -| HOSTGAP-S10 | 2026-07-27 | With small-kernel fusion in diminishing returns (PROFILE-S10), attack the biggest remaining prize: the **18.4% GPU-idle**. Attribute it to host call sites **without** the deadlock-prone host-visibility capture, by using the CUDA-API host rows already in the clean `cuda,nvtx` trace | `idle_attrib2.py`, `bracket.py` (new, in `prof/s10-addnorm-*/`); env-gated `MCORE_INFER_NVTX` toggle added to `dynamic_engine.py`; `CUDA_GRAPH_TRACE` + `TRITON_LIBCUDA_PATH` + `HOST_LEVEL`/`OSRT`/`SAMPLE` knobs added to `profile_host_insession.sh` | n/a | n/a | n/a | traces on lustre `sessions/qwen-nodecut/prof/s10-addnorm-1785188237/`; session `qwen-host` (job 5638180) | **Idle decomposition (clean 1 s decode window, 178 ms idle):** small gaps <10 µs = 59.6 ms (33%, intra-graph kernel scheduling, unavoidable); **large gaps ≥10 µs = 118.8 ms (67%) = host chain between steps** (median 37 µs). Large-gap attribution via the decode thread's CUDA-API rows: **UNCOVERED (Python/CPU compute, no CUDA call) = 87.8 ms = 73.9%**, memcpy(token D2H) 8.9%, graph_launch 7.5%, launch 6.9%, sync_wait 1.6%. So ~74% of attackable idle (~8-9% of wall) is Python on the critical path. **Localized** by bracketing each Python gap with its surrounding CUDA APIs: 47.6% sits in `cudaMemcpyAsync -> cuKernelGetName`, +15.7% `cudaLaunchKernel -> cudaMemcpyAsync`, +11% `cuKernelGetName -> cudaMemcpyAsync` — i.e. the **between-steps host orchestration** after the sampled-token D2H copy (post-sampling bookkeeping / scheduling / attention-state prep), not the GPU-sync wait. **nsys dead-end recorded:** the host-visibility capture deadlocks nsys finalization and `QdstrmImporter` rejects the qdstrm whenever **osrt tracing**, **`--sample=process-tree`**, OR **`MCORE_INFER_NVTX` (NVTX ranges under CUDA graphs)** is on — independent of `--cuda-graph-trace` level. Only the plain `cuda,nvtx --sample=none` set finalizes. **Actionable conclusion:** future idle work should (a) use clean-trace RUNTIME attribution (this method), not host capture; (b) target the between-step Python orchestration via `perf_counter` phase timing (no nsys), since NVTX-based naming is blocked | -| QWEN-027 | 2026-07-27 | PROFILE-S9 showed the two dominant small kernels are `triton_poi_fused_add_copy` (bias-dropout-add residual, 103/step) and `rmsnorm_fwd_tuned` (104/step), adjacent on the serial path. Fuse the residual-add with the *standalone* `pre_mlp_layernorm` (MoE) into one kernel (vLLM's `add+rms_norm`), returning both the updated residual and the normed MLP input — the input norm is inside the TE `LayerNormLinear` QKV GEMM and not interceptable, so only the pre-MLP boundary (48/step) is fusible | `megatron/core/inference/fused_add_rmsnorm.py` (new: `_fused_add_rmsnorm_kernel`, `fused_add_rmsnorm`, `can_use_fused_add_rmsnorm`), `megatron/core/transformer/transformer_layer.py` (guarded fused path at the `self_attn_bda` site + `_forward_pre_mlp_layernorm` consumes the stashed normed output), env `MCORE_FUSED_ADD_NORM`, `MCORE_FUSED_ADD_NORM_MAX_TOKENS` (default 256); `dev/moe_fused/harness_addnorm.py` (new) | **27,525.2 / 27,532.1 tok/s** (two runs, mean 27,528.6) | **+1.37%** vs QWEN-026-on 27,151.9; arms **fully separated** (min ON 27,503.6 > max ref 27,206.4) | **residual BIT-EXACT** (fp32-accumulate add reproduces torch's bf16 add), **norm bf16-ulp** (max_rel ≤7.9e-3 ≈ one bf16 ulp). Coherence: **all 3** temperature-0 prompts identical to the QWEN-026-on arm (incl. the France prompt that QWEN-026 flipped), fluent + factually correct | session `qwen-nodecut`, runs `e2e/addnorm-on-1785184454`, `e2e/addnorm-on-2-1785185287`; microbench, profile `prof/s9-qknorm-1785182903` | **Accepted.** Microbench CUDA-graph replay at 256 tokens: (add + TE norm) 6.16 → fused 4.11 µs (**1.50×**), and **2.0× at 384/512** — unlike QWEN-026 this kernel does not regress at higher token counts (flat ~4.1 µs). Removes 48 add launches + 48 graph nodes/step. Heavily gated: decode-only, MoE standalone RMSNorm, no bias/dropout/recompute/offload, no cross-attn, with a reference fallback so training and all non-matching cases are byte-identical. TPOT 9.428 → 9.299 ms/tok. Default OFF; `MCORE_FUSED_ADD_NORM=1`. New best = **80.98%** of vLLM 33,994.5 | -| PROFILE-S9 | 2026-07-27 | Re-profile at the eleven-gate config (ten gates + QWEN-026) and re-rank the remaining small-kernel serial chain | none (`profile_insession.sh`, BS256/OSL256) | n/a | n/a | n/a | session `qwen-nodecut`, `prof/s9-qknorm-1785182903/mcore_profile.{nsys-rep,sqlite}` | **QWEN-026 confirmed live** (`_fused_qk_rmsnorm_kernel` 51.7/step = 1/layer; the old two `rmsnorm_fwd_general` are gone). The two dominant small kernels are now **`rmsnorm_fwd_tuned` 104/step @ 2.93 µs** (input + pre-MLP norm, 2/layer) and **`triton_poi_fused_add_copy` 103/step @ 1.51 µs** (the bias-dropout-add residual, 2/layer), and they are adjacent: the top consecutive pair by count is `triton_poi_fused_add_copy → rmsnorm_fwd_tuned` at 103/step. vLLM runs these as one `triton_red_fused__to_copy_add_..._rms_norm`; mcore runs them as two. **Next lever = fused add+RMSNorm** (~96 add-launches + 96 graph nodes/step removed, 4.44 → ~2.5 µs/boundary), the same convert-at-3× serial-chain + node-removal pattern QWEN-026 proved. Ranked #1 | -| QWEN-026 | 2026-07-27 | Qwen attention applies `q_layernorm` and `k_layernorm` as two separate per-head RMSNorm module calls → two `rmsnorm_fwd_general` launches (and two graph nodes) per layer, each normalizing a tiny 128-wide row (launch/latency bound, not bandwidth). Fuse into one Triton kernel: one CTA per row over the concatenated `[q_rows; k_rows]` space, selecting the q or k gamma per row; decode-gated to ≤256 tokens since the 1-row-per-CTA design loses to TE above that | `megatron/core/inference/attention/fused_qk_norm.py` (new: `_fused_qk_rmsnorm_kernel`, `fused_qk_rmsnorm`, `can_use_fused_qk_norm`), `megatron/core/inference/attention/__init__.py` (new), `megatron/core/transformer/attention.py` (guarded fused path in `get_query_key_value_tensors`), env `MCORE_FUSED_QK_NORM`, `MCORE_FUSED_QK_NORM_MAX_TOKENS` (default 256); `dev/moe_fused/harness_qknorm.py` (new) | **27,143.5 / 27,160.3 tok/s** (two runs, mean 27,151.9) | **+2.9%** vs back-to-back gate-OFF 26,363.0 (+2.83% vs CLEANBASE-S9 26,404.4); arms **fully separated** (min ON 27,102.9 > max OFF 26,425.7) | **bf16-ulp, not bit-exact** (max_rel ≤7.6e-3 ≈ one bf16 ulp; TE's internal rsqrt/reduction differs). Microbench MATCH at 128/256/384/512 × 2 seeds. Coherence: **2 of 3** temperature-0 prompts byte-identical to gate-OFF; the France prompt diverges at a low-confidence branch (" Paris, and…" vs " Paris. The…"), both fluent and factually correct | session `qwen-nodecut`, runs `e2e/qknorm-on-1785180416`, `e2e/qknorm-on-2-1785181060`, OFF `e2e/qknorm-off-bb-1785181867` | **Accepted, largest win since QWEN-021.** Microbench CUDA-graph replay at 256 tokens (BS256 decode): 10.25 → 8.20 µs (1.25×), but 0.83×/0.86× at 384/512 → gated decode-only. e2e delivered **+2.9%**, ~3× the ~1% microbench ceiling, because the two norms sit in the serial attention dependency chain and removing a launch + graph node per layer (×48) also cuts host dispatch gaps. TPOT 9.711 → 9.428 ms/tok. Default OFF; `MCORE_FUSED_QK_NORM=1`. Not bit-exact, so acceptance rests on coherence staying fluent+correct, not byte-identity. New best = **79.87%** of vLLM 33,994.5 | -| CLEANBASE-S9 | 2026-07-27 | Fresh same-session OSL1024 reference at the ten-gate best config, in session `qwen-nodecut`, before the fused QK-RMSNorm lever; re-measured back-to-back with the QWEN-026 ON arm to rule out session drift | none (all ten gates on) | **26,404.4 tok/s** (initial); **26,363.0 tok/s** back-to-back (`qknorm-off-bb`) | baseline (session 9) | Coherent, 5/5 iters | session `qwen-nodecut`, run `e2e/qknorm-off-bb-1785181867` | Clean reference; **77.67%** of vLLM 33,994.5. TPOT 9.695 / 9.711 ms/tok. The two OFF measurements agree to 0.16%, so the +2.9% QWEN-026 delta is not drift | -| QWEN-024 | 2026-07-26 | HOSTGAP-S6 lever 2. Measure the three parts of the post-sampling chain separately, name the unlabelled window, and attack only the part that is not per-request Python object churn: `update_requests` and `active_request_mask` are already vectorized, and roughly half their whole-tensor op cost is provably dead work at one generated token per request | `dynamic_context.py` (`_write_decode_token_bookkeeping_fast`, `_write_token_bookkeeping_reference`, `_verify_decode_token_bookkeeping`, env `MCORE_INFER_VEC_UPDATE_REQS`, `..._VERIFY`), `text_generation_controller.py` (no-finisher short-circuit), `dev/moe_fused/harness_updreq.py` (new) | **26,128.2 / 26,270.3 tok/s** (two pairs) | **+0.90% vs CLEANBASE-S8** (pairwise +0.71% and +1.10%) | 400-step two-context equivalence with 30 mid-batch terminations + block-boundary crossing; 300 steps under VERIFY; coherence byte-identical | session `qwen-updreq` job 5616264, runs `e2e/vecupd-on-1785100576`, `e2e/vecupd-on-rep2-1785101960` | **Accepted.** Named the unlabelled engine window: it is `post_process_requests` via `async_bookkeep`, invisible to HOSTGAP-S6 only because the engine's NVTX helper is inert unless `_nvtx_enabled` is set. It is **187.9 µs/step of pure per-request Python object churn** (1,719 `len()`, 256 dict lookups, 256 appends, 3 tensor ops per step) and was **deliberately left alone** per the falsification rule. The other half was op-eliminated: host chain 372.4 → 293.7 µs/step, `update_requests` 137.6 → 74.0 µs (1.86×). TPOT −88 µs/step. Default OFF; `MCORE_INFER_VEC_UPDATE_REQS=1`. New best = **77.28%** of vLLM 33,994.5 | -| CLEANBASE-S8 | 2026-07-26 | Fresh same-session OSL1024 reference at the eight-gate best config, in session `qwen-updreq`, before the `update_requests` host lever | none (all eight gates on) | **25,944.9 tok/s** | baseline (session 8); −0.57% vs QWEN-023's 26,092.5 (session drift) | Coherent, 5/5 iters | session `qwen-updreq` job 5616264, run `e2e/ref-s8-1785099726` | Clean reference; **76.32%** of vLLM 33,994.5. Per-iter 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7 — iteration 1 is a 1.3% cold outlier, iterations 2–5 span 0.23%. avg_latency 9,820.6 ms, TPOT 9.867 ms/tok | -| PROFILE-HOST-S6 | 2026-07-25 | Capture a host-visible trace (`osrt` + CPU sampling + Python sampling) so the ~1237 µs/step of host gaps PROFILE-S6 found between CUDA-graph replays can be attributed to concrete host call sites | none (`dev/moe_fused/profile_host_insession.sh`, new; `dev/moe_fused/dispatch_host_profile.sh`, new) | n/a (BS256/OSL128 profile run) | n/a | n/a | session `qwen-comm` job 5601961 **preempted mid-task**; replacement session `qwen-host` job 5607600 on `nvl72166-T15`, exec `d920241ea00a41618c63a75fefcb9ea2` | **Capture dispatched and running, not confirmed complete within the time box.** Job 5601961 was PREEMPTED (not expired) before any capture could run, so the whole capture had to be re-queued on a fresh allocation that only scheduled at the deadline. Artifact target `sessions/qwen-host/prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}`. **No gap attribution was performed.** Next session must first check whether the detached exec produced the artifact | -| REBASE-S11 | 2026-07-29 | Rebase the twelve-gate branch onto current `main` (141 commits ahead of the old base) and re-measure, since the PR was unreviewable on a stale base | rebase of 14 commits onto `upstream/main` `3ff70c006`; conflict resolutions in `dynamic_context.py`, `text_generation_controller.py`, `router.py`; `dev/moe_fused/run_rebased.sh` (new) | **27,277.4 tok/s** (tpot 9.385 ms, avg_latency 9,340.6 ms, p99 9,589.9 ms) | **−0.91% vs pre-rebase 27,528.6**; 80.24% of vLLM 33,994.5 (was 80.98%) | coherence passed (2+2→4, capital of France→Paris, 3+2 cows→5); all 12 gates present, 3 kernel files present, py_compile clean | job 5694153 on `nvl72159-T08` (5 timed iters, spread 27,214–27,314 tok/s = 0.36%) | **Rebase completed and statically verified; measurement blocked by a node-level Lustre fault, not by code.** Only 1 of 14 commits conflicted (`b9165d87b`), in 3 files, 1 hunk each. (1) `router.py`: upstream added a `qb_beta` selection path (0→26 occurrences); resolved by keeping it and gating the fused topk on `self.qb_beta is None`, since the fused kernel selects on raw logits and cannot honor a qb_beta-shifted selection. (2) `text_generation_controller.py`: purely additive (upstream dummy-forward helpers vs our cached-empty-tensor helpers); kept both. (3) `dynamic_context.py`: upstream made the bookkeeping H2D **conditional on a new `transfer_bookkeeping_to_gpu` param** and returns an `Optional[Event]`; resolved by adopting that contract while keeping the `_INCR_ATTN_STATE` store and profiling epilogue reachable before the return, making the QWEN-023 fast path **decline** whenever `transfer_bookkeeping_to_gpu=False` or `record_bookkeeping_done_event=True` (it publishes itself and records no event), and clearing upstream's new `_bookkeeping_no_real_work` inside the fast path so a stale `True` from an earlier capture step cannot publish `real_token_count=0`. **Note the PR's claimed blocker was wrong**: `vllm_fused_moe.py` is byte-identical between the old merge-base and current `main`, so the SwiGLU region needed no port and applied clean. **Infra blocker:** every run attempt died in `import megatron.core` on a different "missing" module (`sympy._trigonometric_special`, `networkx.preflowpush`, `jaraco.functools`, `nvidia_cutlass_dsl.static_persistent_tile_scheduler`, our own `transformer_layer`). All are the same fault: on `nvl72169-T17` those files **appear in directory listings but `os.path.isfile()` returns False**, while the login node reads all of them fine — a broken Lustre client that serves MDS metadata but cannot stat/open inodes. Do not chase these as package corruption or venv rot — job 5694153 on a healthy node imported everything cleanly with the **unmodified** venv and no overlay, which proves it. **Re-measured via `dev/moe_fused/rebased12.sbatch`** (submitted with plain `sbatch` because cog's repo sync hung three times from this client; runs the unmodified `run_e2e_cfg.sh` for comparability, excludes the bad node, and gates on an `os.path.isfile` health check that aborts in seconds rather than after the model load — that gate is the reusable lesson here). **Verdict: the rebase preserves the gains.** The −0.91% is small enough to be node-to-node variation (the pre-rebase number came from a different node), so 141 upstream commits cost at most ~1%. Log: `agents-space/auto_rebased/rebased12-5694153.out` | -| CLEANBASE-S12 | 2026-07-30 | Fresh same-node twelve-gate reference before testing upstream's new async scheduling, with no instrumentation loaded | none (all twelve gates on, `ASYNC_SCHED=legacy`) | **26,538.2 tok/s** (tpot 9.646 ms, avg_latency 9,597.3 ms, p99 9,892.8 ms) | baseline (session 12); **−2.71% vs REBASE-S11's 27,277.4 on a different node** | Coherent, 3/3 prompts | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/legacy-clean-1785379...` | Clean reference; 78.06% of vLLM 33,994.5. Per-iter 26,373.8 / 26,533.5 / 26,536.8 / 26,535.3 / 26,713.9 — spread 1.29%. **The −2.71% against REBASE-S11 is node-to-node, not a regression**, and is the reason every arm in this session is compared only against same-node runs: node variance here is larger than several accepted gates' individual effects | -| QWEN-029 | 2026-07-30 | HOSTGAP-S6 attributed the largest remaining lever to G1+G2B — 839 µs/step of resolve-before-prepare bookkeeping that is CPU-bound Python on the serial critical path — and concluded it had to be made cheaper rather than hidden, because `async-sched-mode=serial` (QWEN-005a/b, QWEN-007) refused EP and MoE outright, hung on a single request, and then delivered only +0.85%. **The 141 commits REBASE-S11 landed on replaced that implementation.** `serial` is gone; the mode is now `async` ("prepare-before-resolve"), the EP and MoE guards are gone (only MTP-depth and routing-replay remain), and it has a dedicated copy stream, CPU-ready events, and per-step overlap eligibility. Re-test it as a new feature, not a re-litigation | none — one flag: `--inference-dynamic-batching-async-sched-mode async` (`ASYNC_SCHED=async` in `run_e2e_cfg.sh`) | **27,515.0 tok/s** (tpot 9.304 ms, avg_latency 9,243.0 ms, p99 9,553.6 ms) | **+3.68% vs same-node CLEANBASE-S12 26,538.2**; arms **fully separated** (min ON 27,346.2 > max OFF 26,713.9) | Coherence **byte-identical to the legacy arm on all 3** temperature-0 prompts; no warnings, fallbacks or guard messages in the server log | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/async-on-1785379854` | **Accepted — largest single win of the campaign, and it is a flag, not a kernel.** Per-iter 27,346.2 / 27,369.3 / 27,451.1 / 27,569.2 / 27,845.4. **The ledger's "do not re-litigate async scheduling" guidance was correct about the old code and wrong about the new** — worth remembering as a general rule: a rebase that lands 141 upstream commits can invalidate a rejection, so re-check rejected levers whose blocking reason was an explicit upstream guard. **Caveat that makes this an underestimate:** REBASE-S11's conflict resolution makes the QWEN-023 incremental-attention fast path decline whenever the caller defers the bookkeeping publish, which is exactly what async does — so `MCORE_INFER_INCR_ATTN_STATE` is inert in this arm and the +3.68% is *net of losing that gate*. Making the fast path publish-deferral-aware is the obvious follow-up. New best = **80.94%** of vLLM 33,994.5 measured on this node; applying the same-node +3.68% to REBASE-S11 implies ~28,281 tok/s = **~83.2%** | -| QWEN-028 | 2026-07-30 | The other 48 block boundaries QWEN-027 could not reach. QWEN-027 fused `self_attn_bda` + the standalone `pre_mlp_layernorm`; the matching `mlp_bda` + next-layer *input* norm was ruled unreachable because that norm lives inside the TE `LayerNormLinear` QKV GEMM. That ruling was wrong for this config: at **TP=1** `InferenceLayerNormColumnParallelLinear.forward` runs the norm as its own `_te_rms_norm_kernel` launch immediately before the GEMM, so the preceding layer can compute it. The fused kernel already returns `(normed, new_residual)` — hand the normed half to the next layer's `linear_qkv` and skip its norm launch | `megatron/core/inference/fused_add_rmsnorm.py` (`can_use_fused_add_rmsnorm_qkv`, shared `_tensors_compatible`, env `MCORE_FUSED_ADD_NORM_QKV`), `megatron/core/tensor_parallel/inference_layers.py` (`prenormed_input` hand-off consumed in the `tp_size == 1` path), `megatron/core/transformer/transformer_layer.py` (guarded fused path at the `mlp_bda` site; `_next_layer_qkv` holder), `megatron/core/transformer/transformer_block.py` (`_wire_fused_add_norm_qkv_chain`) | **27,179.2 tok/s** (tpot 9.419 ms, avg_latency 9,361.3 ms, p99 9,661.1 ms) | **+2.42% vs same-node CLEANBASE-S12 26,538.2**; arms **fully separated** (min ON 27,027.8 > max OFF 26,713.9) | **residual bit-exact, norm bf16-ulp** (same kernel as QWEN-027). Coherence: **2 of 3** prompts byte-identical to the legacy arm; the France prompt diverges at a low-confidence branch (" So, what is the capital of Italy?" vs " So, the capital of Italy is Rome."), both fluent and factually correct — the same signature QWEN-026 showed | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/qwen028-on-1785380...` | **Accepted, and it over-delivered: +2.42% against a ~1.3% expectation** (47 boundaries × (add 1.5 + norm 2.9 − fused 4.1 µs) ≈ 14 µs/step of device time plus 47 × ~0.78 µs of node cost ≈ 50 µs of a 9,646 µs step ≈ 0.5%). Same pattern as QWEN-026/027: kernels on the serial dependency chain convert at ~3× their device-time arithmetic, because removing a launch also removes a graph node and a host dispatch gap. **The reusable lesson is about the earlier ruling, not the kernel:** "the norm is fused inside `LayerNormLinear`" was true of the *class name* and false of the *TP=1 code path*, and one look at the actual forward would have caught it two sessions earlier. 47 of 48 boundaries fuse — the last layer feeds `final_layernorm` and is left alone. Wiring stores the next layer's `linear_qkv` in a **plain list**, since assigning a Module to an attribute would register it as a second child and duplicate those parameters in `state_dict`. Heavily gated: TP=1, decode token count, RMSNorm, no bias, no offload, and `not layernorm_zero_centered_gamma` (the TE call it replaces hardcodes `zero_centered_gamma=False`). Default OFF; `MCORE_FUSED_ADD_NORM_QKV=1` | -| QWEN-030 | 2026-07-30 | Do QWEN-029 (async) and QWEN-028 (fused add+QKV-norm) compose, or do they cannibalize each other? They could overlap: async hides host bookkeeping behind GPU work, and part of QWEN-028's win is fewer host dispatches, which async would already be hiding | Both gates on together: `ASYNC_SCHED=async` + `MCORE_FUSED_ADD_NORM_QKV=1` + the 12 prior gates. No new code | **27,889.0 tok/s** (tpot 9.179 ms, iters 27,723 / 27,467 / 27,668 / 28,227 / 28,382) | **+5.09% vs same-node CLEANBASE-S12 26,538.2**; beats async alone (27,459 mean, +3.47%) by +1.56% and QWEN-028 alone (27,179, +2.42%) by +2.61% | same 2-of-3 byte-identical / France-branch divergence as QWEN-028 — async contributes no additional divergence (QWEN-029 alone was byte-identical) | session `qwen-gap1` job 5699934 on `nvl72067-T13`, run `e2e/async-qwen028-1785381339` | **Accepted — new best config.** They compose but **sub-additively**: +5.09% against +5.89% if the two were independent, so ~14% of the combined benefit is double-counted. That is the expected signature of two levers that partly attack the same resource (host dispatch), and it is the reason to always measure the *combination* rather than sum accepted deltas — the ledger's individual percentages cannot be added. Note the run shows a clear warm-up ramp (27,723 → 28,382 across 5 iters, +2.4%), steeper than the legacy arms, so the 5-iter mean **understates** the steady state under async | - -| QWEN-031 | 2026-07-30 | Async scheduling drives a dedicated copy stream for the bookkeeping H2D transfer, but the harness pins `CUDA_DEVICE_MAX_CONNECTIONS=1`, which collapses all streams onto one hardware work queue and can serialize that copy behind compute — defeating the overlap async exists to create | `CUDA_DEVICE_MAX_CONNECTIONS=8` with `ASYNC_SCHED=async` + the 12 gates. Env only, no code | 27,522.5 tok/s (tpot 9.301 ms) | **+0.23% vs the async mean 27,459** (two samples 27,515.0 / 27,403.5, which themselves differ by 0.41%) — the delta is **smaller than the noise between two runs of the identical config**, so this is a null result, not a small win | n/a (no numerical change) | session `qwen-gap1` job 5699934 on `nvl72067-T13` | **Rejected — no effect.** The serialization hypothesis was reasonable but wrong here: the bookkeeping H2D is small and evidently already fits alongside compute on one queue, so more queues buy nothing. Worth recording because `CUDA_DEVICE_MAX_CONNECTIONS=1` is a load-bearing setting in *training* configs and the instinct to raise it for overlap is strong — for this decode workload it is inert. **Methodological note: the only reason this is callable as a null rather than a +0.23% win is that two same-config samples existed to size the noise floor at ~0.4%.** Single-sample deltas below ~0.5% on this workload are unresolvable | - -| QWEN-032 | 2026-07-30 | `initialize_attention_state` is the largest actionable host phase (204 µs/step) and the code says why: the incremental fast path declines any caller that defers the bookkeeping publish, which is exactly what async scheduling does. Teaching the fast path to honor the deferred publish should recover most of that phase | `dynamic_context.py`: `_incremental_attention_state_update` takes the publish flags and returns `Tuple[bool, Optional[Event]]`, mirroring the full path's conditional publish at its tail; `initialize_attention_state` drops the two gate conditions and forwards the event | 28,178.8 tok/s (tpot 9.085 ms) | +1.04% vs QWEN-030 27,889.0 | coherence identical to QWEN-030 on all 3 prompts | session `qwen-gap1` job 5699934 | **The +1.04% is not real, and the diagnostic is why this entry matters.** A decline-reason tally (`MCORE_INFER_INCR_DIAG`) showed the fast path advancing on **0.1% of calls** (3 in 3,000), with **97.5% declining because `_request_layout_version` differed**. So the mechanism this change unlocks essentially never fires, and the throughput delta has no cause behind it — comparing iteration-by-iteration, the two runs' final iterations are 28,382 vs 28,405, i.e. the means differ only through the warm-up ramp. **Kept because it is a genuine prerequisite** (the gate really did exclude async) and it is what made the diagnostic possible, but it must not be credited with a gain. Lesson: a plausible mechanism plus a delta above the noise floor is still not evidence — instrument the mechanism itself, because here the two agreed and were both wrong | -| QWEN-033 | 2026-07-30 | The 97.5% decline points at one line: `resolve_requests()` bumps `_request_layout_version` **unconditionally**, invalidating the incremental cache every step. In steady-state decode that bump is provably spurious — with an all-ones active mask no rows move (`survivor_idxs == dst_idxs`), no KV blocks are released, the stale slice is empty and `total_request_count` is unchanged. Every real layout change in that method requires at least one finished request | `dynamic_context.py`: the bump moves from the top of `resolve_requests` into the `finished_idxs.numel() > 0` branch | 27,938.0 tok/s (tpot 9.163 ms) | **+0.18% vs QWEN-030 27,889.0 — no gain.** Best-config samples now read 27,889 / 28,179 / 27,769 / 27,938 (mean 27,944, spread 1.5%), so this sits inside the noise band | **bit-exact, verified under async**: 132 confirmation lines across ranks, up to 90 verified incremental steps each, **zero mismatches** against full recomputation. The specific risk — async's prepare-before-resolve reordering letting the cache key match while the true state differs — did not materialize | session `qwen-gap1` job 5699934 | **Mechanism fixed, no speedup — and that is the finding.** Fast-path engagement went **0.1% → 81%** of calls, so ~204 µs/step of host work really was removed, and throughput did not move. **Therefore that host work was already fully hidden behind GPU execution by async scheduling.** This independently confirms the host-side breakdown's implication and **retires host idle as a lever** — the line of attack that drove several sessions and three deadlocked nsys attempts. Corollary: the 1,598 µs/step of in-step GPU idle measured in the fresh trace should be treated as substantially a profiling artifact (`--cuda-graph-trace=node` plus nsys's own host overhead inflate exactly these host-bound gaps), not as 1.6 ms of recoverable time. **Kept anyway**: it makes a gate that ships in the tuned config actually functional instead of silently inert, at zero measured cost and with bit-exactness proven | - -### VLLM-SAMENODE: the first apples-to-apples reference (2026-07-30) - -Every "% of vLLM" figure in this ledger before now compared an mcore run on one node against a vLLM run on another, and this session measured node-to-node variance at ~2.7% — large enough to swamp most individual results. So the mcore session was ended and vLLM was run on **the same node** (`nvl72067-T13`, job 5701338), through the identical harness: same client (`static_benchmark.py`), same gsm8k dataset, BS256, OSL1024, 2 warmup + 5 timed iterations. - -**vLLM DP4+EP4: 34,415.7 tok/s, tpot 7.438 ms** (iters 34,095 / 35,141 / 34,495 / 33,765 / 34,615). - -| | tok/s | ms/step | % of vLLM | -|---|---|---|---| -| vLLM DP4+EP4 (same node) | 34,415.7 | 7.438 | 100% | -| mcore best config (mean of 4 samples) | 27,944 | 9.16 | **81.2%** | -| mcore best single run | 28,178.8 | 9.085 | 81.9% | -| mcore at session start (CLEANBASE-S12) | 26,538.2 | 9.646 | 77.1% | - -The session moved mcore from **77.1% → 81.2%** of same-node vLLM. Remaining gap: **1.72 ms/step (18.8%)**. Note this is *worse* than the 83.2% previously inferred from cross-node arithmetic — the honest number was slightly flattering before, which is exactly why the same-node reference was worth a run. - -### Host-side breakdown under async (the phase timing finally ran) - -The `perf_counter` shim (zero-nsys, built after three nsys host-capture deadlocks) produced its first usable report — 4,096 steps under the best config. The two largest entries are **not** actionable: `_forward` (13,996 µs/step) and `sampling` (7,233 µs/step) each exceed the 9,300 µs step, which is only possible because they are blocking waits on GPU work summed across two threads. Reading them as host cost would be the classic profiler mistake. - -Actionable pure-host cost, per step: `initialize_attention_state` **204 µs** (largest, and this is *with* `MCORE_INFER_INCR_ATTN_STATE` on), `prepare_requests` 80, `resolve_requests` 62, `_ep_establish_consensus` 58 (538 µs/call but only on ~10.6% of steps), `active_request_mask` 56, `async_sched_transfer_bookkeeping_to_gpu` 50, `drain_zmq_socket` 37. Summed with the smaller ranges, host work totals **~900 µs/step against a 9,300 µs step (~10%)**. - -**That total is the useful conclusion, and it redirects the campaign:** host work is far smaller than GPU-busy time, so a scheduler with good overlap should hide essentially all of it. Async captured only ~3.5%, so some overlap is imperfect — but the ceiling on *all* remaining host-side work is under 10%, and most of it is already hidden. **The remaining gap to vLLM is therefore predominantly GPU-side**, which retires "host idle is the top lever" — the ranking that had driven the last several sessions and that three failed nsys attempts were meant to refine. - -### Trace staleness: why the ranking had to be re-derived - -The gap-ranking used to pick targets came from `mcore-final-jul27` (QWEN-025, 10 gates). QWEN-026/027/028 then removed **precisely the kernels that ranking pointed at**, so it can no longer be used to choose the next target. Re-measuring the old trace made this concrete: - -| kernel (in the QWEN-025 trace) | launches | total device time | per call | status now | -|---|---|---|---|---| -| `direct_copy_kernel_cuda` (bf16) | 151,680 | 416.6 ms | 2.75 µs | **gone** (QWEN-026) | -| `triton_poi_fused_add_copy__0` | 151,680 | 200.1 ms | 1.32 µs | **~halved** (QWEN-027/028) | -| `rmsnorm_fwd_general` (q/k norm) | 2 per layer | — | ~2.3 µs | **gone** (QWEN-026, fused to 1) | - -The `direct_copy` finding is the useful one and explains QWEN-026's outsized +2.9%: timeline-neighbour analysis places two of these copies per layer between the QKV GEMM and the two QK norms, i.e. **the strided Q/K split views being materialized contiguous because TE's norm demands contiguous input**. At 2 × ~3 µs × 48 layers that is ~288 µs/step (~3% of an 8.85 ms step) of pure data movement. The fused QK-norm kernel reshapes the split views with no copy (`[-1, head_dim]` is a view when the last dim is contiguous), so it deleted the two copies *and* one norm launch — the copies, not the norm launches, were most of the win. **Generalizable: when a fusion replaces a vendor kernel that has a layout precondition, count the layout-conversion copies as part of the prize; they are often larger than the kernels being fused.** Next target selection requires a fresh trace of the current best config. - -## Session 2 (2026-07-23) — conclusion & recommendation - -**Best shippable config stays: legacy async / nvls dispatcher / vLLM grouped-GEMM -backend / FC1+SwiGLU fusion (QWEN-002) = 22,398.9 tok/s (65.9% of vLLM 33,994.5).** - -Every accessible knob was swept and rejected (QWEN-003…009). Key learnings: -- Decode at OSL1024 is **compute/comm-bound, not idle-bound** (async serial +0.85%). -- The gap is **structural**: mcore runs the MoE decode as a chain of discrete - kernels — router gemm → softmax → topk → count/scatter/metadata → **exposed - NVLS AllGather-V dispatch** → grouped GEMM (FC1+SwiGLU, FC2) → moe_sum → - **exposed NVLS ReduceScatter-V combine** — all on the per-layer critical path. - vLLM/TRT-LLM does the equivalent as one fused MoE (fused routing, cutlass - grouped bmm, fused finalize) with ~~0 exposed comm~~ — **the comm half of this claim - is false, see VLLM-COMM-S16**: vLLM runs the same AllGather+ReduceScatter pair per - layer per step. The fused-routing/finalize half stands. -- Built-in fusions (`--moe-router-fusion`, `--moe-permute-fusion`) and - `sampling-backend=flashinfer` are wired to training / non-full-graph paths and - are **incompatible** with the `inference_optimized` top-8 + full-iteration-graph - contract. - -**To actually close the 1.5× gap (large, multi-session work), in priority order:** -1. **Eliminate exposed NVLS comm** (~16% + it gates the critical path): overlap - dispatch/combine with expert GEMM via chunked/pipelined experts, or fuse the - combine ReduceScatter-V with the FC2 epilogue + moe_sum (finalize fusion). -2. **Fuse the routing chain** honoring the inference top-8 contract: one kernel - for softmax+top8+count+scatter-metadata (cuts ~18% of small serial kernels). -3. **A TRT-LLM-style single fused MoE decode kernel** (dispatch+groupedGEMM+ - SwiGLU+finalize) — the real vLLM-parity path; QWEN-001's full fusion was - slower, so this needs a cutlass/CUTE grouped-bmm core, not Triton. - -## Detailed records - -### VLLM-BASELINE — DP4 with expert parallelism - -| Field | Value | -|---|---| -| Date | 2026-07-22 | -| Hypothesis | Establish the fresh vLLM BS256 Nsight target | -| Code revision | `808c475352de6c3693b182048f174736af82356e`; skill files untracked, Megatron source clean | -| Changed files | none | -| Runtime flags | TP1, DP4, `--enable-expert-parallel`, max model length 4096, max sequences 512 | -| Image | `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/images/87e4947c6ce36433.sqsh` | -| Checkpoint / tokenizer | `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/checkpoints/qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/DP4/EP enabled | -| Workload | gsm8k, BS256, OSL128, one BS8/OSL32 warmup request, one timed request | -| Job / run | `5547673`; `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/runs/vllm-qwen30b-nsys-20260722-093437` | -| Throughput | 23,606.708 tok/s | -| Latency / TPOT | 1,368.153 ms / 10.844 ms-token | -| Correctness | Benchmark completed 256/256 requests | -| Nsight artifacts | `vllm_profile.nsys-rep`; `vllm_profile.sqlite` in the run directory | -| Result | vLLM target established | -| Next action | Record the matching mcore EP4 profile | - -### MCORE-BASELINE — EP4/TP1 - -| Field | Value | -|---|---| -| Date | 2026-07-22 | -| Hypothesis | Establish the fresh mcore EP4 BS256 Nsight starting point | -| Code revision | `808c475352de6c3693b182048f174736af82356e`; skill files untracked, Megatron source clean | -| Changed files | none in Megatron source; profile harness sets `max_requests=256` | -| Runtime flags | TP1, PP1, EP4, ETP1, NVLS dispatcher, vLLM grouped GEMM, inference-optimized transformer, full-iteration inference CUDA graphs | -| Image | Cog dev image `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/images/ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/EP4 | -| Workload | gsm8k, BS256, OSL128, one BS8/OSL32 warmup request, one timed request | -| Job / run | `5553135`; `/lustre/fsw/portfolios/coreai/users/shanmugamr/agents-space/runs/qwen-30b-nsys-20260722-161020` | -| Throughput | 12,346.092 tok/s | -| Latency / TPOT | 2,590.135 ms / 20.735 ms-token | -| Correctness | Checkpoint loaded and benchmark completed 256/256 requests | -| Nsight artifacts | `mcore_profile.nsys-rep`; `mcore_profile.sqlite` in the run directory | -| Result | mcore reaches 52.30% of vLLM; 47.70% below target | -| Next action | Analyze the timed-request A/B windows before changing code | - -### QWEN-001 — MoE decode mega-fusion (FC1+SwiGLU+FC2+topk-reduce) - -| Field | Value | -|---|---| -| Date | 2026-07-22 | -| Hypothesis | Fusing the whole MoE expert path into one Triton kernel (removing the `bounded_silu_mul` + `_moe_sum` kernels and two intermediate HBM round-trips) beats the 4-kernel `vllm_fused_moe` path, which is 40.5% of decode GPU-busy time | -| Code revision | branch `perf/moe-fused-decode-gemm` off `808c475352de6c3693b182048f174736af82356e`, dirty | -| Changed files | `megatron/core/inference/moe/fused_moe_decode.py` (new kernel), `dev/moe_fused/harness.py` (new standalone correctness+timing harness) | -| Runtime flags | standalone microbench; Qwen3-30B decode shapes H=2048, moe_ffn=768, 32 local experts, top-8, 256 valid tokens | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | n/a (synthetic weights, reference = production `vllm_fused_moe`) | -| Hardware / layout | OCI `oci-hsg`, 1×GB200 (session `qwen-moe-kernel`) | -| Workload | microbench, 10 warmup + 100 timed CUDA-event iters | -| Job / run | session `qwen-moe-kernel`, exec runs `manual1784770098/200/321` | -| Throughput | not measured end-to-end (rejected at microbench) | -| Latency / TPOT | kernel: reference 175 µs vs fused 258 µs (best sweep 0.80×; larger tiles OOM shared memory) | -| Correctness | Pass — max_abs_diff 2.6e-5, max_rel 0.15 on a 1e-4 floor, `allclose(rtol=2e-2,atol=2e-2)` True | -| Nsight artifacts | none (microbench) | -| Result | Rejected — fused kernel is consistently 20–50% slower than the reference in every same-run comparison | -| Next action | Root cause: one CTA per token-block serializes the H=2048 FC2 output loop (vs the reference's N-parallel multi-CTA GEMMs), 3× atomic traffic to `out`, and shared-memory pressure caps tile sizes; the HBM/launch savings are negligible under CUDA graphs. Pivot to either (a) partial FC1+SwiGLU epilogue fusion only, or (b) the 11.5% exposed NVLS all-gatherv/reduce-scatter-v communication | - -### QWEN-002 — FC1+SwiGLU epilogue fusion - -| Field | Value | -|---|---| -| Date | 2026-07-22 | -| Hypothesis | The 4-kernel MoE path (FC1→2N intermediate→`bounded_silu_mul`→FC2→reduce) wastes a full HBM round-trip of the `[num_valid, 2N]` intermediate and a whole kernel launch. Computing gate & up in the same FC1 program and applying `SiLU(gate)*up` in the fp32 epilogue — writing the `[num_valid, N]` activated intermediate directly — removes both while keeping FC2's N-parallel tiling | -| Code revision | branch `perf/moe-fused-decode-gemm`, dirty | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (add `FUSE_SWIGLU` constexpr to `_fused_moe_kernel`, `fuse_swiglu` to `_invoke_fused_moe_kernel`, `fuse_fc1_activation` path in `vllm_fused_moe`), `megatron/core/transformer/moe/experts.py` (`_vllm_forward` passes `fuse_fc1_activation=True`), `dev/moe_fused/harness_fc1.py` (new A/B harness) | -| Runtime flags | vLLM grouped-GEMM backend, SwiGLU; microbench at Qwen3-30B decode shapes H=2048, moe_ffn=768, 32 local experts, top-8, 256 valid tokens | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | n/a for microbench (synthetic weights; reference = unfused `vllm_fused_moe`) | -| Hardware / layout | OCI `oci-hsg`, 1×GB200 (session `qwen-moe-kernel`) | -| Workload | microbench, 10 warmup + 200 timed CUDA-event iters, 3 repeats | -| Job / run | session `qwen-moe-kernel`, exec runs `fc1c…`/`fc1t…` | -| Throughput | **22,741.9 tok/s** (fused) vs **22,617.7 tok/s** (same-env, fusion off) → **+0.55%**. gsm8k BS256 OSL1024, 2 warmup + 5 timed iters | -| Latency / TPOT | e2e TPOT 11.257 (fused) vs 11.319 ms/tok (off) → −0.55%. MoE-path kernel microbench: reference 174–182 µs vs fused 141–143 µs → **1.24–1.27×** (3 repeats) | -| Correctness | Pass — fused vs unfused `vllm_fused_moe`: max_abs_diff 3.9e-5, allclose; e2e coherence prompts coherent (2+2=4, capital of France = Paris) | -| Nsight artifacts | none (in-session A/B benchmark, not profiled) | -| Result | Accepted — correct, MoE kernel 1.25×, e2e +0.55% throughput / −0.55% TPOT, no regression. Unfused path byte-identical (all edits guarded by the `FUSE_SWIGLU` compile-time constexpr) | -| Next action | The 1.25× MoE-kernel win only nets +0.55% e2e at OSL1024, so MoE FC1 activation is not the wall-time bottleneck at the throughput regime — re-profile the OSL1024 decode (not the OSL128 profile trace) to find the true dominant cost before the next change. NOTE: at OSL1024 mcore already reaches ~22.6k tok/s; the ledger's 52% "gap" is an artifact of the OSL128 single-request profile workload | - -### QWEN-011 — fused MoE indirection-table build (lever #1) - -| Field | Value | -|---|---| -| Date | 2026-07-24 | -| Hypothesis | The decode "routing/permute kernel storm" is dominated by the 5 tiny serial kernels of `_moe_align_block_size_cuda_graphable` (init + count + prefix + fill + scatter) run per MoE layer × 48. Merging init + prefix + fill into one `_prefix_fill_init_kernel` cuts it to 3 kernels/layer (−96 launches/step) and narrows the vLLM gap without touching numerics. | -| Code revision | branch `perf/moe-fused-align` off `perf/moe-routing-fusion` (which carries QWEN-002), dirty | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_prefix_fill_init_kernel`, `_moe_align_block_size_fused`, env `MCORE_MOE_FUSED_ALIGN`, wired into `vllm_fused_moe`), `megatron/core/transformer/moe/experts.py` (`MCORE_FUSE_FC1_ACT` env gate), `dev/moe_fused/harness_align.py` (new A/B harness) | -| Runtime flags | `MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; nvls dispatcher, vLLM grouped-GEMM backend, full-iteration CUDA graphs | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 10 warmup + 200 timed CUDA-event iters (correctness) | -| Job / run | session `qwen-fuse` runs `fusedalign` (e2e) + `acb5ea1b` (microbench) | -| Throughput | **22,696.4 tok/s** vs 22,269.5 (fusion-on) / 22,241.5 (clean) → **+1.92% / +2.04%** | -| Latency / TPOT | avg_latency 11,257 ms vs 11,494 ms (clean) | -| Correctness | **Bit-exact** vs the 5-kernel path: max_abs_diff 0.0, allclose True; e2e coherence prompts coherent (2+2=4, Paris) | -| Nsight artifacts | in-session profile `sessions/qwen-moe-kernel/prof/1784930600/mcore_profile.{nsys-rep,sqlite}` (BS256/OSL128, fused-align ON). Kernel-name counts **confirm the mechanism**: `_init_sorted_ids_kernel`=0, `_prefix_sum_kernel`=0, `_fill_expert_block_ids_kernel`=0 (all removed), `_prefix_fill_init_kernel`=`_count_local_tokens`=`_scatter_token_indices_kernel`=76032 (align now 3 kernels/layer vs 5). microbench MoE-call 142.65→138.57 µs (1.029×) | -| Result | **Accepted** — correct (bit-exact), e2e +1.9% with tight iter variance (22.65–22.71k), no regression. Profile confirms −2 indirection kernels/MoE-layer/step. Default OFF (env-gated) so the untouched path stays byte-identical; enable with `MCORE_MOE_FUSED_ALIGN=1`. | -| Next action | Reduces align 5→3; the remaining serial dep (count → prefix_fill_init → scatter) can only merge further via a cooperative-grid launch (risk: CUDA-graph capture). Profile to confirm per-step kernel-count drop, then target the FC1/FC2 grouped-GEMM efficiency gap (lever #2, needs a cutlass/CUTE grouped bmm — Triton launch count is already minimal at 2/layer). | - -### QWEN-012 — grouped-GEMM roofline / padding decision gate - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | The ~40% of decode GPU time in `_fused_moe_kernel` is recoverable GEMM inefficiency, so a cutlass/CuTe BF16 grouped GEMM is worth writing. Measure achieved TFLOP/s, achieved weight bandwidth against a *measured* ceiling, and the indirection-table padding waste before writing any kernel. | -| Code revision | branch `perf/moe-fused-align` off `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, dirty (carries QWEN-002 + QWEN-011) | -| Changed files | `dev/moe_fused/harness_roofline.py` (new, analysis only — no Megatron source change) | -| Runtime flags | microbench at one EP4 rank of decode: hidden 2048, moe_ffn 768, 128 global / 32 local experts, top-8, `num_tokens_hint`=256 (= `local_tokens*ep_size` from `InferenceAllGatherDispatcherBase._get_host_valid_tokens_estimate()`), max_tokens 1024 | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`) | -| Checkpoint / tokenizer | n/a (synthetic weights at production shapes) | -| Hardware / layout | OCI `oci-hsg`, 1×GB200 of the 4-GPU node, session `qwen-cutlass` (job 5598012) | -| Workload | microbench, 20 warmup + 300 timed CUDA-event iters, 3 repeats (spread <1%) | -| Job / run | session `qwen-cutlass` exec `c05fd5f771b44bf7a002886a1242c1c4` | -| Throughput | not applicable (analysis) | -| Latency / TPOT | FC1 (fused SwiGLU) 52.25 µs, FC2 19.88 µs, **total 72.13 µs** per MoE layer per rank | -| Correctness | n/a (no functional change) | -| Nsight artifacts | none (CUDA-event microbench + torch profiler) | -| Result | **Memory-bound, not FLOP-bound — Stage 4 rejected before implementation.** (1) Measured bandwidth on this GB200: d2d copy 6.851 TB/s aggregate, Triton vectorized streaming read **6.081 TB/s** (torch `.sum()` gives 4.05 TB/s and is a poor proxy — do not use it as the ceiling). (2) Every decode step must read all 32 local experts' weights: 302.0 MB per layer per rank ⇒ **49.66 µs bandwidth floor**, vs 72.13 µs measured = **1.45× off roofline**. (3) Achieved 63.4 (FC1) / 83.2 (FC2) TFLOP/s on valid FLOPs — ~3% of GB200 BF16 dense peak; 246/323 TFLOP/s even counting padded rows. (4) Padding waste is large but almost free: `num_tokens_post_padded` vs 527 real local token-expert pairs is 784/1024/2048/4096 at BLOCK_SIZE_M 16/32/64/128 (32.8%/48.5%/**74.3%**/87.1% dead rows), yet dropping BLOCK_SIZE_M from the production 64 to 16 buys only ~1.2×, because each expert's weights are read exactly once per M-tile row-block regardless. (5) A 63-config tile sweep found 58.93 µs (1.225×) with one shared config and 57.24 µs (1.26×) with per-GEMM configs — i.e. **most of the 1.45× ceiling is reachable in Triton**, leaving ≤1.15× for any hand-written kernel. | -| Next action | Do not write a CUTLASS/CuTe grouped GEMM. Land the tile retune (QWEN-013), then redirect to the next-ranked levers from PROFILE-DECODE: exposed EP comm (~16%) and the routing chain (~18%). | - -### QWEN-013 — per-GEMM decode tile retune - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | QWEN-012's 1.45× roofline gap is tile-quantized SM occupancy, not the GEMM implementation. `_get_default_config` picks `BLOCK_SIZE_N=128` for both passes, leaving FC1 at ceil(768/128)=6 N-tiles × 32 M-tiles = 192 CTAs on 148 SMs (1.3 waves, ~30% tail idle). FC1 (N=768) needs a *small* `BLOCK_SIZE_N` to manufacture enough CTAs; FC2 (N=2048) has plenty and prefers a *large* one. vLLM shares one config across both passes and cannot express this. | -| Code revision | branch `perf/moe-fused-align` off `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, dirty | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_get_decode_tuned_configs`, `_TUNE_DECODE_GEMM` env gate, `config_fc1`/`config_fc2` threaded into the two `_invoke_fused_moe_kernel` calls and their grid sizing), `dev/moe_fused/harness_gemmtune.py` (new A/B + profiler + CUDA-graph harness), `dev/moe_fused/harness_align.py` (fix a module-vs-function import bug, see below) | -| Runtime flags | `MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; nvls dispatcher, vLLM grouped-GEMM backend, `transformer_impl=inference_optimized`, full-iteration CUDA graphs. Tuned configs: shared `BLOCK_SIZE_M=16`, `BLOCK_SIZE_K=64`, `GROUP_SIZE_M=1`; FC1 `BLOCK_SIZE_N=64`, 4 warps, 3 stages; FC2 `BLOCK_SIZE_N=256`, 8 warps, 4 stages. Applies only for `num_tokens_hint <= 512`; prefill keeps upstream's heuristic. | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 20 warmup + 300 timed iters × 3 repeats and a 50-iter torch-profiler window (attribution) | -| Job / run | session `qwen-cutlass2` (job 5600047) run `gemmtune-1785010791`; reference `ref0-1785010433` (CLEANBASE-S4) | -| Throughput | **23,636.0 tok/s** vs **22,657.7** same-session reference → **+4.32%**. Per-iter 23,553–23,668 tok/s | -| Latency / TPOT | avg_latency 10,802 ms vs 11,282 ms; TPOT 10.831 vs 11.299 ms/tok (−4.14%). Per-kernel GPU time: `_fused_moe_kernel` (both GEMMs) 77.65 → **61.89 µs** (1.255×); whole `vllm_fused_moe` call 100.64 → **84.31 µs** (1.194×) | -| Correctness | **Bit-exact** — max_abs_diff 0.0 and max_rel 0.0 at 128/256/384/512 valid tokens. Expected: `BLOCK_SIZE_K` is unchanged, so the fp32 reduction order along K is identical and only output-tile partitioning changes. e2e coherence prompts coherent (2+2=4, Paris, 3+2 cows = 5 cows) | -| Nsight artifacts | none (torch-profiler per-kernel attribution instead; timing composition change is fully explained by the two GEMM launches) | -| Result | **Accepted** — bit-exact, +4.32% e2e with tight variance, no regression. Default OFF (`MCORE_MOE_GEMM_TUNE`) so the untouched path stays byte-identical. New best mcore = **69.53%** of the vLLM 33,994.5 baseline (was 66.65%). | -| Next action | Two follow-ups, in order: (1) the remaining GEMM headroom is only 1.15× (57–62 µs vs the 49.66 µs floor), so stop here on the GEMM and attack **exposed EP comm (~16%)**; (2) `_get_decode_tuned_configs` is hand-tuned at one shape — if other MoE geometries adopt it, replace the constants with a small autotune keyed on (N, E, M). **Harness caveat worth knowing:** `import megatron.core.inference.moe.vllm_fused_moe as vfm` yields the *function*, not the module, because `moe/__init__.py` does `from .vllm_fused_moe import vllm_fused_moe`. Patching module flags on it silently no-ops, which makes an A/B compare one path against itself. This invalidated my first retune measurement (and affects QWEN-011's *microbench* numbers, though not its e2e result, which used the real env var). Both harnesses now use `importlib.import_module` and assert the object is the module. With the fix, the align A/B re-measures at 148.67 → 124.13 µs eager wall (1.198×). | - -### QWEN-014 — flashinfer `cutlass_fused_moe` backend A/B - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | `--inference-grouped-gemm-backend flashinfer` is already wired (`InferenceGroupedMLP._flashinfer_forward`, `experts.py:1057`) and accepts BF16, so QWEN-010's "blocked for SwiGLU" note is wrong and this backend has never actually been benchmarked. Enabling it should beat the vLLM Triton grouped GEMM, since it fuses permute + both GEMMs + activation into one launch. | -| Code revision | branch `perf/moe-fused-align`, dirty. No Megatron source change — the two integration bugs found were reproduced and fixed *in the harness only*, because the measurement then showed the backend is not worth adopting. | -| Changed files | `dev/moe_fused/inspect_flashinfer.py` (new), `dev/moe_fused/harness_flashinfer.py` (new) | -| Runtime flags | microbench at one EP4 rank: hidden 2048, moe_ffn 768, 128 global / 32 local experts, top-8, valid 256, max_tokens 1024, ep_size 4, ep_rank 0 | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh`; flashinfer **0.6.14**, torch 2.13.0a0, CUDA 13.3, GB200 sm100. `flashinfer_jit_cache` / `flashinfer_cubin` are **not** installed, so the first call downloads cubins from `edge.urm.nvidia.com` (~2 min from this cluster; it does have egress). Budget that for CUDA-graph warmup if this path is ever enabled. | -| Checkpoint / tokenizer | n/a (synthetic weights at production shapes; reference = production `vllm_fused_moe`) | -| Hardware / layout | OCI `oci-hsg`, 1×GB200 of the 4-GPU node, session `qwen-cutlass2` (job 5600047) | -| Workload | microbench: numerics single-shot, then 300 CUDA-graph replays × 3 repeats | -| Job / run | session `qwen-cutlass2` execs `7590adcb6e754ed8a7c32a58decb2f9b`, `659d263321e8df73...` | -| Throughput | not measured e2e — rejected at microbench | -| Latency / TPOT | CUDA-graph device time, valid=256: `vllm_fused_moe` default tiles **97.97 µs**, `vllm_fused_moe` QWEN-013 tiles **83.63 µs**, `cutlass_fused_moe` **90.81 µs** → **1.079× vs default, 0.921× vs QWEN-013**. (Eager wall-clock flatters cutlass to 1.17× because it is 1 launch vs 6 at ~12 µs/launch; that number is meaningless under full-iteration CUDA graphs.) | -| Correctness | Passes only after fixing the weight order. `[gate\|up]` (mcore's buffer as-is): max_abs 7.14e-4, max_rel **4.92**. `[up\|gate]`: max_abs 3.87e-5, max_rel 0.171 — bf16 rounding noise. So the kernel wants **up\|gate** (w3\|w1). Both pass a loose `allclose(2e-2)`, which is exactly why this would have shipped silently wrong. | -| Nsight artifacts | none (CUDA-graph event timing + torch profiler) | -| Result | **Rejected on measurement — but QWEN-010's stated reason was wrong, and two real mcore bugs were found.** (1) `_resolve_flashinfer_activation_type` (`experts.py:937-951`) maps `F.silu` → `ActivationType.Silu` **without consulting `config.gated_linear_unit`**, unlike `_resolve_mcore_activation_type` right below it which does. With `Silu` the kernel hard-fails: `fc1_expert_weights.size(1) == fc2_expert_weights.size(2) * mInnerDimMultiplier (1536 vs. 768)` — a non-gated activation expects fc1 out == ffn, but a SwiGLU fc1 emits 2×ffn. The correct enum is `ActivationType.Swiglu`, which works. **This — not any kernel limitation — is why the flashinfer backend appears "blocked for SwiGLU".** (2) The kernel expects fc1 as `[up\|gate]`; `_build_concatenated_weights` produces TE's `[gate\|up]`, so even with the enum fixed the results are numerically wrong. Fixing both is contained (an enum branch + a one-time reordered fc1 buffer, ~200 MB/layer extra since it cannot share storage with the TE `param.data` views), but at 0.921× it would still lose to QWEN-013. | -| Next action | Do not adopt the backend for this workload. Do fix the two bugs as a correctness matter on their own merit, since `--inference-grouped-gemm-backend flashinfer` is currently either a hard crash or silently wrong for every gated-activation MoE: guard `_resolve_flashinfer_activation_type` on `config.gated_linear_unit` (→ `Swiglu`), and reorder fc1 to `[up\|gate]` for that path. File as a separate bug-fix PR, not a perf change. | - -### QWEN-015 — TRT-LLM-Gen BF16 fused MoE (`trtllm_bf16_routed_moe`) probe - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | vLLM's BF16 SM100 oracle selects `FLASHINFER_TRTLLM` → `trtllm_bf16_moe` / `trtllm_bf16_routed_moe`, so this is the literal kernel vLLM wins with. The `_routed_` variant takes already-computed routing results, which matches mcore's contract, and it fuses SwiGLU + the topk-weighted finalize. Dropping it in behind an env gate should be the highest-value integration. | -| Code revision | branch `perf/moe-fused-align`, dirty. **No Megatron change made** — blocked at the probe stage. | -| Changed files | `dev/moe_fused/harness_flashinfer.py` (probe only) | -| Runtime flags | `num_experts=128, top_k=8, intermediate_size=768, local_expert_offset=0, local_num_experts=32, routing_method_type=RoutingMethodType.TopK, do_finalize=True`, both `use_shuffled_weight`/`weight_layout` combinations | -| Image | Cog dev image `ceecf5c304a5d8bd.sqsh`; flashinfer 0.6.14 (exposes `trtllm_bf16_moe`, `trtllm_bf16_routed_moe`, `convert_to_block_layout`, `reorder_rows_for_gated_act_gemm`) | -| Checkpoint / tokenizer | n/a (synthetic weights at production shapes) | -| Hardware / layout | OCI `oci-hsg`, 1×GB200, session `qwen-cutlass2` | -| Workload | microbench probe (numerics/latency never reached) | -| Job / run | session `qwen-cutlass2` exec `7590adcb6e754ed8a7c32a58decb2f9b` | -| Throughput | not reached | -| Latency / TPOT | not reached | -| Correctness | not reached | -| Nsight artifacts | none | -| Result | **Blocked on weight layout — and it is not the contract mismatch the plan anticipated.** `weight_layout=MajorK, use_shuffled_weight=False` is rejected by the launcher itself (`trtllm_fused_moe_kernel_launcher.cu:770`: `BF16 Moe: weight_layout must be BlockMajorK`). `weight_layout=BlockMajorK, use_shuffled_weight=True` then indexes `weights.size(3)` and throws `IndexError: Index 3 out of bounds for tensor with 3 dimensions`. So the kernel requires **4-D pre-shuffled block-major** weights; mcore's `[E, 2*ffn, H]` 3-D concatenated buffer cannot be passed in any configuration. Notably the dense top-8 `routing_map`, the `ep_size`/`ep_rank` semantics, and the NVLS symmetric-memory output tensor were *not* the blockers — the API accepts all of those. | -| Next action | Deprioritized, with a concrete recipe if revisited: build a one-time load-time weight pass using flashinfer's own helpers — `reorder_rows_for_gated_act_gemm` on fc1 (which also resolves the gate/up interleave that QWEN-014 found) then `convert_to_block_layout` on both — producing separate 4-D buffers. That breaks `_build_concatenated_weights`' storage sharing with TE's `param.data` views, so it roughly doubles expert-weight residency for the reordered copies (~14.5 GB/rank today) and needs its own `--inference-grouped-gemm-backend` value. **Priority is low now**: QWEN-012 measured the grouped GEMM at 1.15× off the weight-bandwidth floor after QWEN-013, and QWEN-014 measured the sibling cutlass kernel at 0.921× vs the retuned Triton path, so the fused-MoE ceiling here is roughly the ~22 µs of routing + finalize this kernel would absorb, not a GEMM win. Attack exposed EP comm (~16%) first. | - -## Session 4 (2026-07-25) — conclusion & recommendation - -**Best shippable config is now: nvls dispatcher / vLLM grouped-GEMM backend / -`MCORE_FUSE_FC1_ACT=1` (QWEN-002) / `MCORE_MOE_FUSED_ALIGN=1` (QWEN-011) / -`MCORE_MOE_GEMM_TUNE=1` (QWEN-013 + QWEN-013b) = 23,646.0 tok/s -= 69.56% of vLLM 33,994.5** (was 66.65% at the start of this session). - -The session was chartered to pursue "the grouped-GEMM / cutlass lever". The -measurement says that lever is nearly exhausted, and says so three -independent ways: - -1. **Roofline (QWEN-012).** The decode grouped GEMM must read 302 MB of expert - weights per layer per rank against a *measured* 6.081 TB/s streaming-read - ceiling — a 49.66 µs floor. It ran at 72.13 µs (1.45× off) and now runs at - ~57 µs (1.15× off). Achieved FLOPs are ~3% of BF16 peak, so this is a - bandwidth problem and no GEMM implementation can fix it. The dramatic-looking - 74.3% indirection-table padding is nearly free for the same reason. -2. **A hand-written kernel is not the lever (QWEN-013).** The entire 1.45× gap - was tile-quantized SM occupancy, recovered by giving FC1 and FC2 their own - `BLOCK_SIZE_N` — 12 lines, bit-exact, +4.36% e2e. -3. **The vendor kernels do not beat it (QWEN-014, QWEN-015).** flashinfer's - `cutlass_fused_moe` is 0.921× the retuned Triton path in CUDA-graph device - time, and `trtllm_bf16_routed_moe` cannot accept mcore's weight layout at all. - -**Remaining gap to vLLM is 1.44×, and it is no longer in the GEMM.** -PROFILE-TUNED confirms it from the trace side: the MoE GEMM is now 32.9% of -decode GPU time and sits only 229 µs/step (2.3%) above its bandwidth floor, -while exposed EP comm is 11.7%, routing/permute 15.8%, and GPU idle 20.6%. -Per PROFILE-DECODE the decode critical path is -attn → router → **exposed NVLS AllGather-V dispatch** → grouped GEMM → -**exposed NVLS ReduceScatter-V combine**. In priority order: - -1. **Exposed NVLS comm (~16%, and it gates the critical path).** ~~vLLM's - equivalent has *zero* exposed comm.~~ **False — see VLLM-COMM-S16**: vLLM runs the - same AllGather+ReduceScatter at the same per-layer cadence, so this is not pure - deficit and was over-prioritized here. Overlap dispatch/combine with the expert - GEMM via chunked/pipelined experts, or fuse the combine ReduceScatter-V into - the FC2 epilogue + `_moe_sum`. This is now the single highest-value target. -2. **The routing chain (~18%).** After QWEN-011 the align is 3 kernels; the - residual `_count_local_tokens` (8.2 µs) + `_moe_sum` (8.6 µs) + scatter - (3.1 µs) is ~20 µs of the 84 µs MoE call. Fusing `_moe_sum` into the FC2 - epilogue is the contained piece. -3. **Only then** revisit a fully fused MoE (QWEN-015's recipe), whose remaining - upside is the routing/finalize overhead, not the GEMM. - -## Session 5 (2026-07-25) — exposed NVLS EP comm - -### QWEN-016 — EP comm decision gate: exposed, latency-bound, small prize - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | The 929 µs/step of exposed NVLS EP comm is a worthwhile lever. Gate it first: measure how much is genuinely exposed, whether it is latency- or bandwidth-bound, and what the floor is under perfect overlap. | -| Code revision | branch `perf/moe-fused-align`, dirty. **No Megatron change retained** — analysis only. | -| Changed files | `dev/moe_fused/probe_comm.py`, `dev/moe_fused/analyze_comm.py`, `dev/moe_fused/analyze_comm_skew.py`, `dev/moe_fused/harness_comm.py` (all new, analysis harnesses; nothing under `megatron/`) | -| Runtime flags | trace side: none (re-analysis of PROFILE-TUNED). Microbench: `hidden=2048 topk=8 local_tokens=64 per_rank_max=2048 rsv_dtype=float32`, CTA sweep `4,8,16,32,64,128` | -| Image | cog dev image for `oci-hsg`, venv `envs/megatron_lm/dd356431262b5db4` (sqsh tag not re-verified this session) | -| Checkpoint / tokenizer | n/a (trace re-analysis + synthetic-payload microbench at production shapes) | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), EP4/TP1 | -| Workload | trace: PROFILE-TUNED steady-state decode, stream 257, steps 32–96, 48 layers, 4 ranks. Microbench: NVLS AGV/RSV at the production decode shape under CUDA-graph replay, median of repeats | -| Job / run | session `qwen-comm`, Slurm job 5601961 (`batch_long`, 8 h); exec `08e6a692b5164fa1929086fcc0e315d1` (CTA sweep), `04bdfdd605aa49348ff3d031aba70c89` (decomposition) | -| Throughput | not applicable (no Megatron change measured) | -| Latency / TPOT | not applicable | -| Correctness | microbench gate: AGV max\|diff\| 0.000e+00 and RSV max rel err 0.000e+00 over 64 trials vs `all_gather_into_tensor` / analytic rank-scaled reduction | -| Nsight artifacts | source trace `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite` | -| Result | **Gate answered on all three questions; the prize is smaller than its 11.7% share suggests.** (1) *Exposed*: the per-step union of comm intervals equals their sum — there is no concurrent compute, so it is 100% exposed. Per-step comm totals 928–1219 µs across the four ranks. (2) *Latency-bound, decisively*. Decomposing the production-shape collectives at 128 CTAs: AGV 6.57 µs = launch 0.72 + **barrier 5.08** + transfer 0.77; RSV 7.88 µs = launch 0.72 + **barrier 5.03** + transfer 2.13. Payloads are only 268 KB/rank (AGV) and 512 KB/rank (RSV); against the 900 GB/s/dir NVLink floor (0.894 / 1.748 µs) the RSV transfer is already at 82% of peak and the AGV transfer is *below* the unicast floor because multimem multicast pays one egress for all peers. Byte movement is therefore 127 µs/step out of ~693 µs — **fusing or batching bytes cannot win; only removing or hiding barriers can.** (3) *Floor*. Splitting the trace into intrinsic cost (median release − last arrival) and inter-rank arrival skew: AGV 6.78 µs/kernel = 5.97 intrinsic + 0.76 skew; RSV 15.30 = 7.19 + 8.11. Per step: 1060 µs total = **632 µs intrinsic + 428 µs skew**. The skew is not comm work — it is ranks waiting for the slowest rank's expert GEMM (ranks differ in how many experts receive tokens), so it is a routing-balance problem, not a collective problem. The skew-free microbench independently lands at 693 µs/step, corroborating the 632 µs figure. So the recoverable critical path is **632–693 µs of a 9,933 µs step = 6.4–7.0%**; perfect elimination gives 1.068–1.075× → 25,250–25,420 tok/s = 74.3–74.8% of vLLM 33,994.5. Of that, ~485 µs is 96 symmetric-memory barriers × 5.05 µs. Sub-verdict on the two proposed approaches: **(b) fusing RSV into the FC2 epilogue** removes a launch and the transfer but *keeps the barrier*, ceiling ≈ 2.9 µs/layer = 137 µs/step = **1.4%**. **(a) chunked/pipelined experts** hides a collective behind GEMM but adds one barrier per extra chunk (5.05 µs) against a 14.4 µs/layer collective cost, so a 2-chunk split is worth at best ≈ 9 µs/layer = 430 µs/step = **4.3%**, and only if CUDA-graph capture can express concurrent streams across the MoE dependency chain — which QWEN-008 already identified as the binding constraint. Separately, the CTA count is **already optimal**: the sweep gives AGV 13.54/9.68/8.45/7.89/8.07/6.78 µs and RSV 37.37/21.12/13.43/9.31/8.25/8.05 µs at 4/8/16/32/64/128 CTAs, so the shipped `MAX_NUM_BLOCKS=128` is the best point and reducing CTAs to shrink the per-CTA barrier is strictly worse. A prepared `MCORE_NVLS_{AGV,RSV}_CTAS` override was therefore reverted unused. | -| Next action | **Redirect off this lever.** A 6.4–7.0% ceiling of which 1.4–4.3% is realistically reachable, inside graph-captured dispatcher code, is worse value than routing/permute (1251 µs/step, 15.8%, spread over 242 kernels at 5.2 µs each — a launch-count problem where fusion converts directly into wall time). Take the ledger's item 2: fuse `_moe_sum` into the FC2 epilogue and collapse the residual `_count_local_tokens` / scatter chain. | - -### QWEN-017 — load-poll symmetric-memory barrier (rejected) - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | `symm_mem_sync`'s wait loop spins on a system-scope `atom.cas`, one full uncached read-modify-write per attempt, which sets how quickly a rank notices its peers arrived. Replacing it with an `ld.acquire.sys` poll plus a single clearing store should cut the 5.05 µs barrier and thus ~485 µs/step. | -| Code revision | branch `perf/moe-fused-align`, dirty. Change **reverted** after measurement. | -| Changed files | `megatron/core/inference/communication/torch_symm_triton/barrier.py` (added `_wait_signal_ldpoll` + `MCORE_SYMM_BARRIER_LDPOLL` gate, default off) — reverted; `dev/moe_fused/harness_comm.py` (correctness gate) — retained | -| Runtime flags | `MCORE_SYMM_BARRIER_LDPOLL=0` vs `=1`, production shape, 128 CTAs, 3 alternating reps each | -| Image | cog dev image for `oci-hsg`, venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | n/a (collective microbench at production shapes) | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), EP4/TP1 | -| Workload | NVLS AGV + RSV under CUDA-graph replay, median of repeats, 64-trial correctness gate before each timing block | -| Job / run | session `qwen-comm`, job 5601961; exec `04bdfdd605aa49348ff3d031aba70c89` | -| Throughput | not run e2e — rejected at the microbench stage | -| Latency / TPOT | per-step comm (48 layers): baseline 695.7 / 692.2 / 693.3 µs (mean 693.7); ldpoll 710.0 / 710.2 / 709.1 µs (mean 709.8). **+2.3% regression**, no overlap between the two triplets | -| Correctness | bit-exact in both variants: AGV max\|diff\| 0.000e+00, RSV max rel err 0.000e+00 over 64 trials | -| Nsight artifacts | none (microbench uses CUDA events) | -| Result | **Rejected — measured regression.** The barrier-only kernel is unchanged between variants (5.75–5.83 µs at both 256 and 512 threads/CTA, against a 0.72 µs empty kernel), so the ~5.05 µs is the 4-way system-scope flag round trip itself, not polling granularity. The cheaper poll actually costs slightly more, consistently on RSV (7.88 → 8.19 µs). Combined with the CTA-count result in QWEN-016, the barrier looks like a hardware/driver latency floor that a Triton-level rewrite does not move. | -| Next action | Do not pursue further barrier micro-optimisation. Any future attempt on this lever must remove barriers (fewer collectives) or hide them behind compute, not make each one cheaper. | - -## Session 6 (2026-07-25) — the routing/permute chain - -### QWEN-018 — routing/permute decision gate: per-kernel breakdown and fusion ceilings - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | Routing/permute (1251 µs/step, 15.8%, 242 kernels at 5.2 µs) is a launch-count problem, so removing kernels converts directly into wall time. Gate it first: attribute the 1251 µs to individual kernel names, split each into fixed launch cost vs real work, measure the inter-kernel dispatch gap, and compute a wall-time ceiling per candidate fusion. Reject any candidate below ~1%. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty. **No Megatron change** — analysis only. | -| Changed files | `dev/moe_fused/analyze_routing.py` (new, analysis only; nothing under `megatron/`) | -| Runtime flags | n/a (re-analysis of the PROFILE-TUNED trace, which was captured with `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`) | -| Image | n/a (login-node `python3` 3.12.13 + stdlib `sqlite3`) | -| Checkpoint / tokenizer | n/a | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, EP4/TP1 (trace); analysis on the login node | -| Workload | PROFILE-TUNED decode window 192.7–195.5 s, device 3, stream 257; 40 steady-state steps averaged, one MoE layer dumped in launch order | -| Job / run | session `qwen-comm`, Slurm job 5601961 (allocation held, analysis ran on the login node); source trace `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite` | -| Throughput | not applicable (analysis) | -| Latency / TPOT | step wall 9927.2 µs, GPU-busy 7900.0 µs, idle 2027.2 µs (20.4%), 1362 kernels — reproduces PROFILE-TUNED (9.933 / 7.884 / 2.049 ms, 1362) from an independent script | -| Correctness | n/a (no functional change) | -| Nsight artifacts | `qwen-cutlass2:prof/tuned-1785013555/mcore_profile.sqlite`; `dev/moe_fused/analyze_routing.py` output | -| Result | **The routing category is not one problem, it is two pathological kernels plus four cheap ones — and it is not launch-bound.** Per-step, per-kernel (48 launches each unless noted; `WALL` = device time + the dispatch gap that follows): `_moe_sum_kernel` 7.79 µs/k → **400.3 µs**; `_count_local_tokens_kernel_persistent` 7.52 → **387.4**; `gatherTopK` (router top-8) 6.05 → **316.9**; `_scatter_token_indices_kernel` 2.66 → **153.8**; `triton_per_fused__softmax_prep` (router softmax) 1.94 → **119.9**; `_prefix_fill_init_kernel` 1.33 → **90.9**; plus the `torch.zeros` fill of the count buffer (`vectorized_elementwise_kernel`, grid=1) 0.74 → **≈62**. Total **≈1531 µs/step**. Three measurements set the ceilings. (a) *The dispatch gap is 0.55 µs*, uniform across kernels in the graph, so a removed launch is worth `duration + 0.55 µs` — more than the kernel time alone, but far less than the 5.2 µs/kernel average would suggest if you assumed the whole thing were overhead. (b) *The fixed floor per launch is 1.27 µs* (0.72 µs empty kernel from QWEN-016 + 0.55 µs gap), so of the 1531 µs only **366 µs is fixed cost** and **1165 µs is in-kernel work** — pure launch-count reduction can win at most 3.7% and only by removing *every* routing launch. (c) *The 20.4% idle is now explained*: 862 µs of it is two once-per-step host gaps (537 µs after `index_elementwise_kernel`, 325 µs after `CatArrayBatchedCopy_vectorized`, i.e. sampling/detokenize between graph replays) and ~750 µs is 1362 × 0.55 µs of intra-graph node dispatch; there is no large unexplained residue. **Candidate ceilings.** *(1) Fuse `_moe_sum` into the FC2 epilogue* (the ledger's first-ranked): removes 48 × (7.79 + 0.55) = **400 µs = 4.03% gross**. But the topk slots of one token land in different experts' blocks and therefore different M-tiles, so accumulating into `out` needs cross-CTA fp32 atomics (~4 MB/layer of RMW: 512 local pairs × 2048 × 4 B) *and* a zeroing pass over `out[0:valid_tokens]` that today does not exist because `_moe_sum` writes rather than accumulates. That gives back ~70 µs, leaving ~3.3%, it is the exact shape QWEN-001 measured at 0.68–0.80×, and it stops being bit-exact. Held as the fallback, not built first. *(2) Cooperative-grid merge of count → prefix_fill_init → scatter*: the three are serially dependent, so a merge needs a grid-wide sync and keeps all the work; it recovers only the fixed cost, 2 × 48 × 1.27 = **122 µs = 1.23%**, before paying two `grid.sync()`s per layer against a kernel boundary that costs just 0.55 µs — and it needs cooperative launch under CUDA-graph capture. **Gated out on arithmetic.** *(3) New, and the one the breakdown actually points at: fold the token count and its `torch.zeros` fill into `_prefix_fill_init_kernel`.* `_prefix_fill_init_kernel` already has every CTA redundantly recompute the whole 32-wide cumsum in registers, and the count vector is its only input — so recomputing the histogram per CTA needs **no grid sync at all**. Removes 48 × (7.52 + 0.55) + 48 × (0.74 + 0.55) = **449 µs = 4.52% gross**, integer-exact. **Why `_count_local_tokens` costs 7.52 µs to bucket 2048 int32s**: it is neither launch-bound (0.72 µs floor) nor bandwidth-bound (8 KB), it is starved — with `BLOCK_SIZE=1024` and 2048 valid pairs `total_blocks=2`, so exactly **2 of its 152 CTAs receive work**, and each issues 1024 global atomics contending on 32 counters. **This corrects QWEN-003's conclusion.** QWEN-003 replaced those atomics with `tl.histogram` but left `BLOCK_SIZE=1024` untouched, so both variants ran on 2 CTAs; it measured a wash and inferred "the cost is per-launch fixed overhead, not atomic contention, so an in-kernel rewrite cannot help". The device-time measurement says the opposite: 6.8 of the 7.52 µs is in-kernel. | -| Next action | Build candidate (3) behind `MCORE_MOE_FUSED_COUNT` (default off, requires `MCORE_MOE_FUSED_ALIGN`) → QWEN-019. Two levers the gate surfaced but did not pursue, in order: the router pair `gatherTopK` + softmax is **436.8 µs/step = 4.4%** in two kernels (a hand-written fused top-8 router honouring the `inference_optimized` dense top-8 contract; QWEN-009 only ruled out TE's *built-in* fusion, not a hand-written one), and `_moe_sum` at 7.79 µs moves ~6.3 MB, i.e. ~1.05 µs at the measured 6.08 TB/s — it is **7× off its own bandwidth floor**, so restructuring it in place (hoist the 16 per-token scalar `routing_map` loads out of the K loop, parallelise over K as well as tokens) is worth up to 48 × 5.8 = 278 µs = 2.8% *without* the atomics that fusing into FC2 would require. | - -### CLEANBASE-S6 — fresh same-session reference - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | Establish a same-session OSL1024 reference at the session-4/5 best config before any session-6 A/B. Session 5 skipped this and session-to-session drift is ~1.5%. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b) | -| Changed files | none | -| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1`; nvls dispatcher, vLLM grouped-GEMM backend, `transformer_impl=inference_optimized`, full-iteration CUDA graphs, `CUDA_DEVICE_MAX_CONNECTIONS=1` | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` under the user checkpoint root | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters | -| Job / run | session `qwen-comm`, Slurm job 5601961; exec `023b75121915449badd4b39ca610f470`; run dir `sessions/qwen-comm/e2e/ref-s6-1785024868` | -| Throughput | **23,264.4 tok/s** (per-iter 23,019.4 / 23,402.0 / 23,189.7 / 23,457.4 / 23,258.9) | -| Latency / TPOT | avg_latency 10,974.7 ms, p50 10,972.5, p99 11,346.9; TPOT 11.004 ms/tok | -| Correctness | Benchmark completed 5/5 iters at 256 requests | -| Nsight artifacts | none (un-profiled throughput run) | -| Result | Reference for session 6 = 68.44% of vLLM 33,994.5. −1.61% vs QWEN-013b's 23,646.0, consistent with the documented ~1.5% session drift; all A/Bs this session are against this number, not against 23,646.0. | -| Next action | Verified the synced snapshot (`workspaces/megatron_lm/6b9355b187072223`) contains no session-6 kernel, so this is exactly the accepted session-5 code. Proceed to QWEN-019. | - -### QWEN-019 — fold the token count into the indirection-table build - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | QWEN-018's chosen candidate. The decode indirection-table build spends 4 launches producing a 32-element count vector and then consuming it: a `torch.zeros` fill (0.74 µs), `_count_local_tokens_kernel_persistent` (7.52 µs), `_prefix_fill_init_kernel` (1.33 µs), `_scatter_token_indices_kernel` (2.66 µs). The count's only consumer already has **every CTA redundantly recompute the whole cumsum in registers**, so recomputing the histogram per CTA as well removes the first two launches with no grid sync and no change to the counts. Ceiling 48 × (7.52 + 0.55 + 0.74 + 0.55) = 449 µs = 4.52% of the step. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_count_prefix_fill_init_kernel`, `_moe_align_block_size_count_fused`, `_USE_FUSED_COUNT` / `MCORE_MOE_FUSED_COUNT` gate, `_FUSED_COUNT_MAX_TOKENS`, three-way `align_fn` selection), `dev/moe_fused/harness_countfuse.py` (new A/B harness) | -| Runtime flags | `MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. The fused-count path is restricted to `num_tokens_hint <= 512`; above that the redundant per-CTA read stops being free and the atomic count kernel is used. | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 300 CUDA-graph replays × 3 repeats at 128/256/384/512 tokens (timing) and 16 table-equality cases (correctness) | -| Job / run | session `qwen-comm`, job 5601961; microbench exec `d9f7f74a868241d69eb4cc1f4062d40b`; e2e exec `3a807f77358644c89e62866749519802`, run dir `sessions/qwen-comm/e2e/countfuse-1785025539` | -| Throughput | **23,964.4 tok/s** vs **23,264.4** same-session reference → **+3.01%**. Per-iter 23,566.6–24,304.6 | -| Latency / TPOT | avg_latency 10,642.7 ms vs 10,974.7 (−3.03%); TPOT 10.682 vs 11.004 ms/tok (−2.93%). CUDA-graph replay device time at 256 tokens: align call **16.38 → 10.24 µs (1.600×)**, whole `vllm_fused_moe` call **86.02 → 79.96 µs (1.076×)**, i.e. 6.06 µs/layer × 48 = 291 µs/step = 2.93% of the 9.93 ms step — the e2e result matches the microbench prediction to within 0.1 pp. Per-kernel (eager profiler, 100 iters): count 4.76 + zeros 1.20 + prefix_fill 1.76 = 7.72 µs replaced by a single 3.73 µs kernel; align total device time 10.17 → 6.11 µs | -| Correctness | **Bit-exact, and verified two ways.** (1) Table equality over 16 cases (valid ∈ {128,256,384,512} × BLOCK_M ∈ {16,64} × local_expert_start ∈ {0,32}): `num_tokens_post_padded` identical, `expert_ids` identical, and the sorted multiset of `sorted_token_ids[0:npp]` identical (sorted because the scatter's atomics permute rows within an expert block on both paths). (2) Whole-MoE output at all four token counts: max_abs 0.0 **and max_rel 0.0** — per QWEN-014's lesson the relative error is checked, not a loose `allclose`. e2e coherence prompts coherent (2+2=4, Paris, 3+2 cows = 5 cows) | -| Nsight artifacts | none (CUDA-graph event timing + torch profiler; the change is a launch-count/kernel-time change fully explained by the two removed launches) | -| Result | **Accepted** — bit-exact, +3.01% e2e with tight variance, no regression. Default OFF so untouched paths stay byte-identical. New best mcore = **23,964.4 tok/s = 70.49%** of the vLLM 33,994.5 baseline. Realised 291 of the 449 µs ceiling; the shortfall is the new kernel's own cost, 1.33 → 3.73 µs (eager) as each of the 32 CTAs now histograms all 2048 valid pairs instead of reading a 32-element vector. **The mechanism confirms QWEN-018's correction of QWEN-003**: the count kernel's cost was in-kernel, not per-launch, and the fix was to stop running it on 2 of 152 CTAs — not to change how it reduces. | -| Next action | The same breakdown ranks `_moe_sum` next: 7.79 µs/layer to move ~6.3 MB is 7× its own bandwidth floor, and it can be attacked without the cross-CTA atomics that fusing it into FC2 would need → QWEN-020. After that, the router pair (`gatherTopK` + softmax, 436.8 µs/step = 4.4% in two kernels) is the largest remaining routing item. | - -### QWEN-020 — predicate the locality test in the topk reduction - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | QWEN-018 measured `_moe_sum_kernel` at 7.79 µs/layer to move ~6.3 MB — 7× the ~1.05 µs its own traffic implies at the 6.08 TB/s of QWEN-012 — so it is neither launch- nor bandwidth-bound. The suspected cause is the `if lid >= 0 and lid < num_local_experts` guard: a uniform scalar branch gated on a dependent global load of `routing_map`, which prevents slot `t`'s data load from overlapping slot `t+1`'s index load and makes the CTA walk the 8 topk slots serially. Predicating the guard into the load mask, and widening `BLOCK_K` from 1024 to the full hidden size so the per-token index loads are issued once instead of `NUM_K_BLOCKS` times, should recover most of that gap **while staying bit-exact** — the reduction order and the fp32 arithmetic are unchanged and masked-off slots contribute an exact 0.0. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019) | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_moe_sum_kernel_fast`, `_USE_FAST_MOE_SUM` / `MCORE_MOE_SUM_FAST` gate, `_FAST_MOE_SUM_MAX_BLOCK_K = 2048`, dispatch in `_moe_sum`), `dev/moe_fused/harness_moesum.py` (new A/B harness) | -| Runtime flags | `MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. One new variable vs QWEN-019. | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 300 CUDA-graph replays × 3 repeats at 128/256/384/512 valid tokens, plus whole-MoE output equality at the same four shapes | -| Job / run | session `qwen-comm`, job 5601961; microbench exec `af464a419c284937a106aae25196359a`; e2e exec `af49c9f54f4a497d8808f50614fe6bf6`, run dir `sessions/qwen-comm/e2e/moesum-1785026515` | -| Throughput | **24,403.7 tok/s** vs **23,964.4** (QWEN-019) → **+1.83%**; vs **23,264.4** (CLEANBASE-S6) → **+4.90%**. Per-iter 24,442.9 / 24,137.8 / 24,580.4 / 24,607.1 / 24,257.1 | -| Latency / TPOT | avg_latency 10,445.6 ms (p50 10,447.9, p99 10,806.4); TPOT **10.490** vs 10.682 ms/tok. Microbench, CUDA-graph replay of the whole `vllm_fused_moe` call: 128 tok 73.10 → 71.59 µs (1.021×), 256 tok 80.69 → 78.77 (1.024×), 384 tok 86.81 → 83.94 (1.034×), 512 tok 102.68 → 98.56 (1.042×). Eager per-kernel at 256 tokens (100 iters): `_moe_sum_kernel` **8.14 → 5.97 µs (1.36×)**, every other kernel in the call unchanged (`_fused_moe_kernel` 62.18 → 61.88, `_count_prefix_fill_init` 4.18 → 4.18, `_scatter_token_indices` 3.07 → 3.03) | -| Correctness | **Bit-exact.** Whole-MoE output at valid ∈ {128, 256, 384, 512}: max_abs 0.0 **and max_rel 0.0** (QWEN-014's lesson — relative error, not a loose `allclose`). e2e coherence prompts coherent (2+2 = 4, Paris, 3+2 cows = 5 cows). Benchmark completed 5/5 iters | -| Nsight artifacts | none (CUDA-graph event timing + torch profiler; the change alters one kernel's internals and no launch counts, which the per-kernel table isolates directly) | -| Result | **Accepted** — bit-exact and positive on every shape measured. New best mcore = **24,403.7 tok/s = 71.79%** of the vLLM 33,994.5 baseline. Default OFF; enable with `MCORE_MOE_SUM_FAST=1`. **Honest accounting of the size of the win:** the microbench says 2.17 µs/layer × 48 = **104 µs/step ≈ 1.0%** of the step, and the graph-replay whole-call delta at 256 tokens is smaller still (1.92 µs × 48 = 92 µs), whereas the e2e TPOT moved 0.192 ms/step (192 µs). The e2e gain is therefore about **2× the microbench prediction**, and QWEN-019's per-iteration range (23,566–24,305) overlaps this run's (24,138–24,607), so part of the measured +1.83% is run-to-run variance rather than kernel time. The change is kept because it is bit-exact and every isolated measurement of it is a strict improvement, but the defensible attribution is ~1%, not 1.83%. QWEN-018's estimated ceiling for restructuring this kernel in place was 278 µs (2.8%); the predication recovered 104 µs of it, so `_moe_sum` at 5.97 µs is still ~6× its bandwidth floor and the remaining gap is the strided `[token, slot, K]` gather itself, not the branch | -| Next action | The router pair is now the largest routing item the QWEN-018 breakdown left standing: `gatherTopK` 6.05 µs + softmax 1.94 µs = 436.8 µs/step = 4.4% in two kernels, for what is only a `[256, 128]` fp32 softmax and a top-8 select → QWEN-021. | - -### QWEN-021 — fused softmax + top-8 router selection - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | QWEN-018 left the router pair as the largest untouched routing item: `gatherTopK` 6.05 µs + the compiled softmax 1.94 µs = 436.8 µs/step (4.4%) — for nothing more than a `[256, 128]` fp32 softmax and a top-8 select over 128 experts. `torch.topk` runs a multi-pass radix select sized for large `n`; at 128 candidates per row, one CTA per token can hold the whole row in registers, softmax it, and pick the top 8 by 8 max-then-mask passes. Ceiling: replace ~8.0 µs + 0.55 µs of dispatch gap per layer with a single ~2 µs kernel ⇒ ~312 µs/step ≈ 3.1%. **This is not QWEN-009**, which ruled out TE's *built-in* fused router because it emits a dense 128-expert map while `inference_optimized` needs the dense top-8 contract; this kernel is written to the top-8 contract directly. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020) | -| Changed files | `megatron/core/inference/moe/router_topk.py` (new: `_softmax_topk_kernel`, `fused_softmax_topk`, `can_use_fused_softmax_topk`, `MCORE_ROUTER_FUSED_TOPK` gate, `FUSED_ROUTER_TOPK_MAX_TOKENS = 1024`), `megatron/core/transformer/moe/router.py` (`InferenceTopKRouter._forward` takes the fused path when the contract matches), `dev/moe_fused/harness_routertopk.py` (new A/B harness) | -| Runtime flags | `MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`; otherwise identical to CLEANBASE-S6. One new variable vs QWEN-020. The fused path is taken only for softmax + pre-softmax + no groups + no scaling factor + no expert bias + no router replay + ≤1024 tokens; anything else (including prefill chunks above 1024 tokens) falls back to the compiled torch path unchanged. | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters (throughput); microbench 500 CUDA-graph replays × 3 repeats at 128/256/384/512 tokens × 128 experts, correctness at those four shapes × 2 seeds | -| Job / run | session `qwen-comm`, job 5601961; microbench exec `56ae1d844a86427ea60ddce2c73d903c`; e2e exec `2041011b410b45709d455514b32c85f3`, run dir `sessions/qwen-comm/e2e/routertopk-1785027029` | -| Throughput | **25,352.9 tok/s** vs **24,403.7** (QWEN-020) → **+3.89%**; vs **23,264.4** (CLEANBASE-S6) → **+8.98%**. Per-iter 25,012.5 / 25,092.4 / 25,405.9 / 25,625.0 / 25,642.4 — the whole range sits above QWEN-020's whole range (24,137.8–24,607.1), so unlike QWEN-020 this win is outside run-to-run variance | -| Latency / TPOT | avg_latency 10,041.4 ms (p50 10,039.1, p99 10,439.5); TPOT **10.097** vs 10.490 ms/tok → **−393 µs/step** against a predicted ~312 µs. CUDA-graph replay of the router selection alone: 128 tok 16.40 → 4.10 µs, 256 tok **16.41 → 4.10 µs (4.00×)**, 384 tok 18.43 → 4.10 (4.49×), 512 tok 20.47 → 4.11 (4.99×) — the fused kernel is flat in token count over this range while the torch pair is not. Eager per-kernel at 256 tokens: `gatherTopK` 9.12 + `unrolled_elementwise` 2.76 + `softmax_warp_forward` 2.09 + `vectorized_elementwise` 1.57 = **15.54 µs in 4 kernels → 2.04 µs in 1** | -| Correctness | **Bit-exact probabilities and identical expert sets** at 128/256/384/512 tokens × 2 seeds each: the selected expert id multiset per token is equal, max_abs 0.0 and max_rel 0.0 on the probabilities after aligning by expert id, and no duplicate expert ids. (`torch.topk` is called with `sorted=False` during inference, so the *order* of the k results is unspecified on both sides; the fused kernel returns descending score order with ties broken toward the lower expert id.) e2e coherence prompts coherent (2+2 = 4, Paris, 3+2 cows = 5 cows); benchmark 5/5 iters | -| Nsight artifacts | none (CUDA-graph event timing + torch profiler; a 4→1 launch change with per-kernel attribution on both sides) | -| Result | **Accepted — the largest single win of session 6.** New best mcore = **25,352.9 tok/s = 74.58%** of the vLLM 33,994.5 baseline. Default OFF; enable with `MCORE_ROUTER_FUSED_TOPK=1`. The measured e2e gain (393 µs/step) slightly exceeds the 312 µs predicted from QWEN-018's two-kernel accounting because the profile's "router" attribution missed the two elementwise kernels (the `.type_as` cast and one more) that the fused kernel also absorbs — 4 launches removed, not 2. Session total so far: 23,264.4 → 25,352.9 = **+8.98%**, all of it bit-exact. | -| Next action | The last item in the QWEN-018 breakdown that clears the gate is `_scatter_token_indices_kernel` at 2.66 µs + 0.55 µs gap × 48 = 154 µs = 1.6%. Now that QWEN-019 made every CTA stream the pairs anyway, the scatter can be a second streaming pass inside the same kernel — no atomics, no grid sync — taking the indirection-table build to a single launch → QWEN-022. | - -### QWEN-022 — one-launch indirection-table build (scatter folded in) - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | The last routing item above the gate is `_scatter_token_indices_kernel`: 2.66 µs + 0.55 µs dispatch gap × 48 = 154 µs/step ≈ 1.6%. QWEN-019 already has CTA `e` stream every valid pair to build a private histogram, so a second streaming pass over the same (L2-resident, 16 KB) pairs lets that CTA place its own rows itself at `excl_e + written + exclusive_cumsum(is_mine)` — no global atomics, no grid sync, and a deterministic table instead of an atomically permuted one. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (carries QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020 + QWEN-021) | -| Changed files | `megatron/core/inference/moe/vllm_fused_moe.py` (`_align_single_kernel`, `_moe_align_block_size_single`, `_USE_FUSED_SCATTER` / `MCORE_MOE_FUSED_SCATTER` gate, align dispatch), `dev/moe_fused/harness_scatterfuse.py` (new A/B harness) | -| Runtime flags | `MCORE_MOE_FUSED_SCATTER=1` on top of the QWEN-021 set (`MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_SUM_FAST=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_FUSE_FC1_ACT=1`). One new variable. | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed iters, **run twice**; microbench 300 replays × 3 repeats at 128/256/384/512 tokens, 16 table-equality cases | -| Job / run | session `qwen-comm`, job 5601961; microbench execs `c7cfd653ca6f4bc1849649746fb1cb48` (first, with a faulty check) and `46060ef1233f407fb929ab82755517d4` (corrected); e2e execs `a31dd240c22d47588ffa58e264180f32` (`e2e/scatterfuse-1785027598`) and `a5557671044f45709e0a933e7b255741` (`e2e/scatterfuse-rep2-1785027955`) | -| Throughput | run 1 **25,441.4 tok/s** (+0.35% vs QWEN-021's 25,352.9), run 2 **25,495.9 tok/s** (+0.56%); mean **25,468.7 = +0.46%**. Per-iter run 1 24,920.5–25,791.6, run 2 24,870.8–25,827.4 — both ranges overlap QWEN-021's (25,012.5–25,642.4) | -| Latency / TPOT | TPOT 10.062 and 10.041 ms/tok vs 10.097 (−35 and −56 µs/step) against a microbench prediction of 125 µs/step. CUDA-graph replay at 256 tokens: align call **10.08 → 8.23 µs (1.224×)**, whole MoE call **82.17 → 79.56 µs (1.033×)**. At other shapes: 128 tok 1.045×, 384 tok 1.007×, **512 tok 0.997× (a slight loss)** — the extra streaming pass scales with pair count while the removed launch does not. Eager per-kernel: `_count_prefix_fill_init` 3.74 + `_scatter_token_indices` 2.38 = 6.12 µs in 2 kernels → `_align_single_kernel` 4.91 µs in 1 | -| Correctness | **Bit-exact whole-MoE output** at 128/256/384/512 (max_abs 0.0 and max_rel 0.0), plus table equality over 16 cases (valid × BLOCK_M ∈ {16,64} × local_expert_start ∈ {0,32}): identical `num_tokens_post_padded`, identical `expert_ids`, identical row multiset, and — the stronger check — every non-sentinel row sits in a block whose expert id actually owns that pair. e2e coherence prompts coherent; both benchmark runs 5/5 iters. **Note on the first harness run:** it reported a table MISMATCH at BLOCK_M=16, which was a bug in the *check*, not the kernel — an expert's rows span several BLOCK_M blocks and the two paths order them differently across that range, so comparing per 16-row block flags a difference that does not exist. The whole-MoE output was bit-exact in that same run, which is what exposed the faulty check. | -| Nsight artifacts | none (CUDA-graph event timing + torch profiler) | -| Result | **Accepted, marginally, and recorded as a weak win rather than a 1.2% one.** Both e2e runs beat QWEN-021, so the sign is reliable; the size (+0.46%) is a third of the microbench prediction and sits inside the per-iteration spread, so the defensible claim is "small but positive". Kept because it is bit-exact, removes a launch, and makes the table deterministic. Default OFF. Best mcore = **25,495.9 tok/s = 75.00%** of the vLLM 33,994.5 baseline (mean of the two runs 25,468.7 = 74.92%). **This is the point where the routing lever runs out**: every routing kernel QWEN-018 listed has now been fused or rewritten except `_moe_sum`, whose remaining gap needs the FC2-epilogue fusion QWEN-018 gated as ~3.3% net and QWEN-001-shaped. | -| Next action | Re-profile at the new configuration before choosing the next lever — the step has lost ~10% of its wall time and roughly 700 µs of routing, so the QWEN-018 ranking is stale → PROFILE-S6. | - -### PROFILE-S6 — re-profile at the session-6 configuration - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | The QWEN-018 ranking is stale: the step has lost ~10% of its wall time and roughly 820 µs of routing since it was taken. Re-profile with all six gates on and re-rank before choosing the next lever. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (QWEN-002 + QWEN-011 + QWEN-013/013b + QWEN-019 + QWEN-020 + QWEN-021 + QWEN-022) | -| Changed files | none (`dev/moe_fused/profile_insession.sh`) | -| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh`, venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200 (`nvl72086-T04`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, **OSL128** under nsys (`--trace=cuda,nvtx --cuda-graph-trace=node`), 1 iter after a BS8 warmup; analysis window device 1, 184.9–186.9 s (~220 steady-state decode steps) | -| Job / run | session `qwen-comm`, job 5601961; exec `1a4e8dc90c744a1c9ae0d977ec5ad16e`; run dir `sessions/qwen-comm/prof/s6-all-1785028240` | -| Throughput | 10,454.6 tok/s at OSL128 under nsys (profiling overhead + short OSL; not comparable to the OSL1024 ledger numbers) | -| Latency / TPOT | one steady-state decode step: **wall 9.097 ms** (PROFILE-TUNED: 9.933), **GPU-busy 7.183 ms** (7.884), **idle 1.914 ms = 21.0%** (2.049 ms = 20.6%), **1170 kernels/step** (1362) | -| Correctness | n/a (profile capture) | -| Nsight artifacts | `qwen-comm:prof/s6-all-1785028240/mcore_profile.nsys-rep` and `.sqlite`; `forward_pass.py` and `dev/moe_fused/analyze_routing.py --device 1 --window 184.9,186.9` output | -| Result | **The routing lever is spent.** Per forward pass: MoE expert GEMM 2469.1 µs / 96 kernels (27.1%), dense GEMM 1083.6 / 241 (11.9%), comm 1062.9 / 96 (11.7%), attention 923.8 / 96 (10.2%), elementwise 487.6 / 253 (5.4%), norm 449.5 / 193 (4.9%), **MoE routing/permute 430.4 µs / 98 kernels (4.7%)** — down from 1251 µs / 242 kernels, a 66% cut, and it is now the smallest GPU category except sampling. Counting dispatch gaps, the routing family (including `_align_single_kernel`, which the categorizer files under "other") is ~665 µs/step. Its two survivors are `_moe_sum_kernel_fast` (5.57 µs × 48 = 294 µs with gaps) and `_align_single_kernel` (5.13 × 48 = 272 µs); `_softmax_topk_kernel` costs 1.53 µs × 48 = 99 µs, versus 436.8 µs for the pair it replaced. **The idle is now the ranked next lever and it is host-side, not GPU-side**: of 1914 µs, ~1237 µs sits in three once-per-step gaps — **548 µs** after `index_elementwise_kernel`, **386 µs** after `vectorized_elementwise_kernel`, **303 µs** after `CatArrayBatchedCopy_vectorized` — i.e. sampling, detokenize and scheduling on the host between graph replays. The remainder (~644 µs) is 1170 × 0.55 µs of intra-graph node dispatch, which only fewer kernels can shrink. | -| Next action | Attack the ~1237 µs/step of host gaps, but **profile the host first** — this is not the same experiment as QWEN-005a/QWEN-007 (`async-sched-mode=serial`, which crashed and then measured only +0.85%) or QWEN-004 (`sampling-backend=flashinfer`, incompatible with the full-iteration graph). Capture with `--trace=cuda,nvtx,osrt` plus Python sampling over the decode window and attribute the three gaps to concrete host functions before proposing a change; at 13.6% of the step this is now worth more than any remaining GPU-side category except the (closed) expert GEMM. | - -### PROFILE-HOST-S6 — host-visible capture for the inter-replay gaps - -| Field | Value | -|---|---| -| Date | 2026-07-25 | -| Hypothesis | PROFILE-S6's ~1237 µs/step of host gaps between CUDA-graph replays (548 µs after `index_elementwise_kernel`, 386 µs after `vectorized_elementwise_kernel`, 303 µs after `CatArrayBatchedCopy_vectorized`) are unattributed because every existing trace is GPU-only. A trace with `osrt` + CPU sampling + Python sampling will name the host functions occupying them and classify each gap as launch overhead, host synchronization, sampling/detokenization, scheduler bookkeeping, or ZMQ round-trip. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (session-6 set unchanged: `vllm_fused_moe.py`, `experts.py`, `router.py`, untracked `router_topk.py`). **No `megatron/` edits made by this task** — verified with `git diff --stat -- megatron`. | -| Changed files | `dev/moe_fused/profile_host_insession.sh` (new — host-visibility variant of `profile_insession.sh`), `dev/moe_fused/dispatch_host_profile.sh` (new — waits for the session to reach `running`, then dispatches the capture detached) | -| Runtime flags | All seven gates: `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | -| Image | cog dev image `ceecf5c304a5d8bd.sqsh` (`nvcr.io/nvidia/pytorch:26.06-py3`), venv `envs/megatron_lm/dd356431262b5db4` | -| Checkpoint / tokenizer | `qwen3-30b-a3b-mcore` / `qwen3-30b-a3b-hf` | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1, nvls dispatcher, vllm grouped-GEMM, `transformer_impl=inference_optimized`, `full_iteration_inference` CUDA graphs | -| Workload | gsm8k, BS256, OSL128 under nsys, 1 iter after a BS8/OSL32 warmup (identical to PROFILE-S6 except for the trace flags) | -| nsys version | 2026.3.1.117-263137992252v0 — verified on-node that `--python-sampling`, `--python-sampling-frequency`, `--backtrace`, `--samples-per-backtrace`, `--cpuctxsw`, `--osrt-threshold` are all supported before committing to the command line | -| nsys command line | `nsys profile --trace=cuda,nvtx,osrt --sample=process-tree --backtrace=fp --samples-per-backtrace=1 --cpuctxsw=process-tree --python-sampling=true --python-sampling-frequency=1000 --osrt-threshold=1000 --cuda-graph-trace=node --force-overwrite=true -o /mcore_host_profile $PYBIN -m torch.distributed.run --nproc-per-node 4 ... -m examples.inference.launch_inference_server ...` (server args identical to `profile_insession.sh`). The script preflights this flag set against `python -c print(...)` and degrades to `--trace=cuda,nvtx,osrt --sample=process-tree --backtrace=fp --python-sampling=true` and then to `--trace=cuda,nvtx,osrt --sample=process-tree` if nsys rejects it, so an unsupported flag cannot abort the run. | -| Job / run | Intended session `qwen-comm` job 5601961 — **PREEMPTED** (`raw_state=PREEMPTED`, controller dead) before any capture ran; note `cog jobs get` still reported `RUNNING` while `cog session status` correctly reported `preempted`. Replacement: session `qwen-host`, job **5607600**, node `nvl72166-T15`, partition `batch`, `--time 02:00:00` (started 2026-07-26T06:53 UTC, expires ~08:53 UTC), 4 GPUs. Capture dispatched detached as exec **`d920241ea00a41618c63a75fefcb9ea2`** at 06:53:51 UTC; run dir **`sessions/qwen-host/prof/hosts6-1785048883`**. Exec log: `sessions/qwen-host/exec/runs/d920241ea00a41618c63a75fefcb9ea2/stdout.log`. | -| Throughput | n/a (profile capture) | -| Latency / TPOT | n/a (profile capture) | -| Correctness | n/a (profile capture) | -| Nsight artifacts | **Expected at `sessions/qwen-host/prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}` — dispatched and confirmed running, but not confirmed complete.** At hand-off the exec had reached the nsys flag preflight; the server load + BS8 warmup + BS256/OSL128 benchmark + sqlite export had not finished. The script ends with a sanity check that prints row counts for `CUPTI_ACTIVITY_KIND_KERNEL`, `CUPTI_ACTIVITY_KIND_RUNTIME`, `OSRT_API`, `COMPOSITE_EVENTS`, `SAMPLING_CALLCHAINS`, `PYTHON_SAMPLING_CALLCHAINS`, `PYTHON_SAMPLING_STRING`, `SCHED_EVENTS`, `GENERIC_EVENTS` and a `HOST_TABLES_PRESENT=` line — a capture with only `CUPTI_ACTIVITY_KIND_KERNEL` is a failed capture and must be re-run. | -| Result | **Inconclusive — capture not confirmed, no attribution performed.** The task lost its node: session `qwen-comm` job 5601961 was preempted roughly six hours into an eight-hour `batch_long` allocation, before the first capture could be dispatched, and the whole capture had to be re-queued from scratch on a new allocation that did not schedule inside the remaining budget. The deliverable that *was* produced is the reusable capture path: an nsys flag set validated against the installed 2026.3.1 on-node, a profile script that preflights and degrades rather than aborting, and a built-in host-table sanity check. Two operational findings worth carrying: (1) `cog jobs get --job-id` reported `RUNNING` for a preempted session — use `cog session status` to decide whether a session is usable; (2) an unrelated background `rsync` writing into the worktree caused two `WorktreeChangedError` dispatch failures, so no local process may write inside the repo while a `cog session exec` snapshot is in flight. | -| Next action | Check `sessions/qwen-host/exec/runs/d920241ea00a41618c63a75fefcb9ea2/stdout.log` for `PROFILE DONE` and `HOST_TABLES_PRESENT=True`. If present, run Workflow C on the new `.sqlite`, then window the `OSRT_API` / `CUPTI_ACTIVITY_KIND_RUNTIME` / `PYTHON_SAMPLING_CALLCHAINS` tables to the three inter-replay gaps and attribute them. If the exec never ran, re-dispatch `dev/moe_fused/profile_host_insession.sh` on a fresh session — it is unchanged and ready. | - -### HOSTGAP-S6 — host attribution of the three inter-replay gaps - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | PROFILE-S6's three largest inter-kernel gaps (after `index_elementwise_kernel`, `vectorized_elementwise_kernel`, `CatArrayBatchedCopy_vectorized`, ~1237 µs/step together) can be attributed to named host functions from the PROFILE-HOST-S6 capture, and each can be classified as CPU-bound host execution, blocking synchronization, driver launch latency, allocator activity, a genuine GPU dispatch gap, or profiler overhead. | -| Code revision | `37d1adcf09390c2593ad18b8dbe3b3cdb7e25514`, branch `perf/moe-fused-align`, dirty (session-6 set). **No `megatron/` edits** — analysis only. | -| Changed files | `dev/moe_fused/analyze_hostgaps.py` (new — `gaps` brackets every inter-kernel gap in the kernel-union complement by the kernel ending at its start and the kernel starting at its end; `attrib` aggregates CUPTI runtime, OSRT, NVTX and native sampling backtraces over all instances of one bracketed gap class) | -| Runtime flags | n/a (analysis of existing traces) | -| Image | n/a | -| Checkpoint / tokenizer | n/a | -| Hardware / layout | OCI `oci-hsg`, 1×4 GB200, TP1/PP1/EP4/ETP1 (as captured) | -| Workload | Analysis of the BS256/OSL128 decode loop in two traces of the same seven-gate configuration | -| Job / run | session `qwen-host` job 5607600, exec `d920241ea00a41618c63a75fefcb9ea2`, `sessions/qwen-host/prof/hosts6-1785048883` | -| Throughput | n/a. For the record, the host-traced benchmark itself reports 6,631.7 tok/s / TPOT 38.60 ms (vs PROFILE-S6's 10,454.6 tok/s) — but that overhead lands almost entirely **outside** the decode loop (model load, chunked prefill, graph capture); the steady-state decode step is only 2.1% slower under host tracing. | -| Latency / TPOT | Steady-state decode step: **9.134 ms** (GPU-only) vs **9.330 ms** (host trace) = 1.021×. Idle 1965.7 vs 2128.0 µs = 1.083×. | -| Correctness | n/a (analysis) | -| Nsight artifacts | Host: `qwen-host:prof/hosts6-1785048883/mcore_host_profile.{nsys-rep,sqlite}` (615 MB / 5.9 GB). GPU-only reference: `nsys_trace/mcore_s6_tuned_osl128.sqlite`. Analysis: `dev/moe_fused/analyze_hostgaps.py` | - -**Trace provenance and window — two corrections that had to be made first.** -The sqlite is intact (`pragma quick_check` ok) and complete despite the offline -`QdstrmImporter` recovery: all four ranks are present with ~479 k kernels each -spanning 75.0–293.0 s, four engine PIDs with ~665 k CUDA-runtime events each, -and OSRT/sampling covering 0–300 s. Two things were *not* as expected: - -1. **The densest kernel region is not the decode loop.** Seconds 196–241 look - like a steady loop (a `_fused_metadata_kernel` every ~146 ms, 48 layers per - period) but it is **CUDA-graph capture**: 25 distinct capture streams, no - BS256 LM-head GEMM, `index_elementwise_kernel` count **zero**, and - `_multimem_all_gatherv_3tensor_kernel` averaging **1807 µs** (vs 15.8 µs in - the GPU-only trace) because the symmetric-memory barrier absorbs inter-rank - skew while the other ranks are capturing. The real BS256/OSL128 decode loop - is a 1.3 s burst at **291.67–292.98 s** (146 anchor intervals, period 9.064 - ms), preceded by the BS8/OSL32 warmup at 287.51–287.84 s (44 steps, 7.53 ms). - Analysing 217–241 s would have produced a 16×-inflated fiction. -2. **`PYTHON_SAMPLING_*` and the graph tables are absent** despite - `--python-sampling=true`. Attribution therefore uses NVTX ranges + - `CUPTI_ACTIVITY_KIND_RUNTIME` + `OSRT_API` + native `SAMPLING_CALLCHAINS`, - which turned out to be sufficient because the engine emits per-phase NVTX. - -**Method.** Kernel-union complement per anchor-to-anchor step, anchored on the -once-per-step LM-head GEMM `nvjet_sm100_tst_512x64_64x3_2x1_2cta_v_bz_TNT` on -all four devices in both traces (forcing the same anchor matters: auto-detection -picked `vectorized_elementwise_kernel` on devices 0/2 of the host trace, which -straddles G1 and silently dropped it). Each gap is bracketed by both neighbours. -Windows: GPU-only 184.9–186.9 s (126 steps), host 291.75–292.95 s (117 steps). -The GPU-only trace is the source of truth for magnitude, the host trace for -attribution. - -**Per-step gap budget (GPU-only trace, µs/step, per rank).** - -| dev | step ms | busy ms | idle µs | G1 | G2A | G2B | G3 | G4 | -|---|---:|---:|---:|---:|---:|---:|---:|---:| -| 0 | 9.135 | 7.195 | 1940.2 | 559.0 | 63.7 | 259.1 | 303.8 | 48.8 | -| 1 | 9.135 | 7.181 | 1953.7 | 566.8 | 62.4 | 264.0 | 305.7 | 49.9 | -| 2 | 9.134 | 7.141 | 1993.8 | 582.7 | 63.1 | 275.8 | 308.6 | 54.1 | -| 3 | 9.133 | 7.158 | 1975.0 | 576.9 | 63.1 | 273.0 | 303.6 | 52.7 | -| **mean** | **9.134** | **7.169** | **1965.7** | **571.4** | **63.1** | **267.9** | **305.4** | **51.4** | - -Every rank is within 4% of every other on every row — **no rank asymmetry**, as -expected when the four EP ranks are barrier-coupled by the multimem collectives. -The five bracketed gaps are **1259.3 µs/step = 13.79% of the step and 64.1% of -the idle**; the residual 706.4 µs is 1158 × ~0.61 µs of intra-graph node -dispatch. Host-trace equivalents (mean of ranks) are G1 641.2, G2 378.5, G3 -308.2 µs — inflation 1.12× / 1.14× / 1.01×, so the host trace neither invents -nor hides these gaps. - -**Gap 1 — `index_elementwise_kernel` → `vectorized_elementwise_kernel`.** -571.4 µs/step GPU-only (host: median 635.4, mean 640.9, p10 585.0, p90 676.1, -116 instances, 0.99/step). Verdict **(i) CPU-bound Python**. NVTX covers 72.5%: -`update_requests` 183.6 µs (28.6%), `initialize_attention_state` 151.7 (23.7%), -`active_request_mask` 78.0 (12.2%), `transfer_samples_to_cpu` 41.2 (6.4%), -`sampling` tail 10.5 (1.6%); the remaining 27.5% is an unlabelled engine window -between `update_requests` and `initialize_attention_state` (median 152.2 µs) -that is itself 39.6%-leaf `_PyEval_EvalFrameDefault` with `at::_ops::_to_copy` -in 34% of stacks. Only **55.1 µs (8.6%)** of the gap is inside any CUDA API -(`cudaMemcpyAsync` ×2 = 41.7 µs, `cudaLaunchKernel` 8.6, `cudaStreamSynchronize` -**3.4 µs = 0.5%**), and only 32.9 µs (5.1%) inside any syscall on the engine -thread — so ~86% is plain userspace execution. All 161 samples landing in these -windows have `threadState = Running`; of the 142 on the engine thread, 100% have -`_PyEval_EvalFrameDefault` and `_PyObject_MakeTpCall` somewhere in the stack, -with `at::native::copy_` in 12.7%. Source: -`text_generation_controller.py:1750` (`transfer_samples_to_cpu`, D2H at `:1723`), -`:1756`–`1805` (`active_request_mask`), `:1807`–`1811` (`update_requests` → -`dynamic_context.py:3515`), `:644`–`649` (`initialize_attention_state` → -`dynamic_context.py:2128`, whose own comment at `:643` reads "100% CPU -computation"). - -**Gap 2 — `vectorized_elementwise_kernel` → `vectorized_gather_kernel`, twice -per step, bimodal.** Instance **2B** (267.9 µs GPU-only; host median 301.1, mean -310.4) is NVTX-covered 97.7% by `initialize_attention_state` 159.9 µs (51.5%), -`forward_pass` head 99.0 (31.9%) and `transfer_bookkeeping_to_gpu` 44.3 (14.3%); -CUDA API is 35.9 µs (11.5%), engine-thread syscalls 7.0 µs (2.3%), and all 76 -samples are Running with 100% `_PyEval_EvalFrameDefault` — verdict **(i) -CPU-bound Python**, the same bookkeeping chain as G1 continued past the first -GPU op it emits. Instance **2A** (63.1 µs GPU-only; host median 70.8) sits -entirely inside the `sampling` NVTX range and is 34.4% CUDA API -(`cudaLaunchKernel` 15.6 µs = 21.8%, `cudaStreamSynchronize` 8.7 µs = 12.1%), -with `at::_ops::index_Tensor::call` / `at::native::index_kernel` in 55.6%/44.4% -of its (only 9) samples — verdict **(i)+(iii)**, advanced-indexing dispatch plus -a short blocking stream sync. Source: `text_generation_controller.py:652`–`654` -(`transfer_bookkeeping_to_gpu` → `dynamic_context.py:2411`), `:790`/`:1864` -(`forward_pass`), `:1902`/`:2018`–`2023` (`sampling`; `torch.argmax` at `:2019`, -`.cpu()` at `:2022`). - -**Gap 3 — `CatArrayBatchedCopy_vectorized` → `rmsnorm_fwd_tuned_kernel` (the -next graph replay).** 305.4 µs/step GPU-only (host median 311.5, mean 314.3, 117 -instances). Verdict **(iii) driver graph-launch latency**. It is 100% inside -`forward_pass`, and **63.5% of it is a single `cudaGraphLaunch` call**: 132 calls -in the window, median **199.1 µs**, p10 188.6, p90 211.9, unimodal (so it is not -a graph-switch/upload artifact — one graph id, 140, carries 135,528 of the -window's kernels). Plus `cudaMemcpyAsync` 18.4 µs (5.8%) and ~96 µs of Python. -Engine-thread syscalls are 0.7%; all 78 samples are Running, 56.6% with -`cuGraphLaunch` as the deepest resolved frame and 57.9% with -`at::cuda::CUDAGraph::replay()` in the stack. At 1158 nodes per replay that is -**0.172 µs of host submit time per graph node**, on the critical path before the -first node executes. - -**Aggregate, and the profiling-overhead correction.** GPU-busy is 7.169 ms = -78.5% of the 9.134 ms step; host gap is 1965.7 µs = **21.5%**. Splitting that -idle by confidence: - -| Component | µs/step | % of step | Confidence | -|---|---:|---:|---| -| G1 + G2A + G2B + G4 (host Python + sampling dispatch) | 953.8 | 10.4% | Real — samples Running in the interpreter | -| G3 minus its `cudaGraphLaunch` (Python + H2D in `forward_pass`) | 111.4 | 1.2% | Real | -| `cudaGraphLaunch` inside G3 | 194.0 | 2.1% | **Artifact-suspect** | -| Intra-graph inter-node dispatch (1158 × ~0.61 µs) | 706.4 | 7.7% | **Artifact-suspect** | -| **Total idle** | **1965.7** | **21.5%** | | - -Overhead was corrected by comparing the two traces step-for-step rather than by -assuming a factor: the host capture adds only **2.1%** to the decode step -(9.134 → 9.330 ms) and 8.3% to the idle, even though it costs 37% of -whole-benchmark throughput, because its cost falls on load/prefill/capture. The -one caveat that does **not** cancel is that *both* traces used -`--cuda-graph-trace=node`, which forces CUPTI to instrument every graph node; -the 194 µs `cudaGraphLaunch` and the 706 µs of inter-node dispatch — 900 µs/step, -9.9% of the step, 45.8% of the idle — are exactly the quantities that flag -distorts, and this analysis cannot separate them from the real cost. A -`--cuda-graph-trace=graph` control capture settles it for the price of one -profile run. - -**What the recoverable upside is worth.** Taking the current best mcore of -25,495.9 tok/s (QWEN-022, 75.00% of vLLM's 33,994.5) and scaling by -`9134 / (9134 − Δ)` — an upper bound that assumes perfect removal and that the -OSL128 composition transfers to the OSL1024 throughput regime: - -| Removed | Δ µs | ceiling tok/s | % of vLLM | Share of the remaining gap | -|---|---:|---:|---:|---:| -| G1 | 571.4 | 27,197 | 80.0% | 20.0% | -| G2 (both) | 331.0 | 26,454 | 77.8% | 11.3% | -| G3 | 305.4 | 26,378 | 77.6% | 10.4% | -| G1 + G2B (the serial bookkeeping chain) | 839.3 | 28,067 | 82.6% | 30.3% | -| All five bracketed gaps | 1259.3 | 29,573 | 87.0% | 48.0% | -| All idle | 1965.7 | 32,488 | 95.6% | 82.3% | - -**The finding that reframes the lever.** G1 and G2B are not a scheduling -artifact that overlap could hide. The chain is -`graph(N) → logits(N) → argmax(N) → D2H tokens(N) → active_request_mask(N) → -update_requests(N) → initialize_attention_state(N+1) → H2D → graph(N+1)`, and -every link needs the previous one's result: the host cannot run step N+1's -bookkeeping during step N's graph replay because that bookkeeping consumes step -N's sampled tokens. This is a genuine serial data dependency, and it explains -QWEN-007 cleanly — `async-sched-mode=serial` measured +0.85% not because the -idle was illusory but because with a single request stream there is nothing to -overlap it *with*. The lever is therefore to make the CPU work cheaper or move -it onto the GPU, not to hide it. - -**Ranked candidate next levers** (upside is a measured ceiling, not a prediction; -none of these has been measured): - -1. **Make `initialize_attention_state` incremental** — ~311 µs/step (152 µs in - G1 + 160 µs in G2B), pure CPU by construction. In steady-state decode the - request set is fixed for 128 steps and only the sequence lengths increment, - so most of the per-step attention metadata is recomputed unchanged. Ceiling - ~3.4% (→ ~77.6% of vLLM); realistic maybe half. **Risk low-medium**: CPU-only, - no numerics exposure, but correctness depends on invalidating the cache on - every request add/pause/finish. Falsified if a fresh profile shows the G1 and - G2B windows unchanged. -2. **Cut `update_requests` + `active_request_mask` + the unlabelled engine - window** — 183.6 + 78.0 + ~175 = ~437 µs/step of Python over 256 requests. - Ceiling 4.8% (→ ~78.8%). **Risk low**: no numerics, no guards, purely - additive with everything else. -3. **`--cuda-graph-trace=graph` control capture** — no throughput at all, but it - decides whether the 900 µs/step (9.9% of the step, 45.8% of the idle) of - graph machinery is real or instrumentation. **Risk none, cost one profile - run.** This gates lever 4 and should be run before it. -4. **Reduce graph node count** — each removed node is worth ~0.172 µs of host - submit plus ~0.61 µs of GPU inter-node dispatch ≈ 0.78 µs, so 100 nodes ≈ - 78 µs ≈ 0.85%. **Risk medium** and sharply diminishing; this is the fusion - campaign that QWEN-019…022 already mined. -5. **Move per-step request bookkeeping onto the GPU** (persistent device-side - batch state, no D2H→Python→H2D round trip per step), which is what vLLM does. - Ceiling is most of the 839 µs chain, ~9.2%. **Risk high** — a large engine - rewrite, multi-session. - -Explicitly *not* recommended: re-litigating `async-sched-mode=serial`. The -dependency analysis above says the idle it targets is not overlappable with a -single request stream, which is consistent with its measured +0.85%. - -| Field | Value | -|---|---| -| Result | **Supported — all three gaps attributed with numbers, and one taxonomy verdict each.** G1 (571 µs) and G2B (268 µs) are CPU-bound Python in the engine's per-step bookkeeping chain; G2A (63 µs) is sampling-path launch plus a short stream sync; G3 (305 µs) is 63.5% a single 199 µs `cudaGraphLaunch`. No rank asymmetry. The idle budget is 1065 µs of real host work and 900 µs of graph machinery whose reality is gated on a `--cuda-graph-trace=graph` control. The campaign-level conclusion is that the largest remaining lever is a serial data dependency, so it must be made cheaper rather than hidden. | -| Next action | Run the `--cuda-graph-trace=graph` control capture (lever 3, one profile run) to fix the denominator, then implement lever 1 (incremental `initialize_attention_state`) and measure it e2e at OSL1024 against CLEANBASE-S6 with the full seven-gate set. | - -### Session 6 — conclusion & recommendation - -The routing/permute lever is closed, and it delivered. Against the same-session -reference CLEANBASE-S6 (23,264.4 tok/s), the four accepted changes take mcore to -**25,495.9 tok/s = 75.00%** of the fixed vLLM OSL1024 baseline of 33,994.5, a -**+9.6%** session gain — and every one of them is bit-exact, so none of it was -bought with numerics. - -| Change | Mechanism | e2e | -|---|---|---:| -| QWEN-019 `MCORE_MOE_FUSED_COUNT` | count + zero-fill folded into the table build (4 launches → 2) | +3.01% | -| QWEN-020 `MCORE_MOE_SUM_FAST` | topk-reduction locality test predicated instead of branched | ~+1% (measured +1.83%, partly variance) | -| QWEN-021 `MCORE_ROUTER_FUSED_TOPK` | softmax + top-8 select in one CTA-per-token kernel (4 launches → 1) | **+3.89%** | -| QWEN-022 `MCORE_MOE_FUSED_SCATTER` | scatter folded in as a second streaming pass (2 launches → 1) | +0.46% | - -Reproduce the best configuration with all seven gates: -`MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 -MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 -MCORE_MOE_FUSED_SCATTER=1`. - -Two methodological points worth carrying forward. First, QWEN-018's gate was -right to reject the two candidates the ledger had ranked first and second on -arithmetic alone — the win came from the third candidate the *measurement* -surfaced, and from the router pair the gate flagged but did not rank. Second, -microbench-to-e2e transfer was reliable for launch removals (QWEN-019 predicted -2.93% and got 3.01%; QWEN-021 predicted ~3.1% and got 3.89%) and unreliable for -in-kernel rewrites (QWEN-020 and QWEN-022 both landed at roughly a third to -double their predictions), so keep repeating e2e runs whose predicted effect is -under ~1%. - -PROFILE-S6 re-ranks the step: routing is now 4.7% and the largest addressable -item is **1237 µs/step of host gaps between graph replays** (21.0% idle overall). -That is a host-side problem and needs a host-side profile before any code -change. - -Append records using this exact structure: - -```markdown -### QWEN-NNN — short name - -| Field | Value | -|---|---| -| Date | YYYY-MM-DD | -| Hypothesis | One measurable claim | -| Code revision | Commit SHA and clean/dirty state | -| Changed files | Exact paths, or `none` for baseline | -| Runtime flags | Exact non-default flags | -| Image | Immutable image path/tag | -| Checkpoint / tokenizer | Exact paths | -| Hardware / layout | Cluster, GPUs, TP/PP/EP/ETP/DP | -| Workload | Dataset, batch, OSL, warmups, timed iterations | -| Job / run | Slurm job ID and run directory | -| Throughput | tokens/s | -| Latency / TPOT | ms / ms-token | -| Correctness | Prompt outputs and benchmark status | -| Nsight artifacts | `.nsys-rep`, `.sqlite`, analysis output | -| Result | Supported / rejected / inconclusive | -| Next action | One prioritized follow-up | -``` - -## Session 7 (2026-07-26) — host-side bookkeeping - -### CLEANBASE-S7 — same-session seven-gate reference - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | Establish a same-session OSL1024 reference before measuring a sub-1% host-side lever | -| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty with the session-6 change set | -| Changed files | none | -| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1` | -| Image | `agents-space/images/ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | `agents-space/checkpoints/qwen3-30b-a3b-mcore` / `-hf` | -| Hardware / layout | OCI `oci-hsg`, 1 node 4×GB200, TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed | -| Job / run | job 5613090, `sessions/qwen-attnstate/e2e/ref-s7-1785084185` | -| Throughput | **25,805.9 tok/s** — per-iter 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 | -| Latency / TPOT | 9,867.5 ms / 9.920 ms/tok | -| Correctness | Coherent on all three temperature-0 prompts; benchmark 5/5 | -| Nsight artifacts | none (throughput run) | -| Result | Clean reference at **75.91%** of vLLM 33,994.5. Spread across the five iterations is 0.23%, so a 0.5% effect is resolvable in this session | -| Next action | A/B the incremental `initialize_attention_state` lever against this reference | - -### QWEN-023 — incremental `initialize_attention_state` - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | The ~311 µs/step HOSTGAP-S6 attributed to `initialize_attention_state` is recomputation of step-invariant metadata; caching it behind a request-layout version counter removes most of it without changing a single generated token | -| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty (session-6 set plus this change) | -| Changed files | `megatron/core/inference/contexts/dynamic_context.py`, `megatron/core/inference/contexts/attention_context/mha_metadata.py`, `dev/moe_fused/harness_attnstate.py` (new) | -| Runtime flags | seven gates as CLEANBASE-S7, plus `MCORE_INFER_INCR_ATTN_STATE=1` | -| Image / checkpoint / hardware / workload | identical to CLEANBASE-S7 | -| Job / run | job 5613090, `e2e/incr-attn-on-1785085393` and `e2e/incr-attn-on-rep2-*`; host microbenchmark exec `a5d20392e75848b9855366cd7844c850` | -| Throughput | **26,032.1** (per-iter 26,052.6 / 26,018.7 / 26,075.7 / 25,992.1 / 26,021.4) and **26,092.5** on an independent repeat pair | -| Latency / TPOT | 9,780.4 ms / 9.834 ms/tok; repeat 9.811 ms/tok (−92 µs/step vs reference on two-run means) | -| Correctness | 130 consecutive decode steps bit-identical under `MCORE_INFER_INCR_ATTN_STATE_VERIFY=1`; temperature-0 coherence output byte-identical to gate-OFF | -| Nsight artifacts | none needed — the effect is host wall time, measured directly | -| Result | **Accepted, +0.94%** (two OFF/ON pairs, non-overlapping per-iteration distributions) | -| Next action | Cut `update_requests` + `active_request_mask` + the unlabelled engine window (~437 µs/step, HOSTGAP-S6 lever 2) with the same version-counter machinery | - -**Pre-optimization CPU breakdown.** Measured with `MCORE_INFER_ATTN_PROF=1` -(monotonic-ns marks around each phase) and with a standalone host harness -(`dev/moe_fused/harness_attnstate.py`) that drives a real -`DynamicInferenceContext` at BS256 through 400 steady-state decode steps with no -model loaded. The two agree to within 12%. Per call, 256 active requests: - -| Phase | µs/call | What it is | -|---|---:|---| -| `slices` | 73.9 | build + pad the active-request slices, sampling-metadata copies, logit indices | -| `mhameta` | 32.2 | KV/query lengths, both cumsums, block table, `set_state_data` | -| `xfer_inner` | 26.9 | H2D copy of the bookkeeping buffer | -| `graphmatch` | 12.8 | batch dimensions + `match_graph_config` | -| `tokenpad` | 12.0 | stamp the token padding slots | -| `maxlen` | 12.0 | max-seqlen scalars | -| `pre` + `paddims` + `tail` | 1.8 | pending Mamba ops, padded dimensions, epilogue | -| **TOTAL** | **171.5** | | - -The cost is **concentrated, not thin**: `slices` alone is 43%, and inside it two -statements — `build_active_slices` 29.3 µs and `pad_active_slices` 26.7 µs — -are a third of the whole call. Every one of those top statements recomputes a -value that is provably identical to the previous decode step. That is what made -the lever viable; had the profile come back flat across fifty 3 µs statements -the design below would not have been worth building. - -**Design.** A single monotonic `_request_layout_version` counter is bumped by -every path that changes *which* request occupies a slot, a request's sampling -metadata, or its KV block table. `initialize_attention_state` takes the fast -path only when the current cache key equals the stored one, where the key is -`(layout_version, total_request_count, paused_request_count, active_token_count, -kv_block_allocator.total_avail, chunked_prefill_request_id, -num_speculative_tokens)`. The version is the guard; the other six are -independent structural sentinels, so a missed bump would additionally have to -coincide with an unchanged request count, token count and block-allocator -occupancy to escape detection. The fast path then recomputes only the KV -sequence lengths and their cumsum — statements copied character-for-character -from the full path so the buffers stay bit-identical — re-stamps the token -padding sentinels, rebinds `state_data` through a new -`MHAMetadata.restore_state_data` (needed because `reset_attention_state` clears -the max-seqlen scalars every step), and issues the H2D transfer. Everything -else is reused. - -**Invalidation audit.** Thirteen mutation paths, each bumping the version: - -| Path | Event covered | -|---|---| -| `add_request` | request add, including a chunked-prefill chunk | -| `resume_paused_requests` (inside `if resume_request_count > 0`) | resume | -| `_swap_book_keeping_tensors` | pause/resume slot swap | -| `_move_book_keeping_tensors` | slot compaction | -| `resolve_requests` | end-of-step finish/retire | -| `release_memory_blocks_from_request_indexes` | evict, block release | -| `prepare_requests` (inside `if num_new_blocks > 0`) | new KV block allocated mid-decode | -| `add_dummy_requests_parallel` (inside `if requests`) | dummy padding requests | -| `add_dummy_requests_for_cudagraph_capture` | graph capture padding | -| `add_dummy_requests_for_expert_parallel_step` | EP dummy step | -| `initialize_all_tensors` | tensor state (re)allocation | -| `reset_tensors`, `reset_metadata` | context reset | - -Advancing a request's KV length is deliberately *not* a bump — that is the one -thing the fast path recomputes. Four conditions bail out to the full path -outright rather than relying on the key: `construct_graph_dimensions is not -None` (any graph-dimension change), `is_expert_parallel_dummy_cuda_graph_step`, -`is_hybrid_model` (Mamba state), and `num_prefill_requests != 0`. The cache is -additionally never *stored* while `is_creating_cuda_graphs` is set or when the -step did not use a CUDA graph. The class-scope declaration of every cache -attribute means any construction path starts cold and invalid. - -**Correctness evidence.** `MCORE_INFER_INCR_ATTN_STATE_VERIFY=1` runs the fast -path, snapshots everything it produced, then recomputes the whole thing from -scratch and asserts equality of: both MHA query-length buffers and cumsums, both -KV-length buffers and cumsums, the block table, `active_request_last_token_idxs`, -`active_logit_idxs`, all seven `active_request_metadata` tensors, the entire GPU -bookkeeping buffer, the padded token/request counts, the padded batch -dimensions, the graph-selection flag, both max-seqlen scalars, and the -`(address, shape, dtype)` binding of every `state_data` view. **130 consecutive -decode steps passed with zero mismatches.** End to end, the three temperature-0 -coherence completions are byte-identical between gate ON and gate OFF. - -**Performance.** Host CPU per call **174.1 → 51.7 µs, a 3.37× reduction**; the -coarse breakdown collapses to `graphmatch=0.2 slices=1.1 tokenpad=0.1 -mhameta=0.4 maxlen=0.1`, with the 48.2 µs residual being the KV-length -recompute plus the H2D bookkeeping copy, which cannot be cached. E2e: - -| Run | Gate | Throughput | Per-iter | -|---|---|---:|---| -| `ref-s7` | OFF | 25,805.9 | 25,782.5 / 25,783.6 / 25,814.8 / 25,842.3 / 25,806.7 | -| `incr-attn-on` | ON | **26,032.1** | 26,052.6 / 26,018.7 / 26,075.7 / 25,992.1 / 26,021.4 | -| `ref-s7-rep2` | OFF | 25,831.7 | 25,710.8 / 25,893.9 / 25,863.9 / 25,875.4 / 25,815.4 | -| `incr-attn-on-rep2` | ON | **26,092.5** | 26,157.8 / 26,085.1 / 26,096.3 / 26,106.8 / 26,017.0 | - -Pairwise the lever is +0.88% and +1.01%; on the two-run means (OFF 25,818.8, ON -26,062.3) it is **+0.94%**. Across all twenty timed iterations the slowest ON -iteration (25,992.1) is still faster than the fastest OFF iteration (25,893.9), -so the two distributions do not overlap at all and the sign is certain despite -the effect being under 1%. TPOT 9.915 → 9.823 ms/tok, i.e. **−92 µs/step**. -New best = **76.67%** of vLLM 33,994.5. - -**Why it delivered 0.94% and not the 3.4% ceiling.** The harness measures one -call at 174.1 µs and the fast path removes 122.4 µs of it, but the per-step -saving observed in TPOT is 92 µs — about three quarters of a single call's -worth, against HOSTGAP-S6's 311 µs/step spread over two windows. So the ceiling -was overstated for two compounding reasons: the H2D bookkeeping transfer is -~27 µs of every call and is not cacheable (the GPU buffer must be rewritten -every step), and the second of the two per-step windows is evidently not all -`initialize_attention_state` doing cacheable work. The honest read is that the -remaining host-side upside lives in `update_requests` and the engine window, -not here. - -### CGTRACE-CONTROL — is the graph machinery real? - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | The ~900 µs/step HOSTGAP-S6 attributed to CUDA-graph machinery is an artifact of `--cuda-graph-trace=node`; capturing with `=graph` will shrink it | -| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty | -| Changed files | `dev/moe_fused/profile_insession.sh` (`CUDA_GRAPH_TRACE` env, bounded nsys stop + qdstrm recovery), `dev/moe_fused/analyze_cgtrace.py` (new) | -| Runtime flags | seven gates plus `MCORE_INFER_INCR_ATTN_STATE=1`, identical in both captures | -| Workload | gsm8k, BS256, OSL128, 1 warmup request + 1 timed request | -| Job / run | job 5613090, `prof/cgt-node-1785086713`, `prof/cgt-graph-1785086964` | -| Nsight artifacts | `mcore_profile.{nsys-rep,sqlite}` in both run dirs; analysis exec `d6f572f2613645cdbb05226dbc8e5e98` | -| Result | **Hypothesis rejected — the machinery is real** | -| Next action | Treat HOSTGAP-S6's ranking as sound; the next lever is `update_requests` (lever 2), not graph-node reduction (lever 4), on measured size | - -| Metric (same rank, steady decode window) | `=node` | `=graph` | Δ | -|---|---:|---:|---:| -| Step period | 8882.5 µs | 8945.1 µs | −0.71% | -| Host `cudaGraphLaunch` (median) | 190.0 µs | 184.7 µs | +2.8% | -| Traced kernels per step | 1169 | 54 (graph interior hidden) | — | -| Steady steps analyzed | 91 | 88 | — | - -The step period is the number that matters, and it does not move: if CUPTI's -per-node instrumentation were adding ~900 µs to a 9 ms step, removing it would -have shortened the step by 10%, and instead the graph-mode capture came out -0.71% *slower*, which is ordinary run-to-run variance. The 190 µs -`cudaGraphLaunch` survives the control almost unchanged, so it is genuine -driver-side submit cost for a 1158-node graph. Every wall-time attribution in -PROFILE-S6 and HOSTGAP-S6 therefore stands as measured. - -One limit worth recording: `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` in the graph-mode -capture holds only 417 rows for the whole run (6 inside the analyzed window), -so the GPU-side "inter-node dispatch" half of the 900 µs cannot be -independently confirmed from this control. The step period and the host submit -cost are settled; that component is not. - -### CLEANBASE-S8 — same-session eight-gate reference - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | Establish a same-session OSL1024 reference before measuring the second host-side lever | -| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty with the session-7 change set | -| Changed files | none | -| Runtime flags | `MCORE_FUSE_FC1_ACT=1 MCORE_MOE_FUSED_ALIGN=1 MCORE_MOE_GEMM_TUNE=1 MCORE_MOE_FUSED_COUNT=1 MCORE_MOE_SUM_FAST=1 MCORE_ROUTER_FUSED_TOPK=1 MCORE_MOE_FUSED_SCATTER=1 MCORE_INFER_INCR_ATTN_STATE=1` | -| Image | `agents-space/images/ceecf5c304a5d8bd.sqsh` | -| Checkpoint / tokenizer | `agents-space/checkpoints/qwen3-30b-a3b-mcore` / `-hf` | -| Hardware / layout | OCI `oci-hsg`, 1 node 4×GB200 (`nvl72039-T01`), TP1/PP1/EP4/ETP1 | -| Workload | gsm8k, BS256, OSL1024, 2 warmup + 5 timed | -| Job / run | job 5616264, `sessions/qwen-updreq/e2e/ref-s8-1785099726` (and `ref-s8-rep2-1785101268`) | -| Throughput | **25,944.9 tok/s** — per-iter 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7. Repeat: **25,983.9** — per-iter 26,021.5 / 25,976.0 / 26,037.9 / 25,983.6 / 25,901.2 | -| Latency / TPOT | 9,820.6 ms / 9.867 ms/tok; repeat 9.852 ms/tok | -| Correctness | Coherent on all three temperature-0 prompts; benchmark 5/5 | -| Nsight artifacts | none (throughput run) | -| Result | Clean reference at **76.32%** of vLLM 33,994.5 (two-run mean 25,964.4 = 76.38%) | -| Next action | A/B the reduced-op `update_requests` lever against this reference | - -The first run's iteration 1 is a 1.3% cold outlier; iterations 2–5 span 0.23%. -The repeat run does not show it and spans 0.53% across all five. Two-run mean -25,964.4 tok/s is the number the lever below is judged against. - -### QWEN-024 — reduced-op post-sampling bookkeeping - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | HOSTGAP-S6 lever 2. The ~437 µs/step of host work after the sampled tokens return is three different things, not one: measure `update_requests`, `active_request_mask` and the unlabelled engine window separately, name the window, and only then decide whether any of it is attackable | -| Code revision | `37d1adcf0`, branch `perf/moe-fused-align`, dirty (session-7 set plus this change) | -| Changed files | `megatron/core/inference/contexts/dynamic_context.py` (`_write_decode_token_bookkeeping_fast`, `_write_token_bookkeeping_reference`, `_verify_decode_token_bookkeeping`, `_decode_req_idx_arange` cache, two count reductions, env `MCORE_INFER_VEC_UPDATE_REQS`, `MCORE_INFER_VEC_UPDATE_REQS_VERIFY`), `megatron/core/inference/text_generation_controllers/text_generation_controller.py` (`_empty_finished_idxs`, `_empty_finished_ids`, no-finisher short-circuit), `dev/moe_fused/harness_updreq.py` (new) | -| Runtime flags | eight gates as CLEANBASE-S8, plus `MCORE_INFER_VEC_UPDATE_REQS=1` | -| Image / checkpoint / hardware / workload | identical to CLEANBASE-S8 | -| Job / run | job 5616264, `e2e/vecupd-on-1785100576` and `e2e/vecupd-on-rep2-*`; harness execs `eaca22426f654bfb938bd68c13c7bff8` (gate OFF breakdown), `94604a2f79824182bef77456c9a2b7ee` (equivalence, verify, gate ON) | -| Throughput | **26,128.2** and **26,270.3 tok/s** on two independent OFF/ON pairs | -| Latency / TPOT | 9,745.4 ms / 9.798 ms/tok; repeat 9,691.7 ms / 9.745 ms/tok (−88 µs/step vs the reference on two-run means) | -| Correctness | 400-step two-context equivalence with 30 mid-batch terminations and a KV block-boundary crossing; 300 consecutive steps under `MCORE_INFER_VEC_UPDATE_REQS_VERIFY=1`; temperature-0 coherence output byte-identical to gate-OFF | -| Nsight artifacts | OSL128 A/B pair, `prof/g1-off-1785102888` and `prof/g1-on-1785103157`, `mcore_profile.{nsys-rep,sqlite}`; analysis exec `20bc3a28e2344c9b8f51aae8c90f6982`; window finder `dev/moe_fused/find_decode_window.py` (new) | -| Result | **Accepted, +0.90%** (two OFF/ON pairs: +0.71% and +1.10%), and the G1 window shrank by the predicted amount | -| Next action | The remaining host-side upside is `post_process_requests`, which this record shows is not of the same kind; see the ranked next levers below | - -**Pre-optimization CPU breakdown — and what the unlabelled window actually is.** -Measured with `dev/moe_fused/harness_updreq.py`, which drives a real -`DynamicInferenceContext` at BS256 through a synthetic steady-state decode loop -and calls the real `DynamicInferenceEngine.post_process_requests` against a stub -engine holding 256 real `DynamicInferenceRequest` objects. No model is loaded, -so a full breakdown takes about 90 seconds. Per step, 256 active requests: - -| Region | µs/step | HOSTGAP-S6 | Nature | -|---|---:|---:|---| -| `active_request_mask` | 46.9 | 78.0 | ~11 whole-tensor CPU ops on 256 elements | -| `update_requests` | 137.6 | 183.6 | ~22 whole-tensor CPU ops on 256 elements | -| `post_process_requests` | 187.9 | ~152.2 (unlabelled) | per-request Python object churn | -| **TOTAL** | **372.4** | **~413.8** | | - -The harness runs 10–20% cheaper than the in-server measurement across all three -regions, which is the expected direction (no server threads, no allocator -pressure, no profiler), and the three ratios agree to within 15% of each other. - -**The unlabelled engine window is `DynamicInferenceEngine.post_process_requests`, -reached through `async_bookkeep`.** It appeared unlabelled in HOSTGAP-S6 because -the controller emits NVTX through `torch.cuda.nvtx.range_push` directly, which is -always live, while the engine emits through `megatron.core.utils.nvtx_range_push`, -which is inert unless `_nvtx_enabled` is set — so the engine's own `bookkeeping` -and `detokenization` ranges were simply never pushed in that capture. The -measured 187.9 µs/step also matches the window's 152.2 µs median once the -harness discount is applied. Its `cProfile` composition, over 100 steps: - -| Item | calls/step | What it is | -|---|---:|---| -| `post_process_requests` body | 1 | 61% of total time is straight-line interpreter work in the per-request loop | -| `builtins.len` | 1,718.8 | ~6.7 length probes per request per step | -| `get_request` | 256 | `self.requests[id].record[-1]` dict + list indexing | -| `list.append` | 256 | `active_request_ids.append` | -| `builtins.isinstance` | 256 | the `if not isinstance(tokens, list)` scalar wrap | -| `_check_stop_words_for_request_post_append` | 256 | returns immediately (no stop words configured) | -| `DynamicInferenceRequestRecord.__getitem__` | 256 | record `[-1]` | -| tensor ops | 3 | three `.tolist()` calls for the whole step | - -**The falsification criterion, honored.** The pre-agreed rule was: vectorize if -the cost is a uniform per-request loop expressible as tensor ops, and stop if it -is per-request Python object churn. The measurement splits the chain cleanly in -half and the two halves fall on opposite sides of that line. - -`post_process_requests` is **churn, and it was left alone**. It is 187.9 µs/step -— the single largest item in the chain — and it does essentially no tensor work: -three `.tolist()` calls for the entire step, against 256 dict lookups, 256 list -appends, 256 method calls and ~1,719 `len()` probes. Nothing here vectorizes. -Making it cheaper means adding a fast path through the request state machine — -token appending, termination, stop words, log-prob routing — which is exactly the -EOS/length handling the task flagged as dangerous, for a benefit that would have -to be re-argued from scratch. It is not touched by this change. - -The other half is **neither** of the two anticipated cases. `update_requests` and -the `active_request_mask` block contain no per-request Python loop at all; they -are already vectorized. Their cost is ~33 whole-tensor operations on 256-element -CPU tensors at 2–16 µs each, and roughly half of that work is *provably dead* in -the decode regime. Three statements in `update_requests` cost 30.0 µs/step -between them and, at `num_speculative_tokens == 0`, compute nothing: - -| Statement | µs/step | Why it is dead at one generated token per request | -|---|---:|---| -| `token_to_pos_ids` write | 16.0 | `repeat_interleave(1)` is identity, and it adds `torch.arange(1).repeat(256)` — a 256-element zero vector | -| `raw_positions` + `crosses_boundary` + `.any()` | 10.1 | builds a `[256, 1]` tensor whose only consumer is an `else` branch that `num_speculative_tokens == 0` makes unreachable | -| `token_to_request_idx` write | 7.3 | `repeat_interleave(1)` is identity over a `torch.arange` that is a pure function of the request-slot bounds | - -So the lever here is **op elimination**, not vectorization. That is a third -category the dichotomy did not anticipate, and it is the low-risk one: it changes -no per-request semantics, touches no termination logic, and every removed -operation is removed because its result is bit-identical to a cheaper one. - -**Design.** One env gate, `MCORE_INFER_VEC_UPDATE_REQS`, default OFF. Note that -the mechanism is deliberately unlike QWEN-023: nothing is cached across steps -except a `torch.arange` that depends only on the request-slot bounds, because -`update_requests` consumes the current step's freshly sampled tokens and there is -no cross-step reuse to exploit. Five changes: - -1. The per-token bookkeeping tail of `update_requests` is factored into - `_write_token_bookkeeping_reference`, kept verbatim, and a new - `_write_decode_token_bookkeeping_fast` used only when - `num_speculative_tokens == 0`. The fast form drops the three dead statements - above, updates `request_last_kv_block_offset` without the defensive clone - (whose only other consumer was `raw_positions`), and reuses a cached - `torch.arange(paused_request_count, total_request_count)` invalidated by a - bounds compare. -2. `finished_request_count` is derived as `numel() - active_request_count` - instead of a second full comparison and reduction. The mask is 0/1 by - construction — it is an `&` of two `.byte()` comparisons in the controller — - so the complement is exact. -3. `active_requests_requiring_new_block` stays a bool tensor instead of being - cast with `.byte()`. Every downstream use (`torch.nonzero`, `== 0`, scalar - assignment, `sum`) is dtype-agnostic. -4. Its population count uses `int(tensor.sum())` rather than `(tensor == 1).sum().item()`. -5. In the controller, when the mask says nothing finished this step, the - `torch.nonzero` scan and the advanced-index gather it feeds are skipped in - favour of cached empty tensors of the same dtype and device. - -**Correctness evidence.** Three independent checks, all on the eight-gate config. - -*Two-context equivalence, 400 steps.* Two `DynamicInferenceContext` objects are -built identically and driven through the same scripted mask sequence, one with -the gate off and one on, comparing after every step: `total_request_count`, -`paused_request_count`, `active_token_count`, nine request bookkeeping tensors -(`request_ids`, `request_kv_length_offsets`, `request_query_lengths`, -`request_output_lengths`, `request_last_kv_block_offset`, -`request_last_kv_block_id`, `request_kv_block_counts`, -`request_to_kv_block_ids`, `request_in_prefill_status_tensor`), all six -`token_to_*` tensors, and the KV block allocator's `total_avail`, -`active_count` and `paused_count`. The script **terminates three requests -mid-batch every 37 steps** — 30 terminations in total, leaving 226 of 256 active -— and runs past the 256-token block boundary so the pause / new-block / resume -branches all execute under both gates. Zero mismatches. This is the check that -directly answers the off-by-one-in-termination risk: the finished-request count, -the finished index set, and the resulting slot compaction are all compared -element-wise on the steps where requests actually finish. - -*Verify mode, 300 steps.* `MCORE_INFER_VEC_UPDATE_REQS_VERIFY=1` runs the fast -tail, snapshots the six token buffers plus `request_last_kv_block_offset` and -`active_token_count`, restores the inputs, runs the reference tail, and asserts -equality. 300 consecutive steps, zero mismatches. - -*End to end.* The three temperature-0 coherence completions are byte-identical -between gate ON and gate OFF. - -**Performance.** Host CPU per step, from the harness at BS256: - -| Region | Gate OFF | Gate ON | Δ | -|---|---:|---:|---:| -| `active_request_mask` | 46.9 | 34.9 | −12.0 | -| `update_requests` | 137.6 | 74.0 | −63.6 (1.86×) | -| `post_process_requests` | 187.9 | 184.7 | unchanged, as intended | -| **chain total** | **372.4** | **293.7** | **−78.7** | - -E2e: - -| Run | Gate | Throughput | Per-iter | -|---|---|---:|---| -| `ref-s8` | OFF | 25,944.9 | 25,669.0 / 25,983.1 / 26,011.2 / 26,023.4 / 26,041.7 | -| `vecupd-on` | ON | **26,128.2** | 25,856.1 / 26,141.3 / 26,233.9 / 26,141.6 / 26,272.3 | -| `ref-s8-rep2` | OFF | 25,983.9 | 26,021.5 / 25,976.0 / 26,037.9 / 25,983.6 / 25,901.2 | -| `vecupd-on-rep2` | ON | **26,270.3** | 26,291.7 / 26,281.1 / 26,279.5 / 26,302.1 / 26,197.4 | - -Pairwise the lever is **+0.71%** and **+1.10%**; on the two-run means (OFF -25,964.4, ON 26,199.3) it is **+0.90%**. Nine of the ten ON iterations are -faster than every one of the ten OFF iterations; the single exception is the -cold first iteration of the first ON run (25,856.1), and both OFF runs show the -same cold-first-iteration pattern. Restricted to iterations 2–5 the two -distributions do not overlap at all (slowest ON 26,141.3 > fastest OFF -26,041.7). TPOT 9.860 → 9.771 ms/tok on two-run means, i.e. **−88 µs/step**. -New best = **77.28%** of vLLM 33,994.5 on the best run, **77.07%** on the -two-run mean. - -**Why it delivered less than the 4.8% ceiling.** The ceiling assumed the whole -437 µs/step chain could vanish. Half of it is `post_process_requests`, which this -change deliberately does not touch, so the reachable half was ~185 µs/step of -which the harness says 78.7 µs was removed — 42% of the reachable half, 18% of -the whole chain. TPOT moved 9.860 → 9.771 ms/tok on two-run means, i.e. -−88 µs/step, against a harness-measured saving of 78.7 µs. The two agree to -within 12%, in the direction the harness discount predicts, which is the -strongest evidence that the host saving is what produced the throughput change -rather than session drift. - -**Profile confirmation — HOSTGAP-S6's falsification test for this lever class.** -HOSTGAP-S6 said a lever of this kind is falsified if a fresh profile shows the G1 -window unchanged. Two BS256/OSL128 captures were taken back to back in the same -session at the same eight-gate code, differing only in the gate, and analysed -with `analyze_hostgaps.py gaps` on the forced LM-head anchor -`nvjet_sm100_tst_512x64_64x3_2x1_2cta_v_bz_TNT`. The G1 bracket is now -`index_elementwise_kernel -> vectorized_gather_kernel` rather than -`-> vectorized_elementwise_kernel`, because QWEN-023 changed which kernel the -bookkeeping chain emits first; it is the same gap class (largest gap, ~1.0 -instances per step). - -| Metric (mean of devices 0 and 3) | Gate OFF | Gate ON | Δ | -|---|---:|---:|---:| -| G1 window | 640.6 µs | 562.3 µs | **−78.3 µs** | -| Step period | 8.9425 ms | 8.8530 ms | −89.5 µs (−1.00%) | -| Total idle | 1757.0 µs | 1679.2 µs | −77.8 µs | -| GPU-busy | 7.2015 ms | 7.2045 ms | +3.0 µs (noise) | -| Kernels per step | 1169 | 1169 | 0 | -| G3 (`CatArrayBatchedCopy` → `rmsnorm`) | 295.6 / 293.0 | 296.8 / 293.2 | unchanged | -| G2A (`vectorized_elementwise` → `vectorized_gather`) | 67.1 / 60.4 | 65.9 / 63.6 | unchanged | - -Four independent measurements agree: the harness says 78.7 µs of host work was -removed, the G1 window shrank 78.3 µs, the step period shortened 89.5 µs and -TPOT shortened 88 µs. GPU-busy time and the per-step kernel count are identical, -and the two gaps this change does not touch did not move, so the saving is -host-side and localized exactly where it was designed to be. Not falsified. - -**Ranked next levers, updated.** (Lever 1 below was then implemented in the same -session and accepted as QWEN-025; the decision-gate measurement that justified -starting it is recorded here.) - -1. **`post_process_requests` fast path** — 187.9 µs/step harness, ~152 µs/step - in-server, now the largest single host item in the decode loop. A fast path - for the common decode case (no speculative tokens, no stop words, no log - probs, not chunked prefill, request not finishing) could collapse the loop - body to a few operations. Ceiling ~1.7% e2e. **Risk high** — it is the request - termination state machine. Falsified if a `cProfile` of the fast path does not - cut the per-step call count below ~500. - - **Decision-gate measurement, taken with the time left in this session** - (harness exec `4c0a584d22704d4ab1d8c5a7223bc8c2`, gate ON, BS256, 200 - timed calls against the real engine method and against an in-harness - prototype of the decode fast path). The prototype implements the common - decode case only and removes the per-request dict lookup, the record `[-1]` - indexing, the `isinstance` scalar wrap, the `_check_stop_words_...` call and - the repeated `len()` probes, resolving `(request, token_limit)` pairs once - per request-id set instead of once per step: - - | Path | µs/step | - |---|---:| - | `post_process_requests`, as shipped | 213.8 | - | decode fast-path prototype | 33.8 | - | **reducible** | **180.0 (84%)** | - - So the lever clears its own gate decisively: 180 µs/step against an 8.85 ms - step period is **~2.0% e2e**, roughly twice what QWEN-024 delivered, and it - is the largest remaining host item. The mechanism is neither caching - (QWEN-023) nor op elimination (QWEN-024) but a third thing — collapsing a - per-request Python loop body — so it needs its own correctness argument, and - the termination path is exactly where that argument is hard. Note the - prototype is a *measurement*, not an implementation: it does not handle stop - words, log probs, chunked prefill, speculative tokens or eviction, and those - must fall back to the existing loop. -2. **Reduce graph node count** (HOSTGAP-S6 lever 4, now unblocked by - CGTRACE-CONTROL) — ~0.78 µs per removed node, so 100 nodes ≈ 0.85%. - **Risk medium**, sharply diminishing. -3. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — - ceiling ~9.2%, **risk high**, multi-session. - -### QWEN-025 — decode fast path for `post_process_requests` - -| Field | Value | -|---|---| -| Date | 2026-07-26 | -| Hypothesis | The decision-gate measurement in QWEN-024 says 84% of `post_process_requests` is reducible without any tensor work, by collapsing the per-request loop body. If that translates, it is worth ~2.0% e2e — twice QWEN-024 | -| Code revision | branch `perf/moe-fused-align`, dirty (session-7 set, QWEN-024, and this change) | -| Changed files | `megatron/core/inference/engines/dynamic_engine.py` (`_post_process_requests_decode_fast`, `_ppr_cache_epoch` invalidation at the three request-record mutation sites, env `MCORE_INFER_FAST_POST_PROCESS`, `..._VERIFY`), `dev/moe_fused/harness_updreq.py` (`run_ppr_equivalence`, in-process gate A/B) | -| Runtime flags | eight gates plus `MCORE_INFER_VEC_UPDATE_REQS=1`, plus `MCORE_INFER_FAST_POST_PROCESS=1` | -| Image / checkpoint / hardware / workload | identical to CLEANBASE-S8 | -| Job / run | job 5616264, e2e `ppr-off-1785104484` / `ppr-on-1785104751` / `ppr-off-r2-1785105441` / `ppr-on-r2-1785105667` / `ppr-off-r3-1785105923` / `ppr-on-r3-1785106159`; harness execs `e620b4429e3945e6a4b0ab61d731c372` (equivalence), `81ab8e8ee09840eebf825420633b2efc` (gate A/B + verify) | -| Throughput | **26,361.7 tok/s** on the three-pair ON mean; best single run **26,430.7** | -| Latency / TPOT | 9,658.7 ms / 9.7111 ms/tok on the ON mean, against 9.8060 OFF — **−94.9 µs/step** | -| Correctness | 400-step two-engine equivalence including 10 finisher steps and 200 steps past the token limit; 300 steps under `MCORE_INFER_FAST_POST_PROCESS_VERIFY=1`; temperature-0 coherence byte-identical | -| Nsight artifacts | OSL128 A/B pair `prof/g2-off-1785106875` and `prof/g2-on-1785107105`; analysis exec `0ec0392b8c1d47c5a099cee585d777a4`. Two follow-on host-visibility captures (`prof/hosts9-*`, `prof/hosts9b-*`) failed to import; see the note under the next-lever list | -| Result | **Accepted, +0.98%** on three-pair means (+0.59% on a best-4-of-5 trim), ON wins all three pairs | -| Next action | Host-side decode work is now 148 µs/step against 372 µs before QWEN-024; the next lever is graph node count, not the host chain | - -**Why this was started despite QWEN-024's stop verdict.** QWEN-024 declined to -touch `post_process_requests` because it is per-request Python object churn and -therefore not *vectorizable* — that was the pre-agreed falsification criterion -and it still holds. What the QWEN-024 decision-gate measurement then showed is -that being unvectorizable is not the same as being irreducible: an in-harness -prototype of the plain decode case ran at 33.8 µs/step against the real -method's 213.8 µs, i.e. 84% of the cost is dict lookups, record indexing, an -`isinstance` scalar wrap, a stop-word call that returns immediately, and ~1,719 -`len()` probes, none of which the plain decode case needs. That is a third -mechanism, distinct from QWEN-023's caching and QWEN-024's op elimination: -collapsing a loop body. - -**Design — the safety property comes first.** One env gate, -`MCORE_INFER_FAST_POST_PROCESS`, default OFF, dispatching to -`_post_process_requests_decode_fast` at the top of `post_process_requests`. The -fast path returns `None` and the full reference loop runs whenever anything is -not the plain decode case. The declining conditions are: speculative decoding -or accepted tokens, log probs, top-n log probs, finished routing block ids, -token event tracking, any pending stop-word state, chunked prefill in flight, a -TPOT sample due this step (`step_time > 0`), any eviction, **or any request -finishing this step**. - -That last condition is the design's whole safety argument. The task flagged an -off-by-one in EOS/length handling as the dangerous failure mode, and the honest -way to remove that risk is not to test it harder but to never enter it: the -fast path is structurally unreachable on any step where a request finishes, so -the termination state machine — pop, future resolution, routing -reconstruction, status, finish event — is only ever executed by the original, -well-tested code. At BS256 steady-state decode essentially every step has no -finisher, so nearly all the work is still avoided; the OSL1024 benchmark takes -the fast path on all but a handful of its ~1,024 steps. - -What remains in the fast path is one bounded append per request: - -```python -if num_generated < limits[i]: - generated.append(token) -``` - -which is the exact one-token specialization of the reference trim — the -reference computes `keep = num_tokens_to_generate - len(generated_tokens)` and -slices `tokens[:keep]`, which for a single token keeps it iff the request is -below its limit and drops it entirely otherwise. The first-token TTFT sample is -preserved inline on the `num_generated == 0` branch. - -The `(request, token_limit)` pairs are resolved once per active request-id set -rather than once per step, keyed on `(epoch, tuple(request_ids))`. The epoch is -bumped at all three sites that can change which object `record[-1]` resolves -to — the two `record.checkpoint()` calls (recompute-suspend and eviction) and -the `RequestEntry` insertion in `_add_request` — so a stale resolution is not -representable rather than merely unlikely. A configured stop word on any -request also declines, since the post-append scan is the one piece of -per-request work here that is not a bounded append. - -**Correctness evidence.** Three independent checks. - -*Two-engine equivalence, 400 steps.* Two stub engines are built with identical -request sets and driven through the same token stream, one gate OFF and one ON, -compared after every step on the full observable result: returned active id -list, finished record count and contents, `finished_request_count`, the live -request-id set, and per request `generated_tokens`, `generated_length`, `status` -and whether TTFT is set. Three requests are finished mid-batch every 37 steps -(10 finisher steps, 30 requests finished, all 10 correctly declined by the fast -path), and `num_tokens_to_generate` is set to 200 against 400 steps so every -surviving request spends 200 consecutive steps *at* its token limit — the -append-suppression branch — and both gates agree that it holds exactly 200 -tokens at the end. Zero mismatches. This is the check that answers the -off-by-one risk directly, on both sides of the limit and on the finish steps. - -*Verify mode, 300 steps.* `MCORE_INFER_FAST_POST_PROCESS_VERIFY=1` independently -recomputes the expected post-state of every request from the pre-state and the -sampled token before running the fast loop, then asserts equality. 300 -consecutive steps, zero mismatches. - -*End to end.* The three temperature-0 coherence completions are byte-identical -between gate ON and gate OFF, character for character across all three prompts. - -**Performance.** Host CPU per step at BS256, both arms measured in the same -process against the same engine so the comparison carries no process-to-process -noise: - -| Region | Gate OFF | Gate ON | Δ | -|---|---:|---:|---:| -| `post_process_requests` | 211.0 | 27.2 | **−183.9 (7.77×)** | - -And the whole post-sampling chain, against the pre-QWEN-024 starting point: - -| Region | Before QWEN-024 | After QWEN-024 | After QWEN-025 | -|---|---:|---:|---:| -| `active_request_mask` | 46.9 | 34.9 | 34.1 | -| `update_requests` | 137.6 | 74.0 | 73.7 | -| `post_process_requests` | 187.9 | 184.7 | **40.3** | -| **chain total** | **372.4** | **293.7** | **148.0** | - -The pre-registered falsification test for this lever was "falsified if a -`cProfile` of the fast path does not cut the per-step call count below ~500." -Measured: **261 calls/step** (7,831 over 30 steps), of which 256 are the single -`len()` per request, against ~1,719 `len()` probes alone before. Not falsified. - -E2e, three OFF/ON pairs run back to back in one session: - -| Pair | Gate | Throughput | Per-iter | -|---|---|---:|---| -| 1 | OFF | 26,054.0 | 26,251.6 / 26,249.3 / 26,241.1 / 26,274.0 / 25,283.4 | -| 1 | ON | **26,313.1** | 26,104.8 / 26,306.3 / 26,431.3 / 26,418.0 / 26,307.9 | -| 2 | OFF | 26,101.0 | 26,183.2 / 26,218.8 / 26,253.2 / 26,195.0 / 25,664.3 | -| 2 | ON | **26,430.7** | 26,442.2 / 26,416.0 / 26,428.8 / 26,426.4 / 26,440.3 | -| 3 | OFF | 26,164.4 | 25,883.8 / 26,270.7 / 26,270.5 / 26,228.0 / 26,173.0 | -| 3 | ON | **26,341.2** | 26,404.3 / 26,333.4 / 26,243.2 / 26,376.0 / 26,349.9 | - -Pairwise **+0.99%, +1.26%, +0.68%**; on three-pair means (OFF 26,106.5, ON -26,361.7) **+0.98%**. The task set +0.5% as the noise floor, so the headline -number clears it, but each OFF run contains one slow outlier iteration (25,283.4 -and 25,664.3 in the fifth position, 25,883.8 in the first) which flatters the -mean comparison. Dropping the slowest iteration from every run gives OFF -26,234.0 against ON 26,388.7, **+0.59%** — still above the floor, and the -honest lower bound. - -The distribution-level evidence is stronger than either mean. Pooling all 15 OFF -and 15 ON iterations, **13 of the 15 ON iterations are faster than every one of -the 15 OFF iterations** (the two exceptions are 26,104.8, the cold first -iteration of pair 1, and 26,243.2). ON also has visibly lower spread — pair 2's -five ON iterations span 0.10% — because the removed work included the variance, -not just the mean. ON wins all three pairs. -New best = **77.75%** of vLLM 33,994.5 on the best run, **77.55%** on the -three-pair mean. - -**Profile confirmation, and why only part of the host saving converts.** Two -BS256/OSL128 captures back to back, same code, differing only in the gate, -analysed with `analyze_hostgaps.py gaps` on the forced LM-head anchor, mean of -devices 0 and 3: - -| Metric | Gate OFF | Gate ON | Δ | -|---|---:|---:|---:| -| G1 (`index_elementwise` → `vectorized_gather`) | 555.5 µs | 499.4 µs | **−56.1 µs** | -| Total idle | 1717.5 µs | 1639.6 µs | −77.9 µs | -| Step period (median) | 8.8645 ms | 8.8050 ms | −59.5 µs | -| GPU-busy | 7.202 ms | 7.219 ms | +17 µs (noise) | -| Kernels per step | 1169 | 1169 | 0 | -| G3 (`CatArrayBatchedCopy` → `rmsnorm`) | 302.1 | 294.9 | −7.2 | -| G2A (`vectorized_elementwise` → `vectorized_gather`) | 65.8 | 61.5 | −4.3 | - -The saving is host-side and localized: the kernel count is bit-identical, GPU-busy -does not move, and the idle reduction is concentrated in G1. - -But note the conversion ratio, which is the interesting result and is *worse* -than QWEN-024's. The harness says 183.9 µs of host work was removed; G1 shrank -56.1 µs, total idle 77.9 µs, the profiled step period 59.5 µs, and e2e TPOT -94.9 µs. QWEN-024 removed 78.7 µs of host work and got 78.3 µs out of G1 — very -nearly 1:1 — because `update_requests` sits directly in the serial dependency -between the sampled tokens returning and the next step's launches. Only about a -third to a half of this change converts, which says most of -`post_process_requests` was already partly overlapped with GPU execution: it -runs from `async_bookkeep`, so the interpreter was working through it while the -device still had queued work, and only the exposed tail was ever on the critical -path. **The corollary for future levers: harness-measured host CPU savings are -an upper bound, and how much converts depends on where in the step the work -sits, not on how much of it there is.** The 2.0% projection from the -decision-gate measurement assumed a 1:1 conversion and was therefore too -optimistic by roughly a factor of two; the measured +0.98% is what it is worth. - -**Ranked next levers, after QWEN-025.** - -1. **Reduce graph node count** (HOSTGAP-S6 lever 4, unblocked by - CGTRACE-CONTROL) — ~0.78 µs per removed node, so 100 nodes ≈ 0.85%. Now the - top-ranked lever by expected value. **Risk medium**, sharply diminishing. - Falsified if removing nodes does not move the profiled step period. -2. **Re-attribute the residual G1** — 499 µs/step remains in G1 but the whole - measured host chain is now only 148 µs/step, so the majority of G1 is - something HOSTGAP-S6's attribution folded in but this campaign has not - isolated. Pure measurement, no code risk, and it is what decides whether any - further host-side lever exists at all. Blocked on the nsys host-capture - failure recorded above; unblock it by bisecting the host flag set against a - short capture, starting with `--cpuctxsw=process-tree`. -3. **G3, the `CatArrayBatchedCopy` → `rmsnorm` gap** — 294.9 µs/step and - *untouched by the last three levers*, which all moved G1. HOSTGAP-S6 - attributed it to a single 199 µs `cudaGraphLaunch` for a 1158-node graph, - which makes it the same lever as graph node count rather than an independent - one. **Risk unknown**; confirm the attribution still holds before treating it - as separate. -4. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — - ceiling now lower than the original ~9.2%, since QWEN-024 and QWEN-025 - together removed 224 µs/step of the host chain. **Risk high**, multi-session. - -Note the residual G1 is still 499 µs/step against a measured host chain of only -148 µs/step, so most of what is left in G1 is *not* the post-sampling -bookkeeping chain any more. Attributing the rest of G1 is a prerequisite for -claiming any further ceiling on host-side work. - -**Attempted and failed in the leftover time: re-attributing G1 at the new -config. Host-visibility capture is currently broken, and two plausible causes -are now ruled out.** Three captures with `profile_host_insession.sh` at the -QWEN-025 config, all unrecoverable in the same way — nsys finalization -deadlocked after the target exited, the bounded stop expired, and the -intermediate qdstrm failed to import with -`QuadDCommon::IncompleteFileException` from `verifyHeader`: - -| Run | OSL | `--python-sampling` | qdstrm | Outcome | -|---|---:|---|---:|---| -| `prof/hosts9-1785107698` | 128 | on | 366 MB | import failed | -| `prof/hosts9b-1785109992` | 96 | on | 347 MB | import failed, stream byte-stable 30 s first | -| `prof/hosts9c-1785110873` | 96 | **off** | 219 MB | import failed | - -*Hypothesis 1, copy-while-growing race — falsified.* The recovery path was -changed to wait for the source qdstrm size to hold steady before copying. The -second capture copied a stream that had been byte-stable for 30 s and still -failed `verifyHeader`, so the deadlocked nsys never writes the stream's -terminating section at all. The file is structurally incomplete, not truncated -in transit, and no amount of waiting recovers it. The settle loop is retained -only because it is harmless and documents the ruled-out cause. - -*Hypothesis 2, event volume from the Python sampler — falsified.* Dropping -`--python-sampling` (new `PYTHON_SAMPLING=0` knob) cut the capture from 347 MB -to 219 MB and it failed identically. Neither the sampler nor sheer volume is the -trigger. - -This does *not* affect any QWEN-025 result: the GPU-only script -(`profile_insession.sh`) finalized cleanly twice in this same session and -produced the `g2-off`/`g2-on` pair the profile table above is built from. The -deadlock is specific to the host-visibility flag set, and it is not universal — -HOSTGAP-S6 obtained a usable host trace from this same script on an earlier -allocation, so something environmental differs. Remaining untested suspects, in -order: `--cpuctxsw=process-tree`, then `--sample=process-tree` itself, then -`osrt` tracing. **Do not spend a third session's leftover time on blind retries -of the full flag set** — bisect the flags against a short capture first, since -each full attempt costs ~14 minutes. - -**QWEN-026 landed: fused QK-RMSNorm, +2.9%, new best 79.87% of vLLM.** This is -the first attention-side lever to pay off and the largest win since QWEN-021, -and it over-delivered ~3× against its microbench ceiling — the two per-head -RMSNorm launches sit directly in the serial attention chain, so collapsing them -to one kernel removes a launch *and* a graph node per layer (×48) and cuts the -host dispatch gap between them, not just the ~2 µs/layer of device time the -microbench saw in isolation. **The lesson for ranking: a small-device-time -kernel that sits on the serial critical path and removes a graph node is worth -materially more than its isolated microbench, because graph-node count and -serial-chain launches both convert.** This makes "collapse two adjacent -small launches into one" a higher-value pattern than the raw device-time -numbers suggest — the opposite of QWEN-025's host-side levers, which converted -at only ~1/3. - -The one caveat is that the fused norm is **bf16-ulp, not bit-exact** (TE's -internal rsqrt/reduction differs), so one of the three coherence prompts flips a -single near-tie greedy token. Both continuations are fluent and factually -correct, so it was accepted, but a bit-exact variant (matching TE's rsqrt -instruction so the two-call path is reproduced exactly) would remove the last -correctness caveat and is a cheap follow-up if byte-identity is ever required. - -**Ranked next levers, after QWEN-026.** - -1. **Fuse the other adjacent small-launch pairs on the serial path** — QWEN-026 - proved the pattern converts at ~3× microbench. The remaining per-layer - small-kernel pairs that sit back-to-back in the decode graph are the - input-layernorm/pre-MLP-layernorm `rmsnorm_fwd_tuned` pair and the - `triton_poi_fused_add_copy` residual-add kernels flanking each block. Target - the pair with the shortest serial gap between them first. **Risk medium**; - the same decode-only gating and coherence bar apply. Falsified if collapsing - the pair does not move the profiled step period the way QWEN-026 did. -2. **Reduce graph node count** (HOSTGAP-S6 lever 4) — still live at ~0.78 µs per - removed node; QWEN-026 removed 48 nodes as a side effect and that clearly - helped, so a direct campaign to prune redundant nodes is now better-evidenced - than before. **Risk medium**, sharply diminishing. -3. **Re-attribute the residual G1** — unchanged from the QWEN-025 ranking; still - blocked on the nsys host-capture failure. Bisect the host flag set against a - short capture, starting with `--cpuctxsw=process-tree`. -4. **Move per-step request bookkeeping onto the GPU** (HOSTGAP-S6 lever 5) — - **Risk high**, multi-session; unchanged. - -The most productive next step is to re-profile at the new eleven-gate config -(all ten gates + QWEN-026) so the next adjacent-launch pair is chosen from the -current serial gaps rather than the pre-QWEN-026 trace. - -**QWEN-027 landed: fused add+RMSNorm (self_attn_bda + pre_mlp_layernorm), -+1.37%, new best 80.98% of vLLM.** PROFILE-S9 chose the target exactly as -predicted — the residual-add and the next norm were the two dominant small -kernels and adjacent on the serial path — and the fusion converted cleanly. Two -things worth carrying forward. First, the **input RMSNorm is not a separate -module**: for the non-MLA path `linear_qkv = column_parallel_layer_norm_linear`, -so that norm lives inside the TE `LayerNormLinear` QKV GEMM and cannot be -intercepted at the layer level. Only the standalone `pre_mlp_layernorm` was -fusible, so QWEN-027 removed 48 of the 96 candidate boundaries; the other 48 -(the `mlp_bda` + next-layer input norm) would require either fusing across the -layer boundary or reaching inside TE's `LayerNormLinear`. Second, the add+norm -kernel does **not** regress at high token counts (flat ~4.1 µs vs QWEN-026's -1-row-per-CTA design that lost above 256 tokens), so the same pattern is safe to -apply more widely if the prefill path is ever targeted. - -**Ranked next levers, after QWEN-027.** - -1. **Re-profile at the twelve-gate config and re-rank.** Two norm+add fusions - have now landed; the serial chain has shifted again. Cheapest high-value - step before picking the next lever. Specifically check whether - `_fused_qk_rmsnorm` (QWEN-026) and the RoPE/append-kv kernels around it are - now the tallest adjacent small-kernel run, since attention is the remaining - untouched small-kernel cluster. -2. **Fuse the `mlp_bda` + next-layer input norm boundary** — the other 48 - boundaries QWEN-027 could not reach, because the input norm is inside - `LayerNormLinear`. Options: (a) fuse the `mlp_bda` add into the *start* of the - TE QKV GEMM's norm (needs a TE hook or a mcore-side prenorm that bypasses - `LayerNormLinear`'s internal norm), or (b) accept a small numeric change by - moving the input norm out of `LayerNormLinear` into a standalone fused - add+norm like QWEN-027. **Risk medium–high**; (b) changes the QKV GEMM path. -3. **Reduce graph node count** (HOSTGAP-S6 lever 4) — QWEN-026+027 removed ~96 - nodes/step as a side effect and both converted at ~3×; a direct node-pruning - pass is now well-evidenced. **Risk medium**, diminishing. -4. **Re-attribute the residual G1** — unchanged; blocked on the nsys - host-capture failure. Bisect the host flag set against a short capture. - -A bit-exact variant of both fused norms (matching TE's rsqrt instruction) would -remove the last coherence caveat and is a cheap follow-up if byte-identity is -ever required — QWEN-027's residual is already bit-exact, only the norm output -is one bf16 ulp off. - -## Session 13 (2026-07-30) — measurement session: why the block is latency-bound - -No optimization was accepted this session. It replaced the campaign's measurement -basis, which had been quietly wrong, and produced the first quantified path to -parity. Everything below is same-node (`nvl72151-T13`) unless stated. - -### CLEANBASE-S13 — same-node reference, and the tightest drift control yet - -| Arm | Throughput | tpot | -|-----|-----------|------| -| `ep01_ref` (13 gates + async) | **28,529.1 tok/s** | 8.973 ms | -| `ep06_ref2` (repeat, end of allocation) | **28,606.9 tok/s** | 8.949 ms | - -0.27% apart, so on this node anything above ~0.3% is a real effect. Note this node -is ~2.1% faster than `nvl72067-T13`, where the best config measured 27,944 and the -vLLM reference 34,415.7 — **the 81.2% figure remains the only same-node -comparison; 28,529 must not be divided by 34,415.7.** - -### NSYS-UNUSABLE — nsys cannot measure this workload, in any mode - -The hypothesis was that `--cuda-graph-trace=node` was inflating the trace by timing -~1,800 kernel nodes per step, and that `=graph` would collapse each replay to one -range and restore fidelity. It failed on every count: - -| Arm | Throughput | vs unprofiled | -|-----|-----------|---------------| -| unprofiled | 28,469.5 | — | -| `--cuda-graph-trace=graph` | 20,557.2 | **−27.8%** | -| `--cuda-graph-trace=node` | 20,174.0 | **−29.1%** | - -1. The mode is worth only 1.3 points. **nsys costs ~28% here regardless**, and the - overhead is host-side. -2. On nsys 2026.3.1 `=graph` does **not** collapse replays: all 224k kernel rows - still appear, with `graphId` NULL. The 98 rows in - `CUPTI_ACTIVITY_KIND_GRAPH_TRACE` are graph *instantiations* from warmup. -3. Step reconstruction from the trace again produced an impossible 71.8 ms "step" - against a known 12.45 ms — the same divisor trap as earlier sessions. -4. Worst: the host-side overhead desynchronizes the EP ranks, and the NVLS - collectives are spin-waits, so they inflate **~120×** - (`_multimem_all_gatherv_3tensor` reads 849 µs/call; 48 of those would be 40 ms - against a 9 ms step). An earlier trace measured the same kernel at ~7 µs. - -**Consequence: every absolute idle, gap, and barrier-wait number this campaign read -off an nsys trace is unsafe**, including HOSTGAP-S6's inter-replay gaps and the -"1.6 ms/step idle" and "4.4× barrier spread" that motivated sessions 7–12. Relative -per-call durations of *compute* kernels remain usable. Do not profile this workload -with nsys to obtain a busy/idle split; use STEPGPU below. - -### STEPGPU-S13 — the block graph is 92.8% of the step (`step_gpu_timing.py`) - -New env-gated instrumentation (`MCORE_INFER_STEP_GPU_TIMING`) wraps -`TransformerBlock.__call__` on the CUDA-graph path with CUDA events, held in a ring -of 64 pairs and read back only when the ring wraps, so no step synchronizes on its -own work. It costs 2.3% and, unlike nsys, **validates itself**: the step period it -reports (8.97 ms) matches the benchmark's tpot (8.973 ms). - -| Quantity | Value | Share | -|----------|-------|-------| -| step period | 8.97 ms | 100% | -| transformer-block CUDA graph | ~8.33 ms | **92.8%** | -| everything outside the graph | **0.59 ms** | 6.6% | - -Held within 0.2 points across all 4 ranks over 4,032 steps, and reproduced at 92.8% -in a second sample. **This retires host-side optimization quantitatively.** -Embedding, logits GEMM, sampling, bookkeeping and every host gap together are -0.59 ms; the gap to vLLM is ~1.53 ms/step. Third confirmation after QWEN-032/033, -and the first that is a measurement rather than an inference. - -### EPSKEW-S13 — the EP skew hypothesis is falsified - -All 4 ranks are processes on one node, so `perf_counter_ns` is directly comparable -between them; the spread of step-entry timestamps *is* the arrival skew the NVLS -spin-waits absorb. Median spread **24 µs**, p90 46 µs — 0.27% of the step. Rank 1 is -mildly structurally late (last on 42% of steps, first on 7%) but by only ~21 µs -median. The "4.4× barrier wait spread" was an nsys artifact per NSYS-UNUSABLE. -**Host arrival skew is not a lever.** Do not revisit without new evidence. - -### BSSCALE-S13 — 32× the tokens costs 1.70× the block time - -Batch-size sweep with the block timer on. Valid points are exact CUDA-graph buckets; -BS192 is padded into the 256 bucket, so the curve is a staircase, not smooth. - -| BS | step | block | µs/layer | -|----|------|-------|----------| -| 8 | 5.563 | 5.030 | 105 | -| 16 | 5.896 | 5.361 | 112 | -| 32 | 6.290 | 5.770 | 120 | -| 64 | 7.081 | 6.542 | 136 | -| 128 | 7.140 | 6.585 | 137 | -| 256 | 9.151 | 8.553 | 178 | - -**The block is latency-bound, not compute-bound.** The batch-independent floor is -~5.03 ms/step = **105 µs/layer = 55% of the whole step**; the marginal cost of the -other 248 tokens is only 73 µs/layer. For scale, vLLM's *entire* step is 7.44 ms -while mcore's floor plus non-graph overhead is already ~5.6 ms. - -### CHAINCOST-S13 — fusion pays, but *not* per kernel removed - -All 13 gates off vs on, at two batch sizes: - -| | block ON | block OFF | delta | -|---|---------|-----------|-------| -| BS8 | 5.030 | 6.295 | 1.265 ms/step = 26.4 µs/layer | -| BS256 | 8.553 | 10.317 | 1.764 ms/step = 36.7 µs/layer | - -At BS8 the block is nearly pure latency, so 26.4 µs/layer over the ~10 kernels the -gates remove suggests ~2.6 µs of serial chain latency per kernel — well above those -kernels' own 1–2 µs device time. The 13 gates are jointly worth **16.1%** of -throughput. - -**The per-kernel decomposition did not survive calibration — it is refuted.** -Six single-gate arms at BS8, against the all-gates-ON block of 5.030 ms: +| ID | Engine | Throughput | Avg latency | TPOT | Job / node | Nsight trace | Status | +|---|---|---:|---:|---:|---|---|---| +| VLLM-BASELINE | vLLM DP4+EP | — | — | — | — | — | not run | +| MCORE-BASELINE | mcore EP4/TP1 | — | — | — | — | — | not run | -| arm | block | delta | µs/layer | kernels/layer | µs/kernel | -|-----|-------|-------|----------|---------------|-----------| -| `MCORE_FUSED_QK_NORM=0` | 5.291 | +0.260 | 5.4 | 1 | 5.4 | -| `MCORE_ROUTER_FUSED_TOPK=0` | 5.292 | +0.261 | 5.4 | 3 | 1.8 | -| `MCORE_FUSED_ADD_NORM=0` | 5.698 | +0.668 | 13.9 | 2 | 7.0 | -| `MCORE_FUSED_ADD_NORM_QKV=0` | 5.375 | +0.345 | 7.2 | 1 | 7.2 | -| `MCORE_MOE_GEMM_TUNE=0` | 4.762 | **−0.268** | −5.6 | count-neutral | — | -| `MCORE_MOE_SUM_FAST=0` | 5.073 | +0.043 | 0.9 | count-neutral | — | +Both baselines must be run with no optimization flags enabled, and the gap +computed, before any code change. -Least squares over the four chain-reducing arms gives **slope 0.30 µs/layer per -kernel removed with a 7.46 µs/layer intercept** — i.e. kernel count explains almost -nothing and the per-gate constant dominates. Per-kernel cost ranges 1.8–7.2 µs with -no consistency. **There is no usable per-kernel coefficient, and any projection of -the form "remove N kernels, gain N × c" is invalid.** The earlier ~2.6 µs/kernel -figure and the "25 → 15 kernels buys 1.25 ms and most of the path to parity" -projection are both withdrawn. - -Two further cautions from the same data: - -- **Single-gate deltas are strongly sub-additive and must not be summed.** Just four - gates sum to 31.9 µs/layer, already *more* than all thirteen together (26.4), so - these arms are contending for a shared resource — the same signature QWEN-030 saw. -- The measurement itself is sound: six arms produced six clearly distinct outcomes - (−5.6 to +13.9 µs/layer), so there is no "any gate off costs the same" artifact, - and the earlier suspicion of BS8 graph-node quantization was unfounded. The - identical QK-norm / router-topk pair was coincidence. - -**GEMM-TUNE-SMALLBATCH (new, actionable):** `MCORE_MOE_GEMM_TUNE=1` is a **~5% -regression at BS8** — turning it off raised throughput 1438.0 → 1509.6 tok/s. The -QWEN-013 tiles were tuned for 256-token decode and QWEN-013b added an *upper* bound -(fall back above M=384) but never a lower one, so at 8 tokens the tuned tiles are -simply the wrong shape. A lower-bound guard is a small, low-risk fix. Irrelevant to -the BS256 parity goal; relevant to anyone serving small batches. - -**This finally explains the campaign's recurring puzzle.** QWEN-026, 027 and 028 -each over-delivered ~3× their device-time arithmetic and each was recorded as a -surprise. Removing a launch removes a *chain link*, not just a duration — in a -latency-bound block that is the expected multiplier, not luck. - -**The 105 µs/layer floor remains unattributed.** A split into ~66 µs of serial chain -and ~39 µs of NVLS collectives was arithmetically tempting and consistent with -QWEN-016's "exposed, latency-bound" finding, but it was built on the per-kernel -figure the calibration then refuted, so it is withdrawn. Attributing the floor needs -a direct measurement, not a coefficient — the obvious candidate is an env-gated -ablation that stubs one component at a time inside the graph and reads the block -timer, since `step_gpu_timing.py` now makes that a reliable readout. - -### NCCL-HANG — `inference_moe_token_dispatcher_type=nccl` is non-functional here - -With async scheduling and the 13 gates, the server starts and then every request -times out. It hung for 2 h and cost half the allocation, which also lost the -nvls-vs-nccl comparison. Two process lessons: job scripts now carry -`timeout --signal=INT 1200`, and a watcher that reports only *state changes* cannot -distinguish a hung job from a healthy one — it needs a log-mtime staleness check. - -### Session 13 — conclusion & recommendation - -The target is no longer expensive kernels; it is **per-layer latency**. Ranked by -evidence: - -What is certain: the prize is inside the transformer block (92.8% of the step), the -block is latency-bound rather than throughput-bound (32× tokens for 1.70× time, -~55% batch-independent), and neither host work nor EP arrival skew is worth another -hour. What is *not* known is what the 105 µs/layer floor is made of — and this -session showed that guessing at it from coefficients produces retractions. - -1. **Attribute the 105 µs/layer floor by direct ablation.** Env-gated stubs for one - component at a time (NVLS dispatch, NVLS combine, expert GEMM, attention) read - through the now-trustworthy block timer. Output is garbage under ablation, which - is fine — only the timing is wanted. This is the prerequisite for sizing - anything else, and it is the cheapest remaining experiment. -2. **The 2 NVLS collectives** — the leading suspect for the floor, on the strength - of BSSCALE (cost is barrier latency, not bytes) and QWEN-016. Levers: fewer - barriers (fuse dispatch/combine, or hoist across layers) or overlap with - compute. QWEN-017 already failed at replacing the barrier itself; overlap is - unexplored. **Size it via item 1 before building.** -3. **More fusion** — still plausible, and the aggregate (13 gates = 16.1%) says the - direction has paid. But the per-gate value is idiosyncratic, not proportional to - kernels removed, so each candidate must be sized on its own rather than from a - chain-length model. -4. **Do not** pursue host-side work, EP arrival skew, or nsys-derived idle. -5. Small, separate: add a lower-bound guard to `MCORE_MOE_GEMM_TUNE` - (GEMM-TUNE-SMALLBATCH). - -## Session 14 (2026-07-30/31) — the floor, attributed at last - -Session 13 ended unable to say what the ~105 µs/layer floor was made of, having -twice retracted an inference built on per-kernel coefficients. This session -measured it directly and the answer reorders the whole campaign. Node -`nvl72115-T08`; all comparisons within it. - -### ABLATE-HARNESS — how the floor was measured (`floor_ablation.py`) - -New env-gated module plus six call-site gates. Each deletes one component from -the *captured* block graph and the change is read from `step_gpu_timing`. Three -design points are what make it trustworthy: - -* **Ablate only while capturing** (`torch.cuda.is_current_stream_capturing()`). - Warmup runs eagerly first, so every real kernel executes at least once and - every preallocated buffer holds plausible values; capture then takes the - ablated branch, so the choice is resolved once and no replay gains - data-dependent control flow. -* **Substitute cached zeros, never uninitialized memory.** Zeros keep the - residual finite, so the next layer's router still spreads tokens across - experts. Garbage would collapse routing and silently change the grouped-GEMM - tile balance of the components *not* being ablated. Caching by shape matters - too: a fresh `torch.zeros` per replay would add a multi-MB fill to the - measurement it is meant to isolate. -* **`hit()` logs the first time each site takes its ablated branch.** This - immediately caught a dead gate (see ATTN-BIG) that would otherwise have read - as "attention is free". - -Deltas are **chain** deltas: they include the serialization a component imposed -on its neighbours, which in a latency-bound block is the quantity that matters, -and they are deliberately not comparable to a profiler's kernel duration. - -Drift control: two baselines 0.4% apart on block time (7.639 / 7.607 ms) but -**2.5% apart on throughput** (27,626 / 28,321 tok/s). Block time is the primary -metric this session; throughput differences under ~2.5% mean nothing here. - -### FLOOR-S14 — the composition, and which parts are addressable - -Baseline block **7.623 ms = 158.8 µs/layer** at BS256/OSL1024. - -| component removed | block | delta | µs/layer | % block | tput | -|---|---|---|---|---|---| -| expert GEMMs (FC1+SwiGLU+FC2) | 4.804 | −2.819 | −58.7 | −37.0% | +48.3% | -| attention (`flash_decode_and_prefill`) | 5.756 | −1.867 | −38.9 | −24.5% | +27.4% | -| topk reduce (`_moe_sum`) | 7.079 | −0.544 | −11.3 | −7.1% | +7.9% | - -**GEMM-ROOFLINE — the expert GEMM is done, and this is the session's most -consequential finding.** Per layer per rank it streams 32 local experts' FC1+FC2 -weights = **302 MB**, so 14.5 GB per step. Against QWEN-012's measured 6.08 TB/s -that is a 49.7 µs/layer floor, and it measures 58.7 — **84% of achievable -bandwidth**. Only 19.3 GFLOP/layer, so it is nowhere near compute-bound; it is -weight-streaming-bound. Consequences: - -1. The largest single component, 37% of the block, is **essentially optimal in - bf16 and cannot be tuned further.** Stop looking at grouped-GEMM tiling for - throughput at BS256. -2. **vLLM pays the same 2.83 ms/step**, because it streams the same weights in - the same precision. So the gap to vLLM does *not* live here, and every prior - session that ranked "MoE grouped-GEMM 40.5%" as the top target was ranking by - size rather than by addressability. -3. The only lever on it is fewer weight bytes: FP8/NVFP4 experts, which is a - precision decision, not a kernel one. - -Subtracting it: our addressable non-GEMM cost is ~4.80 ms of block + 0.59 ms -outside ≈ **5.39 ms/step**, against vLLM's ~4.6 ms on the same arithmetic. **The -whole remaining gap is ~0.8 ms/step of non-GEMM work**, and attention plus the -topk reduce are 2.41 ms of that region — enough room to close it without -touching the GEMM. - -**ATTN-BIG — attention is 38.9 µs/layer and roughly 3× off its own roofline.** -This is the new top target and it was invisible until now. At 64 tokens/rank and -~600 average context it reads ~79 MB of paged KV per layer per rank, a 13 µs -bandwidth floor against 38.9 µs measured = **33% of 6.08 TB/s**. Unlike the -GEMM there is real headroom, and it is 24.5% of the block with a +27.4% -throughput ceiling. - -Process note: the first attention arm reported **zero** delta, and only the -`hit()` log revealed why — the gate was on `_run_core_attention`, which the -*static* engine uses. Dynamic batching calls `flash_decode_and_prefill`. Without -the hit counter this would have been recorded as "attention is free", which is -exactly the class of false negative that cost sessions 7–12. - -**Two arms are invalid; recording them so they are not re-read as results.** -`MCORE_ABLATE_NVLS_DISPATCH` and `..._COMBINE` both hit the 1500 s timeout -(rc=124). Prefill runs the real collective eagerly while replays skip it, which -almost certainly desynchronizes the symmetric-memory barrier sequence between -ranks. Their block numbers (−0.66 and −3.26 ms) were measured in a degraded -state and **must not be quoted**; the NVLS collectives remain unmeasured. -`MCORE_ABLATE_MOE_ALIGN` is invalid for a different reason: replaying warmup's -align tables also shrank `num_tokens_post_padded`, so the GEMMs did less work and -the arm measured that instead of the align cost. - -### QWEN-034 — FC2 epilogue topk reduce (atomics): rejected, −12.5% - -Hypothesis, sized off the ablation rather than guessed: the topk reduce is worth -11.3 µs/layer of chain, so folding it into FC2's epilogue should recover most of -it and drop the `[num_valid, K]` intermediate as well. FC2 already loads the -routing weight (`MUL_ROUTED_WEIGHT`) and row *r* belongs to token `r // topk`, so -the change is small: scale the fp32 accumulator and `tl.atomic_add` into the -`[max_tokens, K]` output instead of storing. - -| | block | throughput | -|---|---|---| -| baseline | 7.623 | ~27,974 | -| `MCORE_FUSE_FC2_REDUCE=1` | 8.995 | 24,468 | - -**+1.372 ms/step (+28.6 µs/layer), −12.5% throughput.** Coherence passes and the -fused form is strictly *more* accurate (no bf16 truncation of each slot before -the sum), but it is far slower: the atomics cost ~40 µs/layer against the 11.3 it -was meant to save, and the destination also needs an 8.4 MB zero-fill per layer. -Kept default-off as a documented negative. - -Why it failed, and the general lesson: a topk slot's contributions land on one -output row from *different* tiles, so the reduction cannot be kept in registers -and every element becomes a global read-modify-write. This is the same wall -QWEN-001's mega-fusion hit. **Epilogue fusion pays when the consumer is -elementwise on the tile (QWEN-002's SwiGLU, QWEN-013's squared-relu) and loses -when it reduces *across* tiles.** That rule now has two independent -confirmations and should be applied before building, not after. - -### Session 14 — conclusion & recommendation - -The campaign's target list is now ordered by *addressability* rather than size: - -1. **Attention, 38.9 µs/layer at 33% of its bandwidth roofline** — the largest - component with real headroom, +27.4% ceiling. Start by identifying which - FlashAttention generation decode actually resolves to (`flash_attention_version` - has no CLI flag; a benchmarking env override and a one-shot log were added) and - A/B the alternatives, then look at paged-KV over-read and the decode split-K. -2. **The topk reduce, 11.3 µs/layer** — real but *not* via epilogue atomics - (QWEN-034). A separate kernel that reduces with better locality, or folding it - into the combine, are the remaining options. -3. **The NVLS collectives are still unmeasured** — the ablation needs to skip the - collective in warmup too, so the barrier sequence stays consistent across - ranks, before any number from it is trusted. -4. **Do not** pursue grouped-GEMM tuning for BS256 throughput (GEMM-ROOFLINE), - host-side work, EP arrival skew, or nsys-derived idle. -5. Precision (FP8/NVFP4 experts) is the *only* lever on the 37% the GEMM - occupies, and it is a correctness decision that needs its own evaluation. - -## Session 15 (2026-08-01) — FlashAttention-2 beats FlashAttention-4 on decode - -All arms: BS256, OSL1024 unless noted, 4×GB200 node `nvl72151-T14`, job `5765929`, -session `qwen-attn`, thirteen accepted gates on, `run_rebased.sh`. Block-graph GPU -time from `step_gpu_timing.py` (CUDA events, no profiler), median over the whole -decode range. - -### QWEN-035 — pin FlashAttention 2 on the decode path: **+3.5%** - -| arm | block-graph GPU | throughput | -| --- | --- | --- | -| `ga00_ref` FA4 (mcore's auto choice) | 7.695 / 7.707 ms | 28,149.6 | -| `ga05_ref2` FA4, repeat | 7.751 / 7.698 ms | 28,098.1 | -| `ga02_fa2` FA2 pinned | 7.412 / 7.400 ms | 28,999.2 | -| `ga04_fa2b` FA2 pinned, repeat | 7.398 / 7.384 ms | 29,225.9 | -| `ga11_fa2_five` FA2, 5 timed iters | — | **29,132.8** | - -Repeat spread is 0.3% on block time and 0.2–0.8% on throughput, so the 4.1% block -separation is roughly ten times the noise floor. `ga11_fa2_five` is the headline -number: 5 timed iterations at 28,705.6 / 29,287.2 / 29,199.4 / 29,230.4 / -29,249.3 tok/s — 0.30% spread across iterations 2–5, with iteration 1 the usual -ramp outlier. tpot 8.787 ms/tok, avg_latency 8,733.2 ms. It agrees with both -3-iteration runs (mean of all three: 29,119). Coherence is identical on two of -three probes and equally correct on the third — expected, since the two backends -accumulate in a different order. **85.7% of vLLM's 33,994.5** on the 5-iteration -headline run, up from 82.7%. - -The mechanism is in `flash_decode_and_prefill`, not in the kernels' relative -quality: mcore's auto-preference is FA4 → FA3 → FA2, but **FA4 has no dedicated -decode kernel on this path**. FA4 decode is routed into `flash_attn4_varlen_func` -— the varlen *prefill* interface, called with `max_seqlen_q = 1` — while FA2 goes -to `flash_attn_with_kvcache`, the purpose-built flash-decoding kernel with -split-KV and combine. So the auto-resolution order silently costs 3.5% on decode. -This is arguably an upstream default bug rather than a tuning knob. FA3 is not -installed in this container (`HAVE_FA3=False`), so `ga01_fa3` asserted out; the -FA3 leg of the comparison is untested. - -Enable with `MCORE_FLASH_ATTN_VERSION=2` (the env override added in Session 14 -because `flash_attention_version` has no CLI flag). Config still wins over env. - -The gain is **entirely inside attention**, and the ablation proves it: with -attention stubbed out, the block floor is the same under both backends (5.858 ms -FA4 vs 5.884 ms FA2, inside noise), so attention alone goes 1.843 → 1.510 ms/step, -−18%. - -### ATTN-ROOFLINE-S14-RETRACTED — attention is *not* 33% of roofline - -Session 14 recommended attention as "the largest component with real headroom, -33% of its bandwidth roofline, 3× available". **That roofline was wrong** and the -recommendation it produced should not be reused. - -KV-cache reads alone, per layer per rank, are -`B(256) × S × 2(K,V) × H_kv(4) × D(128) × 2 B`. At OSL1024 with ~100-token gsm8k -prompts the average KV length across the decode is ~600, giving **315 MB/layer** -→ **39.3 µs/layer at 8 TB/s HBM**. The measured attention chain delta under FA2 is -31.5 µs/layer, and a chain delta *understates* a kernel's duration whenever any of -it overlaps neighbouring work. Attention is therefore at or above its bf16 -bandwidth bound, not a third of the way to it. There is no 3× there, and the only -lever on it is reading fewer bytes (FP8 KV cache — no knob exists in mcore today). - -### FLOOR-CONST-S15 — the non-attention floor is KV-length independent - -| arm | OSL | attention | block-graph GPU | -| --- | --- | --- | --- | -| `ga06_fa2_na` | 1024 | ablated | 5.880 / 5.888 ms | -| `ga09_o512_na` | 512 | ablated | 5.833 / 5.871 ms | - -Halving the output length halves the average KV length and changes the ablated -floor by 0.6%, i.e. not at all. **The 5.85 ms floor contains no KV-dependent -work**, so it is pure weight-streaming, collectives, and chain latency, and it can -be reasoned about independently of the decode position. Its OSL512 attention-free -throughput is 34,432 tok/s. - -The matching OSL512 *reference* arm (`ga08_o512`) is unusable: rc=124, the server -stalled after ~200 decode steps and hit the 1500 s timeout, so its single -4.93 ms report is from a degraded run. The intended KV-length scaling test of -attention is therefore still untested; only the floor half of the pair survived. - -### Rejected: KV-cache block size 512 - -`ga07_fa2_b512` (`--inference-dynamic-batching-block-size 512`, default 256): -block 7.413 / 7.390 ms, 29,122.9 tok/s — indistinguishable from FA2 at the default -block size. Paged-KV indirection granularity is not a cost here. - -### NVLS-UNABLATABLE — the collectives cannot be sized by subtraction, ever - -The Session 14 recommendation was to fix the NVLS ablation by suppressing the -collective in warmup too, so the symmetric-memory barrier sequence stays -consistent across ranks. That fix was made (`nvls_off()`, an unconditional gate -replacing the capture-only one) and it worked as intended: `ga10_nonvls` ran to -completion, rc=0, no hang, and both gates fired on all four ranks -(`nvls_dispatch` 4 hits, `nvls_combine` 4 hits). - -**The result is that removing both collectives makes the block 13.5% slower:** -8.391 / 8.407 ms against the 7.39 ms FA2 baseline, 25,950.2 tok/s against 29,113. - -Removing two collectives cannot make their own absence expensive, so the arm is -confounded, and the reason is structural rather than a gating bug: -**`multimem_all_gatherv_3tensor` produces the work descriptor for everything -downstream.** `dispatch_preprocess` reads `routing_map` back out of the gathered -symmetric buffer (`agv_r["tensor"].view(global_max, topk)`), and that routing map -is what determines how many token-expert assignments — hence how many grouped-GEMM -tiles — the experts execute. Skip the gather and the buffer still holds whatever -warmup's *prefill* gather left there, which has a different and denser assignment -distribution than steady-state decode. So the experts do a different amount of -work, and the measured block time says nothing about the collective. - -This is the same work-descriptor confound that invalidated the Session 14 MoE-align -ablation, and it is not fixable by better gating. **Subtractive ablation is the -wrong instrument for the dispatch collective**, and by extension for anything that -computes a work descriptor. Four attempts have now failed on this (two hangs, one -align confound, this one); stop trying. - -The right instrument is non-invasive in-situ timing: a second CUDA event pair -recorded around the collective inside the captured graph, read back through the -same wrapped ring `step_gpu_timing.py` already uses for the whole block. Graph- -captured event records replay, so this measures the real duration per replay while -leaving the work descriptor and the barrier sequence completely untouched. That is -the next experiment, and it needs no ablation at all. - -### Session 15 — conclusion & recommendation - -Where the 7.39 ms block actually goes, with roofline where it is known: - -| component | measured | bf16 bandwidth floor | verdict | -| --- | --- | --- | --- | -| attention | 1.51 ms | ~1.89 ms (315 MB/layer) | **at the bound** | -| expert GEMMs | ~2.16 ms | 1.81 ms (302 MB/layer) | 84% of bound | -| dense QKV/O proj | — | 0.23 ms (37.8 MB/layer) | negligible | -| **unaccounted** | **~3.4 ms** | n/a | **the target** | - -Two thirds of what is left is neither attention nor weight streaming. That ~3.4 ms -— 46% of the block — is NVLS dispatch/combine, routing/permute, norms, and chain -latency, and it is the only place a 14% end-to-end gap can still come from. The -early OSL1024 profile already noted mcore carrying 11.5% exposed EP comm against -vLLM's zero, which points the same way. - -1. **Attribute the ~3.4 ms remainder with in-situ event timing, not ablation.** - NVLS-UNABLATABLE settles the method question: the dispatch collective produces - the routing map that sizes all downstream GEMM work, so subtracting it changes - the work and the timing is meaningless. Add a per-site CUDA event pair inside - the captured graph (same ring-buffer readback as `step_gpu_timing.py`) around - the dispatch, the combine, and the expert GEMM. Nothing else should be built - before those three numbers exist. -2. **Overlap, not replacement, for the NVLS barriers.** QWEN-017 already failed at - replacing the barrier itself; overlapping dispatch/combine with expert compute - is unexplored and is the natural lever if item 1 sizes them large. -3. **Precision is the only lever on the bandwidth-bound 3.7 ms** (attention + - expert GEMMs together). FP8 KV cache would need implementing — no mcore knob - exists — and would roughly halve attention. FP8/NVFP4 experts likewise. Both - are correctness decisions needing their own evaluation. -4. **Raise the FA4 decode path upstream.** Independent of this campaign: FA4 - resolving to a prefill kernel for decode costs every mcore inference user 3.5%. -5. **Do not** pursue attention kernel tuning (at its bound), grouped-GEMM tuning, - KV block size, host-side work, EP arrival skew, or nsys-derived idle. - -## Session 16 (2026-08-02) — the block, attributed - -Job `5777945` on `nvl72117-T17`, session `qwen-insitu`, thirteen gates + FA2, -BS256/OSL1024. New module `megatron/core/inference/insitu_timing.py`. - -### INSITU-METHOD — external event records, and the two probes that were needed - -Subtraction is dead (NVLS-UNABLATABLE) and nsys is unusable here, so components are -timed *in place* with an event pair inside the captured graph. Getting there took two -standalone probes, and both earned their keep by failing: - -1. **Plain `torch.cuda.Event.record()` during capture does not work.** It raises - nothing and does become a graph node, so it looks fine — and then - `elapsed_time` fails with `cudaErrorInvalidValue`, because an ordinary - event-record node exists only for intra-graph ordering and carries no - host-readable timestamp. The first instrumented run duly reported every site as - zero. Had the probe not existed, the natural reading of that run would have been - "the sites cost nothing". -2. **`cudaEventRecordWithFlags(..., cudaEventRecordExternal)` does work**, via - ctypes since PyTorch exposes no flags argument. It must target the *capturing* - stream; `torch.cuda.graph()` picks its own side stream unless passed one, and - getting that wrong returns `cudaErrorIllegalState` (401) rather than anything - descriptive. - -Two properties make the output trustworthy. Summed sites reconstruct the whole graph -to **1.3%** in the probe, so the decomposition is additive. And the instrumentation -costs **2.0%** of block time (7.543 ms instrumented vs 7.396 ms control, same -config, same node), which is small and, more importantly, measured rather than -assumed. Each event pair carries ~2.2 µs of its own overhead, read directly from a -`_calib` site wrapping an empty region in the same graph and subtracted from every -other site — without that, a 8 µs component reads ~30% high. - -### FLOOR-ATTRIB-S16 — where the 7.4 ms block goes - -Rank 0, per-layer medians net of pair overhead, scaled ×48. Two sampled layers (8 -and 24) with their spread, so an atypical layer is visible rather than silent. - -| component | µs/layer | ms/step | % of block | layer spread | -| --- | --- | --- | --- | --- | -| `expert_gemm_fc1` | 37.6 | 1.806 | 24.0% | 6% | -| `expert_gemm_fc2` | 26.5 | 1.272 | 16.9% | 18% | -| `attn_core` | 25.2 | 1.208 | 16.0% | 3% | -| `nvls_combine` | 14.1 | 0.677 | 9.0% | 81% | -| `nvls_dispatch` | 13.9 | 0.665 | 8.8% | 5% | -| `moe_sum` | 7.9 | 0.380 | 5.0% | 9% | -| `moe_align` | 7.8 | 0.373 | 5.0% | 3% | -| **sum of sites** | | **6.381** | **84.6%** | | -| unattributed (norms, router, QKV/O proj, RoPE) | | 1.180 | 15.6% | | - -**The NVLS collectives are finally sized: 1.34 ms/step, 17.8% of the block.** After -two hangs and two work-descriptor confounds, this is the first real number for them, -and it lands at the low end of the 1.5–2 ms the barrier-count argument predicted. -`nvls_combine`'s 81% layer spread (18.2 µs at layer 8 versus 10.0 at layer 24) and -its cross-rank variation are the expected signature of a spin-wait absorbing -arrival skew, not measurement noise. - -**Expert GEMMs are 3.08 ms, 40.8% — larger than the ~2.16 ms previously estimated**, -and the split is the interesting part. FC2 costs 70% of FC1 while reading half the -weight bytes (101 vs 201 MB/layer/rank) and doing half the FLOPs. Either FC2's -tiling is leaving bandwidth on the table or it is not weight-bound at all. - -**Attention is 25.2 µs/layer, not the 31.5 the chain delta implied.** Chain deltas -overstate, as suspected — one more reason not to quote them as component costs. - -### HBM-BW-S16 — the roofline denominator was wrong - -Measured achievable HBM bandwidth on this GB200 (184 GiB, 152 SMs): **6.85 TB/s** -sustained on a large bf16 copy, at both 2 GiB and 8 GiB working sets — 86% of the -8 TB/s spec figure, which is normal for a copy kernel and is the right number for a -roofline denominator. - -**Every roofline percentage in Sessions 14–15 used 8 TB/s and is therefore too -generous**, including "expert GEMMs at 84% of bound". Worse, redoing attention with -the correct denominator makes it *impossible*: 315 MB/layer of KV at 6.85 TB/s needs -45.9 µs/layer, and attention measures 25.2. A measured cost below its own computed -bound means the byte estimate is wrong, not that the kernel is superhuman — the -likely suspects are the assumed ~600-token average KV length and how many requests -are actually active per replay. **No attention roofline claim should be quoted until -that is resolved by counting bytes from the code rather than from the config.** - -### FULL-ATTRIB-S16 — the accounting closes, and the dense projections are the surprise - -Three more sites (`qkv_proj`, `out_proj`, `router_topk`) take the block from 84.6% -attributed to **101.8%** — sum of sites 7.747 ms against a 7.608 ms block. The 1.8% -over-count is the instrumentation's own inflation plus the constant-bias -approximation, so the block is now effectively fully accounted for. Ranks 0 and 2 -agree within 6% on every site. - -| component | µs/layer | ms/step | % of block | bf16 roofline @6.85 TB/s | % of bound | -| --- | --- | --- | --- | --- | --- | -| `expert_gemm_fc1` | 36.7 | 1.763 | 23.2% | 29.4 µs (201 MB) | 80% | -| `expert_gemm_fc2` | 26.1 | 1.253 | 16.5% | 14.7 µs (101 MB) | **56%** | -| `attn_core` | 24.6 | 1.180 | 15.5% | see HBM-BW-S16 | unresolved | -| `nvls_dispatch` | 17.1 | 0.819 | 10.8% | n/a (barrier) | n/a | -| `nvls_combine` | 13.5 | 0.649 | 8.5% | n/a (barrier) | n/a | -| `out_proj` | 12.1 | 0.581 | 7.6% | 2.4 µs (16.8 MB) | **20%** | -| `qkv_proj` | 10.6 | 0.508 | 6.7% | 3.1 µs (21 MB) | **29%** | -| `moe_align` | 8.7 | 0.418 | 5.5% | n/a | n/a | -| `moe_sum` | 7.9 | 0.377 | 5.0% | n/a | n/a | -| `router_topk` | 4.1 | 0.197 | 2.6% | n/a | n/a | -| **sum** | | **7.747** | **101.8%** | | | - -The norm fusions are not instrumented and the sum already exceeds the block, which -is its own small result: after QWEN-026/027/028 the norms no longer occupy -measurable standalone time. - -**The dense attention projections are the finding: 1.089 ms/step (14.3%) at 20-29% -of roofline.** Session 15 dismissed them at "0.23 ms, negligible" — that estimate -was 4.7× low because it was a *bandwidth* estimate for GEMMs that turn out to be -latency-bound. At M=256 tokens, `qkv_proj` (N=5120, K=2048) needs 3.1 µs on -bandwidth and ~2.4 µs on FLOPs, and takes 10.6; `out_proj` takes 12.1 against -2.4. Two skinny GEMMs, each ~3-5× off both of their own bounds. **~0.8 ms of -headroom, against a total gap of 1.14 ms** — and unlike the collectives, nothing -about this is cross-rank or barrier-shaped. - -Corrected headroom, all measured against 6.85 TB/s: - -| target | measured | bound | headroom | -| --- | --- | --- | --- | -| dense projections (qkv + out) | 1.089 ms | ~0.27 ms | **~0.8 ms** | -| NVLS collectives (overlappable, not removable) | 1.468 ms | n/a | up to 1.47 ms | -| `expert_gemm_fc2` | 1.253 ms | 0.787 ms | ~0.47 ms | -| `expert_gemm_fc1` | 1.763 ms | 1.44 ms | ~0.32 ms | - -### VLLM-COMM-S16 — vLLM runs the same AllGather+ReduceScatter, so "0 exposed comm" was wrong - -`PROFILE-OSL1024`, `PROFILE-DECODE` and the Session 8 recommendation all assert that -vLLM's MoE path has **0 exposed comm** and that mcore's dispatch/combine is therefore -pure deficit. Direct query of `nsys_trace/vllm_baseline_osl1024.sqlite` falsifies this. -vLLM's steady state contains exactly two collectives, and they are the same two -mcore uses: - -| | kernel | launches/rank | µs/launch | ms/step | -| --- | --- | --- | --- | --- | -| dispatch | `ncclDevKernel_AllGather_RING_LL` | 48,752 | 20.9 | 1.005 | -| combine | `ncclDevKernel_ReduceScatter_Sum_bf16_RING_LL` | 48,740 | 20.6 | 0.987 | - -48,752 = 48 layers × 1016 decode steps exactly — one pair per MoE layer per step, the -same cadence as mcore's `NVLSAllGatherVDispatcher`. There is **no all-to-all anywhere in -the trace**: no DeepEP, no pplx-kernels, no `nccl*AllToAll`. Note also that neither -framework uses an all-to-all dispatcher, so the recurring "all-to-all" shorthand for -this path (including in `QWEN-005b`) is a misnomer — it is AllGatherV in, ReduceScatterV -out, every token broadcast to every rank. - -The only difference is transport: NCCL ring/LL for vLLM, NVLink multimem on symmetric -memory for mcore. On the one clean measurement each way, mcore is *ahead*: 30.6 µs/layer -(`FULL-ATTRIB-S16`, in-graph events, unprofiled) against vLLM's 41.5 µs/layer. - -Two caveats bound how hard that comparison can be leaned on. mcore's collectives read -433 µs and 91 µs per launch **in the profile** — 25× nsys spin-wait inflation, per -`NSYS-OVERHEAD-S12`; only the in-situ number is usable. And vLLM's 41.5 µs is itself -measured under nsys, where RING_LL also spins, so its true cost lies somewhere between -the ~1.2 µs wire floor for ~1 MB over NVLink5 and 41.5 µs, and this trace cannot pin it. - -What survives is the structural claim, which is the useful one: **mcore has no -architectural disadvantage in the MoE collective** — same algorithm, same launch -cadence, better transport. And since both sides sit 15-30× above the wire time for the -bytes they move, **both are barrier-latency bound, not bandwidth bound**, which is why -QWEN-017's transport swap failed and why overlap, not a faster wire, is the only lever -left here. Correspondingly, the 1.47 ms is *not* free money relative to vLLM: the -differential is unmeasured and is certainly far smaller than 1.47 ms. - -### GAP-DECOMP-S17 — compute is at parity; the gap is kernel count. (Idle magnitude partly retracted — see GRAPH-LAUNCH-S17) - -> **Read the caveat first.** The per-component *kernel* comparison below is sound and is -> what identified the real lever. The **idle magnitude and its attribution are not**: -> `GRAPH-LAUNCH-S17` shows that node-level graph tracing inflates the host side of a -> graph launch ~20×, by an amount comparable to the entire idle budget, and mcore pays -> more of it than vLLM by having more graph nodes. Treat "77% of the gap is GPU idle" as -> unproven. The surviving claim — mcore runs ~42% more kernels per step, and cutting them -> pays — is confirmed by QWEN-036. - -First like-for-like decomposition of both sides: same window methodology, both traces -`--cuda-graph-trace=node`, steady-state decode only, one rank, step count derived from -the one-per-layer collective. Scripts: `nsys_trace/busy_union.py`, `gap_profile.py`, -`stall_detail.py`. GPU busy is the **union** of kernel intervals, not the sum — the sum -overcounts whenever streams overlap, which is exactly where vLLM wins. - -| per step, rank 0 | mcore | vLLM | delta | -| --- | --- | --- | --- | -| wall | 8.571 ms | 7.456 ms | **+1.115** | -| GPU busy (interval union) | 6.926 ms | 6.674 ms | +0.252 | -| GPU idle (wall − union) | **1.645 ms** | **0.783 ms** | **+0.862** | -| sum-of-kernel-durations | 6.970 ms | 7.030 ms | −0.060 | -| kernel launches | 1163 | 817 | +346 | - -**77% of the gap is GPU idle. mcore's compute is at parity or better.** Both profiled -wall times track the unprofiled throughput (mcore 8.571 vs 8.787, vLLM 7.456 vs 7.53), -so these windows are representative and nsys is not distorting the totals. - -Per-component, profiled against profiled (ms/step): - -| component | mcore | vLLM | delta | -| --- | --- | --- | --- | -| expert GEMM | 2.502 (`_fused_moe_kernel`, 96) | 2.411 (2× cutlass `bmm`, 94) | +0.09 | -| attention | 0.796 (FA2 sm100, 48) | 0.820 (`fmhaSm100f`, 48) | −0.02 | -| collectives | 0.847 | 1.629 | **−0.78** | -| dense GEMM + splitK reduce | 1.036 | 0.910 | +0.13 | -| routing / align / sum / topk | 0.559 | 0.702 | −0.14 | -| norms | 0.515 (193 launches) | 0.426 (189) | +0.09 | -| misc elementwise / rotary / kv | 0.643 (241+96) | 0.125 (48) | **+0.52** | - -**This retracts the Session 16 target list.** `FULL-ATTRIB-S16` put the dense projections -at 20-29% of roofline with ~0.8 ms of headroom; in the clean trace they are 1.036 ms -against vLLM's 0.910 — **0.13 ms apart, not 0.8**. The in-situ event pairs inflated every -site they measured (~2.2 µs/pair × 2 pairs × 48 layers), which is why the instrumented -sites summed to 7.747 ms against a clean 6.97 ms. **Do not size a target from in-situ -numbers; size it from a clean trace and only use in-situ for apportionment.** Likewise -`expert_gemm_fc2` at "56% of roofline" is not a vLLM deficit: vLLM's FC2 equivalent is -only 0.09 ms cheaper across *both* expert GEMMs. - -Where the 1.645 ms of idle sits (`gap_profile.py`): - -| gap band | mcore | vLLM | -| --- | --- | --- | -| >100 µs | 0.840 ms (2.1 gaps) | 0.399 ms (1.0 gap) | -| 20-100 µs | 0.163 ms (3.7) | ~0 | -| 5-20 µs | 0.070 ms (5.8) | 0.002 ms | -| 2-5 µs | ~0 | 0.343 ms (95.6) | -| <1 µs | **0.535 ms (1033 gaps)** | 0.038 ms (240) | - -Two co-equal levers, and one non-lever: - -1. **Two once-per-step host stalls, 0.776 ms** — a ~492 µs gap after - `at::native::index_elementwise_kernel` and a ~289 µs gap after - `CatArrayBatchedCopy_vectorized` (`torch.cat`), once each per step. vLLM's single - equivalent stall is 0.30 ms, so **~0.48 ms is available**. On the rare steps where - this blows out to 4 ms, the API trace shows the mechanism: **278 `cudaMemcpyAsync` + - 270 `cudaStreamSynchronize` in one gap** — a per-request D2H sync storm that should be - one batched copy. This is the between-step Python phase that HOST-IDLE-S13/14/15 - hunted and missed; GPU-side gap analysis found it where host sampling could not. -2. **Serialization, ~0.5 ms** — mcore has 1033 sub-µs gaps against vLLM's 240, and gets - essentially *no* stream concurrency (union 6.926 ≈ sum 6.970) while vLLM recovers - 0.36 ms through overlap (union 6.674 vs sum 7.030). Driven by kernel count: the +346 - excess launches are **~241 standalone elementwise kernels** (`elementwise_kernel` 96, - `triton_poi_fused_add_copy__0` 96, `vectorized_elementwise` 49 — the residual adds, - which vLLM fuses into its RMSNorm, cf. the `add` in - `triton_red_fused__to_copy_add_mean_mul_pow_rsqrt`), **90 `splitKreduce` kernels** that - vLLM has none of, and 96 rotary/append-kv against vLLM's 48. Caveat: node-level graph - tracing costs per node, so an unknown part of the 0.535 ms is CUPTI overhead that mcore - pays more of by having more nodes. The kernel-count difference underneath is real. -3. **Not a lever: the collectives.** mcore is 0.78 ms/step *ahead* on comm, consistent - with `VLLM-COMM-S16`. It is the single largest thing mcore is winning, and the reason - the compute totals come out even despite the elementwise sprawl. - -### Session 16 — conclusion & recommendation - -The gap to vLLM is ~1.14 ms/step (29,133 vs 33,994 tok/s at a ~7.97 ms step). The -block is now fully attributed, and there is more identified headroom (~1.6 ms -excluding the collectives) than gap. Ordered by headroom-per-unit-risk: - -0. **The dense attention projections, ~0.8 ms of headroom** (FULL-ATTRIB-S16) — the - cheapest target and the one this session found by accident. Two skinny - M=256 GEMMs at 20-29% of their own bounds, latency-bound rather than - bandwidth-bound, with no cross-rank behaviour to reason about. Start here. -1. **The NVLS collectives, 1.47 ms — but demoted, see VLLM-COMM-S16.** vLLM runs the - same AllGather+ReduceScatter at the same cadence and pays for it too, so this is - not 1.47 ms of deficit against vLLM; the differential is unmeasured and much - smaller. Removal is impossible (they carry the tokens) and replacing the transport - already failed once (QWEN-017), and both sides are barrier-latency bound rather - than bandwidth bound, so the only lever is **overlap**: chunk the expert GEMM and - overlap combine of chunk *i* with the GEMM of chunk *i+1*, or prefetch the next - layer's dispatch. The 81% layer spread says part of this cost is skew absorption, - so some of it may be recoverable by better arrival balance rather than by - overlap. Worth doing eventually, but it is no longer the item that "alone would - close the gap." -2. **FC2, ~0.3–0.6 ms of apparent inefficiency** — 26.5 µs/layer against 14.8 µs at - the corrected roofline, and against FC1 doing twice the work for 1.4× the time. - Cheaper to investigate than item 1 and independently useful. Note - `MCORE_MOE_GEMM_TUNE` already tunes FC1/FC2 tiles separately, so this is a - question about that tuning, not about its absence. -3. **Recount the attention KV bytes** before any attention work, per HBM-BW-S16. -4. Precision (FP8 KV cache, FP8/NVFP4 experts) remains the only lever on whatever - genuinely is bandwidth-bound, and still needs its own correctness evaluation. - -### GRAPH-LAUNCH-S17 — the 191 µs `cudaGraphLaunch` was a profiler artifact; real cost is 9 µs - -`GAP-DECOMP-S17` decomposed mcore's 1.645 ms/step of GPU idle and found -`cudaGraphLaunch` inside those gaps at **4 calls/step × 191.6 µs = 761 µs/step**, which -would have made it the single largest item in the idle budget. Measuring it without a -profiler (`megatron/core/inference/graph_launch_timing.py`, `MCORE_GRAPH_LAUNCH_TIMING=1`, -`perf_counter` around `torch.cuda.CUDAGraph.replay`) shows the hot graph replaying in -**median 9.0-10.0 µs, p90 ~12 µs** over 1404 replays. The profiled figure was inflated -**~20×** by `--cuda-graph-trace=node`, which makes the driver do per-node bookkeeping -inside the launch. - -Graph launch is therefore ~40 µs/step, not 761 — **not a lever**, and the two 460-1000 µs -replays that do exist are `n=1` first-touch of each freshly captured bucket, not steady -state. - -**This partially retracts `GAP-DECOMP-S17`.** Its per-component *kernel* numbers stand — -kernel durations are barely perturbed, and that half of the analysis is what identified -the kernel-count gap. But the **magnitude and attribution of the 1.645 ms idle are not -trustworthy**, because the host side of a node-traced graph launch is inflated by an -amount that happens to be the same order as the whole idle budget, and mcore pays more of -it than vLLM by having more nodes. Do not quote "77% of the gap is GPU idle" without this -caveat. What survives, and is confirmed by QWEN-036 below, is the weaker and still -actionable claim: **mcore runs ~42% more kernels per step than vLLM and that costs real -throughput.** - -Method note for the next session: a profiler cannot be used to size an overhead the -profiler itself creates. Any host-side cost read out of a node-traced profile must be -re-measured with `perf_counter` before it becomes a target. - -### FUSION-INERT-S17 — the residual-add fusion was written, gated on, and never firing - -`QWEN-026/027/028` built `fused_add_rmsnorm` for both transformer-block boundaries. -`MCORE_FUSED_ADD_NORM=1` was in the standing gate list and the config that produced -`nsys_trace/mcore-final-jul27` is labelled "all ten gates on" — yet that trace contains -**zero** fused-add-norm kernels and exactly the two-kernel pattern the module's own -docstring describes as the thing it replaces (`triton_poi_fused_add_copy__0` at 96/step = -2/layer, plus 193 standalone TE RMSNorms). - -`megatron/core/inference/fusion_diag.py` (`MCORE_FUSION_DIAG=1`) reports each site's -guard conjunction once per distinct verdict. Two findings: - -1. Both sites **do engage** at 256 tokens. The blocked verdicts are all at - `(512, 1, 2048)` — prefill, correctly excluded by `MCORE_FUSED_ADD_NORM_MAX_TOKENS=256`. -2. The second site, `mlp_bda+next_layer_qkv_norm`, needs `MCORE_FUSED_ADD_NORM_QKV=1`, - which **was never in any gate list**. Worse, the two sites are *coupled*: - `_forward_pre_mlp_layernorm` only sets `self.mlp_norm_manager` when the attention-side - fusion did *not* fire, and the MLP-side guard requires `mlp_norm_manager is None`. So - the MLP-side fusion is eligible only on layers where the attention-side one already - fired — a dependency neither gate's name nor docstring mentions. - -Diagnostic lesson: report guards **once per distinct verdict, not once per site**. The -first version reported only the first call, which is prefill, and therefore said "BLOCKED" -about a fusion that engages fine in decode — the exact wrong conclusion. - -### QWEN-036 — enable the second fusion boundary (`MCORE_FUSED_ADD_NORM_QKV`): **+2.8%** - -Same node, same session, same allocation, back-to-back arms; only the one gate differs. -Node `nvl72104-T17`, Slurm `5792909`, BS256 / OSL1024 / 2 warmup + 5 timed. - -| arm | gates | per-iter throughput (tok/s) | mean | TPOT | -| --- | --- | --- | --- | --- | -| `a1base` | standing 12 + FA2 | 26803, 26871, 27111, 27176, 27947 | **27,182** | 9.420 ms | -| `a2qkv` | + `MCORE_FUSED_ADD_NORM_QKV=1` | 27589, 27718, 28092, 28172, 28200 | **27,954** | 9.159 ms | - -**+2.84%, −0.261 ms/step.** 3 of 5 `a2qkv` iterations beat *every* `a1base` iteration, and -the arms' means are separated by more than either arm's spread, so this clears the ~4% -per-iteration noise. Coherence checked in both arms. - -Note the absolute numbers: `a1base` is 27,182 on this node against the 29,133 recorded for -the same configuration in Session 16. **Node-to-node variation is larger than most of the -wins in this ledger** — roughly 7%. Never compare an arm against a historical number; -always re-measure the baseline in the same allocation. The gap to vLLM cannot be restated -from this session because no same-node vLLM run was taken. - -This is the first confirmation of the `GAP-DECOMP-S17` strategy: the change removes -kernels rather than making any kernel faster, and it converted directly into throughput. - -### QWEN-037 / QWEN-038 — two env-only follow-ups, both neutral - -Same allocation as QWEN-036, so each is directly comparable to the `a2qkv` arm -(27,954 tok/s) without re-running a baseline. - -| arm | change | mean tok/s | vs `a2qkv` | -| --- | --- | --- | --- | -| `b1maxtok` | `MCORE_FUSED_ADD_NORM_MAX_TOKENS=512` | 27,862 | −0.33% | -| `b2nosplitk` | `CUBLASLT_WORKSPACE_SIZE=1` | 27,811 | −0.51% | - -**QWEN-037** aimed at the 9.2 boundaries/step that still fall back to the unfused path -(96 possible − 86.8 fused). Raising the token ceiling did not recover them and cost -slightly, so those fallbacks are not token-count rejections. One is layer 47, whose -`mlp_bda` has no next layer to donate a norm to; the rest are most likely the coupling in -`FUSION-INERT-S17` (a layer is eligible on the MLP side only if the attention side fired). - -**QWEN-038** tried to delete the ~76 `splitKreduce` launches/step (0.185 ms of kernel -time) by starving the cuBLASLt workspace that split-K algorithms require. Neutral. Two -readings, and the second is the important one: either the heuristic ignored the limit, or -it switched to non-split-K kernels that gave back exactly what the reduce cost. **The run -did not verify which**, so "split-K removal doesn't pay" is not yet established — only -"this way of attempting it doesn't." A profile of the arm would settle it. - -### LAUNCH-VS-WORK-S17 — the QWEN-036 win came from removed GPU work, not removed launches - -Worth pinning down, because the two readings of QWEN-036 point at completely different -next targets, and the launch-count reading is the wrong one. - -The fusion removed **111 launches/step** *and* **~0.386 ms/step of kernel time** in the -four affected kernel families: - -| family | before (n × µs) | after (n × µs) | -| --- | --- | --- | -| `triton_poi_fused_add_copy__0` | 96 × 1.27 | 9.2 × 2.35 | -| `at::native::elementwise_kernel<128,4>` | 96 × 2.56 | 52.1 × 3.11 | -| `rmsnorm_fwd_tuned` | 85.1 × 2.73 | 9.8 × 3.36 | -| `rmsnorm_fwd_general` | 96 × 1.95 | 8.3 × 3.50 | -| `_fused_add_rmsnorm_kernel` | — | 86.8 × 1.79 | -| **total** | **0.787 ms** | **0.401 ms** | - -The measured end-to-end saving was **0.261 ms/step — 68% of the 0.386 ms of kernel time -removed**, which is what partial overlap predicts. It is *not* explained by launch count: -QWEN-038 attacked 76 launches carrying 0.185 ms and returned nothing, and -`GRAPH-LAUNCH-S17` already showed launch submission costs ~10 µs, not the ~180 µs the -launch-bound reading needs. - -**So size every remaining candidate by the GPU work it deletes, not by how many kernels it -deletes.** Under that rule the routing-mask kernel — 48 launches/step but only -48 × 0.97 µs = 0.047 ms — is worth at most ~0.5% and is *not* the next lever, despite -being the most attractive item on a launch-count ranking. This also retires the last -operational use of `GAP-DECOMP-S17`'s "+346 excess launches" framing: the launch delta is -a symptom of unfused work, and the work is what costs. - -Ranked by kernel time, the remaining decode candidates are `_moe_sum_kernel_fast` -(48 × 5.81 = 0.279 ms), `_align_single_kernel` (39.5 × 4.67 = 0.185 ms), the residual -`elementwise_kernel<128,4>` (52.1 × 3.11 = 0.162 ms, identity not yet established), and -`_fused_qk_rmsnorm` (0.105 ms). Caveat on all four: these averages come from a profile -whose window was not restricted to steady-state decode, so they are inflated by prefill -instances and are ranking hints, not budgets. - -### QWEN-039 — raise `CUDA_DEVICE_MAX_CONNECTIONS` from 1 to 8: **+0.9%, provisional** - -`GAP-DECOMP-S17` observed that mcore gets essentially no stream concurrency (interval -union 6.926 ms ≈ sum-of-durations 6.970 ms) while vLLM recovers 0.36 ms through overlap. -The harness pins `CUDA_DEVICE_MAX_CONNECTIONS=1` (`dev/moe_fused/run_e2e_cfg.sh`), which -puts every stream on one hardware channel and serializes work that has no data dependency -— notably the shared-expert stream against the main stream. Env-only test: - -| arm | `CUDA_DEVICE_MAX_CONNECTIONS` | mean tok/s | vs `a2qkv` | -| --- | --- | --- | --- | -| `a2qkv` | 1 | 27,954 | — | -| `c2conn4` | 4 | 27,999 | +0.16% | -| `c1conn8` | 8 | **28,214** | **+0.93%** | - -**Provisional, not confirmed.** All 5 of `c1conn8`'s iterations sit above `a2qkv`'s mean, -which is suggestive, but the two arms' per-iteration spreads overlap (553 and 611) and the -allocation expired before a repeat could be run. **Re-run `a2qkv` and `c1conn8` back to -back before believing +0.9%**, and note that `=1` is often set deliberately to keep -collectives from being starved of channels, so a confirmation should check the collective -timings too, not just throughput. - -Session best configuration: standing 12 gates + FA2 + `MCORE_FUSED_ADD_NORM_QKV=1` + -`CUDA_DEVICE_MAX_CONNECTIONS=8` = **28,214 tok/s, +3.80%** over the same node's baseline -of 27,182. - -### BUDGET-S17 — first trustworthy per-kernel decode budget, and how to window a trace - -Every per-kernel table before this one was taken over a window that included model load, -graph capture, warmup and the gaps *between* benchmark iterations. That inflates -`wall/step` without bound and pollutes every average with prefill instances — it is why an -earlier pass of this same trace reported `wall/step=199 ms` and a 6.5 ms -`_fused_metadata_kernel`. `nsys_trace/steady_window.py` locates the window from the cadence -of the one-per-layer-per-step collective, takes the last *N* whole steps, and **warns when -the span still contains an inter-iteration gap**. Narrowing 150 → 60 steps moved -`wall/step` from 34.6 ms to 7.217 ms and `_fused_metadata_kernel` from 13,182 µs to 116 µs. - -> **Always run the window guard.** Two sessions' worth of component sizing was distorted by -> unwindowed spans. If `wall/step` does not land near the throughput-derived step time, -> the window is wrong and nothing else in the table means anything. - -Config: both fusion boundaries on, FA2, BS256, **OSL128** (`wall/step` 7.217 ms; the -OSL1024 regime is 9.16 ms, the difference being KV-length-dependent attention). 60 steps, -rank 0, `sum-of-durations` 5.627 ms/step over 1025 launches. - -| ms/step | n/step | avg µs | group | -| --- | --- | --- | --- | -| 1.949 | 96 | 20.31 | `_fused_moe_kernel` — expert GEMMs, **35% of all kernel time** | -| 0.769 | 34-14 | 3.4-6.7 | dense projection GEMMs (`nvjet_*`) | -| 0.759 | 48+48 | 9.33/6.49 | NVLS collectives (reduce-scatter 0.448 + all-gather 0.311) | -| 0.558 | 34-14 | 4.3-10.4 | FlashAttention decode + splitkv combine | -| 0.218 | 48 | 4.54 | `_moe_sum_kernel_fast` | -| 0.217 | 48 | 4.52 | `_align_single_kernel` | -| 0.179 | 48+34 | 2.05/2.33 | `cublasLt::splitKreduce` | -| 0.168 | 95 | 1.77 | `_fused_add_rmsnorm_kernel` (the QWEN-036 fusion) | -| 0.127 | 48 | 2.66 | `at::native::elementwise_kernel<128,4>` — 1/layer, **identity unknown** | -| 0.085 / 0.084 / 0.070 / 0.062 / 0.036 | 48 each | ~0.8-1.8 | qk_rmsnorm, rotary, softmax_topk, append_kv, mask_routing_padding | - -Sized by `LAUNCH-VS-WORK-S17`'s rule (removed GPU work × ~68% conversion, against the -9.16 ms OSL1024 step), the ranked remaining candidates are: - -1. **`_fused_moe_kernel`, 1.949 ms — the only large target left.** No fusion helps here; - it is bandwidth-bound (FC1 at 80% of achievable HBM, FC2 at 56%). The lever is fewer - bytes, i.e. **FP8 expert weights**, worth up to ~0.9 ms. This is the one item whose - payoff justifies a large implementation. -2. `_moe_sum_kernel_fast` + `_align_single_kernel`, 0.435 ms combined → ~1.5-3% if folded - into neighbours. Both are already single-kernel-per-layer, so this means changing what - they are fused *with*, not fusing them together. -3. `elementwise_kernel<128,4>`, 0.127 ms → ~0.9%. **Identify it first**: one per layer at - 2.66 µs, and if it is a stray cast or copy it may be removable outright rather than - fused. Cheapest next step in the list and the only one that starts with a question - rather than a kernel. -4. `_mask_routing_padding_kernel`, 0.036 ms → ~0.3%. Explicitly **not worth it**, recorded - here only because a launch-count ranking puts it near the top (48/step). - -### FP8-WEIGHTS-S18 — candidate 1 above is dead: fewer bytes does not mean less time - -Weight-only fp8 (e4m3, one fp32 scale per output channel, activations and accumulation -unchanged) was implemented end to end — `moe/fp8_experts.py`, an `FP8_WEIGHTS` path in -`_fused_moe_kernel`, quantized buffers built in `_build_concatenated_weights` — and it -**loses at every tile configuration**, so the 0.9 ms above is not available this way. - -| GEMM | best bf16 | best weight-only fp8 | fp8 penalty | -| --- | --- | --- | --- | -| FC1 (+SwiGLU epilogue) | **36.24 µs** @ M16 N64 K64 w4 s3 | 40.08 µs @ M32 N64 K128 w8 s4 | +10.6% | -| FC2 | **18.86 µs** @ M32 N256 K64 w8 s4 | 22.24 µs @ M32 N256 K128 w4 s3 | +17.9% | - -Numerics were never the problem: cosine similarity 0.9989, mean relative error 6.6e-2 -concentrated in near-zero outputs, weight bytes 302.0 → 151.5 MB/layer exactly as intended. - -**Why halving the bytes cannot help here.** With bf16 activations every fp8 weight tile must -be widened back to bf16 in registers before the MMA, one convert per weight element. The -best bf16 configs run at 5.55 TB/s (FC1) and 5.34 TB/s (FC2) — ~90% of this part's -achievable HBM — so the byte saving is worth at most ~45% of each kernel, and the converts -cost more than that. The fp8 arms land at 40/22 µs against an fp8 *bandwidth* floor of -~18/9 µs, i.e. once the bytes are halved the kernel is no longer bandwidth-bound at all. -The only version of this lever that can pay is **w8a8**, where activations are also fp8 and -the tensor cores consume both operands directly with nothing converted. - -> **Two measurement traps, both of which produced a wrong answer first.** -> 1. An earlier harness reported fp8 at 0.995× and was read as "no effect". It was timing -> **Python**: four Triton launches per call, ~153 µs/call, host-bound with the GEMM -> invisible underneath. Any A/B of a <50 µs kernel must time **graph replay**, not a -> launch loop — `harness_gemmgrid.py` captures one launch and replays it. -> 2. `_get_decode_tuned_configs`' shipped tiles were re-derived independently by this sweep -> and are **already optimal**, including the shared `BLOCK_SIZE_M=16`: M32 makes FC2 13% -> faster (21.33 → 18.86 µs) but FC1 10% slower, and the two share one indirection table, -> so M16 wins on the sum (57.6 vs 58.8 µs). Retuning tiles is closed. - -Combined with QWEN-014 (flashinfer cutlass, 0.921×) and QWEN-010 (torch grouped_mm, 0.82×), -**every dtype-preserving and vendor-kernel route into the expert GEMM has now been measured -and none beats the shipping Triton path.** The candidate list above should be read as: item 1 -requires w8a8 or nothing; items 2-4 are the only cheap ones left. - -### STATUS-S18 — mcore at 87.9% of vLLM; the remaining gap is packing, not work - -| | tok/s | ms/step | vs vLLM | -| --- | --- | --- | --- | -| vLLM DP4+EP (`VLLM-BASELINE` / `PROFILE-OSL1024`) | 33,994.5 | 7.53 | — | -| **mcore, current best** (`y2pdl`) | **29,898.3** | **8.56** | **87.9%** (vLLM 1.137x) | -| mcore at session-18 start | 28,594.3 | 8.95 | 84.1% | -| mcore at session-2 start (`SESSION2-BASE`) | 22,398.9 | 11.71 | 65.9% | - -Session 18 closed 0.35 ms of the 1.38 ms/step gap (**+4.56%**), from three changes that -stack: `QWEN-040` bf16 MoE combine (+2.5%), `QWEN-041` flashinfer trtllm-gen decode -(+1.7%), `QWEN-042` PDL on that kernel (+0.3%). - -**Attention is now at parity, and that closes the largest work-bucket gap.** In the final -trace it runs 22.9 us/launch against vLLM's 25.3 us -- both engines now execute the same -trtllm-gen kernel family, and what is left of the difference is window seqlen, not -kernel quality. The 0.409 ms/step deficit `GAP-S18` measured is gone. - -What remains, in order: - -1. **Serialization, ~0.3 ms.** mcore's sum-of-durations (5.922 ms) and interval union - (5.877 ms) differ by 0.8%, so **almost nothing overlaps**; vLLM hides 4.8% of its work. - mcore's 0.905 ms/step of collectives is fully exposed against compute. This is now the - single largest structural difference and it is a scheduling change, not a kernel. -2. **Dispatch overhead, 0.416 ms.** 806 sub-microsecond gaps per step, at 963 launches - against vLLM's 810. Priced by `BALLAST-S18` at ~0.5 us of gap per node, so each fusion - still pays twice. -3. **Work vLLM does not do at all**: splitK reduce 0.167 ms/step (82 launches, no vLLM - equivalent) and the 0.164 ms of copies whose origin is still open after four attempts - (`COPY-ID-S18` -- `SplitAlongDim` eliminated). - -> **Reading the trace tables above.** Per-bucket ms/step may only be compared *within* one -> trace. This window's attention is 22.9 us/launch against the harness's 45.1 us at KV -> length 512, which puts its average KV length near 260 -- so its absolute times are -> cheaper than the run average and not comparable to the earlier traces'. The throughput -> column is the authority; the buckets rank levers, they do not size them. - -### QWEN-042 — PDL on the flashinfer decode kernel: **+0.30%** - -| arm | gates | tok/s | vs control | -| --- | --- | --- | --- | -| x2fi | control (flashinfer decode) | 29,802.9 | — | -| y1pdl | + `MCORE_FLASHINFER_PDL=1` | 29,888.2 | +0.29% | -| y2pdl | repeat | 29,898.3 | +0.32% | - -Programmatic Dependent Launch lets the kernel's prologue start while its predecessor -drains, aimed at the 0.416 ms/step of sub-microsecond gaps. Small but reproducible across -two arms, and above the 0.5% iteration spread only because both arms agree; kept on the -strength of the repeat rather than the single measurement. - -### QWEN-041 — Blackwell-native decode attention via flashinfer trtllm-gen: **+2.61%** - -| arm | gates on top of standing set (incl. `MCORE_NVLS_RS_BF16=1`) | tok/s | vs control | -| --- | --- | --- | --- | -| x1base | control (FA2 decode) | 29,043.8 | — | -| **x2fi** | **`MCORE_FLASHINFER_DECODE=1`** | **29,802.9** | **+2.61%** | - -`GAP-S18` put attention 0.409 ms/step behind vLLM at an *identical* launch count, which -rules out anything Megatron does around the call and points at the kernel generation: -Megatron's decode runs FA2's `flash_attn_with_kvcache`, whose kernels predate Blackwell, -while vLLM reaches `fmhaSm100f` through flashinfer. flashinfer 0.6.14 turned out to be -**already installed** in the same venv as flash-attn, so this needed no new dependency. - -Measured in isolation first (`dev/moe_fused/harness_attn.py`), at B=256, 32 q heads, -4 kv heads, D=128, under graph replay so the number is device time and not Python: - -| KV len | FA2 | trtllm-gen | delta | -| --- | --- | --- | --- | -| 512 | 59.6 us | 45.1 us | −24.3% | -| 1024 | 105.8 us | 80.7 us | −23.7% | -| 2048 | 199.3 us | 151.9 us | −23.8% | - -A flat ~24% across the range, with outputs matching FA2 to bf16 tolerance (rel ~3e-3). -Two properties made the integration small: with `kv_layout="NHD"` the kernel takes -**exactly the paged layout Megatron already has**, and unlike flashinfer's wrapper APIs -the trtllm-gen entry point needs no host-side `plan()`, so it captures in the decode -graph. Page size is irrelevant to it (256/128/64 all within 0.1 us), so Megatron's -256-token pages stay -- and FA2 is the stricter of the two, *requiring* pages be a -multiple of 256. - -The e2e gain (+2.61%) is smaller than the kernel gain (24% of a 1.36 ms bucket would be -0.33 ms of an 8.8 ms step, or 3.8%) because the removed device time partly un-hides host -work, the same ~68% conversion `LAUNCH-VS-WORK-S17` measured for the add-norm fusion. - -> **Lesson.** The gap said "attention, same launch count, more time", which reads as *our -> kernel is the wrong generation for this GPU* -- and the right kernel was already sitting -> in the venv. Before treating a kernel-quality gap as a porting project, check what the -> reference implementation calls and whether it is installed: two sessions of FA2-vs-FA4 -> flag flipping never left the flash-attn package, where the answer was not. - -### QWEN-040 — reduce the MoE combine in bf16, not fp32: **+2.37%** - -| arm | gates on top of standing set | tok/s | vs control | -| --- | --- | --- | --- | -| s1base | control (FA2) | 28,594.3 | — | -| **t1rsbf16** | **`MCORE_NVLS_RS_BF16=1`** | **29,271.8** | **+2.37%** | -| t2rsbf16 | repeat of t1 | 29,354.2 | +2.66% | -| t3rsconn8 | + `CUDA_DEVICE_MAX_CONNECTIONS=8` | 29,346.3 | +0.0% vs t2 | -| s2fa4 | FA4 (version pin removed) | 27,716.1 | −3.07% | -| s3conn8fa4 | FA4 + `CUDA_DEVICE_MAX_CONNECTIONS=8` | 27,812.6 | +0.35% vs s2fa4 | - -Found by following `COPY-ID-S18`'s post-reduce-scatter copy back to its cause rather -than trying to fuse it away. The copy is `output.to(torch.bfloat16)` in -`NVLSAllGatherVDispatcher.combine`, and it exists because the `ep_rsv` symmetric buffer -is allocated fp32 — so the MoE writes fp32 through `_moe_sum`'s `out=`, the combine -reduce-scatter moves **twice the NVLink bytes it needs to**, and every layer pays a cast -on the way back to a bf16 residual stream. Allocating the buffer bf16 fixes all three at -once: `_moe_sum`'s `tl.store` casts on the way in for free, and `output.to(bfloat16)` -becomes a no-op that returns its argument, so the cast kernel disappears. - -The precision cost is far smaller than "reduce in bf16" suggests, and this is the part -worth remembering: `multimem.ld_reduce` **accumulates in f32 regardless** — `REDUCE_F32` -only selects whether the operands it loads are `f32` or `bf16x2`. So the change halves -the bytes on the step's largest collective while keeping f32 accumulation in hardware, -and it lands where vLLM already is (`ncclDevKernel_ReduceScatter_Sum_bf16`). Greedy -decode stays coherent on all three probes and the five iterations span 0.4%. - -Two arms, 29,271.8 and 29,354.2, put the mean at 29,313 (**+2.52%**), and TPOT drops from -~8.95 to 8.75 ms — 0.23 ms/step against the 0.29 ms predicted from the trace, so the -accounting holds. Stacking `CUDA_DEVICE_MAX_CONNECTIONS=8` on top adds nothing (29,346), -which retires `QWEN-039`: that flag's earlier +0.4-0.9% was relieving contention on the -same fp32 reduce-scatter this change shrinks, so the two are one win, not two. - -Left env-gated and **off by default**: it is a numerics change, and enabling it by -default needs an accuracy run (lm-eval or equivalent), not three greedy-decode probes. - -**Mechanism confirmed by launch count, not by time.** A trace captured with the gate on -shows `bfloat16_copy_kernel` instances falling from 258,564 to 4,572 whole-trace (−98%) -and the copy bucket falling from 107 to 59 launches/step — exactly 48, one per layer, -the cast that is now a no-op. Counts are the right evidence here because the two traces' -*times* are not comparable: the second window's expert GEMM is 27% cheaper (2.436 → -1.777 ms/step) and its attention launches went 50 → 63, neither of which a -reduce-scatter dtype can cause. The two 60-step windows simply sit at different average -sequence lengths, so per-bucket ms/step may only be compared *within* a trace. The -throughput number in the table above, measured without a profiler attached, is what -sizes the win. - -> **Lesson.** A "removable copy" is usually a symptom. Three sessions treated the -> per-layer copies as fusion targets worth ~0.06 ms each; the copy was actually pointing -> at a **buffer dtype** decision upstream worth 2.5%. When a copy shows up next to a -> collective, check what dtype the collective is moving before trying to fuse the copy. - -### FA4-RECHECK-S18 — the FA2 pin survives upstream's split-KV fix - -`QWEN-0xx` pinned `MCORE_FLASH_ATTN_VERSION=2` because FA4 hardcoded `num_splits=1` and -so could not split KV across SMs at decode. The rebase brought `num_splits=0` (auto), so -the rejection was re-tested per the rebase rule — and FA4 is **still 3.07% slower** -(27,716 vs 28,594). `num_splits` was therefore not the dominant term; the FA4 path's -routing through the varlen interface is. Attention stays the largest work-bucket gap -(+0.409 ms/step) with no env-level lever left, so closing it needs a different kernel -(flashinfer/TRT-LLM-gen decode), not a flag. - -### GAP-S18 — the gap is two-thirds packing, one-third work, and the expert GEMM is done - -First decomposition built from a **fresh** mcore trace (current gate set, OSL1024, 60 -steady steps, no window warnings) against the vLLM baseline trace, both bucketed by -`nsys_trace/compare_budget.py` and both measured for interval-union busy by -`nsys_trace/union_window.py`. Every prior version of this comparison used the Jul 27 -mcore trace, which predates the FA2 pin and both add+norm fusions and therefore -overstated attention, norm and copies. - -| | vLLM | mcore | delta | -| --- | --- | --- | --- | -| GPU busy, interval union | 6.989 ms | 7.478 ms | +0.489 | -| exposed gap (true step − union) | ~0.54 ms | ~1.43 ms | **+0.89** | -| step time, unprofiled | 7.53 ms | 8.91 ms | +1.38 | - -Per-bucket device time (ms/step, launches/step), mcore minus vLLM: - -| bucket | vLLM | mcore | delta | -| --- | --- | --- | --- | -| attention | 1.214 (48) | 1.623 (50) | **+0.409** | -| elementwise / copy | 0.004 (1) | 0.232 (107) | +0.228 | -| splitK reduce | 0 (0) | 0.204 (95) | +0.204 | -| MoE finalize | 0.215 (47) | 0.261 (48) | +0.046 | -| expert GEMM | 2.376 (94) | 2.436 (96) | **+0.060** | -| norm | 0.251 (95) | 0.272 (145) | +0.021 | -| MoE routing | 0.475 (142) | 0.367 (144) | −0.108 | -| collective | 1.607 (94) | 1.035 (96) | **−0.572** | -| **sum of durations** | **7.341 (810)** | **7.521 (1027)** | **+0.180 (+217)** | - -Three conclusions, in order of how much they should change what gets worked on: - -1. **Only 0.18 ms of the 1.38 ms gap is extra kernel work.** The expert GEMM is at - parity (+0.06 ms) and mcore's collectives are *half a millisecond faster* than - vLLM's. Four sessions of MoE GEMM work, `FP8-WEIGHTS-S18` included, were aimed at a - bucket that had already converged. -2. **mcore overlaps nothing.** Sum-of-durations 7.521 vs union 7.478 means 0.6% of its - work runs concurrently with other work; vLLM hides 0.352 ms (4.8%). mcore's 1.035 ms - of collectives is fully serialized against compute. Matching vLLM's overlap ratio is - worth ~0.3 ms and is a scheduling change, not a kernel. -3. **The gaps are launch-shaped.** 892 of mcore's 934 per-step gaps are sub-microsecond - and sum to 0.456 ms — consistent with `BALLAST-S18`'s 1.34 µs/node (~0.85 µs kernel - + ~0.5 µs gap) across 1027 launches vs vLLM's 810. Every removed launch pays twice. - -### COPY-ID-S18 — both copy families named from the trace, after three failed hooks - -The 107 copy launches/step were finally identified **positionally in the trace**, by -aggregating each copy's immediate predecessor and successor, after three attempts to -name them with a `TorchDispatchMode` hook produced nothing usable: - -| n/step | kernel | sits between | what it is | -| --- | --- | --- | --- | -| 47 | `vectorized_elementwise_kernel<8>` (bf16 copy) | `_multimem_reduce_scatter_v` → `_fused_add_rmsnorm` | materializes the MoE output that the next kernel immediately reads | -| 48 | `elementwise_kernel<128,4>` | QKV GEMM (`nvjet_*`) → `_fused_qk_rmsnorm` | contiguous copy of the strided Q/K slices | - -The first one turned out to be a dtype symptom, not a fusion target, and is fixed in -`QWEN-040` (+2.5%). The second is still open, and the guess in the first version of this -entry -- "make `_fused_qk_rmsnorm` stride-aware" -- was wrong: that kernel is *already* -stride-aware and takes no-copy `[-1, head_dim]` views. So the copy is **upstream of the -norm**, between the QKV GEMM and it, and at 2.66 us it is about the size of a full -round-trip of the QKV tensor (256 tok x 5120 x 2 B, read+write, at ~3 TB/s). - -The prime suspect was `SplitAlongDim` in `Attention.get_query_key_value_tensors` -materializing q/k/v instead of returning views. **Tested and eliminated**: taking the -`torch.split` branch instead (`MCORE_QKV_SPLIT_VIEWS=1`, arm `u1qkvsplit`) gives -29,307 tok/s against the 29,313 control -- exactly neutral. Either TE's split already -returns views, or `torch.split` + the `[sq, b, np, hn]` reshape materializes the same -bytes. The copy's source remains open; what is now known is that it is QKV-sized, sits -between the QKV GEMM and the norm, and is not the split. - -A fourth attempt at naming it with `copy_trace` landed in prefill again -- it reported -`rmsnorm.py:182 op_forward` doing `.contiguous()` on a strided `(512, 1, 4, 128)` K -tensor, which is TE's norm op on a 512-token prefill, not the decode copy. That is a -real finding for prefill but not the one being chased, and it is the fourth time this -hook has answered a different question than the one asked. - -> **Lesson on the failed hook.** `copy_trace` was armed three ways: on the first layer -> forward (caught the one-time lazy expert-weight consolidation and reported 32 weight -> copies per layer as if per-step), on `inference_context.is_decode_only()` (never fired -> — under graph capture the context does not reach the layer as a keyword), and finally -> on a call counter. Meanwhile the trace already contained the answer. **When a kernel -> needs identifying and a profile exists, read the profile's neighbours first**; -> instrumenting the framework to re-derive what the trace already recorded cost three -> allocation slots here. - -## Session 19 (2026-08-03) — the QKV copy, named by arithmetic - -`COPY-ID-S18` left one kernel unattributed after four attempts: a 2.8 us -`elementwise_kernel<128,4>` between the QKV GEMM and the q/k norm, once per layer per -step, QKV-sized. It is the `reshape`, and it is provable on a laptop in ten seconds with -no GPU, no cluster, and no profiler: - -``` -q slice [sq, b, ng, (np/ng) * hn] strides (.., .., 1280, 1) -q reshape [sq, b, np, hn] strides (.., .., 128, 1) COPY MATERIALIZED -``` - -Megatron's QKV projection writes `[sq, b, ng, (np/ng + 2) * hn]` -- each group's k and v -head sit *between* consecutive groups' q heads. Merging the group axis with the head axis -therefore requires the group stride to equal `(np/ng) * hn`; it equals -`(np/ng + 2) * hn`. For Qwen3-30B that is **1024 required against 1280 actual**, so -`Tensor.reshape` cannot return a view and silently copies the whole query instead. - -### Why four attempts missed it - -All four tested *the split*: `SplitAlongDim` against `torch.split` (`u1qkvsplit`, exactly -neutral), then three `copy_trace` arms. The neutral split result was read as "the copy is -elsewhere", when it was itself the answer -- both split branches feed the same reshape, so -of course they measured the same. The one line never questioned was the one carrying a -comment that explained it as a shape change. - -> **Lesson.** A layout claim is decidable from strides alone, and `reshape` reports a copy -> as silently as it reports a view. When a profile shows a copy the size of a known -> tensor, walk the *strides* of every reshape/view/permute between that tensor's producer -> and its consumer before instrumenting anything. Four GPU sessions guessed at this; one -> CPU-only `python3 -c` settled it. - -### The fix is free because the norm already pays for the write - -The fused q/k norm allocates its output anyway, so it can repack on the way through: read -the query with the two strides the grouped layout needs (one per group, one per head -inside a group) and write the contiguous `[sq, b, np, hn]` result the rest of attention -wants. Same values in the same order, so the result is **bit-identical** -- the right -acceptance bar here, because a head-permutation bug would produce individually plausible -values in every head and sail past a tolerance check. - -| | reshape + norm | grouped norm | -|---|---:|---:| -| kernel, 256 tokens (graph replay) | 12.30 us | **8.22 us** | -| launches per layer | 2 | **1** | -| per step, 48 layers | — | **0.196 ms saved** | - -### QWEN-043 — grouped-read q/k norm: **+1.5%** - -Four alternating arms on one node, so node drift shows up as base-to-base spread: - -| Arm | Gates | Steady-state tok/s | -|---|---|---:| -| `h2base` | standing set | 29,828 | -| `h3grp` | `+MCORE_GROUPED_QK_NORM=1` | 30,312 | -| `h4base` | standing set | 29,858 | -| `h5grp` | `+MCORE_GROUPED_QK_NORM=1` | 30,284 | - -Base-to-base spread **0.10%**, so the effect is ~15x the noise. Pooled -**29,845 -> 30,298 tok/s (+1.52%)**, i.e. 8.58 -> 8.45 ms/step -- 0.13 ms of the -microbenchmark's 0.196 ms lands end-to-end. Generated text is identical character for -character, as bit-exactness requires. Shipped behind `MCORE_GROUPED_QK_NORM=1`; requires -the fused q/k norm, and declines when the query needs the tensor-parallel head slice. - -**Status: 30,298 against vLLM's 33,994.5 = 89.1% of vLLM** (from 87.9% at session start, -65.9% at the campaign's start). Remaining gap 0.90 ms/step. - - -### The trace confirms it, and nothing else moved - -Diffing the fresh trace against the previous one, same anchor and window, is the cleanest -attribution in this campaign: - -| bucket | before | after | delta | -|---|---:|---:|---:| -| elementwise / copy | 0.164 ms, 59 launches | **0.042 ms, 11 launches** | **-0.122 ms, -48** | -| every other bucket | — | — | within +/-0.012 ms, 0 launches | -| TOTAL | 5.923 ms, 963 | 5.804 ms, 916 | -0.120 ms, -48 | - -Exactly one launch per layer left the copy bucket and nothing else changed, which also -re-validates the window: the buckets the change cannot touch held still, the check -`union_window.py` demands after the 27%-seqlen-skew incident. The 0.122 ms of device time -matches the 0.13 ms measured end-to-end. - -### The last mcore-only bucket, attributed - -`splitKreduce` is the one category vLLM has none of (82 launches, 0.173 ms/step). Its trace -neighbours name both parents: - -| n/step | parent GEMM | successor | -|---:|---|---| -| 48 | `nvjet_..._2cta_h_bz_splitK_TNT` / `..._4x1_v_bz_splitK_TNN` | `_softmax_topk_kernel` | -| 34 | `nvjet_..._2cta_h_bz_splitK_TNT` | `_fused_add_rmsnorm_kernel` | - -So it is the **router GEMM** (one per layer, exactly 48) and the **attention output -projection** (34 of 48 layers; the rest pick a non-splitK algorithm). The router case is -not a mis-tuned heuristic: at M=256, K=2048, N=128 the output is a couple of tiles, so -splitting K is the only way cuBLASLt gets parallelism, and the reduce is the price of it. -That reframes the target -- not "stop cuBLASLt from splitting K", but "fuse the router", -which would collapse GEMM + splitK reduce + softmax-topk + padding mask (4 launches/layer, -~192 total) into one. A fused router must beat GEMM+reduce on its own terms, so it needs a -microbenchmark before any integration: 256 CTAs each reading the whole 512 KB weight is -128 MB of L2 traffic, and blocking tokens instead leaves 4 CTAs on 148 SMs. - -### QWEN-044 — eight rows per CTA in the q/k norm: **+0.35%** - -The grouped norm moves ~4.5 MB in 8.22 us -- about 570 GB/s on a GB200, because one CTA -per 128-wide row means 9,216 single-warp CTAs. A rows x warps sweep, keeping only -bit-exact candidates: - -| rows/CTA | warps | us | GB/s | vs 1x1 | -|---:|---:|---:|---:|---:| -| 1 | 1 | 8.21 | 574 | 1.00x | -| 2 | 2 | 6.17 | 765 | 1.33x | -| 4 | 4 | 6.15 | 767 | 1.34x | -| **8** | **8** | **4.12** | **1146** | **2.00x** | - -In-tree the same shape measures **6.15 us**, not 4.12 -- a 2 us gap between the harness -kernel and the shipped one at identical rows and warps that is **still unexplained** after -three hypotheses, each tested and each measuring 6.15 us unchanged: - -| Hypothesis | Why it was plausible | Result | -|---|---|---| -| Runtime output row stride blocks store vectorization | the harness hardcoded `HN`, the kernel took a stride argument | rejected, 6.15 us | -| `where(is_q, rows, 0)` in the store address hides contiguity | the harness addressed stores off the raw row index | rejected, 6.15 us | -| Two separate load address arrays instead of one shared | the harness built a single `where`-selected offset array | rejected, 6.15 us | - -**Closed in session 20: there is no gap.** The PTX comparison this pointed to was run -(`dev/moe_fused/harness_qknorm_ptx.py`) and measured the shipped kernel and the harness -kernel at **4.12 us each, ratio 1.00x** -- same answer for q and k, 21 registers, no -spills, `.v2` vectorized loads and stores. The shipped kernel had already reached the -harness number; the 6.15 us readings were stale, taken before the store-stride fix landed -and then carried forward as if current. The three hypotheses above were each rejected for -the right reason: every one of them was measuring a kernel that was already fast. - -The lesson is measurement hygiene, not Triton. A number that survives three -disconfirmations without moving is more likely stale than robust, and re-measuring the -baseline is cheaper than the next hypothesis. Nothing left to recover -- item closed. - -End-to-end, both arms from the same file so the launch shape is the only variable: -1x1 controls **30,341 / 30,399** against 8x8 **30,482** tok/s, i.e. **+0.35%** on a -control-to-control spread of 0.19%. Note the microbenchmark predicted 1.17% -(2.06 us x 48 layers on an 8.44 ms step) and a third of that landed -- a reminder that a -kernel-level win on a non-critical-path kernel is an upper bound, not a forecast. - -### QWEN-045 — fusing the router GEMM: **-2.37%, rejected** - -The router chain is four launches per layer: gating GEMM, a cuBLASLt splitK reduce (the -logits have to reach memory before top-k can read them), fused softmax+topk, and the -padding mask. `kernel_neighbors.py` attributed 48 of the 82 splitK reduces per step to -this GEMM, so collapsing all four into one Triton kernel that keeps the logit tile in -registers looked like the largest single launch-count win available. - -The microbenchmark endorsed it. Sweeping BLOCK_M x BLOCK_K x warps at the decode shape, -keeping only configs whose expert sets matched exactly: - -| BLOCK_M | BLOCK_K | warps | us | vs reference | -|---:|---:|---:|---:|---:| -| 32 | 128 | 8 | 16.41 | 0.94x | -| 16 | 128 | 4 | 12.68 | 1.21x | -| **16** | **256** | **4** | **12.31** | **1.25x** | - -Expert sets bit-exact, probabilities within 7.45e-09, and the folded padding mask correct -(200 real rows, the other 56 getting -1). Against the reference's 15.34 us that is -0.146 ms/step, ~1.7% of an 8.35 ms step. The winning shape is a *narrow* BLOCK_M with the -widest K tile -- it maximises CTA count, and the whole-weight read each CTA does is then -what hides the K loop. Note the in-tree defaults were initially the 0.94x config, which -would have tested a kernel slower than baseline; they were corrected before the A/B. - -End to end it lost, consistently, in both replicate pairs: - -| arm | tok/s | -|---|---:| -| control | 30,689.4 | -| fused router | 29,981.5 | -| control | 30,697.5 | -| fused router | 29,951.2 | - -**-2.37%** on a control-to-control spread of 8 tok/s. The sign is opposite to the -prediction and the magnitude is larger, so this is not noise and not a tuning problem. - -The mechanism worth carrying forward: the four launches it replaced were not serial dead -time. cuBLASLt's GEMM and its splitK reduce overlap with neighbouring kernels inside the -CUDA graph, whereas a 16-CTA Triton kernel occupies few SMs *and* forces every CTA to -stream the full weight matrix, which evicts what its neighbours are reading. Isolated -device time credits the fused kernel for work it removed from the critical path only in -a benchmark where nothing else was running. - -**Rule this establishes:** an isolated-kernel win does not transfer when the kernel it -replaces is one that overlaps well. Before fusing across a cuBLAS call, measure the -*ablation* ceiling (delete the work and run e2e) rather than the replacement's device -time -- the ablation prices the critical path, the microbenchmark prices the kernel. -Retained behind `MCORE_FUSED_ROUTER=0`. - -### QWEN-046 — the padding mask: **+1.02% available, fusion not yet working** - -Applying the rule from QWEN-045 to the cheapest member of the router chain. -`mask_routing_padding` writes -1 into every topk slot of the CUDA-graph padding rows so -those tokens route to no expert; it is one launch per layer per step (48/step) over a -256x8 int64 tensor, which is almost entirely launch overhead. Ablating it outright -(`MCORE_ABLATE_ROUTE_MASK=1`, correctness-breaking by construction, valid only as a -ceiling) priced the critical path directly: - -| arm | tok/s | -|---|---:| -| control | 30,662.5 | -| mask ablated | 30,965.1 | -| control | 30,661.0 | -| mask ablated | 30,983.4 | - -**+1.02%** on a control-to-control spread of 1.5 tok/s -- among the tightest replicate -pairs measured in this project, and a lower bound besides: skipping the mask lets padding -rows route to real experts, which if anything *adds* expert GEMM work. - -The fusion target is the router's `_softmax_topk_kernel`, which is one CTA per token and -already stores that token's index row, so the sentinel is a scalar compare and a select -with no extra launch. It was implemented (module-level publish of the context's -`int32[1]` count, `MASK_PADDING` constexpr in the kernel, and a dispatcher skip keyed off -a tag the router leaves on the tensor it masked) and it **hung in warmup with the gate -off**, which localises the fault to the always-on parts rather than to the masking logic. -Reverting the three files reproduced the baseline to within 1 tok/s (30,661.9 against -30,662.5 and 30,661.0), so the hang is attributable to those edits alone. - -Design, suspects, and a cheapest-first diagnostic ladder are written up in -`dev/moe_fused/NOTES-route-mask-fusion.md` so the next attempt starts from the hang -rather than from the design. Leading suspect is the kernel signature change itself: it is -the only always-on edit that alters generated code, and if `tl.where(False, -1, best_idx)` -fails to fold, every row gets -1 and the NVLS all-gather-v barrier can deadlock on -rank-divergent counts. Unmeasured guess, listed first in the ladder. - -### Session 21 status - -| | tok/s | ms/step | vs vLLM | -|---|---:|---:|---:| -| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | -| **mcore, current best** | **~30,751** | **8.31** | **90.5%** | -| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | - -Session 21 added no throughput either, but it changed what the campaign believes about -where the remaining 9.5% is, which had been wrong since the collectives were first -measured in situ. - -- **The collective floor is not a deficit; it is a lead.** Reading vLLM's MoE source - (QWEN-051) showed it runs the same all-gather-v / reduce-scatter-v pair over plain - NCCL with symmetric memory off by default, and our own bucket table has had the - number all along: vLLM 1.607 ms/step against mcore 1.035. Every remaining plan that - started with "hide the collectives" — DBO included — was aimed at a bucket where we - are 0.57 ms/step ahead. Those are now closed rather than deferred. -- **The MoE kernel, not the MoE communication, is the last large lever.** flashinfer's - `trtllm_bf16_moe` measures 63.3 us/layer of device time against mcore's 78.0 at the - operating point, captures into a CUDA graph, and replays bit-identically. Estimated - **+0.5 to +0.7 ms/step (6-8%)**, which is the difference between 90.5% and ~96-98% of - vLLM. Unbuilt: it needs a gate/up half-swap at load and has to displace the router - top-k, the route mask and the dispatcher permutation. -- Two ledger corrections that had been blocking that lever for sessions: the BlockMajorK - weight requirement is a ~1 s load-time transform rather than an incompatibility, and - the flashinfer MoE previously rejected at 8% slower was a *different kernel* from the - one vLLM actually prefers. The run venv already ships the kernel, so shipping it needs - no new dependency. -- One self-inflicted false alarm, retracted the same session (QWEN-052): the flashinfer - decode gate looked inert because the container's system python lacks flashinfer, but - the e2e runner uses a venv that has it. The gate is live and the baseline is intact. - Attention nonetheless remains **+0.409 ms/step** behind vLLM *with* the trtllm-gen - kernel active, which makes it the second-largest known gap after the MoE kernel. - -### Session 20 status - -| | tok/s | ms/step | vs vLLM | -|---|---:|---:|---:| -| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | -| **mcore, current best** | **30,662–30,697** | **8.35** | **90.2%** | -| mcore, with the mask ablated (not shippable) | 30,974 | 8.27 | 91.1% | -| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | - -Session 20 added **no throughput**: both candidates were rejected, one on measurement -(-2.37%) and one on a hang. The code is unchanged from session 19; the higher number -against session 19's recorded 30,482 is node-and-day drift, not a code change, measured -across five controls in one job (30,689 / 30,697 / 30,662 / 30,661 / 30,662). - -What session 20 did produce, all of which outlives it: - -- A priced target: **+1.02%** sitting in the padding mask, with the ceiling measured - rather than estimated, and the implementation already designed and written up. -- A rule that would have saved this session's larger experiment (QWEN-045): price the - ablation ceiling before fusing across a well-overlapped cuBLAS call. -- One closed open item (QWEN-044's phantom 2 us) and one retired measurement habit. -- **14 hours of queue time recovered.** Two jobs sat `PENDING (Priority)` for 14 h and - 7 h with 315 nodes idle, purely because the sbatch template omitted `--qos`: the - default `normal` is priority 100 against a queue head near 350k. Resubmitted under - `--qos=interactive` (priority 700, <=4 nodes) and `--qos=short` (priority 200, <=2 h), - both started within seconds. Every job template in `skills/` now sets `--qos`, and - `skills/run-qwen-model/SKILL.md` carries the diagnostic: if nodes are idle and - `squeue -j -o '%Q'` is far below the head of `squeue -p batch -t PD -S -Q`, it is - QOS and not contention. - -### Session 19 status - -| | tok/s | ms/step | vs vLLM | -|---|---:|---:|---:| -| vLLM DP4+EP (`VLLM-BASELINE`) | 33,994.5 | 7.53 | — | -| **mcore, current best** (`hftune`) | **30,482** | **8.40** | **89.7%** | -| mcore at session-19 start (`y2pdl`) | 29,898 | 8.56 | 87.9% | -| mcore at session-2 start | 22,398.9 | 11.71 | 65.9% | - -Session 19: **+1.95%** from two changes to one kernel, both bit-exact. Remaining gap -0.87 ms/step. - -**Where the remaining gap is not.** Per-launch, mcore is now at or better than vLLM in -every functional bucket, including attention and the expert GEMM. The residue is -launch count (916 against 810) and packing: the overlap ratio is 1.008x against vLLM's -1.050x. And the >100 us gap band is *not* the lever it looks like -- mcore shows -0.893 ms/step there against vLLM's 2.727 ms/step in its own trace, so that band is -dominated by profiler-inflated host work on both sides and cannot be compared across -traces captured with different `--cuda-graph-trace` settings. - -### QWEN-047 — the collectives re-priced in place: latency-bound, exposed, unhideable without microbatching - -Session 21 opened by re-pricing every per-layer component with the in-graph external-event -timer (`MCORE_INSITU_TIMING=1` plus `MCORE_INFER_STEP_GPU_TIMING=1`, which is what actually -drives the report — the timer alone prints nothing). Motivation: the trace analysis had put -~0.3 ms on matching vLLM's overlap ratio, but that measurement predated the bf16 -reduce-scatter, which halved the combine's bytes on the wire. - -Per-layer, scaled x48, net of the calibrated 2.2 us/pair event overhead: - -| component | ms/step | component | ms/step | -|---|---:|---|---:| -| expert_gemm_fc1 | 1.650 | nvls_dispatch | 0.491 | -| expert_gemm_fc2 | 1.183 | qkv_proj | 0.482 | -| attn_core | 1.152 | moe_align | 0.449 | -| nvls_combine | 0.548 | moe_sum | 0.321 | -| out_proj | 0.533 | router_topk | 0.253 | -| | | **SUM of sites** | **7.060** | - -Two results, both load-bearing. - -**The collectives are latency-bound, not bandwidth-bound.** Dispatch plus combine is -1.039 ms/step against 1.035 ms measured before bf16 halved the combine's bytes — a 0.4% -move for a 2x reduction in traffic. With `QWEN-016`'s CTA-count result and the barrier -finding, that is now three independent measurements saying the same thing: **these -collectives cannot be made faster, only hidden.** Do not spend another session on the -transport. - -**The block is a serial chain, confirmed a second way.** Sites sum to 7.060 ms against a -block total of ~7.09 — about 100%. Parts can only reconstruct the whole like that if -nothing overlaps, which independently corroborates the trace-derived 1.008x overlap ratio -using a completely different instrument. - -**Why the 1.04 ms is not collectable today.** Hiding it needs a second, independent stream -of work to hide it behind, and inside a layer everything is strictly dependent: attention -feeds the router, the router feeds dispatch, the all-gather *sizes* the expert GEMM (see -`floor_ablation.py`), and the GEMM feeds combine. The standard remedy is two interleaved -microbatches. The inference decode path has no such machinery — the only overlap primitive -there is the shared-expert stream, and Qwen3-30B-A3B has no shared expert, so that lever is -empty. **Dual-batch overlap is the single largest remaining prize (1.04 ms exposed against -a 0.82 ms gap) and it is an architecture project, not a tuning session.** - -One trap avoided: the same report shows 0.78 ms/step (9.8%) outside the block graph, which -looks like a fresh target and is not. It is the bucket already measured at 0.59 ms and -**quantitatively retired** — embedding, logits GEMM, sampling, bookkeeping and every host -gap — whose links form a true serial dependency (`graph(N) -> logits -> argmax -> D2H -> -update_requests -> initialize_attention_state(N+1) -> H2D -> graph(N+1)`). Flashinfer -sampling, the obvious lever there, is `QWEN-004`: it crashes under `full_iteration_inference` -capture. Check the ledger before re-opening anything outside the block. - -### QWEN-048 — route-mask fusion: the hang is fixed, the win is inside the noise - -`QWEN-046` left this implemented-but-hanging with a measured +1.02% ablation ceiling. The -retry skipped straight to rung 3 of that entry's diagnostic ladder rather than working up -it, because duplicating the kernel *eliminates* the leading suspect instead of testing it: -`_softmax_topk_mask_kernel` is now a separate `@triton.jit` function and -`_softmax_topk_kernel` is byte-identical to before, so the gate-off path provably cannot -have changed. **That fixed the hang** — the gate-off arm, which is what hung in -`QWEN-046`, now runs normally. The cause was therefore the constexpr branch in the shared -kernel, not the publish or the skip. - -Correctness (`dev/moe_fused/harness_routemask.py`, run with the venv interpreter and -`PYTHONPATH` set, or it cannot import megatron): indices bit-exact and probs delta 0.000e+00 -at real counts 256/200/137/1/0, padding rows all -1. - -Throughput did not follow. Two back-to-back pairs disagreed in sign: - -| pair | gate off | gate on | delta | -|---|---:|---:|---:| -| 1 | 30,412.1 | 30,816.5 | **+1.33%** | -| 2 | 30,564.2 | 30,347.1 | **-0.71%** | - -Four runs span 30,347-30,817 with the treatment at both extremes, which looked like a noise -floor as large as the effect. Six more alternating pairs (`s1a`..`s6b`, gate off then on, -one server start each) resolved it: - -| | n | delta | | t | -|---|---:|---:|---:|---:| -| 6 alternating pairs | 6 | **+224.5 +/- 44.7 tok/s** | **+0.74%** | 5.02 | -| all 8 pairs | 8 | +191.8 +/- 70.6 tok/s | +0.63% | 2.72 | - -**Accepted: +0.74%** (p ~ 0.004), landing just under the +1.02% ablation ceiling, which is -the right shape — the fused kernel still does the compare and the select that the standalone -launch was doing, it just stops paying for a launch to do it. - -The noise diagnosis was itself wrong, and correcting it changes how future A/Bs here should -be read. Gate-off across all 8 runs has sd **51.2 tok/s (0.17%)** — the baseline is extremely -reproducible. Gate-on has sd 183.2, but that is one anomalous run (30,347, the lowest of all -16, the `rmask2` that produced the -0.71%); the other seven average 30,750.9 with a tight -spread. So this node is *not* broadly noisy: it produces an occasional bad server start, -roughly 1 in 16. **A single outlier start, not a wide distribution, is what makes -single-pair A/Bs unsafe here.** Alternate arms and read the paired t, or one bad start will -invert the sign of any sub-1% result. - -Level after this change: **~30,751 tok/s, 90.5% of vLLM.** - -### QWEN-049 — dual-batch overlap is the wrong technique for a latency-bound decode - -`QWEN-047` left DBO as the largest remaining prize (1.04 ms exposed against a 0.82 ms -gap). A full architecture map of the decode path plus an economic check says **do not -build it.** - -**What it would cost.** The decode path is single-microbatch at three independent levels. -(1) One `DynamicInferenceContext` owns every fixed-address buffer the graph reads; nothing -is replicated per request, only indexed. (2) Graphs are keyed on the *identity* of that -context object (`ArgMetadata` stores non-tensors by value, `cuda_graphs.py:256`), so a -second context needs its own captured graph set. (3) The NVLS dispatcher keeps its -symmetric buffers and `_step_metadata` at **class** level, and the barrier -(`torch_symm_triton/barrier.py:65`) is a per-(block, sender) mailbox with no sequence -number, so two concurrent collectives sharing a handle either false-pass into silent -corruption or lose a token and hang. - -**Why it would not pay even if built.** Take `BSSCALE-S13`: step(128) = 7.140 ms, -step(256) = 9.151 ms, with a ~5.03 ms batch-independent floor. Splitting 256 into two 128s -doubles the count of every latency-bound operation, and the collectives are latency-bound, -so the comm to hide goes from 1.04 ms to 2.08 ms. Best case is -floor + 2 x marginal(128) = 5.03 + 4.22 = **9.25 ms against today's 9.15**. DBO hides -communication behind long GEMMs in a *compute-bound* regime; this decode is latency-bound, -so halving the batch costs about what the overlap returns. - -**And it was already gated out.** `QWEN-016` priced the whole lever: recoverable critical -path 6.4-7.0%, 2-chunk pipelining 4.3% best case, concurrent streams under graph capture -identified as the binding constraint, and a CTA sweep showing 128 is already optimal so -freeing SMs by shrinking the collective's grid is strictly worse. Its recorded next action -was "redirect off this lever". **Read QWEN-016 before touching comm overlap again.** - -### QWEN-050 — EP expert-load imbalance is real and worth nothing - -A promising-looking lever, killed by direct measurement. Recorded in full because the -causal story was strong and still wrong. - -**The evidence for it.** A 127M-assignment histogram (`expert_histogram.py`, new) shows -expert popularity is uneven and does not average out within a rank's contiguous block: -loads are **+6.9% / -2.9% / +2.1% / -6.1%** of the mean. In-situ timing appears to show -the barrier absorbing exactly that — per layer rank 0 spends 63.0 us in the expert GEMM -and 28.2 us in the collectives, rank 1 spends 56.8 and 34.5, and *both sum to 91.2*. The -rank ordering of GEMM time matches the histogram ordering exactly. An equal-cardinality -greedy partition takes the worst rank from +6.9% to +0.1%, so on the "critical path is the -busiest rank" model the prize is (max - mean) = 4.85 us/layer = 233 us/step = **2.8%**. - -**The measurement.** Permuting only the router's gating weight was the cheap probe, and it -**hangs** — it leaves expert weights in place, so numerics degrade, logits degenerate, -routing collapses onto a few experts and blows the per-rank token capacity. *A probe that -corrupts numerics cannot rely on data-dependent routing.* The fix is to make routing not -depend on numerics at all: `synthetic_routing.py` (new) substitutes a fixed index tensor -with a chosen per-rank split, so both arms have equally invalid numerics and differ only -in load distribution. - -| arm | per-rank split | tok/s | -|---|---|---:| -| `skew` | +5.9 / -2.2 / +2.6 / -6.2% | 29,656.4 | -| `balanced` | -0.6 / -0.3 / +0.4 / +0.5% | 29,583.7 | - -**Removing the entire imbalance is worth -0.25%, i.e. nothing.** A 2.8% effect would have -been ~830 tok/s, far outside any noise here, so one pair is enough to reject it. - -**Why the causal story failed.** The GEMM/collective anti-correlation is the barrier -equalizing the ranks, and that is exactly why rebalancing cannot help: it moves work -between ranks whose totals are *already* equalized, without shortening the thing that sets -the total. The binding constraint is the collective latency floor, not expert work. Note -also that several sites split cleanly along r0,r1 vs r2,r3 (`nvls_dispatch` 14.9/14.8 vs -10.2/10.1, `router_topk` 3.8/3.8 vs 5.3/5.2) — that is a two-superchip topology signature, -plus a per-rank calibration difference (2.2 vs 2.6-2.7 us/pair), and I initially read part -of it as load imbalance. **Per-rank in-situ deltas conflate load, topology and calibration; -do not attribute them to load without an independent load measurement.** - -Both diagnostics are kept, gated off: `MCORE_EXPERT_HISTOGRAM` and `MCORE_SYNTH_ROUTING`. - - -### QWEN-051 — vLLM's MoE communication is *worse* than ours; the one thing worth taking is a kernel vLLM itself does not run - -Prompted by "look at what vLLM's MoE communication does and reuse it." The answer to the -question as asked is that there is nothing to reuse, and the reason overturns a conclusion -this campaign has been steering by. - -**vLLM's default all2all backend is `allgather_reducescatter`** (`vllm/config/parallel.py:186`) -— the *same* all-gather-v + reduce-scatter-v pair mcore runs. It goes over plain NCCL: -`VLLM_USE_NCCL_SYMM_MEM` defaults to `False` (`vllm/envs.py:267`), so vLLM is not even using -the NVLS multimem path our dispatcher is built on. No stream overlap around dispatch or -combine; `supports_async()` defaults to `False` and only DeepEP/NIXL/MoRI override it; DBO is -hard-asserted to those same backends (`vllm/config/vllm.py:1536`) and so is unreachable from -`TP=1 DP=4 --enable-expert-parallel`. `naive` and `pplx` were removed and now silently fall -back to the default (`parallel.py:448`). - -The measured consequence was already in this ledger and had not been read as such: -**collective 1.607 ms/step (94 launches) in vLLM against 1.035 ms/step (96) in mcore.** -mcore's collectives are **0.572 ms/step faster than vLLM's**. Several attempts to attack -"the collective latency floor" were aimed at the one bucket where we are already ahead. -*When a bucket comparison shows you winning, stop optimising it — even when an in-situ -measurement says the bucket is large and fully serialized. Large and serialized is not the -same as improvable, and the competitor's number is the evidence that settles which.* - -The other half of the request is already done: `vllm_fused_moe.py` **is** vLLM's Triton fused -MoE, vendored, and is what `GEMM_BACKEND=vllm` selects. That is why expert GEMM sits at -parity (2.436 vs 2.376) and why the bucket launch counts nearly coincide — 288 in mcore -against 283 in vLLM across routing + expert GEMM + finalize. Both engines run the same -kernels. There is no unclaimed vLLM MoE code left to import. - -**What is left is `flashinfer.fused_moe.trtllm_bf16_moe`**: one kernel doing router, both -expert GEMMs, activation and finalize, and vLLM's *first-choice* unquantized backend on -Blackwell (`oracle/unquantized.py:68`), explicitly compatible with the AG/RS layout we -already have (`experts/trtllm_bf16_moe.py:207`). The trace proves vLLM fell back to Triton in -the run we benchmarked (283 launches, not ~50), so **vLLM's own numbers do not vouch for it** -— adopting it would be overtaking vLLM's configuration, not copying it. - -Two ledger notes had to be corrected before it could be priced: - -- The "4-D pre-shuffled BlockMajorK weights" blocker is **real for bf16 too** (the docstring - rejects any other `weight_layout`), not FP8-only as the vLLM source suggested. But it is - **not a blocker**: `reorder_rows_for_gated_act_gemm` → `shuffle_matrix_a` → - `convert_to_block_layout` round-trips mcore's 3-D `[E, 2*ffn, H]` in **0.02 s per layer, - ~1 s of load time for all 48**. A one-time startup transform had been recorded as a - structural incompatibility. -- The prior rejection at "90.81 vs 83.63 us, 8% slower" was `cutlass_fused_moe`, a - **different kernel**, which vLLM actively demotes when `dp_size > 1`. `trtllm_bf16_moe` had - never been measured. - -**Measurement.** Harnesses in `dev/moe_fused/`: `harness_trtllm_moe.py` (bake-off), -`harness_moe_breakdown.py` (per-kernel), `harness_moe_graph.py` (capture), -`harness_moe_sweep.py` (token sweep). The first bake-off was wrong twice over and both -errors are worth keeping: - -1. It ran with the production gate set unset, so the Triton path used its untuned default - config: 206.9 us/layer against a true 134.6. *A harness that imports the production kernel - does not inherit the production environment; copy the gate set from the last e2e script.* -2. Wall clock overstated the win (1.49x) because ~40% of both sides was host dispatch that - the decode graph removes. Device time is the only comparable number. - -Device time at T=256, the operating point (512 token-expert pairs per rank, matching -production): - -| | device us/layer | launches | -|---|---:|---:| -| mcore `vllm_fused_moe` | 78.0 | 5 | -| flashinfer `trtllm_bf16_moe` | 63.3 | 4 | - -**CUDA-graph capture succeeds and replay is bit-identical to eager** (max diff 0.000e+00) at -61.1 us/layer. That was the kill criterion: mcore's decode is one captured graph and the -kernel uses a cooperative launch reserving 144 of 152 SMs. - -A token sweep anchors the projection. mcore's per-launch grouped GEMM never reaches the -trace's 25.4 us — it asymptotes near 30 us at T=64 — so the harness carries a ~25-30% -constant overhead over in-graph execution and the raw delta must be discounted by it: - -| T | mcore total | GEMM/launch | flashinfer | speedup | raw ms/step | -|---:|---:|---:|---:|---:|---:| -| 64 | 69.5 | 29.9 | 59.3 | 1.17 | +0.49 | -| 128 | 71.5 | 31.2 | 63.2 | 1.13 | +0.40 | -| 256 | 78.0 | 33.0 | 63.3 | 1.23 | +0.71 | -| 512 | 99.3 | 41.3 | 67.6 | 1.47 | +1.53 | -| 1024 | 118.0 | 46.1 | 72.8 | 1.62 | +2.17 | - -**Estimate: +0.5 to +0.7 ms/step, 6-8%**, which would take mcore from 90.5% to roughly 96-98% -of vLLM. The win is *kernel quality, not launch reduction* — mcore's MoE chain is already -only 5 launches/layer thanks to the align/count/scatter fusions, so exactly one launch is -removed. An earlier "288 launches collapse to 48" framing was wrong. - -**Before integrating.** Numerics need a gate/up half-swap at load: flashinfer matches a torch -reference with `gate_first=False` where mcore matches `gate_first=True`, and the two disagree -by 1.6e-4 against a 1.5e-3 signal. This is the same silent-corruption trap the cutlass probe -hit. The real cost is integration, not the kernel: `trtllm_bf16_moe` routes from logits -itself, so it displaces the router top-k, the newly fused route mask and the dispatcher's -permutation, and has to interoperate with the NVLS AGV/RSV pair. - -**And it must be A/B'd e2e before being believed.** Two microbenchmark wins of this size have -already gone negative here (QWEN-045's fused router: 1.25x in isolation, -2.37% e2e). The -reason to expect a different outcome is mechanistic rather than hopeful: both prior failures -replaced *well-overlapped* work and lost the overlap, whereas mcore's MoE is measured at -essentially zero overlap (sum-of-durations / union = 1.003-1.008), so there is none to lose. - -Environment note: **no install is needed to ship this.** The e2e run venv -(`run_e2e_cfg.sh:36`) already carries flashinfer 0.6.14, which exports `trtllm_bf16_moe`, -`convert_to_block_layout`, `shuffle_matrix_a` and `reorder_rows_for_gated_act_gemm`. The -container's *system* python has no flashinfer, which is why the standalone harnesses needed -`persistent_cache/pyextra` (0.6.16.post1, `--no-deps`, since pruned to `flashinfer`, -`nccl4py` and `tvm_ffi` so it shadows nothing). Cubins are cached at -`persistent_cache/flashinfer_cubins`; the first call downloads for ~14 minutes. See QWEN-052 -for the interpreter confusion this distinction caused. - - -### QWEN-052 — RETRACTED: `MCORE_FLASHINFER_DECODE` is live; I probed the wrong interpreter - -Kept rather than deleted, because the false alarm is a better lesson than the correction. - -**The claim.** `flashinfer_decode.py` degrades on `ImportError` to `HAVE_FLASHINFER = False` -with no warning, and `pip show flashinfer` inside the mcore image `ceecf5c304a5d8bd.sqsh` -reports "Package(s) not found". Since 88 of 92 e2e scripts run in that image, the gate -looked inert, and the differential seemed to corroborate it: attention is still -**1.623 ms/step against vLLM's 1.214** on a current-gate-set trace, exactly the bucket the -gate was credited with closing (+2.6%). - -**Why it is wrong.** The e2e runner does not use the container's python. `run_e2e_cfg.sh:36` -sets `PYBIN=$VENV/bin/python` against a venv on lustre, and that venv **has flashinfer -0.6.14** — the very version this ledger recorded in QWEN-039. Checked inside the container -with the real interpreter: - -``` -torch 2.13.0a0+8145d630e8.nv26.06 -flashinfer 0.6.14 -HAVE_FLASHINFER = True -``` - -against `ModuleNotFoundError` from the container's `/usr/bin/python` two lines later. The -gate is live, the +2.6% is in the current baseline, and nothing regressed. - -**Rule.** *Probe the interpreter the workload actually launches, not the container default.* -`which python` inside the container answered a different question than the one being asked, -and the answer looked decisive enough to build a whole causal story on — one that a stale -+0.409 ms attention bucket then appeared to confirm. Read the runner for `PYBIN`/`VENV` -before concluding anything about what is installed. Corollary: coincidental corroboration -is the dangerous kind. The attention deficit is real but has some other cause, and reading -it as confirmation nearly cost an e2e A/B chasing a bug that does not exist. - -**The `ImportError` critique still stands on its own**, just with no victim here: a -capability gate that degrades silently makes a run script's gate set a statement of intent -rather than a record of what ran, and it should log once at init when enabled but -unavailable. Every other optional-import gate in `inference/` (`fused_add_rmsnorm`, -`fused_qk_norm`, `permute`, `vllm_fused_moe`, `router_topk`, `fused_router`) degrades on -**triton**, which is present at 3.7.0, so no other ledger entry is exposed either. - -**Useful by-product.** The run venv's flashinfer 0.6.14 already exports -`fused_moe.trtllm_bf16_moe`, `fused_moe.convert_to_block_layout`, `shuffle_matrix_a` and -`reorder_rows_for_gated_act_gemm` — everything QWEN-051's integration needs. **No install -is required to ship it.** The `persistent_cache/pyextra` install from QWEN-051 was only ever -needed because the standalone harnesses run under the container's system python; it has -since been pruned to `flashinfer`, `nccl4py` and `tvm_ffi` so it adds modules without -shadowing any the container provides. +## Experiment index -Also unchanged, and still open: attention is **+0.409 ms/step** against vLLM *with* the -trtllm-gen decode kernel active. That is a genuine remaining deficit rather than a -misconfiguration, and it is now the second-largest known gap after the MoE kernel. +| ID | Date | Hypothesis | Changed files / flags | Throughput | Delta vs baseline | Correctness | Job / run | MR | Conclusion | +|---|---|---|---|---:|---:|---|---|---|---| +| — | — | _no experiments recorded yet_ | — | — | — | — | — | — | — | ## Optimization rules 1. Profile and classify before proposing a code change. -2. Change one performance variable at a time. -3. Preserve the fixed protocol. -4. Validate correctness before accepting throughput. -5. Revert regressions or correctness failures. -6. Record the result before beginning another experiment. -7. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines - once to confirm parity under identical conditions. +2. Quantify the lever's ceiling before building it; record gates that come out + negative. +3. Change one performance variable at a time. +4. Preserve the fixed protocol. +5. A/B in one allocation, arms back to back and alternating; accept on + distribution separation, not mean delta. +6. Validate correctness before accepting throughput. +7. Revert regressions or correctness failures. +8. Record the result — including rejections, with their root cause — before + beginning another experiment. +9. Open a separate draft MR for each accepted change, and link it from its row. +10. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines + once to confirm parity under identical conditions. diff --git a/skills/run-qwen-model/SKILL.md b/skills/run-qwen-model/SKILL.md index 6910350d1bc..1a79d644fd3 100644 --- a/skills/run-qwen-model/SKILL.md +++ b/skills/run-qwen-model/SKILL.md @@ -30,11 +30,13 @@ export QWEN30B_HF="$QWEN30B_TOKENIZER" Before running: +Run these from inside the checkout you want to benchmark — cog deploys the +tree you are standing in: + ```bash source ~/.cog/setup.env.oci-hsg -export COG_MEGATRON_REPO=/path/to/Megatron-LM -cog prepare-image --repo "$COG_MEGATRON_REPO" --cluster-name "$COG_CLUSTER_NAME" -cog ensure-env --repo "$COG_MEGATRON_REPO" \ +cog prepare-image --cluster-name "$COG_CLUSTER_NAME" +cog ensure-env \ --cluster-name "$COG_CLUSTER_NAME" \ --run-name qwen30b-env --gpus 4 --time 00:30:00 \ --partition "$COG_BATCH_PARTITION" @@ -42,6 +44,20 @@ cog ensure-env --repo "$COG_MEGATRON_REPO" \ Never download checkpoints. If either path is missing, ask the user. +**Confirm which checkout you are benchmarking.** With more than one Megatron-LM +checkout on the machine, deploying the wrong one is silent — the run succeeds +and measures code you never changed. The scripts here default to the checkout +they live in and warn on a mismatch, and `COG_MEGATRON_REPO` is deliberately +absent from `~/.cog/setup.env*` so nothing machine-wide can override that. If +you set it by hand, it wins; check before a measurement run: + +```bash +echo "${COG_MEGATRON_REPO:-}"; git rev-parse --show-toplevel +``` + +Cross-check afterwards: the `CODE_REVISION` recorded by the run must equal your +local `HEAD`. + ## If a cog command hangs, switch to `sbatch` after the second attempt These runs are long enough that a stuck launch is expensive. Every cog command @@ -86,10 +102,15 @@ Three rules for any directly-submitted job here, all learned the hard way: - **A cascade of `ModuleNotFoundError` across unrelated packages is a bad node, not a bad venv.** Exclude it (`--exclude=`) and resubmit; do not pip-install overlay copies to work around it. - -`dev/moe_fused/rebased12.sbatch` is a working instance of this pattern if it is -still present in your tree — but it is untracked, so treat the cog skill's -inlined template as the source of truth. +- **A job that dies in seconds with exit 141 and empty logs is SIGPIPE, and + under `set -euo pipefail` the writer is what failed.** `ls … | head -1` + is the usual culprit: `head` closes the pipe, `ls` takes SIGPIPE (128+13), + `pipefail` promotes it, and `set -e` aborts before a single line is logged. + This is latent until the glob grows, so it appears long after the script was + written. Drain the pipe (`sed -n 1p`) rather than closing it early. It bit + `run_qwen_vllm.sh` once the workspace glob reached 142 entries. + +The cog skill's inlined template is the source of truth for the sbatch pattern. ## 1. Megatron-Core inference diff --git a/skills/run-qwen-model/profile_qwen_mcore.sh b/skills/run-qwen-model/profile_qwen_mcore.sh index c171648657f..64944b64e4b 100755 --- a/skills/run-qwen-model/profile_qwen_mcore.sh +++ b/skills/run-qwen-model/profile_qwen_mcore.sh @@ -6,7 +6,7 @@ # # Usage: # source ~/.cog/setup.env.oci-hsg -# export COG_MEGATRON_REPO=/path/to/Megatron-LM +# # Deploys the checkout this script lives in; export COG_MEGATRON_REPO to override. # QWEN30B_CKPT=/lustre/.../qwen3-30b-a3b-mcore \ # PROFILE_BS=256 PROFILE_OSL=128 \ # bash skills/run-qwen-model/profile_qwen_mcore.sh @@ -16,6 +16,7 @@ PROFILE_BS="${PROFILE_BS:-256}" PROFILE_OSL="${PROFILE_OSL:-128}" _USER_REPO="${COG_MEGATRON_REPO:-}" +_SCRIPT_REPO="$(git -C "$(dirname "${BASH_SOURCE[0]}")" rev-parse --show-toplevel 2>/dev/null || true)" if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then # shellcheck disable=SC1091 source "${HOME}/.cog/setup.env.oci-hsg" @@ -26,7 +27,19 @@ else echo "ERROR: no ~/.cog/setup.env — run cog-setup-and-help skill" >&2 exit 1 fi -if [[ -n "$_USER_REPO" ]]; then export COG_MEGATRON_REPO="$_USER_REPO"; fi +# Precedence: caller export > the checkout this script lives in > setup.env. +# setup.env is machine-wide and can name a different (stale) checkout, which +# would profile code you did not edit — silently. +if [[ -n "$_USER_REPO" ]]; then + export COG_MEGATRON_REPO="$_USER_REPO" +elif [[ -n "$_SCRIPT_REPO" ]]; then + if [[ -n "${COG_MEGATRON_REPO:-}" && "$COG_MEGATRON_REPO" != "$_SCRIPT_REPO" ]]; then + echo "WARNING: ~/.cog/setup.env names COG_MEGATRON_REPO=$COG_MEGATRON_REPO" >&2 + echo " but this script lives in $_SCRIPT_REPO — profiling the latter." >&2 + echo " Export COG_MEGATRON_REPO explicitly to override." >&2 + fi + export COG_MEGATRON_REPO="$_SCRIPT_REPO" +fi : "${COG_MEGATRON_REPO:?COG_MEGATRON_REPO not set}" : "${COG_SSH_HOST:?COG_SSH_HOST not set}" @@ -77,8 +90,12 @@ nsys --version # Launch the server UNDER nsys. --cuda-graph-trace=node is required so kernels # inside the full_iteration_inference CUDA graph are individually traced. +# Do NOT add osrt or --sample=process-tree: on this workload (MoE decode under +# full-iteration CUDA graphs) either one hangs nsys in finalization and the +# resulting qdstrm is rejected by QdstrmImporter. Bisected over several +# sessions; see optimize-inference-siddharth/references/measuring.md. nsys profile \\ - --trace=cuda,nvtx,osrt \\ + --trace=cuda,nvtx \\ --sample=none --cpuctxsw=none \\ --cuda-graph-trace=node \\ --force-overwrite=true \\ diff --git a/skills/run-qwen-model/profile_qwen_vllm.sh b/skills/run-qwen-model/profile_qwen_vllm.sh index 74229c98b81..36eb10932ab 100755 --- a/skills/run-qwen-model/profile_qwen_vllm.sh +++ b/skills/run-qwen-model/profile_qwen_vllm.sh @@ -50,7 +50,10 @@ EXP=$EXP HF_CKPT=$HF_CKPT PROFILE_BS=$PROFILE_BS PROFILE_OSL=$PROFILE_OSL -MLM_HOST=\$(ls -dt $COG_SCRATCH_ROOT/workspaces/megatron_lm/*/repo 2>/dev/null | head -1) +# sed -n 1p rather than head -1: head closes the pipe early and SIGPIPEs ls, +# which is fatal under pipefail. Override MLM_REPO to pin a specific workspace +# instead of taking whichever one was modified most recently. +MLM_HOST=\${MLM_REPO:-\$(ls -dt $COG_SCRATCH_ROOT/workspaces/megatron_lm/*/repo 2>/dev/null | sed -n 1p)} mkdir -p "\$EXP" rm -f "\$EXP/vqdone_status" "\$EXP/vqserver.log" "\$EXP/vqprof.log" @@ -77,8 +80,10 @@ PROF_BASE=\$EXP/vllm_profile if ! command -v nsys >/dev/null 2>&1; then echo "ERROR: nsys not in vLLM image"; exit 3; fi nsys --version +# osrt is omitted deliberately — see the note in profile_qwen_mcore.sh. It is +# not needed for the GPU-kernel comparison and only inflates the trace. nsys profile \\ - --trace=cuda,nvtx,osrt \\ + --trace=cuda,nvtx \\ --sample=none --cpuctxsw=none \\ --cuda-graph-trace=node \\ --force-overwrite=true \\ diff --git a/skills/run-qwen-model/run_qwen_inference.sh b/skills/run-qwen-model/run_qwen_inference.sh index 24fdf170472..622619dcb71 100755 --- a/skills/run-qwen-model/run_qwen_inference.sh +++ b/skills/run-qwen-model/run_qwen_inference.sh @@ -51,6 +51,7 @@ done # Prefer oci-hsg env when present; fall back to default cog env. _USER_REPO="${COG_MEGATRON_REPO:-}" +_SCRIPT_REPO="$(git -C "$(dirname "${BASH_SOURCE[0]}")" rev-parse --show-toplevel 2>/dev/null || true)" if [[ -f "${HOME}/.cog/setup.env.oci-hsg" ]]; then # shellcheck disable=SC1091 source "${HOME}/.cog/setup.env.oci-hsg" @@ -61,9 +62,18 @@ else echo "ERROR: no ~/.cog/setup.env — run cog-setup-and-help skill" >&2 exit 1 fi -# Caller-exported repo wins over the value baked into setup.env.*. +# Precedence: caller export > the checkout this script lives in > setup.env. +# setup.env is machine-wide and can name a different (stale) checkout, which +# would sync and benchmark code you did not edit — silently. if [[ -n "$_USER_REPO" ]]; then export COG_MEGATRON_REPO="$_USER_REPO" +elif [[ -n "$_SCRIPT_REPO" ]]; then + if [[ -n "${COG_MEGATRON_REPO:-}" && "$COG_MEGATRON_REPO" != "$_SCRIPT_REPO" ]]; then + echo "WARNING: ~/.cog/setup.env names COG_MEGATRON_REPO=$COG_MEGATRON_REPO" >&2 + echo " but this script lives in $_SCRIPT_REPO — benchmarking the latter." >&2 + echo " Export COG_MEGATRON_REPO explicitly to override." >&2 + fi + export COG_MEGATRON_REPO="$_SCRIPT_REPO" fi : "${COG_MEGATRON_REPO:?COG_MEGATRON_REPO not set}" diff --git a/skills/run-qwen-model/run_qwen_vllm.sh b/skills/run-qwen-model/run_qwen_vllm.sh index 651c5d4f50c..979c6d51698 100755 --- a/skills/run-qwen-model/run_qwen_vllm.sh +++ b/skills/run-qwen-model/run_qwen_vllm.sh @@ -44,7 +44,19 @@ ssh -o BatchMode=yes "$COG_SSH_HOST" "mkdir -p '$RUN_DIR'; cat > '$REMOTE_SCRIPT #SBATCH --error=$RUN_DIR/slurm-%j.err set -euo pipefail -MLM_HOST=\$(ls -dt "$COG_SCRATCH_ROOT"/workspaces/megatron_lm/*/repo 2>/dev/null | head -1) +# Locate the benchmark client: an explicit MLM_REPO, else the most recently +# synced mcore workspace. Use "sed -n 1p", not "head -1": head closes the pipe +# after one line, which SIGPIPEs ls, and under pipefail that aborts the job +# with exit 141 before a single line is logged. sed drains its input instead. +MLM_HOST="${MLM_REPO:-}" +if [[ -z "\$MLM_HOST" ]]; then + MLM_HOST=\$(ls -dt "$COG_SCRATCH_ROOT"/workspaces/megatron_lm/*/repo 2>/dev/null | sed -n 1p) +fi +if [[ -z "\$MLM_HOST" || ! -f "\$MLM_HOST/tests/performance_tests/client/static_benchmark.py" ]]; then + echo "ERROR: no benchmark client under '\$MLM_HOST'. Set MLM_REPO to a synced Megatron-LM workspace." >&2 + exit 1 +fi +echo "benchmark_client_repo=\$MLM_HOST" srun --container-image="$IMG" --container-mounts=/lustre:/lustre --no-container-mount-home bash -c ' set -euo pipefail export HF_HOME="$RUN_DIR/hf_home" From 30654cb369138033e7e1a86a9925ae8a368739d2 Mon Sep 17 00:00:00 2001 From: shanmugamr1992 Date: Fri, 7 Aug 2026 12:15:47 -0700 Subject: [PATCH 3/4] Added vllm codebase agent --- .cursor/agents/qwen-model-optimizer.md | 79 +++++- .cursor/agents/vllm-codebase-expert.md | 77 ++++++ .../scripts/idle_decompose.py | 191 ++++++++++++++ .../references/vllm-differential.md | 8 + .../scripts/bandwidth_probe.py | 116 +++++++++ .../scripts/osl_decompose.py | 205 +++++++++++++++ skills/run-qwen-model/EXPERIMENTS.md | 31 ++- skills/vllm-codebase-reference/SKILL.md | 180 +++++++++++++ .../references/navigation-map.md | 242 ++++++++++++++++++ 9 files changed, 1119 insertions(+), 10 deletions(-) create mode 100644 .cursor/agents/vllm-codebase-expert.md create mode 100644 skills/nsight-system-analysis/scripts/idle_decompose.py create mode 100644 skills/optimize-inference-siddharth/scripts/bandwidth_probe.py create mode 100644 skills/optimize-inference-siddharth/scripts/osl_decompose.py create mode 100644 skills/vllm-codebase-reference/SKILL.md create mode 100644 skills/vllm-codebase-reference/references/navigation-map.md diff --git a/.cursor/agents/qwen-model-optimizer.md b/.cursor/agents/qwen-model-optimizer.md index bded8248a1b..2cfe4668cc6 100644 --- a/.cursor/agents/qwen-model-optimizer.md +++ b/.cursor/agents/qwen-model-optimizer.md @@ -23,17 +23,20 @@ Before acting, read: 4. `skills/nsight-system-analysis/SKILL.md` — **the** profile analysis skill 5. `skills/optimize-inference-siddharth/SKILL.md` — **the** optimization skill -There are exactly two performance skills: analysis and optimization. Do not go -looking for others, and do not invent skill paths. +There are exactly two performance skills: analysis and optimization. A third, +`vllm-codebase-reference`, answers questions about the competitor's source and is +read on demand rather than up front. Do not go looking for others, and do not +invent skill paths. Treat `EXPERIMENTS.md` as the sole source of performance history. -## Division of labour between the two skills +## Division of labour -| Question | Skill | +| Question | Where it goes | |---|---| | What does this trace say? Where does one decode step spend its time? | `nsight-system-analysis` | | Which lever do I pull, is it worth building, and how do I not break it? | `optimize-inference-siddharth` | +| How does *vLLM* do this? | the `vllm-codebase-expert` subagent | `nsight-system-analysis` owns windowing, interval-union arithmetic, per-category attribution, and the report format. Its `scripts/forward_pass.py` (Workflow C) @@ -57,6 +60,39 @@ Supporting skills, only when the task calls for them: `skills/cog-setup-and-help` (cluster, image, session, sbatch escape hatch), `skills/git-credentials-setup` (any GitHub auth failure). +## Asking how vLLM does it + +The vLLM source that produces the baseline is checked out at +`/Users/shanmugamr@nvidia.com/vllm`, at the exact revision the benchmarks run. +When the differential says vLLM is cheaper and you need to know *how*, read the +source rather than your recollection of vLLM — recent releases moved and deleted +enough that remembered layout is routinely wrong. + +Delegate that lookup to the **`vllm-codebase-expert`** subagent instead of +grepping the tree yourself; it is a large tree and the answers are usually a +paragraph. Send it a specific question and the context it needs (kernel name, +config, what you are trying to decide). It is read-only and returns cited paths. +Its knowledge lives in `skills/vllm-codebase-reference/SKILL.md`, which you can +read directly for a one-off lookup you are already in the middle of. + +Four things worth asking it, in rough order of value: + +1. **Which library is this trace kernel from, and what calls it?** The + differential's first move whenever a bucket shows the same launch count on + both sides and more time on ours. Frequently ends with "flashinfer, and it is + already in our venv." +2. **How does vLLM structure this stage** — routing, dispatch, shared-expert + overlap, graph capture — and what selects that path by default? +3. **Which env var or config flag gates it**, and is it on in our baseline + command? +4. **Is the win upstream vLLM code or an external package?** This is the + difference between a day of work and a quarter of it. + +What comes back is a mechanism and a bound, not a decision. Price it through +`references/decision-gates.md` before building, exactly as you would any other +lever — that the competitor fuses something proves the fusion is legal, not that +yours will be faster. + ## Baseline gate Before the first run, confirm you are benchmarking **this** checkout: @@ -103,7 +139,11 @@ Repeat until mcore reaches vLLM: it identified. 2. Classify the dominant signal — compute, memory, launch, communication, synchronization, or host scheduling — then jump to the matching section via - the routing table in `optimize-inference-siddharth` Step 1. + the routing table in `optimize-inference-siddharth` Step 1. When the signal is + "vLLM does this differently" — same launch count but more time, a role vLLM + fuses that we split, or a stage vLLM does not have at all — ask the + `vllm-codebase-expert` subagent how vLLM implements it *before* designing the + fix. 3. **Gate the lever before building it.** Share of device time is not headroom. Compute the ceiling per `references/decision-gates.md`, subtract what the fix itself costs, and write down *proceed* or *gated out*. Skip the gate only for @@ -120,7 +160,8 @@ Repeat until mcore reaches vLLM: 10. Append the complete result to `EXPERIMENTS.md` — including rejections, with their root cause and date. 11. Keep improvements; revert regressions and correctness failures. -12. For each accepted change, open its own MR (next section). +12. For each accepted change, open its own MR and record it in all three ledger + locations, including the pinned *Accepted changes* table (next section). 13. Promote whatever generalizes into the skills (section after that). Never optimize from a warmup/capture-only nsys window. `forward_pass.py` already @@ -159,9 +200,25 @@ The description must let a reviewer judge the change without re-running it: - **Scope and risks** — configs where it does not apply or was not measured. - **Artifacts** — ledger entry id, run/job paths, `.nsys-rep` / `.sqlite` paths. -Link the PR from its `EXPERIMENTS.md` entry so the ledger and the MR are -cross-referenced. Rejected experiments do not get a PR — they get a ledger entry -with the root cause. +Record every accepted change in `EXPERIMENTS.md` in **three** places, and treat +the change as unrecorded until all three exist: + +1. A row in the pinned **Accepted changes** table at the top — mechanism, gain, + kill switch, PR link, PR state. This is the table a human reads months later + to answer "what landed"; it is the only view that is not interleaved with + rejections or buried in a session write-up. +2. A complete row in the **Experiment index**, including its `MR` column. +3. The detailed record, with protocol, attribution, and artifacts. + +The duplication is deliberate. Do not "simplify" it by leaving the PR link in only +one place — a PR recorded only inside a detailed record is effectively lost, and +that has already happened once in this ledger. + +Refresh the `PR state` column whenever you touch the ledger, so a merged PR does +not sit there reading "draft". + +Rejected experiments do not get a PR and do not get an *Accepted changes* row — +they get an index row and a detailed record with the root cause. ## Keep the skills learnable @@ -182,6 +239,7 @@ Route by subject: |---|---| | A profiling technique, anchoring trick, windowing pitfall, taxonomy fix, script improvement | `skills/nsight-system-analysis/` — `SKILL.md` for a workflow or hard rule, `references/pitfalls.md` for a trap, `references/sql_recipes.md` for a query, `references/taxonomy_template.yml` for categories, `scripts/` for tooling | | An optimization pattern, decision gate, flag behavior, invariant, A/B methodology, competitor-diff insight | `skills/optimize-inference-siddharth/` — route via the table in `references/updating-this-skill.md`; only new invariants, flag behavior, and routing lines go in its `SKILL.md` | +| A vLLM path that moved, a stale-layout trap that cost a search, a resolved kernel-to-package mapping | `skills/vllm-codebase-reference/` — `references/navigation-map.md` for paths, the traps section of its `SKILL.md` for stale knowledge. Record the vLLM HEAD sha you verified against. | | Cluster, queue, image, or launch failure that cost real time | `skills/run-qwen-model/SKILL.md` or `skills/cog-setup-and-help/SKILL.md` | | Everything else about this campaign | `skills/run-qwen-model/EXPERIMENTS.md` | @@ -200,6 +258,9 @@ to its revision log. Apply the same discipline to ## Guardrails - Do not change the vLLM baseline configuration. +- The vLLM checkout at `/Users/shanmugamr@nvidia.com/vllm` is read-only. Never + edit, commit, checkout, or clean it — it defines the baseline, and changing it + invalidates every number in `EXPERIMENTS.md` with no error. - Do not change batch size, OSL, checkpoint, hardware, or parallelism to claim a speedup. - Do not stack unmeasured changes; one mechanism per experiment and per PR. diff --git a/.cursor/agents/vllm-codebase-expert.md b/.cursor/agents/vllm-codebase-expert.md new file mode 100644 index 00000000000..163be27ae08 --- /dev/null +++ b/.cursor/agents/vllm-codebase-expert.md @@ -0,0 +1,77 @@ +--- +name: vllm-codebase-expert +description: Answers questions about how vLLM implements something by reading the vLLM source checkout at /Users/shanmugamr@nvidia.com/vllm. Use when you need vLLM's mechanism for MoE routing, expert dispatch, attention, CUDA graph capture, kernel/backend selection, or its host-side decode loop — usually to judge whether Megatron-Core can do the same. Also use to resolve a kernel name from a vLLM trace back to its call site, to find which env var gates a path, or to check whether an optimization is upstream code or an external package. Read-only; never modifies either repo. +model: inherit +readonly: true +is_background: false +--- + +You answer questions about how vLLM implements things, by reading source. + +Your caller is optimizing Megatron-Core inference against a vLLM baseline. It +delegates to you so it does not have to spend its own context reading a +~2200-file tree. Your value is a grounded, cited answer it can act on — not a +tour of the codebase. + +## Mandatory context + +Read `skills/vllm-codebase-reference/SKILL.md` before answering anything. It +defines the checkout, the search discipline, the four question types, the answer +format, and the traps in this revision. Its +`references/navigation-map.md` is a verified path index — consult it before +searching, since it usually removes the need to search at all. + +Read nothing else by default. If the question concerns the Qwen3-30B-A3B +campaign's own history, `skills/run-qwen-model/EXPERIMENTS.md` is the ledger, but +only open it when the question actually depends on what was already tried. + +## The checkout + +`/Users/shanmugamr@nvidia.com/vllm`, the exact revision that runs in the +benchmarks. **Treat it as strictly read-only.** Do not edit, stage, commit, +checkout, stash, or clean anything there. Modifying it silently invalidates every +baseline in the ledger. You may run read-only git commands (`log`, `describe`, +`status`, `show`, `blame`). + +You also do not modify the Megatron-LM repo. You are a research agent: you read +and you report. If the answer implies a code change, describe it; do not make it. + +## How to work + +1. Pin the revision (`git -C ... log --oneline -1`) and confirm the tree is + clean. Report the sha in your answer; flag a dirty tree loudly. +2. Consult the navigation map. Grep for symbols, not concepts. +3. Read the definition and the dispatch point that selects it. +4. Verify before asserting. Every path you cite must exist; every claim gets a + `file.py:line`. If you did not open it, do not cite it. + +## Answer with + +- The mechanism, in two or three sentences, first. +- `vllm/path/file.py:123` for every claim. +- A short verbatim snippet of the load-bearing code. +- **What selects this path and its default.** vLLM usually has several + implementations behind a selector, and only one is on the benchmarked path. + Describing a path the benchmark never executes is worse than saying nothing, + because it is actionable and wrong. +- Whether a kernel or library is upstream vLLM or an external package + (`flashinfer`, `deep_ep`, `deep_gemm` are imported, not vendored). +- **What you did not verify**, stated plainly. +- The HEAD sha, once, so the caller can cite it in `EXPERIMENTS.md`. + +## Stay in your lane + +- Report what vLLM does and what its approach costs. Do not recommend whether + Megatron-Core should adopt it — pricing a lever belongs to the caller and to + `skills/optimize-inference-siddharth/references/decision-gates.md`. +- Do not interpret nsys traces; that is `nsight-system-analysis`. You may resolve + a kernel *name* the caller gives you back to its call site — that is question + type B in the skill and squarely your job. +- Do not explain Megatron-Core's implementation. If a comparison is wanted, state + the vLLM side precisely and say the mcore side is out of scope. +- Never guess to fill a gap. "Not found in this revision, I searched X and Y" is + a correct and useful answer. Inventing a plausible path costs the caller an + experiment. + +If the question is ambiguous, answer the most likely reading and say which +reading you took, rather than returning without an answer. diff --git a/skills/nsight-system-analysis/scripts/idle_decompose.py b/skills/nsight-system-analysis/scripts/idle_decompose.py new file mode 100644 index 00000000000..b8a177cba5a --- /dev/null +++ b/skills/nsight-system-analysis/scripts/idle_decompose.py @@ -0,0 +1,191 @@ +"""Decompose GPU idle inside a steady-state decode window into attackable vs not. + +Total idle is not an actionable number: a large share of it is intra-graph kernel +scheduling that no code change reaches. This splits idle two ways, in the order +that decides whether a host-path lever exists at all. + +1. **By gap size.** Small gaps (default < 10 us) are launch/scheduling granularity + between kernels that are already enqueued. Large gaps are the host chain between + steps, and only those are attackable. +2. **By what occupies each large gap**, read from the CUDA API rows + (`CUPTI_ACTIVITY_KIND_RUNTIME`) that a clean `--trace=cuda,nvtx` capture already + contains. A gap containing no CUDA API call at all is host compute -- Python on + the critical path. A gap covered by a memcpy or a sync is a different lever with + a different fix. + +This is the recommended method rather than a fallback: `osrt`, process-tree +sampling, and NVTX under graph capture all deadlock nsys finalization on MoE decode +workloads, so the trace that can be captured is GPU-only, and the CUDA API rows are +the only host visibility in it. See references/measuring.md. + +Usage: + python idle_decompose.py trace.sqlite '' \ + [--steps N] [--skip N] [--threshold-us F] [--device N] + +The anchor should be a kernel that fires a known number of times per step; the +window is taken between anchor firings so it contains whole steps only. +""" + +import argparse +import sqlite3 +from collections import defaultdict + + +def union_len(intervals): + """Total length covered by a list of (start, end), merging overlaps.""" + if not intervals: + return 0, [] + intervals = sorted(intervals) + merged = [list(intervals[0])] + for s, e in intervals[1:]: + if s <= merged[-1][1]: + merged[-1][1] = max(merged[-1][1], e) + else: + merged.append([s, e]) + return sum(e - s for s, e in merged), merged + + +def main(): + ap = argparse.ArgumentParser() + ap.add_argument("db") + ap.add_argument("anchor", help="kernel name LIKE pattern marking step boundaries") + ap.add_argument("--steps", type=int, default=20, help="steps in the window") + ap.add_argument("--skip", type=int, default=200, help="anchor firings to skip (warmup)") + ap.add_argument("--threshold-us", type=float, default=10.0) + ap.add_argument("--device", type=int, default=0) + args = ap.parse_args() + + con = sqlite3.connect(args.db) + thr = args.threshold_us * 1000.0 # ns + + anchors = [ + r[0] + for r in con.execute( + """SELECT k.start FROM CUPTI_ACTIVITY_KIND_KERNEL k + JOIN StringIds s ON k.demangledName = s.id + WHERE k.deviceId = ? AND s.value LIKE ? + ORDER BY k.start""", + (args.device, args.anchor), + ) + ] + if len(anchors) < args.skip + args.steps + 2: + raise SystemExit( + f"anchor '{args.anchor}' fired {len(anchors)} times; " + f"need > {args.skip + args.steps + 2} for skip={args.skip} steps={args.steps}" + ) + + t0 = anchors[args.skip] + t1 = anchors[args.skip + args.steps] + wall = t1 - t0 + print(f"window: {wall/1e6:.3f} ms over {args.steps} anchor periods " + f"({wall/args.steps/1e6:.4f} ms per period)") + + # --- GPU busy as an interval union across all streams ------------------- + kern = con.execute( + """SELECT k.start, k.end FROM CUPTI_ACTIVITY_KIND_KERNEL k + WHERE k.deviceId = ? AND k.end > ? AND k.start < ?""", + (args.device, t0, t1), + ).fetchall() + # clip to window + kern = [(max(s, t0), min(e, t1)) for s, e in kern] + busy, merged = union_len(kern) + idle = wall - busy + print(f"kernels in window : {len(kern)}") + print(f"GPU busy (union) : {busy/1e6:.3f} ms ({100*busy/wall:.1f}%)") + print(f"GPU idle : {idle/1e6:.3f} ms ({100*idle/wall:.1f}%)") + + # --- idle gaps, split by size ------------------------------------------ + gaps = [] + prev = t0 + for s, e in merged: + if s > prev: + gaps.append((prev, s)) + prev = max(prev, e) + if prev < t1: + gaps.append((prev, t1)) + + small = [(s, e) for s, e in gaps if e - s < thr] + large = [(s, e) for s, e in gaps if e - s >= thr] + small_t = sum(e - s for s, e in small) + large_t = sum(e - s for s, e in large) + + print() + print(f"idle gaps: {len(gaps)} (threshold {args.threshold_us} us)") + print(f" < {args.threshold_us:>5.1f} us : {len(small):6d} gaps {small_t/1e6:8.3f} ms " + f"{100*small_t/idle:5.1f}% of idle {100*small_t/wall:5.1f}% of wall " + f"-- intra-graph scheduling, not attackable") + print(f" >={args.threshold_us:>5.1f} us : {len(large):6d} gaps {large_t/1e6:8.3f} ms " + f"{100*large_t/idle:5.1f}% of idle {100*large_t/wall:5.1f}% of wall " + f"-- host chain, attackable") + if large: + ls = sorted(e - s for s, e in large) + med = ls[len(ls) // 2] + print(f" large-gap size: median {med/1000:.1f} us, " + f"max {ls[-1]/1000:.1f} us, mean {large_t/len(large)/1000:.1f} us") + print(f" large gaps per anchor period: {len(large)/args.steps:.2f}") + + if not large: + print("\nno attackable idle at this threshold") + return + + # --- what occupies the large gaps, from the CUDA API rows --------------- + api = con.execute( + """SELECT r.start, r.end, s.value FROM CUPTI_ACTIVITY_KIND_RUNTIME r + JOIN StringIds s ON r.nameId = s.id + WHERE r.end > ? AND r.start < ?""", + (t0, t1), + ).fetchall() + print(f"\nCUDA API calls in window: {len(api)}") + + by_name = defaultdict(float) + covered_total = 0.0 + uncovered_total = 0.0 + uncovered_n = 0 + api_sorted = sorted(api) + starts = [a[0] for a in api_sorted] + import bisect + + for gs, ge in large: + # API calls overlapping this gap + i = bisect.bisect_left(starts, gs) + j = i + while j > 0 and api_sorted[j - 1][1] > gs: + j -= 1 + hit = [] + k = j + while k < len(api_sorted) and api_sorted[k][0] < ge: + s, e, name = api_sorted[k] + if e > gs: + hit.append((max(s, gs), min(e, ge), name)) + k += 1 + if not hit: + uncovered_total += ge - gs + uncovered_n += 1 + continue + cov, _ = union_len([(h[0], h[1]) for h in hit]) + covered_total += cov + uncovered_total += (ge - gs) - cov + for s, e, name in hit: + by_name[name] += e - s + + print(f"\nlarge-gap idle ({large_t/1e6:.3f} ms) attribution:") + print(f" {'covered by a CUDA API call':46s} {covered_total/1e6:8.3f} ms " + f"{100*covered_total/large_t:5.1f}%") + print(f" {'NO CUDA call -- host/Python on critical path':46s} " + f"{uncovered_total/1e6:8.3f} ms {100*uncovered_total/large_t:5.1f}% " + f"({uncovered_n} fully-empty gaps)") + print(f"\n top CUDA APIs inside large gaps:") + for name, t in sorted(by_name.items(), key=lambda x: -x[1])[:12]: + print(f" {name[:52]:52s} {t/1e6:8.3f} ms {100*t/large_t:5.1f}% of large-gap idle") + + print(f"\nper-anchor-period budget ({args.steps} periods):") + for label, t in ( + ("wall", wall), ("busy", busy), ("idle", idle), + (" idle small (fixed)", small_t), (" idle large (attackable)", large_t), + (" of which host/Python", uncovered_total), + ): + print(f" {label:28s} {t/args.steps/1000:9.2f} us") + + +if __name__ == "__main__": + main() diff --git a/skills/optimize-inference-siddharth/references/vllm-differential.md b/skills/optimize-inference-siddharth/references/vllm-differential.md index c7a4963fea4..e9cc925e431 100644 --- a/skills/optimize-inference-siddharth/references/vllm-differential.md +++ b/skills/optimize-inference-siddharth/references/vllm-differential.md @@ -149,6 +149,14 @@ microbenchmark it at *your* shapes under graph replay, verify the memory layout demands, and only then consider tuning your own. Steps 1 and 2 are minutes; the two sessions spent skipping them were not. +Step 1 does not have to be guesswork. The vLLM source that produced the baseline +is checked out locally, and the `vllm-codebase-expert` subagent resolves a trace +kernel name to its call site, package, and gating flag against that exact +revision — see [vllm-codebase-reference](../../vllm-codebase-reference/SKILL.md). +Ask it rather than inferring the call site from the kernel name; symbol names +identify the library but not which of vLLM's several backends selected it, and +that distinction is the whole of step 2. + An important asymmetry: this works for **leaf kernels with a narrow contract** (attention, norms, sampling) and fails for **fused mega-kernels**, which come with layout and weight-preparation demands — the two traps below. diff --git a/skills/optimize-inference-siddharth/scripts/bandwidth_probe.py b/skills/optimize-inference-siddharth/scripts/bandwidth_probe.py new file mode 100644 index 00000000000..753438e57ff --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/bandwidth_probe.py @@ -0,0 +1,116 @@ +"""Measure the achievable HBM streaming ceiling on this GPU. + +Every memory-bound decision gate in this campaign divides a byte count by a +bandwidth constant. The constant in use (6.081 TB/s) predates this session and is +the denominator behind the MoE grouped-GEMM floor of 49.66 us/layer/rank that +vLLM appears to beat -- so it needs to be a measured machine constant, not a +datasheet number or an inherited one. + +Three probes, in increasing fidelity to the thing being modelled: + 1. pure read -- reduction over one large buffer + 2. read+write -- copy between two buffers + 3. MoE-shaped -- read the actual per-layer expert weight footprint as the + grouped GEMM sees it: `local_experts` separate tensors of + [ffn, hidden] and [ffn, hidden] gate/up plus [hidden, ffn] down +""" + +import torch + +GiB = 1024**3 + + +def timed(fn, warmup=5, iters=20): + for _ in range(warmup): + fn() + torch.cuda.synchronize() + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + start.record() + for _ in range(iters): + fn() + end.record() + torch.cuda.synchronize() + return start.elapsed_time(end) / iters / 1e3 # seconds + + +def report(label, seconds, nbytes): + tbps = nbytes / seconds / 1e12 + print(f" {label:44s} {seconds*1e6:9.2f} us {nbytes/1e6:9.2f} MB {tbps:6.3f} TB/s") + return tbps + + +def main(): + dev = torch.device("cuda") + print(f"device: {torch.cuda.get_device_name(0)}") + props = torch.cuda.get_device_properties(0) + print(f" SMs={props.multi_processor_count} total_mem={props.total_memory/GiB:.1f} GiB") + print() + + results = {} + + print("1. pure read (reduction over one buffer)") + for gib in (1, 2, 4): + n = int(gib * GiB // 2) + x = torch.empty(n, dtype=torch.bfloat16, device=dev) + x.normal_() + t = timed(lambda: x.sum()) + results[f"read_{gib}GiB"] = report(f"read {gib} GiB bf16", t, n * 2) + del x + torch.cuda.empty_cache() + print() + + print("2. read+write (copy between buffers)") + for gib in (1, 2): + n = int(gib * GiB // 2) + a = torch.empty(n, dtype=torch.bfloat16, device=dev) + a.normal_() + b = torch.empty_like(a) + t = timed(lambda: b.copy_(a)) + results[f"copy_{gib}GiB"] = report(f"copy {gib} GiB bf16 (r+w)", t, n * 2 * 2) + del a, b + torch.cuda.empty_cache() + print() + + # 3. The shape the MoE gate actually models: Qwen3-30B-A3B, EP4. + hidden, ffn, n_local = 2048, 768, 32 + print(f"3. MoE per-layer weight footprint (hidden={hidden} ffn={ffn} " + f"local_experts={n_local})") + # gate+up fused = [2*ffn, hidden]; down = [hidden, ffn] + gate_up = [torch.empty(2 * ffn, hidden, dtype=torch.bfloat16, device=dev) for _ in range(n_local)] + down = [torch.empty(hidden, ffn, dtype=torch.bfloat16, device=dev) for _ in range(n_local)] + for t_ in gate_up + down: + t_.normal_() + nbytes = sum(t_.numel() * 2 for t_ in gate_up + down) + + def read_all(): + s = 0 + for t_ in gate_up: + s = s + t_.sum() + for t_ in down: + s = s + t_.sum() + return s + + t = timed(read_all) + tbps = report("read all local expert weights (per layer)", t, nbytes) + + # single contiguous buffer of the same size, as an upper bound on how fast + # that footprint could possibly be streamed + flat = torch.empty(nbytes // 2, dtype=torch.bfloat16, device=dev) + flat.normal_() + t2 = timed(lambda: flat.sum()) + tbps2 = report("same bytes, one contiguous buffer", t2, nbytes) + + print() + print("=== summary ===") + peak = max(results.values()) + print(f" best pure-read bandwidth : {peak:6.3f} TB/s") + print(f" MoE-shaped read (32 tensors x2) : {tbps:6.3f} TB/s") + print(f" MoE-shaped, contiguous : {tbps2:6.3f} TB/s") + print(f" per-layer weight bytes : {nbytes/1e6:.2f} MB") + for name, bw in (("best pure read", peak), ("MoE-shaped", tbps), + ("MoE contiguous", tbps2)): + print(f" floor at {name:16s}: {nbytes/(bw*1e12)*1e6:7.2f} us/layer/rank") + + +if __name__ == "__main__": + main() diff --git a/skills/optimize-inference-siddharth/scripts/osl_decompose.py b/skills/optimize-inference-siddharth/scripts/osl_decompose.py new file mode 100644 index 00000000000..345f86f33f2 --- /dev/null +++ b/skills/optimize-inference-siddharth/scripts/osl_decompose.py @@ -0,0 +1,205 @@ +"""Three-way decomposition: prefill, decode step vs context, for both engines. + +The client reports tpot = wall / num_output_tokens, which folds prefill into every +token. With a fixed prompt set and batch, prefill is the *same constant* in every +wall time, so it cancels in a difference between two OSLs: + + wall(N) = prefill + sum_{i=1..N-1} step(c0 + i) + (wall(N2) - wall(N1)) / (N2 - N1) = mean step over contexts [c0+N1, c0+N2-1] + +Each consecutive pair therefore yields one clean measurement of the decode step at a +known mean context, with no profiler and no prefill contamination. Fitting +step(c) = a + b*c to those points gives the context-scaling slope directly, and +back-solving prefill from every OSL independently checks the linear model: if the +model is wrong, the recovered prefill will not be constant. + +The slope is also checked against the KV-cache bandwidth floor, since re-reading the +KV cache is decode's only meaningful per-context-token cost. A slope already at that +floor has no headroom regardless of how it compares to a competitor. + +Input: one directory per engine containing `osl_.log` files, each the stdout of a +benchmark client that prints `wall=` per iteration. Generate them by running the +same client against the same server at several OSLs, warming up once at the longest +OSL first so graph capture and any autotuning land outside every timed point. + +Usage: + python osl_decompose.py [--label-a A] [--label-b B] + [--prompt-tokens F] [--seqs-per-gpu-a N] [--seqs-per-gpu-b N] + [--bandwidth-tbs F] [--kv-bytes-per-ctx-token-per-seq F] + +Defaults describe Qwen3-30B-A3B on 4xGB200 (mcore EP4/TP1 vs vLLM DP4); override +them for any other model or sharding. +""" + +import argparse +import re +import statistics + +OSLS = [64, 128, 256, 512, 768, 1024, 1152] + +ap = argparse.ArgumentParser() +ap.add_argument("dir_a") +ap.add_argument("dir_b") +ap.add_argument("--label-a", default="engine A") +ap.add_argument("--label-b", default="engine B") +ap.add_argument("--prompt-tokens", type=float, default=60.86328125, + help="mean prompt length; must be identical across all OSL points") +ap.add_argument("--seqs-per-gpu-a", type=int, default=64, + help="sequences whose KV this GPU reads each step (A)") +ap.add_argument("--seqs-per-gpu-b", type=int, default=72) +ap.add_argument("--bandwidth-tbs", type=float, default=7.0, + help="achievable HBM streaming ceiling, TB/s") +ap.add_argument("--kv-bytes-per-ctx-token-per-seq", type=float, default=96.0 * 1024, + help="2*kv_heads*head_dim*dtype_bytes*layers") +ap.add_argument("--osls", type=int, nargs="+", default=OSLS) +args = ap.parse_args() + +C0 = args.prompt_tokens +BW_TBS = args.bandwidth_tbs +OSLS = args.osls + + +def load(run_dir, osls=OSLS): + out = {} + for osl in osls: + try: + txt = open(f"{run_dir}/osl_{osl}.log").read() + except FileNotFoundError: + continue + walls = [float(m) for m in re.findall(r"wall=(\d+(?:\.\d+)?) ms", txt)] + if walls: + out[osl] = walls + return out + + +def fit(points): + """Least squares a + b*c, with the standard error on the slope. + + The slope is the whole argument, so it needs an error bar: the finite + differences are individually noisy and a slope difference inside the error + bars is not a finding. + """ + n = len(points) + sx = sum(p[0] for p in points) + sy = sum(p[1] for p in points) + sxx = sum(p[0] * p[0] for p in points) + sxy = sum(p[0] * p[1] for p in points) + den = n * sxx - sx * sx + b = (n * sxy - sx * sy) / den + a = (sy - b * sx) / n + resid = [p[1] - (a + b * p[0]) for p in points] + s2 = sum(r * r for r in resid) / (n - 2) + se_b = (s2 * n / den) ** 0.5 + ybar = sy / n + sstot = sum((p[1] - ybar) ** 2 for p in points) + r2 = 1 - sum(r * r for r in resid) / sstot if sstot else float("nan") + return a, b, se_b, r2 + + +def analyse(label, run_dir, seqs_per_gpu): + data = load(run_dir) + if not data: + print(f"\n### {label}: no data at {run_dir}") + return None + print(f"\n### {label}") + print(f"{'OSL':>6} {'mean wall ms':>13} {'sd':>7} {'spread':>8} {'n':>3} {'tpot ms':>9}") + means = {} + for osl, walls in sorted(data.items()): + m = statistics.mean(walls) + sd = statistics.stdev(walls) if len(walls) > 1 else 0.0 + means[osl] = m + print(f"{osl:6d} {m:13.1f} {sd:7.1f} " + f"{(max(walls)-min(walls))/m*100:7.2f}% {len(walls):3d} {m/osl:9.4f}") + + print(f"\n finite differences (prefill cancels):") + print(f" {'interval':>14} {'mean context':>13} {'decode step ms':>15}") + pts = [] + ks = sorted(means) + for n1, n2 in zip(ks, ks[1:]): + step = (means[n2] - means[n1]) / (n2 - n1) + ctx = C0 + (n1 + n2 - 1) / 2.0 + pts.append((ctx, step)) + print(f" {f'{n1}->{n2}':>14} {ctx:13.1f} {step:15.4f}") + + a, b, se_b, r2 = fit(pts) + print(f"\n fit: step(c) = {a:.4f} + {b:.6f} * c ms (R2 {r2:.3f})") + print(f" context-free term {a:.4f} ms") + print(f" context slope {b*1000:.3f} +/- {se_b*1000:.3f} us per context token") + # The decode KV read is the only per-context-token cost of consequence: + # 2(K,V) * 4 kv-heads * 128 dim * 2 bytes * 48 layers = 96 KiB per context + # token per sequence, streamed once per step by whichever sequences this GPU + # owns. That sets a hard floor on the slope. + gb_per_ctx_tok = (args.kv_bytes_per_ctx_token_per_seq * seqs_per_gpu + / 1024 ** 3) + floor_us = gb_per_ctx_tok / BW_TBS * 1e3 + print(f" KV floor at {seqs_per_gpu} seqs/GPU: {gb_per_ctx_tok*1024:.2f} MB " + f"per context token => {floor_us:.3f} us at {BW_TBS} TB/s") + print(f" measured / floor = {b*1000/floor_us:.2f}x") + + def mean_step(n): + return a + b * (C0 + n / 2.0) + + print(f"\n {'OSL':>6} {'mean-context step':>18} {'recovered prefill ms':>21}") + prefills = [] + for osl in ks: + pf = means[osl] - sum(a + b * (C0 + i) for i in range(1, osl)) + prefills.append(pf) + print(f" {osl:6d} {mean_step(osl):18.4f} {pf:21.1f}") + print(f" recovered prefill: mean {statistics.mean(prefills):.1f} ms, " + f"sd {statistics.stdev(prefills):.1f} ms " + f"(constant => the linear step model holds)") + return { + "a": a, "b": b, "se_b": se_b, "means": means, + "step128": mean_step(128), "step1024": mean_step(1024), + "prefill": statistics.mean(prefills), + "prefill_sd": statistics.stdev(prefills), + } + + +# seqs-per-gpu is the count whose KV this rank reads each step, and it is easy to get +# wrong: read it off a per-step kernel's grid dimension (e.g. the KV-write kernel) +# rather than assuming batch/world_size. +m = analyse(args.label_a, args.dir_a, args.seqs_per_gpu_a) +v = analyse(args.label_b, args.dir_b, args.seqs_per_gpu_b) + +if m and v: + print("\n\n########## THREE-WAY DECOMPOSITION ##########") + print(f"{'quantity':<34} {'mcore':>11} {'vLLM':>11} {'ratio':>8}") + print("-" * 68) + rows = [ + ("prefill (ms, BS256)", m["prefill"], v["prefill"]), + ("pure decode step @OSL128 (ms)", m["step128"], v["step128"]), + ("pure decode step @OSL1024 (ms)", m["step1024"], v["step1024"]), + ("context slope (us/context token)", m["b"] * 1000, v["b"] * 1000), + ("context-free term (ms)", m["a"], v["a"]), + ("client tpot @OSL1024 (ms)", m["means"].get(1024, 0) / 1024, + v["means"].get(1024, 0) / 1024), + ] + for name, mv, vv in rows: + print(f"{name:<34} {mv:11.4f} {vv:11.4f} {mv/vv:8.3f}x") + + print("\n--- is the context-scaling lever real? ---") + gm = m["step1024"] / m["step128"] + gv = v["step1024"] / v["step128"] + print(f" pure-decode growth OSL128 -> OSL1024: mcore {gm:.4f}x, vLLM {gv:.4f}x") + if_v = m["step128"] * gv + print(f" mcore @OSL1024 if it scaled at vLLM's rate: {if_v:.4f} ms") + print(f" excess attributable to context scaling: {m['step1024']-if_v:+.4f} ms/step") + deficit = m["step1024"] - v["step1024"] + print(f" pure-decode deficit at OSL1024: {deficit:+.4f} ms/step") + if deficit: + print(f" context scaling is {(m['step1024']-if_v)/deficit*100:.1f}% of that deficit") + print("\n--- where the OSL1024 deficit actually lives ---") + cm = C0 + 1024 / 2.0 + print(f" context-free term: mcore {m['a']:8.4f} vLLM {v['a']:8.4f} " + f"delta {m['a']-v['a']:+8.4f} ms") + print(f" context term @c={cm:.0f}: mcore {m['b']*cm:8.4f} vLLM {v['b']*cm:8.4f} " + f"delta {(m['b']-v['b'])*cm:+8.4f} ms") + tot = (m['a']-v['a']) + (m['b']-v['b'])*cm + print(f" total delta {tot:+8.4f} ms") + print(f" => {(m['a']-v['a'])/tot*100:.1f}% of the deficit is context-FREE per-step work") + + print(f"\n parity at the OSL1024 pure decode step needs " + f"{(1 - v['step1024']/m['step1024'])*100:.1f}% off mcore's step") + print(f" parity at the client throughput metric needs " + f"{(m['means'][1024]/v['means'][1024] - 1)*100:.1f}% more throughput") diff --git a/skills/run-qwen-model/EXPERIMENTS.md b/skills/run-qwen-model/EXPERIMENTS.md index 97f1bd6a5aa..af24d04d105 100644 --- a/skills/run-qwen-model/EXPERIMENTS.md +++ b/skills/run-qwen-model/EXPERIMENTS.md @@ -7,6 +7,32 @@ This ledger starts empty. Append every experiment, including failures and regressions. Never edit an earlier result after it is recorded — supersede it with a new row. +## Accepted changes + +**The at-a-glance answer to "what actually landed."** Every accepted experiment +gets a row here *and* a full row in the experiment index below. This table is a +duplicate on purpose: the index interleaves accepts with rejections and gated-out +levers, and detailed records bury PR links hundreds of lines down, so neither is +readable as a summary. One row per accepted mechanism, newest last. + +| ID | Date | Mechanism | Gain | Kill switch | PR | PR state | +|---|---|---|---:|---|---|---| +| — | — | _nothing accepted yet_ | — | — | — | — | + +Rules for this table: + +- A row appears here **only** after the change passed its A/B on distribution + separation and its correctness gate. Nothing provisional. +- `Gain` is the percentage against the same-session OFF arm, not against a + historical baseline. +- `PR` is a markdown link to the draft MR. If a PR is superseded, replace the + link and note the superseded number in the detailed record, not here. +- `PR state` is draft, open, merged, or closed. Refresh it when you touch the + ledger; a stale "draft" against a merged PR is the failure mode this column + exists to catch. +- Reverted-after-the-fact changes are struck through here with a pointer to the + regression's entry. Do not delete the row. + ## Fixed protocol | Setting | Value | @@ -62,6 +88,9 @@ computed, before any code change. 7. Revert regressions or correctness failures. 8. Record the result — including rejections, with their root cause — before beginning another experiment. -9. Open a separate draft MR for each accepted change, and link it from its row. +9. Open a separate draft MR for each accepted change, and link it from **both** + the *Accepted changes* table and its experiment-index row. An accepted change + that is missing from the *Accepted changes* table is not recorded, however + thorough its detailed write-up is. 10. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines once to confirm parity under identical conditions. diff --git a/skills/vllm-codebase-reference/SKILL.md b/skills/vllm-codebase-reference/SKILL.md new file mode 100644 index 00000000000..50f60a0cffd --- /dev/null +++ b/skills/vllm-codebase-reference/SKILL.md @@ -0,0 +1,180 @@ +--- +name: vllm-codebase-reference +description: Answers questions about how vLLM implements something, by reading the vLLM source checkout at /Users/shanmugamr@nvidia.com/vllm. Use when you need to know how vLLM does MoE routing, expert dispatch, attention, CUDA graph capture, kernel or backend selection, sampling, scheduling, or its host-side decode loop — typically to decide whether Megatron-Core can do the same thing. Also use to resolve a kernel name seen in a vLLM nsys trace back to its call site, to find which env var or config flag gates a code path, or to check whether a vLLM optimization is upstream code or an external package (flashinfer, deep_ep, deep_gemm). Do not use for questions about Megatron-Core's own implementation, for running vLLM, or for interpreting a trace — those belong to the mcore skills, run-qwen-model, and nsight-system-analysis. +license: Apache-2.0 +metadata: + author: NVIDIA Corporation +--- + +# vLLM Codebase Reference + +vLLM is the performance target for the Qwen3-30B-A3B campaign. When a +differential says vLLM does something cheaper, the next question is always *how*, +and the answer is in source, not in recollection. This skill turns that question +into a grounded lookup. + +## The checkout + +``` +/Users/shanmugamr@nvidia.com/vllm +``` + +This is the **exact revision that runs in the benchmarks**, so its answers are +authoritative for this campaign in a way that general vLLM knowledge is not. +Recent vLLM releases moved and deleted things aggressively; anything you +remember about vLLM's layout from training data should be treated as a +hypothesis to verify, never as an answer. + +**Read-only.** Never edit, stage, commit, checkout, or `git clean` anything under +that path. Changing it silently invalidates every baseline in `EXPERIMENTS.md`. + +Start every session by pinning the revision, and quote it in your answer: + +```bash +git -C /Users/shanmugamr@nvidia.com/vllm log --oneline -1 +git -C /Users/shanmugamr@nvidia.com/vllm describe --tags 2>/dev/null +git -C /Users/shanmugamr@nvidia.com/vllm status --short # must be clean +``` + +If the tree is dirty, say so in your answer — someone has modified the baseline. +Note that `vllm/_version.py` is generated at build time and is absent here, so +`vllm.__version__` reports `"dev"`. Use the `git describe` string instead. + +## Search discipline + +The tree is ~2200 files (~1600 Python) under `vllm/`. Browsing directories to +orient yourself wastes a large amount of context for a small amount of signal. + +1. **Grep for the symbol, not the concept.** `class FusedMoE`, `def + forward_impl`, `SharedExpertsOrder`, `VLLM_ALL2ALL_BACKEND`. Concept + greps ("expert dispatch") return prose in docstrings and tests. +2. **Consult [references/navigation-map.md](references/navigation-map.md) + first** for anything in the MoE, decode-loop, CUDA-graph, attention, + communication, or config areas. It is a verified path index; it will usually + save you the search entirely. +3. **Read the definition, then the single call site that matters.** vLLM has + many abstract bases and registries. The dispatch point (which concrete class + gets built, and under what condition) is nearly always more informative than + the implementation. +4. **Prefer `vllm/` over `tests/`** for behavior, but `tests/` is the fastest way + to see an API's intended calling convention when a signature is unclear. +5. Do not read `configs/*.json` tile files in bulk. There are hundreds. Grep for + the specific `E=`/`N=`/`device_name=` combination. + +## Four question types + +### A. "How does vLLM implement X?" + +The default. Find the mechanism and, critically, the **condition under which it +is used** — vLLM almost always has several implementations of X behind a +selector, and only one of them is on the benchmarked path. An answer that +describes a code path the benchmark never executes is worse than no answer, +because it is actionable and wrong. + +So always report both the implementation *and* what selects it. For MoE that +means `oracle/` and the all2all manager choice; for attention, the backend +selector; for graphs, the cudagraph mode resolution. + +### B. "Which call site produces kernel symbol `Y` in the trace?" + +The highest-value query in this campaign, and the one the differential workflow +depends on. Names like `fmhaSm100...`, `nvjet...`, `trtllm_gen...`, or an +Inductor-generated `triton_red_fused__to_copy_add_rms_norm_0` need to be +resolved to a package and an entry point before anyone tunes anything. + +Work the name for its provenance first: + +| Name shape | Origin | +|---|---| +| `triton_{poi,red}_fused_*` | `torch.compile` Inductor output — the fusion boundary is the finding, and the source is the Python module that got compiled, not a kernel file | +| `trtllm_gen*`, `fmha*Sm100*` | flashinfer / TRT-LLM-gen, an **external package** | +| `nvjet*`, `sm100_xmma*` | cuBLAS / CUTLASS via an external library | +| `deep_ep*`, `deep_gemm*` | external DeepEP / DeepGEMM packages | +| plain snake_case C++ symbols | usually vLLM's own `csrc/` | + +Then confirm by grepping the Python call site for the launching function. Report +the package, the vLLM file and line that calls it, and whether it is upstream +vLLM code or an external dependency — see question type D. + +### C. "What flag or env var gates path Z?" + +Check `vllm/envs.py` for the `VLLM_*` definition and its default, then grep for +the constant's uses, then check whether a config field in `vllm/config/` +overrides it. Report the default, because "vLLM does X" is only interesting if X +is on by default in the benchmarked configuration. When the answer changes +behavior for our runs, cross-check against the launch command in +`skills/run-qwen-model/` rather than assuming the default applies. + +### D. "Is this upstream vLLM, or an external package?" + +Decisive for whether Megatron-Core can adopt the same thing, and the difference +between a day of work and a quarter of it. `flashinfer`, `deep_ep`, `deep_gemm`, +and `pplx_kernels` are **imported, not vendored** in this checkout. If the win +lives in an external package, the follow-up question is whether that package is +already installed in our environment — per the differential workflow, it +frequently is. + +Genuinely vendored in-tree: `vllm/vllm_flash_attn/` (a FlashAttention interface +shim) and `vllm/third_party/`. + +## Answer format + +Optimize the answer for a caller who will not read the source themselves. Be +complete on mechanism and exact on location; skip everything else. + +- **Direct answer first** — the mechanism, in two or three sentences. +- **Paths with line numbers**, `vllm/path/file.py:123`, for every claim. +- **A short verbatim snippet** for the load-bearing part. Quote it; do not + paraphrase code. +- **What selects this path**, and its default. If it is off by default, say so + prominently. +- **Upstream or external package**, when a kernel or library is involved. +- **What you did not check.** Say "I did not verify whether this is on the + benchmarked path" rather than implying you did. An unverified claim presented + as verified is the one failure mode that costs real experiment time here. +- **The HEAD sha**, once, so the answer can be cited in `EXPERIMENTS.md`. + +Do not editorialize about whether Megatron-Core should adopt something. Report +what vLLM does and what it costs to do it that way; the pricing decision belongs +to `optimize-inference-siddharth/references/decision-gates.md`. + +## Traps in this revision + +Three pieces of common vLLM knowledge are stale here, and each has already sent +a search in the wrong direction: + +1. **`vllm/attention/` does not exist.** Attention moved to `vllm/v1/attention/` + (backends, selector, metadata builders) and + `vllm/model_executor/layers/attention/` (the `Attention` nn.Module). +2. **`fused_moe/` is now subpackages**, not flat files: `prepare_finalize/`, + `experts/`, `router/`, `runner/`, `oracle/`. Paths like + `fused_moe/cutlass_moe.py` have moved. +3. **The pplx and naive all2all backends were removed.** Both strings are still + accepted and are silently rewritten to `allgather_reducescatter` in + `vllm/config/parallel.py`. There is no pplx code in the tree. A config that + *says* pplx is not running pplx. + +Also: there are **two** GPU model runners. `vllm/v1/worker/gpu_model_runner.py` +is production; `vllm/v1/worker/gpu/model_runner.py` is an experimental "V2" +rewrite with its own separate `InputBatch`. Read the former unless you have +confirmed the run uses the latter. + +## Related skills + +| Question | Skill | +|---|---| +| How does vLLM do this? | this skill | +| What does the trace say? | `nsight-system-analysis` | +| Should we build it, and how? | `optimize-inference-siddharth` | +| How do I run vLLM or mcore? | `run-qwen-model` | + +`optimize-inference-siddharth/references/vllm-differential.md` is the consumer of +this skill: it defines the comparison that generates these questions, and records +what past answers led to. + +## Keeping this skill current + +Update [references/navigation-map.md](references/navigation-map.md) when a lookup +finds that a path has moved, and add a numbered entry to the traps section when a +piece of stale knowledge costs a real search. Record the HEAD sha you verified +against. Paths are the perishable part of this skill; the discipline above is not. diff --git a/skills/vllm-codebase-reference/references/navigation-map.md b/skills/vllm-codebase-reference/references/navigation-map.md new file mode 100644 index 00000000000..4cd3c19dfca --- /dev/null +++ b/skills/vllm-codebase-reference/references/navigation-map.md @@ -0,0 +1,242 @@ +# vLLM Navigation Map + +Verified path index for `/Users/shanmugamr@nvidia.com/vllm`. + +**Verified at HEAD `b8160878f`** (`v0.19.2rc0-219-gb8160878f`). Every path below +was confirmed to exist at that revision. Re-verify before quoting if the checkout +has advanced; vLLM moves files frequently. Line numbers are approximate anchors — +grep the symbol rather than trusting the number. + +Paths are relative to the checkout root. + +## MoE and expert parallelism + +The area this campaign asks about most. + +### Layer and dispatch abstraction + +| Path | What it is | +|---|---| +| `vllm/model_executor/layers/fused_moe/layer.py` | `FusedMoE` module (~219), `forward` (~1545) | +| `vllm/model_executor/layers/fused_moe/modular_kernel.py` | The modular-kernel abstraction: `FusedMoEPrepareAndFinalize`, `FusedMoEExperts`, `FusedMoEKernel`, `TopKWeightAndReduce`, `ExpertTokensMetadata` | +| `vllm/model_executor/layers/fused_moe/config.py` | MoE parallel and quant config resolution | +| `vllm/model_executor/layers/fused_moe/fused_moe_method_base.py` | Base quant-method interface for MoE | +| `vllm/model_executor/layers/fused_moe/fused_moe_modular_method.py` | Modular-kernel quant-method wiring | +| `vllm/model_executor/layers/fused_moe/unquantized_fused_moe_method.py` | The bf16/fp16 path — **this is the one Qwen3-30B-A3B bf16 takes** | +| `vllm/model_executor/layers/fused_moe/oracle/` | Backend *selection* per quant type: `unquantized.py`, `fp8.py`, `nvfp4.py`, `mxfp4.py`, `mxfp8.py`, `int8.py`, `int_wna16.py` | + +`oracle/` answers "which kernel does vLLM actually pick", which is usually the +question behind "how does vLLM do MoE". + +### Prepare/finalize — dispatch and combine + +`vllm/model_executor/layers/fused_moe/prepare_finalize/` + +| File | Backend | +|---|---| +| `deepep_ht.py` | DeepEP high-throughput | +| `deepep_ll.py` | DeepEP low-latency | +| `no_dp_ep.py` | Single-rank, no DP/EP | +| `naive_dp_ep.py` | Allgather-reducescatter DP/EP | +| `batched.py` | Batched per-expert padded layout | +| `flashinfer_nvlink_one_sided.py`, `flashinfer_nvlink_two_sided.py` | FlashInfer NVLink all2all | +| `mori.py`, `nixl_ep.py` | MORI (ROCm) and NIXL transports | + +### All2all manager selection + +| Path | What it is | +|---|---| +| `vllm/distributed/device_communicators/all2all.py` | All manager classes: `AgRsAll2AllManager` (~41), `DeepEPHTAll2AllManager` (~197), `DeepEPLLAll2AllManager` (~261), `FlashInferNVLinkTwoSidedManager` (~449), `FlashInferNVLinkOneSidedManager` (~556), `MoriAll2AllManager` (~672) | +| `vllm/distributed/device_communicators/cuda_communicator.py` (~118-166) | The if/elif chain mapping the `all2all_backend` string to a manager | +| `vllm/config/parallel.py` (~40-51) | `All2AllBackend` Literal — the valid names | +| `vllm/config/parallel.py` (~418) | Where `pplx` and `naive` are warned about and rewritten to `allgather_reducescatter` | +| `vllm/model_executor/layers/fused_moe/all2all_utils.py` | Shared helpers | + +### Triton kernel and tuned configs + +| Path | What it is | +|---|---| +| `vllm/model_executor/layers/fused_moe/fused_moe.py` | The Triton fused MoE kernel. `get_moe_configs` (~1035), `fused_experts` (~1605), `fused_experts_impl` (~1682) | +| `vllm/model_executor/layers/fused_moe/configs/` | ~316 tuned tile-config JSONs, `E=,N=,device_name=[,dtype=...].json`. Grep for the specific shape; do not list. | +| `moe_align_block_size.py`, `moe_permute_unpermute.py`, `topk_weight_and_reduce.py` | Scatter/gather/reduce support ops (same directory) | + +### Expert implementations + +`vllm/model_executor/layers/fused_moe/experts/` + +| File | Implementation | +|---|---| +| `cutlass_moe.py` | CUTLASS grouped GEMM | +| `deep_gemm_moe.py`, `batched_deep_gemm_moe.py` | DeepGEMM contiguous and masked-batched | +| `flashinfer_cutedsl_moe.py`, `flashinfer_cutedsl_batched_moe.py` | FlashInfer CuteDSL | +| `trtllm_bf16_moe.py`, `trtllm_fp8_moe.py`, `trtllm_nvfp4_moe.py`, `trtllm_mxfp4_moe.py` | TRT-LLM-gen. `trtllm_bf16_moe.py` is the one probed and rejected for mcore over its 4-D pre-shuffled weight-layout requirement. | +| `gpt_oss_triton_kernels_moe.py`, `nvfp4_emulation_moe.py`, `ocp_mx_emulation_moe.py`, `xpu_moe.py` | Others | + +Still flat at the `fused_moe/` root: `fused_batched_moe.py`, +`flashinfer_cutlass_moe.py`, `triton_cutlass_moe.py`, `triton_deep_gemm_moe.py`, +`fused_marlin_moe.py`, `rocm_aiter_fused_moe.py`, `fallback.py`. + +### Routing and top-k + +`vllm/model_executor/layers/fused_moe/router/` + +| Path | What it is | +|---|---| +| `grouped_topk_router.py` | `fused_grouped_topk` (~29), `grouped_topk` (~81), `GroupedTopk` CustomOp (~167), `GroupedTopKRouter` (~247) | +| `router_factory.py` | `create_fused_moe_router` (~35) — the dispatch point | +| `fused_moe_router.py` | `FusedMoERouter` ABC | +| `fused_topk_router.py`, `fused_topk_bias_router.py`, `custom_routing_router.py`, `zero_expert_router.py`, `routing_simulator_router.py`, `gate_linear.py` | Other routers | +| `csrc/moe/topk_softmax_kernels.cu`, `csrc/moe/grouped_topk_kernels.cu`, `csrc/moe/moeTopKFuncs.cuh` | The CUDA routing kernels | + +This is where vLLM's one-kernel router lives — the counterpart to mcore's +four-to-five-kernel routing chain in the differential. + +### Shared experts and overlap + +| Path | What it is | +|---|---| +| `vllm/model_executor/layers/fused_moe/runner/shared_experts.py` | `SharedExpertsOrder` IntEnum (~27): `NONE`, `NO_OVERLAP`, `MK_INTERNAL_OVERLAPPED` (overlapped with dispatch/combine inside the modular kernel), `MULTI_STREAM_OVERLAPPED` (aux stream, overlapped with gate/router/experts). `SharedExperts` (~41), `_disable_shared_experts_overlap` heuristic (~82) | +| `vllm/model_executor/layers/fused_moe/runner/moe_runner.py`, `moe_runner_interface.py` | The runner invoking experts | + +### EPLB — expert load balancing + +`vllm/distributed/eplb/`: `eplb_state.py` (`EplbState` ~210), +`rebalance_execute.py` (physical weight movement), `eplb_communicator.py`, +`async_worker.py`, `eplb_utils.py`, `policy/`. Runner-side hooks in +`vllm/v1/worker/gpu/eplb_utils.py`. + +## Qwen3 models + +| Path | What it is | +|---|---| +| `vllm/model_executor/models/qwen3_moe.py` | `Qwen3MoeMLP` (~92), `Qwen3MoeSparseMoeBlock` (~137), `Qwen3MoeAttention` (~261), `Qwen3MoeDecoderLayer` (~364), `Qwen3MoeModel` (~440), `Qwen3MoeForCausalLM` (~675) — **the benchmarked model** | +| `vllm/model_executor/models/qwen3.py` | Dense Qwen3 | +| `qwen3_next.py`, `qwen3_next_mtp.py`, `qwen3_vl_moe.py`, `qwen3_omni_moe_thinker.py` | Hybrid/SSM, MTP, and multimodal variants | + +## V1 decode loop and host path + +| Path | What it is | +|---|---| +| `vllm/v1/worker/gpu_model_runner.py` | **The production GPU model runner** | +| `vllm/v1/worker/gpu/model_runner.py` | Experimental "Model Runner V2" — see `vllm/v1/worker/gpu/README.md`. Its `vllm/v1/worker/gpu/input_batch.py` defines a *different* `InputBatch`. Do not confuse the two. | +| `vllm/v1/worker/gpu_worker.py` | GPU worker: device init, memory profiling, KV cache allocation | +| `vllm/v1/worker/gpu_input_batch.py` | The persistent batch: `CachedRequestState` (~30), `InputBatch` (~81) | +| `vllm/v1/worker/block_table.py` | Per-request device block table across steps | +| `vllm/v1/core/sched/scheduler.py` | The scheduler | +| `vllm/v1/core/sched/async_scheduler.py` | `AsyncScheduler` (~12) — schedules step N+1 while N runs on GPU | +| `vllm/v1/core/sched/output.py`, `interface.py`, `request_queue.py` | Scheduler output struct, ABC, queue policies | +| `vllm/v1/engine/core.py` | `EngineCore` (~91), `EngineCoreProc` (~806), `run_busy_loop` (~1164, and a DP variant ~1731) | +| `vllm/v1/engine/async_llm.py` | Async front end | +| `vllm/forward_context.py` | Per-forward context threaded into the model: attention metadata, DP metadata, cudagraph runtime mode | + +CPU/GPU overlap beyond async scheduling: `vllm/v1/worker/ubatching.py` and +`vllm/v1/worker/gpu_ubatch_wrapper.py` (DBO-style micro-batch overlap), +`vllm/v1/worker/gpu/async_utils.py` (async output copy, V2 runner). + +## CUDA graphs and compilation + +| Path | What it is | +|---|---| +| `vllm/v1/cudagraph_dispatcher.py` | `CudagraphDispatcher` — picks runtime mode and key per step | +| `vllm/compilation/cuda_graph.py` | `CUDAGraphWrapper` (~145) which captures and replays; `CUDAGraphEntry`, `CUDAGraphOptions` | +| `vllm/config/compilation.py` | `CUDAGraphMode` (~53) with `FULL_AND_PIECEWISE` (~63, the v1 default), `cudagraph_capture_sizes` (~622), `resolve_cudagraph_mode_and_sizes` (~1301) | +| `vllm/config/vllm.py` (~1423) | `_set_cudagraph_sizes` — **where the default bucket list is generated**: `[1, 2, 4] + list(range(8, 256, 8)) + ...` (~1432). Read this for bucket policy, not `compilation.py`. | +| `vllm/compilation/backends.py` | `VllmBackend` (~798), `PiecewiseCompileInterpreter` (~680), `CompilerManager` (~123) | +| `vllm/compilation/piecewise_backend.py` | `PiecewiseBackend` (~86) | +| `vllm/compilation/decorators.py` | `@support_torch_compile`, dynamic-dim marking | +| `vllm/compilation/partition_rules.py`, `passes/` | Graph-partition rules and fusion passes | + +## Attention + +`vllm/attention/` does not exist at this revision. + +| Path | What it is | +|---|---| +| `vllm/v1/attention/selector.py` | `get_attn_backend` (~53) — the selection entry point | +| `vllm/v1/attention/backends/registry.py` | `AttentionBackendEnum` (~34), `register_backend` (~211) | +| `vllm/platforms/cuda.py` | Platform-level `get_attn_backend` override applying CUDA eligibility rules | +| `vllm/v1/attention/backend.py` | The abstract base file: `AttentionBackend`, `AttentionMetadataBuilder`, `CommonAttentionMetadata` all live here — **not** in `backends/utils.py`, which is helpers only | +| `vllm/v1/attention/backends/flash_attn.py` | FlashAttention | +| `vllm/v1/attention/backends/flashinfer.py` | FlashInfer — the backend behind the `trtllm-gen` Blackwell decode kernel in the differential | +| `vllm/v1/attention/backends/triton_attn.py` | Triton unified attention | +| `vllm/v1/attention/backends/mla/` | MLA variants | +| `vllm/v1/attention/ops/` | Attention op kernels | +| `vllm/model_executor/layers/attention/attention.py` | The `Attention` nn.Module used inside models | +| `vllm/vllm_flash_attn/flash_attn_interface.py` | Vendored FA Python shim | + +## KV cache and paging + +`vllm/v1/kv_cache_interface.py` (specs and layout), `vllm/v1/core/kv_cache_manager.py` +(paged manager), `vllm/v1/core/block_pool.py` (allocation and prefix-cache hash +table), `vllm/v1/core/kv_cache_utils.py` (block hashing), +`vllm/v1/core/kv_cache_coordinator.py` and `single_type_kv_cache_manager.py` +(multi-group, hybrid models), `csrc/cache_kernels.cu` and +`csrc/cache_kernels_fused.cu` (reshape-and-cache). + +## Distributed and communication + +| Path | What it is | +|---|---| +| `vllm/distributed/parallel_state.py` | Process group setup and accessors for TP/PP/DP/EP | +| `vllm/distributed/device_communicators/base_device_communicator.py` | `DeviceCommunicatorBase` | +| `vllm/distributed/device_communicators/cuda_communicator.py` | CUDA communicator; also selects the all2all manager | +| `vllm/distributed/device_communicators/custom_all_reduce.py` | Custom all-reduce; kernel in `csrc/custom_all_reduce.cu` | +| `vllm/distributed/device_communicators/symm_mem.py` | `SymmMemCommunicator` (~25) — PyTorch symmetric memory, the closest thing to an NVLS/multicast path. Multicast pointer check ~105 falls back when `multicast_ptr == 0`. | +| `vllm/distributed/device_communicators/flashinfer_all_reduce.py`, `mnnvl_compat.py` | FlashInfer all-reduce, MNNVL compat | +| `vllm/distributed/device_communicators/pynccl.py`, `pynccl_wrapper.py`, `pynccl_allocator.py` | NCCL bindings | +| `vllm/distributed/communication_op.py` | High-level TP collectives | + +## Config and environment + +| Path | What it is | +|---|---| +| `vllm/envs.py` | Every `VLLM_*` env var and its default | +| `vllm/config/vllm.py` | `VllmConfig` — stitches sub-configs and runs cross-config resolution, including cudagraph sizes | +| `vllm/config/parallel.py` | `ParallelConfig`: TP/PP/DP/EP, all2all backend | +| `vllm/config/compilation.py` | `CompilationConfig`, `CUDAGraphMode` | +| `vllm/config/cache.py` | `CacheConfig`: block size, KV dtype, GPU memory utilization | +| `vllm/config/model.py`, `scheduler.py` | `ModelConfig`; `SchedulerConfig` (max num seqs, chunked prefill, async scheduling) | +| `vllm/config/attention.py`, `kernel.py`, `quantization.py`, `speculative.py`, `mamba.py`, `load.py` | The rest | + +## Benchmarks + +Thin CLI scripts over a library. `benchmarks/benchmark_throughput.py`, +`benchmark_serving.py`, `benchmark_latency.py` are entry points; +`vllm/benchmarks/throughput.py`, `serve.py`, `latency.py` are the implementations +behind `vllm bench {throughput,serve,latency}`. Microbenchmarks worth knowing +about for kernel comparison: `benchmarks/kernels/`, +`benchmarks/attention_benchmarks/`, `benchmarks/fused_kernels/`, +`benchmarks/cutlass_benchmarks/`. + +## csrc + +88 `.cu` files. Notable: `csrc/moe/` (topk softmax, grouped topk, +permute/unpermute, marlin wna16, mxfp8, DSv3 router GEMM), `csrc/attention/` +including `csrc/attention/mla/`, `csrc/quantization/` (machete, marlin, w8a8, +awq, gptq, gguf, hadamard), `csrc/mamba/mamba_ssm/`, `csrc/cutlass_extensions/`, +`csrc/quickreduce/`, `csrc/rocm/`, `csrc/cpu/`. + +## External vs vendored + +All imported, none vendored: + +| Package | Where imported | +|---|---| +| `flashinfer` | 13 files, incl. `vllm/v1/attention/backends/flashinfer.py`, `experts/trtllm_{bf16,fp8,nvfp4}_moe.py`, `flashinfer_all_reduce.py` | +| `deep_ep` | Exactly 3 files: `device_communicators/all2all.py`, `prepare_finalize/deepep_{ht,ll}.py` | +| `deep_gemm` | Wrapped behind `vllm/utils/deep_gemm.py`; direct import only in `vllm/model_executor/warmup/kernel_warmup.py`. MoE code imports the wrapper (`fused_moe/deep_gemm_utils.py`). | +| `pplx_kernels` | **Not present.** Removed from the tree. | + +Vendored in-tree: `vllm/vllm_flash_attn/` and `vllm/third_party/` +(`flashmla`, `pynvml.py`). + +## Known absent + +Do not send a search after these; they do not exist at this revision. + +- `vllm/attention/` — moved, see the attention section. +- `pplx_kernels` and any pplx all2all backend — removed. The `pplx` config string + is accepted and silently rewritten. +- A naive all2all implementation — same treatment. The nearest real code is + `AgRsAll2AllManager` in `all2all.py` and `prepare_finalize/naive_dp_ep.py`. +- `vllm/_version.py` — generated at build time, absent in this source tree. From 3d769a5395e32a0e5ff74a0d8ef227f3e74bf787 Mon Sep 17 00:00:00 2001 From: shanmugamr1992 Date: Fri, 7 Aug 2026 13:26:05 -0700 Subject: [PATCH 4/4] Added staced pr --- .cursor/agents/qwen-model-optimizer.md | 167 ++++++++++++++---- skills/optimize-inference-siddharth/SKILL.md | 15 +- .../references/updating-this-skill.md | 2 +- .../references/vllm-differential.md | 10 +- skills/run-qwen-model/EXPERIMENTS.md | 120 ++++++++----- 5 files changed, 224 insertions(+), 90 deletions(-) diff --git a/.cursor/agents/qwen-model-optimizer.md b/.cursor/agents/qwen-model-optimizer.md index 2cfe4668cc6..8ee425f6032 100644 --- a/.cursor/agents/qwen-model-optimizer.md +++ b/.cursor/agents/qwen-model-optimizer.md @@ -28,7 +28,12 @@ There are exactly two performance skills: analysis and optimization. A third, read on demand rather than up front. Do not go looking for others, and do not invent skill paths. -Treat `EXPERIMENTS.md` as the sole source of performance history. +`EXPERIMENTS.md` is the sole record of **this campaign's** results — every number +you measure goes there and nowhere else. It is not the only source of prior +knowledge: the two performance skills carry measured findings from earlier work on +this same model, including flags and levers already tried and rejected with their +mechanism. Trust those to avoid re-deriving dead ends, but never treat a number +quoted in a skill as a baseline. Re-measure before you compare. ## Division of labour @@ -115,6 +120,9 @@ Do not modify Megatron-Core until both fresh baselines are recorded: Use the commands and fixed protocol in `skills/run-qwen-model/SKILL.md`. +These two baselines are the **only** measurements taken at `main`. From the first +experiment onward every run happens on the stack described below. + ## Optimization loop Repeat until mcore reaches vLLM: @@ -127,8 +135,9 @@ Repeat until mcore reaches vLLM: .sqlite .sqlite --label-a mcore --label-b vllm ``` - The current baseline trace paths are recorded in `EXPERIMENTS.md`; re-capture - rather than hunt for a trace you cannot open. `forward_pass.py` auto-isolates + Baseline trace paths live in `EXPERIMENTS.md` once recorded. At the start of a + campaign the ledger has none — capture them per the baseline gate rather than + hunting for an older trace. `forward_pass.py` auto-isolates one decode step per engine and prints wall time, GPU-busy vs idle, launch counts, and a per-category Δ table. Same µs/kernel with more launches ⇒ the lever is @@ -149,19 +158,29 @@ Repeat until mcore reaches vLLM: itself costs, and write down *proceed* or *gated out*. Skip the gate only for cheap reversible changes (flag flips, tile retunes, backend swaps). 4. State one measurable hypothesis. -5. Back up or capture the current diff before editing. +5. Cut a branch from the **stack tip** — every accepted change applied — and never + from `main`. Capture the current diff before editing. 6. Implement one change in the bottleneck's source path, honoring the optimization skill's hard rules, and add a kill switch. 7. Run focused correctness tests. 8. Run the fixed BS256 mcore benchmark using the same-session, back-to-back, alternating-arms A/B protocol (hard rule 10). Cross-session comparisons drift - more than most individual wins. -9. Capture a new profile when the timing composition could have changed. + more than most individual wins. **Both arms run on the stack tip**, with every + accepted change below this one enabled in both; the only thing that flips + between arms is this change's kill switch. The result is a *marginal* gain on + top of the stack, which is what the reviewer of this PR will get — not the gain + the change would have shown against `MCORE-BASELINE`. +9. Capture a new profile when the timing composition could have changed. Re-anchor + against the current stack tip, not the baseline trace: earlier accepted changes + have already moved the composition, so a category that looked dominant at + baseline may no longer be. 10. Append the complete result to `EXPERIMENTS.md` — including rejections, with - their root cause and date. -11. Keep improvements; revert regressions and correctness failures. -12. For each accepted change, open its own MR and record it in all three ledger - locations, including the pinned *Accepted changes* table (next section). + their root cause, date, and the stack position they were measured on. +11. Keep improvements; revert regressions and correctness failures. A reverted + change is removed from the stack, and the tip moves back to its parent. +12. For each accepted change, open its stacked MR and record it in all three ledger + locations, including *The merge stack* table (next section). The new branch + becomes the stack tip for the next iteration. 13. Promote whatever generalizes into the skills (section after that). Never optimize from a warmup/capture-only nsys window. `forward_pass.py` already @@ -170,25 +189,88 @@ its reported forward-pass period against measured TPOT before acting. Always reason in terms of one forward pass: GPU-busy interval union, GPU idle, per-category time, launch counts, and critical communication or kernel tails. -## Ship each accepted experiment as its own MR +## Ship accepted experiments as a stacked chain of MRs -Every experiment that succeeded gets a **separate** draft PR against -`main` in `NVIDIA/Megatron-LM`. One mechanism per PR, so the measurement -attributes cleanly and CODEOWNERS review stays narrow (see +Every experiment that succeeded gets its own draft PR, one mechanism per PR, so +the measurement attributes cleanly and CODEOWNERS review stays narrow (see `skills/mcore-split-pr/SKILL.md` if a change spans several owner groups). -Before opening it, walk `skills/optimize-inference-siddharth/assets/review-checklist.md`. - -Mechanics, per `CLAUDE.md`: branch off `main`, commit with both `-s` and `-S`, -push to your **personal fork** — never to `NVIDIA/Megatron-LM` — then -`gh pr create --draft`. `origin` here points at the upstream repo, so confirm a -fork remote exists and ask the user for it if it does not. +**The PRs are stacked, not independent.** `main` is used exactly once — to cut the +branch for stack position 1. Every accepted change after that branches off the +previous accepted change, **whether or not the two are logically related**. Do not +return to `main` between experiments. + +``` +main ──● (baselines only) + └──● 01 ──● 02 ──● 03 ──● … +``` + +The reviewer merges in ascending order, rebasing each row as its parent lands, so +each review sees one mechanism even though the work was cumulative. Your job is to +make that sequence unambiguous. + +**Maintain one local branch that is always the stack tip** — every accepted change +applied, nothing rejected. Cut each new experiment from the tip, and when the +experiment is accepted, the tip becomes that new branch. This is also what makes +the next experiment's measurement correct: you are always measuring on top of what +has already been accepted, so the marginal gain you report is the marginal gain a +reviewer will actually get. + +Before opening a PR, walk +`skills/optimize-inference-siddharth/assets/review-checklist.md`. + +### Mechanics + +Per `CLAUDE.md`: commit with both `-s` and `-S`, push to your **personal fork** — +never to `NVIDIA/Megatron-LM` — then `gh pr create --draft`. `origin` here points +at the upstream repo, so confirm a fork remote exists and ask the user for it if it +does not. + +Stack-specific mechanics on top of that: + +- **Branch names carry the stack position**: `perf/qwen-NN-`, zero-padded, + ascending. The merge order must be readable from the branch name alone. +- **Cut from the predecessor**, not `main`: + `git checkout -b perf/qwen-02- perf/qwen-01-`. +- **The PR base is still `main`.** Branches cannot be pushed to + `NVIDIA/Megatron-LM`, so a PR cannot target a parent branch that only exists on + your fork. Consequence: until the parent merges, the PR diff shows the parent's + changes too. Say so in the description — see the stack block below — so a + reviewer does not read it as scope creep. +- **When the parent merges, rebase past it** and force-push. A plain + `git rebase main` reapplies commits already in `main` and conflicts, especially + when the parent was squash-merged: + + ```bash + git fetch origin main + git rebase --onto origin/main perf/qwen-01- perf/qwen-02- + git push --force-with-lease perf/qwen-02- + ``` + + Force-push only to your own fork's stack branches. +- **If a row is rejected upstream** rather than merged, rebase the rows above it + past it with the same `--onto` form, strike its ledger row through with the + reason, and retire its number. Re-measure the rows that were sitting on it if + they touched the same code path — their marginal gains were measured against a + baseline that no longer exists. + +### Every PR description states its stack position + +Add this block at the top of the description, above the usual sections: + +``` +Stack position 2 of N. Base: main. Parent: # (perf/qwen-01-). +Merge after #. Until then this PR's diff also contains the parent's +changes; the commit for this mechanism is . +``` The description must let a reviewer judge the change without re-running it: - **What changed and why** — a brief summary of the mechanism, not a file list. -- **Measured gain** — percentage over baseline, with the baseline it is measured - against, and the absolute throughput/latency/TPOT numbers. +- **Measured gain** — the marginal percentage over the stack level below this + change, named explicitly, plus the cumulative absolute + throughput/latency/TPOT with the whole stack applied. Never quote a marginal + gain against `MCORE-BASELINE`; on a stack that overstates the change. - **Protocol** — hardware, model, batch, OSL, parallelism, warmup/timed counts, and that the arms ran back to back in the same allocation. State arm separation (`min(ON) > max(OFF)`), not just the mean delta. @@ -200,25 +282,28 @@ The description must let a reviewer judge the change without re-running it: - **Scope and risks** — configs where it does not apply or was not measured. - **Artifacts** — ledger entry id, run/job paths, `.nsys-rep` / `.sqlite` paths. +### Record it in the ledger + Record every accepted change in `EXPERIMENTS.md` in **three** places, and treat the change as unrecorded until all three exist: -1. A row in the pinned **Accepted changes** table at the top — mechanism, gain, - kill switch, PR link, PR state. This is the table a human reads months later - to answer "what landed"; it is the only view that is not interleaved with - rejections or buried in a session write-up. -2. A complete row in the **Experiment index**, including its `MR` column. +1. A row in **The merge stack** table at the top, with its `#`, mechanism, branch, + `Cut from`, PR link, PR state, marginal gain, cumulative throughput, and kill + switch. This table is the merge order; it is the one view a human reads later + to answer "what landed, and in what sequence". +2. A complete row in the **Experiment index**, including `Measured on` and `PR`. 3. The detailed record, with protocol, attribution, and artifacts. -The duplication is deliberate. Do not "simplify" it by leaving the PR link in only -one place — a PR recorded only inside a detailed record is effectively lost, and -that has already happened once in this ledger. +The duplication is deliberate — do not collapse it. A PR link that exists only +inside a detailed record is unreachable in practice, and the `#`/`Cut from` +columns are the only machine-readable record of the stack's shape once the +branches are gone. -Refresh the `PR state` column whenever you touch the ledger, so a merged PR does -not sit there reading "draft". +Assign `#` at creation and never reuse or reorder it. Refresh `PR state` whenever +you touch the ledger, so a merged PR does not sit there reading "draft". -Rejected experiments do not get a PR and do not get an *Accepted changes* row — -they get an index row and a detailed record with the root cause. +Rejected experiments do not get a PR and do not get a merge-stack row — they get +an index row and a detailed record with the root cause. ## Keep the skills learnable @@ -263,7 +348,13 @@ to its revision log. Apply the same discipline to invalidates every number in `EXPERIMENTS.md` with no error. - Do not change batch size, OSL, checkpoint, hardware, or parallelism to claim a speedup. -- Do not stack unmeasured changes; one mechanism per experiment and per PR. +- One mechanism per experiment and per PR. Accepted changes *are* stacked, each on + the last — that is the required workflow. What is forbidden is carrying more than + one unmeasured mechanism at a time: never have two changes in flight on the tip + whose effects have not been separated by their own A/B. +- Never measure on `main` after the baselines, and never cut an experiment branch + from `main`. A number measured off the stack is not comparable to anything in the + ledger. - Never invent expected or observed metrics. - Do not retain a change that only improves throughput by breaking coherence. - Do not use `git reset --hard` or `git clean`. @@ -284,5 +375,7 @@ Finish only when one of these is true: concrete next step. Return a concise summary containing baseline, best mcore result, remaining gap, -accepted changes with their MR links, rejected experiments with root causes, -skill updates made, and artifact paths. +the merge stack in order with each row's PR link and marginal gain, rejected +experiments with root causes, skill updates made, and artifact paths. State the +merge sequence explicitly — the reviewer's first question is which PR to merge +first. diff --git a/skills/optimize-inference-siddharth/SKILL.md b/skills/optimize-inference-siddharth/SKILL.md index c36c6d33aef..019e268990f 100644 --- a/skills/optimize-inference-siddharth/SKILL.md +++ b/skills/optimize-inference-siddharth/SKILL.md @@ -67,8 +67,8 @@ often they paid off: identical launch count, the usual cause is that the competitor is calling a different *library*, not a better-tuned version of yours — and on a shared image that library is often already installed. Switching the decode attention to - flashinfer's `trtllm-gen` was +2.6%; the two sessions spent tuning inside the wrong - package were worth nothing. See + flashinfer's `trtllm-gen` was +2.6%, while the preceding effort spent tuning + inside the wrong package was worth nothing. See [references/vllm-differential.md](references/vllm-differential.md). Everything else in this skill is a consequence of these six. @@ -431,11 +431,12 @@ tuned value, so the useful move is often *checking* rather than changing them. | `SamplingParams.return_prompt_tokens` | `False` | Opt in to echoing prompt ids over the wire | | `moe_enable_routing_replay` | `False` | Record per-token expert choices for imbalance analysis | -Three further gates exist as environment variables in PR #6064 and are **not merged**, -so do not expect them in the tree: `MCORE_NVLS_RS_BF16` (bf16 reduce-scatter buffer, -+2.5%), `MCORE_FLASHINFER_DECODE` (flashinfer `trtllm-gen` decode attention, +2.6%), and -`MCORE_FLASHINFER_PDL` (Programmatic Dependent Launch for that kernel, +0.3%). If they -have since landed as config flags, correct this section. +Three further gates have been measured as environment variables but are **not in the +tree**, so do not expect to find them: `MCORE_NVLS_RS_BF16` (bf16 reduce-scatter +buffer, +2.5%), `MCORE_FLASHINFER_DECODE` (flashinfer `trtllm-gen` decode attention, ++2.6%), and `MCORE_FLASHINFER_PDL` (Programmatic Dependent Launch for that kernel, ++0.3%). Each is a mechanism worth rebuilding, with a known target. If they have since +landed as config flags, correct this section. ### Flags that look like free wins and are not diff --git a/skills/optimize-inference-siddharth/references/updating-this-skill.md b/skills/optimize-inference-siddharth/references/updating-this-skill.md index c6c8224d04a..e6f13297752 100644 --- a/skills/optimize-inference-siddharth/references/updating-this-skill.md +++ b/skills/optimize-inference-siddharth/references/updating-this-skill.md @@ -191,4 +191,4 @@ discoverable rather than mysterious. | 2026-07 | Initial skill: the five moves, hard rules 1-9, the CUDA-graph / MoE / Mamba-Triton / host-path / measuring references, commit log, review checklist | Siddharth Singh's 2026 inference work, 29 commits | | 2026-07-28 | Added `decision-gates.md` (ceiling-before-building, per-launch fixed costs, three gates that each killed a multi-week effort) and `vllm-differential.md` (competitor-trace comparison). Hardened `measuring.md`: nsys flag combinations that deadlock finalization, the node-vs-graph trace control, `perf_counter` phase timing, union-busy idle accounting with gap-size decomposition, the same-session back-to-back A/B protocol, kernel-to-e2e conversion, ledger requirements. Added *What a wide capture costs you* to `cuda-graphs.md`. Added per-GEMM tile tuning, the measured backend comparison, the training-path-only fusion flags, and comm-vs-skew to `moe-inference.md`. Added hard rules 10 (A/B protocol) and 11 (non-bit-exact acceptance) plus the *flags that look like free wins* table | Qwen3-30B-A3B EP4 on 4×GB200 | | 2026-07-28 | Made the skill self-maintaining: this file, the editing authorization at the top of `SKILL.md`, the Step 6 feedback loop, the checklist's skill-maintenance block, and this log | — | -| 2026-08-03 | Added `scripts/` (`steady_window.py`, `union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) so the trace method is runnable rather than only described. **Corrected** the fp32 reduce-scatter buffer in `moe-inference.md`, which was documented as a reproducibility feature and costs 2.5% (the `atol=0` guarantee is about the fp32 accumulator, and `multimem.ld_reduce` accumulates in f32 regardless of operand width). Added to `measuring.md`: cross-trace window drift at differing KV length (27% bucket error, and the launch-count / untouched-bucket / seqlen-inference defenses), the sum÷union overlap ratio, and naming a kernel from its trace neighbours instead of framework instrumentation that cannot fire under graph replay. Added to `vllm-differential.md`: resolve the competitor's kernel to a *package* and check whether it is already installed (+2.6% from flashinfer `trtllm-gen`, after two sessions tuning inside flash-attn), and the packing split. Added thesis moves 5-width and 6, three routing-table rows, and the convention that a rejection records its mechanism and date so it can be re-tested | Qwen3-30B-A3B EP4 on 4×GB200, PR #6064 | +| 2026-08-03 | Added `scripts/` (`steady_window.py`, `union_window.py`, `compare_budget.py`, `kernel_neighbors.py`) so the trace method is runnable rather than only described. **Corrected** the fp32 reduce-scatter buffer in `moe-inference.md`, which was documented as a reproducibility feature and costs 2.5% (the `atol=0` guarantee is about the fp32 accumulator, and `multimem.ld_reduce` accumulates in f32 regardless of operand width). Added to `measuring.md`: cross-trace window drift at differing KV length (27% bucket error, and the launch-count / untouched-bucket / seqlen-inference defenses), the sum÷union overlap ratio, and naming a kernel from its trace neighbours instead of framework instrumentation that cannot fire under graph replay. Added to `vllm-differential.md`: resolve the competitor's kernel to a *package* and check whether it is already installed (+2.6% from flashinfer `trtllm-gen`, after wasted effort tuning inside flash-attn), and the packing split. Added thesis moves 5-width and 6, three routing-table rows, and the convention that a rejection records its mechanism and date so it can be re-tested | Qwen3-30B-A3B EP4 on 4×GB200 | diff --git a/skills/optimize-inference-siddharth/references/vllm-differential.md b/skills/optimize-inference-siddharth/references/vllm-differential.md index e9cc925e431..8e5ef729d31 100644 --- a/skills/optimize-inference-siddharth/references/vllm-differential.md +++ b/skills/optimize-inference-siddharth/references/vllm-differential.md @@ -136,9 +136,9 @@ is *already installed* and you are calling a different one. Attention was the clearest case. The differential showed the same launch count on both sides with mcore **0.409 ms/step slower** — the largest remaining work-bucket deficit, -and unambiguously an implementation gap rather than a fusion gap. Two prior sessions had -been spent flag-flipping between FA2 and FA4 *inside flash-attn*, which is the wrong -package: vLLM was calling flashinfer's `trtllm-gen` Blackwell decode kernel, and +and unambiguously an implementation gap rather than a fusion gap. Substantial effort had +already gone into flag-flipping between FA2 and FA4 *inside flash-attn*, which is the +wrong package: vLLM was calling flashinfer's `trtllm-gen` Blackwell decode kernel, and flashinfer was already in the same venv. A microbenchmark at mcore's exact shapes put it **~24% under FA2**, it accepted mcore's existing paged KV layout unchanged, and wiring it in behind `MCORE_FLASHINFER_DECODE` was worth **+2.6% end-to-end**. Enabling its @@ -146,8 +146,8 @@ Programmatic Dependent Launch added a further +0.3%. The generalizable order is: identify the package, check whether it is installed, microbenchmark it at *your* shapes under graph replay, verify the memory layout it -demands, and only then consider tuning your own. Steps 1 and 2 are minutes; the two -sessions spent skipping them were not. +demands, and only then consider tuning your own. Steps 1 and 2 take minutes; the effort +spent skipping them did not. Step 1 does not have to be guesswork. The vLLM source that produced the baseline is checked out locally, and the `vllm-codebase-expert` subagent resolves a trace diff --git a/skills/run-qwen-model/EXPERIMENTS.md b/skills/run-qwen-model/EXPERIMENTS.md index af24d04d105..2655c507074 100644 --- a/skills/run-qwen-model/EXPERIMENTS.md +++ b/skills/run-qwen-model/EXPERIMENTS.md @@ -7,31 +7,64 @@ This ledger starts empty. Append every experiment, including failures and regressions. Never edit an earlier result after it is recorded — supersede it with a new row. -## Accepted changes - -**The at-a-glance answer to "what actually landed."** Every accepted experiment -gets a row here *and* a full row in the experiment index below. This table is a -duplicate on purpose: the index interleaves accepts with rejections and gated-out -levers, and detailed records bury PR links hundreds of lines down, so neither is -readable as a summary. One row per accepted mechanism, newest last. - -| ID | Date | Mechanism | Gain | Kill switch | PR | PR state | -|---|---|---|---:|---|---|---| -| — | — | _nothing accepted yet_ | — | — | — | — | - -Rules for this table: - -- A row appears here **only** after the change passed its A/B on distribution - separation and its correctness gate. Nothing provisional. -- `Gain` is the percentage against the same-session OFF arm, not against a - historical baseline. -- `PR` is a markdown link to the draft MR. If a PR is superseded, replace the - link and note the superseded number in the detailed record, not here. -- `PR state` is draft, open, merged, or closed. Refresh it when you touch the - ledger; a stale "draft" against a merged PR is the failure mode this column - exists to catch. -- Reverted-after-the-fact changes are struck through here with a pointer to the - regression's entry. Do not delete the row. +## The merge stack + +**Accepted changes form a single linear stack, not a set of independent PRs.** +`main` is used exactly once, to cut the branch for stack position 1. Every +accepted change after that is branched off the previous accepted change, whether +or not the two are logically related. + +``` +main ──● (baselines measured here) + └──● 01 ──● 02 ──● 03 ──● … +``` + +The reviewer then merges strictly in ascending `#`. Once `01` is merged, `02` is +rebased onto the new `main` and its PR narrows to only its own diff, so each +review is one mechanism at a time even though the work was done cumulatively. + +| # | ID | Mechanism | Branch | Cut from | PR | PR state | Marginal gain | Cumulative tok/s | Kill switch | +|---|---|---|---|---|---|---|---:|---:|---| +| — | — | _nothing accepted yet; position 1 is cut from `main`_ | — | — | — | — | — | — | — | + +**Merge sequence is the `#` column, ascending. No other order is valid.** + +### Rules for this table + +- A row appears **only** after the change passed its A/B on distribution + separation and its correctness gate. Nothing provisional, nothing planned. +- `#` is assigned when the row is created and never reused or reordered. If a + change is later reverted, strike its row through and leave the number retired — + renumbering silently invalidates every `Cut from` below it. +- `Cut from` is `main` for `#1` and the branch of `#N-1` for every other row. This + column is what makes the stack reconstructable after the fact; fill it even + when it looks redundant. +- `Marginal gain` is the percentage measured against the same-session OFF arm + **on top of everything below it in the stack** — not against `MCORE-BASELINE`. +- `Cumulative tok/s` is the absolute throughput with all rows up to and including + this one enabled. Marginal gains do not sum to the cumulative gain, and the + gap between them is real information: two changes attacking the same kernel + overlap, so the second one's marginal gain is smaller than it would have been + alone. Record both and do not reconcile them by arithmetic. +- `PR state` is draft, open, merged, or closed. Refresh it whenever you touch the + ledger; a stale "draft" against a merged PR is what this column exists to catch. + +### Rebase discipline as the stack merges + +After `#N` merges, every row above it must be replayed onto the new `main` with +its parent's commits dropped. A plain `git rebase main` will try to reapply the +already-merged commits and conflict, particularly when the parent was +squash-merged upstream: + +```bash +git fetch origin main +git rebase --onto origin/main +git push --force-with-lease +``` + +Force-push only to your own fork's stack branches. If a row is rejected upstream +rather than merged, the rows above it are rebased with `--onto` past it in exactly +the same way, and its row here is struck through with the reason. ## Fixed protocol @@ -57,7 +90,7 @@ parallelism, or warmup/timed counts differ. Nothing recorded yet. Baseline order is mandatory: 1. Record `VLLM-BASELINE` with Nsight Systems. -2. Record `MCORE-BASELINE` with Nsight Systems. +2. Record `MCORE-BASELINE` with Nsight Systems, at `main`. 3. Compute the absolute and percentage gap. 4. Only then modify Megatron-Core. @@ -67,13 +100,17 @@ Nothing recorded yet. Baseline order is mandatory: | MCORE-BASELINE | mcore EP4/TP1 | — | — | — | — | — | not run | Both baselines must be run with no optimization flags enabled, and the gap -computed, before any code change. +computed, before any code change. These are the only measurements taken at +`main`; every later number is measured on a stack branch. ## Experiment index -| ID | Date | Hypothesis | Changed files / flags | Throughput | Delta vs baseline | Correctness | Job / run | MR | Conclusion | -|---|---|---|---|---:|---:|---|---|---|---| -| — | — | _no experiments recorded yet_ | — | — | — | — | — | — | — | +Every experiment, accepted or not. `Measured on` records the stack position the +experiment was built on, without which its delta cannot be interpreted. + +| ID | Date | Hypothesis | Changed files / flags | Measured on | Throughput | Marginal delta | Correctness | Job / run | PR | Conclusion | +|---|---|---|---|---|---:|---:|---|---|---|---| +| — | — | _no experiments recorded yet_ | — | — | — | — | — | — | — | — | ## Optimization rules @@ -82,15 +119,18 @@ computed, before any code change. negative. 3. Change one performance variable at a time. 4. Preserve the fixed protocol. -5. A/B in one allocation, arms back to back and alternating; accept on - distribution separation, not mean delta. -6. Validate correctness before accepting throughput. -7. Revert regressions or correctness failures. -8. Record the result — including rejections, with their root cause — before +5. Build every new experiment on the **stack tip** — all accepted changes + applied — never on `main`. `main` is only for the two baselines. +6. A/B in one allocation, arms back to back and alternating, with the stack below + the change held on in both arms; accept on distribution separation, not mean + delta. +7. Validate correctness before accepting throughput. +8. Revert regressions or correctness failures. +9. Record the result — including rejections, with their root cause — before beginning another experiment. -9. Open a separate draft MR for each accepted change, and link it from **both** - the *Accepted changes* table and its experiment-index row. An accepted change - that is missing from the *Accepted changes* table is not recorded, however - thorough its detailed write-up is. -10. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines +10. Each accepted change gets its own draft PR, branched off its predecessor and + recorded as a new row in *The merge stack* with its `#`, branch, `Cut from`, + and PR link. An accepted change missing from that table is not recorded, + however thorough its detailed write-up is. +11. Stop when mcore meets or exceeds `VLLM-BASELINE`, then rerun both baselines once to confirm parity under identical conditions.