Skip to content

feat(k3): kernel mining + capsule substrate — vendored vLLM KDA/top-k cubins (+7%/+20% decode) - #982

Open
xiaguan wants to merge 10 commits into
mainfrom
feat/graph-dump-json
Open

feat(k3): kernel mining + capsule substrate — vendored vLLM KDA/top-k cubins (+7%/+20% decode)#982
xiaguan wants to merge 10 commits into
mainfrom
feat/graph-dump-json

Conversation

@xiaguan

@xiaguan xiaguan commented Aug 28, 2026

Copy link
Copy Markdown
Collaborator

What

Two layers of work on one branch, tooling first, then the first production use of it:

Kernel mining tooling

  • --dump-graph-png PATH now also writes a machine-readable PATH.json sibling (per-kernel symbol, grid/block, function attrs, every staged param classified against the allocation map).
  • tools/kernel-capture/: provider-agnostic CUPTI injection lib (CUDA_INJECTION64_PATH) that lifts every launched kernel's cubin + full call ABI out of any CUDA process (vLLM/sglang/ourselves). Multi-rank-hardened: per-pid output subdirs, and the module/ABI tables are mutex-locked — the unlocked g_abi realloc segfaulted our 4-ranks-in-one-process K3 engine on concurrent lazy module load.

Capsule substrate, live for K3 (PEGAINFER_K3_CAPSULE)

  • The serving kernel becomes an external cubin artifact; the framework keeps the op contract and the native reference twin. Two vendored vLLM v0.28.0 kernels (pegainfer-kernels/cubin/k3/, Apache-2.0, image digest + sha256 + offline-build recipe in its README), embedded at build time, bound in csrc/k3/k3_capsule.cu via cuModuleLoadData + a fail-closed cuFuncGetParamInfo ABI walk. No vLLM source enters the build; build.rs compiles nothing new for them.
  • topk: single_group_topk warp kernel replaces the native router top-k at the step call site.
  • kda: the fused KDA decode kernel replaces conv_silu×3 + kda_core, collapsing the four wbig band GEMMs into one full GEMM + two landings. Same formulas/constants as our TileLang generator (lower bound −5, scale 128^-0.5, eps 1e-5, tap order, [head, v, k] state) — zero new weights. It updates conv + recurrent state in place: capsule mode adds a packed conv slab filled by adopt_row at the prefill→decode handover, pins recurrent parity slab 0, and refuses to combine with dspark / CP prefill.
  • Flag unset ⇒ byte-identical native serving (default).

Production invariant

K3 EP4 decode step composition: the KDA chain and router top-k are served by vendored external cubins when PEGAINFER_K3_CAPSULE selects them; everything else (MegaMoE, MLA, norms, GEMMs) unchanged. Not duplicated work: no existing capsule/loader path served K3, and the native kernels remain the reference twin.

Evidence

  • Per-kernel numeric gates (GPU, run on GB300 sm_103): k3_capsule_topk_gate — expert sets equal, weights ≤1e-5, b ∈ {1,3,8,32}; k3_capsule_kda_gate — conv windows bitwise-equal, state ≤1.6% / out ≤6.3% rel err (bf16-rounding-chain difference: native lands intermediates in bf16 where vLLM stays f32).
  • Production E2E A/B (pruned-224 checkpoint, EP4, tray03, 128-token completions, same serve config both arms):
    • 4-concurrent: 96.3 → 103.0 tok/s (+7%)
    • 32-concurrent: 495 → 594 tok/s (+20%)
    • Greedy 100-token texts: 4/8 byte-identical, rest diverge at a near-tie token with comparable quality.
  • Clean vLLM v0.28.0 baseline on the same phases (DP4×EP4, cudagraphs on): 139 / 724–788 tok/s — still ~25% ahead; the remaining gap is dominated by graphs-vs-forced-eager (~4.3k launches/rank-step) plus the unported skinny GEMM, both recorded as next levers in docs/models/k3/vllm-kernel-ab.md.
  • Workspace lib tests 22/22 green (includes fixing a stale chunk-bucket assertion that perf(k3): raise the MegaMoE protocol max to 16896 tokens per rank #962 left behind); cargo fmt --check clean.

Not ported (documented in the A/B doc)

CuTe skinny GEMM/dotprod (captured instantiations only cover M ≤ 2; highest-complexity ABI), attn-res (NB=3 online kernel vs our 8-block walk — structural mismatch for ≤2% of step), MoE/MLA (ours wins / negligible). Mined cubins + ABI manifest stay in the tray-local capture archive.

🤖 Generated with Claude Code

https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 39abd17f74

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

Comment on lines +724 to +725
if capsule_flags().topk {
k3_capsule_router_topk_launch(

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P1 Badge Validate the capsule router against the expert count

For the supported 896-expert K3 checkpoint, selecting PEGAINFER_K3_CAPSULE=topk or the advertised all path reaches this call with experts == 896, but k3_capsule_router_topk_launch rejects any count above 512. Under EP this forward error is passed to ep_fatal, so the fully loaded service terminates on its first request. Validate the flag against the checkpoint during startup or retain the native router for the 896-expert model.

Useful? React with 👍 / 👎.

Comment on lines +24 to +27
flags.kda = true;
}
"topk" => flags.topk = true,
"kda" => flags.kda = true,

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Badge Require a model eval before enabling capsule KDA

When PEGAINFER_K3_CAPSULE=kda or all is set, this exposes a serving path whose retained evidence reports capsule output error up to 6.3% and divergent greedy text in 4/8 prompts, but the change includes no model-level accuracy evaluation. A synthetic operator tolerance and a handful of qualitative generations cannot establish that the changed recurrent-state rounding preserves model quality, so keep this path experimental/non-serving or add a representative model eval gate before exposing it.

AGENTS.md reference: AGENTS.md:L110-L112

Useful? React with 👍 / 👎.

xiaguan and others added 10 commits August 28, 2026 12:04
--dump-graph-png now also writes a machine-readable .json next to the PNG/DOT:
per-kernel function attributes (regs, static/const/local sizes, ptx/binary
version) and every staged parameter's bytes via cuFuncGetParamInfo, with
8-byte values resolved against the CUDA allocation map (cuPointerGetAttribute)
into device/host pointer ranges. Driver floor moves 12.3 -> 12.4.

Verified on GB300 (sm_103) Qwen3-4B bs=1 decode: 543 kernel nodes, 2785
params extracted, 1086 device pointers classified; only cuBLASLt nvjet
kernels launch with packed extra buffers and report params: null.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_017gqqDEArVgaujRMYffr2Jt
Signed-off-by: JinYan Su <751080330@qq.com>
…ngines

Provider-agnostic capture of every CUDA module and kernel launch from any
host process (vLLM/sglang/PegaInfer) via CUDA_INJECTION64_PATH. Module-loaded
resource callback dumps each cubin image straight from memory (Triton,
cuBLAS/cuBLASLt, CUTLASS, FlashInfer, hand-written CUDA alike -- no on-disk
Triton-cache archaeology); cuLaunchKernel(Ex) driver callback records symbol,
launch geometry, function attributes, and staged parameter bytes via
cuFuncGetParamInfo, mirroring the in-tree graph JSON dumper.

Smoke-tested on GB300 (sm_103) against PegaInfer Qwen3-4B serving one decode:
100 cubins dumped (all ELF, 3.4KB-8.3MB), 14758 launches (5742 with staged
params, 9016 cuBLASLt nvjet extra-buffer launches correctly reported null).

Signed-off-by: JinYan Su <751080330@qq.com>
…pointers

Kernels launched through the PyTorch runtime <<<>>> path (vllm::rms_norm,
rotary_embedding, act_and_mul, reshape_and_cache, all at::native ops) hand
cuLaunchKernel a CUfunction that cuFuncGetParamInfo/cuFuncGetAttribute reject,
so the first cut captured launch config but empty params + zero attributes for
~66% of launches (only driver-API kernels like the FMHA attention answered).

Fix: at module-load, self-load a private copy of the cubin and enumerate its
functions (cuModuleEnumerateFunctions), caching each kernel's parameter layout
and attributes from the driver's own parse; a reentrancy guard skips the
recursive MODULE_LOADED. record_launch resolves layout from the live handle,
else the cache. Also classify each 8-byte value against the allocation map
(cuPointerGetAttribute), mirroring the in-tree graph dumper.

Re-verified on GB300 vLLM 0.26.0 Qwen3-4B bs=1 decode: 11418/11418 launches
now carry attributes; 8374 full param records (3044 cuBLASLt nvjet extra-buffer
correctly null); 7467 pointer args classified with owning ranges. rms_norm bf16
decodes cleanly -- out/input/weight pointers tagged device, eps=1e-6 scalar.

Signed-off-by: JinYan Su <751080330@qq.com>
Capture layer under the roadmap's kernel ledger: the graph JSON dumper and the
CUPTI injection lib, the vLLM-vs-PegaInfer bs=1 decode JSON A/B (attention +
GEMM-dispatch levers, norm/rope/silu/kv-append cleared), the settled capsule/
manifest design, storage+license discipline per the cubin/glm52 precedent, and
the next action (nsys attention span before porting fmhaSm100fKernel).

Signed-off-by: JinYan Su <751080330@qq.com>
Per-process pid<N>/ output subdirs so TP/EP workers (vLLM one-process-per-
rank) and single-process multi-rank engines don't clobber each other's
launches.jsonl and cubin dumps.

Lock the module-load path: cache_module_abi grew the shared g_abi table
(realloc + count++) with no lock, and record_launch's fallback abi_lookup
read it unlocked. Fine under vLLM's process-per-rank shape; PegaInfer's K3
engine hosts 4 ranks in one process, and the first inference step's
concurrent lazy module loads corrupted the heap and killed the server with
no log. The recursive MODULE_LOADED from our own cuModuleLoadData is
filtered by the thread-local guard before it can reach the lock.

Verified on GB300: vLLM kimi-k3 EP4 (4 processes) and PegaInfer K3
--k3-ep-size 4 (1 process, 4 ranks) both complete a bs=1/8/32 ladder.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
CUPTI-mined bs=1/8/32 ladder on the pruned-224 checkpoint, PegaInfer EP4 vs
vLLM TP4+EP, same GB300 tray. Cross-referenced with the existing EP4 decode
step profile to rank ports: vLLM's CuTe-DSL skinny GEMM (their answer to the
52% B=1 cuBLASLt splitK cliff), the E=224 routing kernel, then KDA/attn-res
fusion. Also records where our structure already wins (MegaMoE launch count,
zero collectives) and the per-bs dispatch evidence behind the (op, bucket)
manifest axis.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
#962 added the 4224/8448 prefill buckets but left the test asserting
chunk_bucket(4096) rounds to the ladder max — true only before 4224
existed. Assert the actual neighbors instead.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
First live use of the capsule design: the serving kernel is an external
cubin artifact, the framework keeps the op contract and the native
reference twin. Two vendored vLLM v0.28.0 kernels (Apache-2.0,
provenance + sha256 in cubin/k3/README.md), embedded at build time and
bound in csrc/k3/k3_capsule.cu via cuModuleLoadData + a fail-closed
cuFuncGetParamInfo walk against the capture-recorded ABI. No vLLM source
enters the build; nvcc compiles nothing new for them.

PEGAINFER_K3_CAPSULE (unset = byte-identical native serving; `all` or
csv of topk,kda) flips the ops per launch site:

- topk: single_group_topk warp kernel replaces k3_router_topk at the
  step.rs router site (needs the routed scale as a host scalar; pairs
  come back sort-ordered, consumers treat them as unordered).
- kda: the fused decode kernel replaces conv_silu x3 + kda_core, and
  the four wbig band GEMMs collapse into one full GEMM + two landings
  (out_gate + a new packed q|k|v land config). Same formulas and
  constants as our generator (lower bound -5, scale 128^-0.5, eps 1e-5,
  tap order, [head, v, k] state); no new weights. The kernel updates
  conv + recurrent state in place: capsule mode allocates a packed
  [rows, 3 taps, q|k|v, 12288] conv slab filled by adopt_row at the
  prefill->decode handover, pins recurrent parity slab 0, and refuses
  to combine with dspark or CP prefill.

Numeric gates (GPU, --ignored-style manual): k3_capsule_topk_gate
(expert sets equal, weights <=1e-5, b in {1,3,8,32}) and
k3_capsule_kda_gate (conv windows bitwise-equal; state <=1.6% / out
<=6.3% rel err — bf16-rounding-chain difference, native lands
intermediates in bf16 where vLLM stays f32).

E2E A/B on the pruned-224 EP4 serve (same tray, 128-token completions):
4-concurrent 96.3 -> 103.0 tok/s (+7%), 32-concurrent 495 -> 594 tok/s
(+20%); greedy 100-token texts 4/8 byte-identical, rest diverge at a
near-tie token with comparable quality.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
vllm-kernel-ab.md gains the port log (what shipped, layout deltas,
offline-build recipe pointer, why skinny GEMM and attn-res stay
mined-but-unported); TL;DRs in it, kernel-mining.md and the index row
now lead with the measured +7%/+20% decode throughput.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01PvpKoX2Wh868p9UqpqNeFP
Signed-off-by: JinYan Su <751080330@qq.com>
@xiaguan
xiaguan force-pushed the feat/graph-dump-json branch from 39abd17 to 0581707 Compare August 28, 2026 12:05
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant