perf(megakernel): MK-GEMM v2 -- 비상주 GEMM 레인 (VLLM_GLM53_MK_GEMM2, 기본 off); 원장 30차: mk_gemm 이 deep_gemm 과 같은 14 ms 인 이유 - #305
Merged
Conversation
…nsion, two blocks per SM (VLLM_GLM53_MK_GEMM2, default off) The 09-04 18:42 armed trace read against the offline bench: the shared expert's down [4096 x 512] is 18 us alone and 135 us in the step (42 launches, 5.7 of the lane's 14.1 ms) because the persistent 48 x 69.6 KB grid cannot share an SM with the routed MoE kernel (90 KB) and the publish barrier holds every landed block for the last one; deep_gemm's independent blocks ran the same GEMM inside the MoE tail (09-01 stock trace, 36 us). On the main stream the prologue and the static first-unit assignment cost 15-35 us a launch (dense gate_up 119 us for 62 us of bytes). mk_gemm2_kernel: grid = (n/128) x ksr blocks, no barrier, no shared A tiles (each block quantizes its slice from x, same arithmetic), the e2m1 expansion in registers into the mma fragments (same LUT), 37 KB smem so two blocks sit on an SM, deterministic last-arriver fold for the slices. Bench: --gemm2 both (side by side, outputs diffed) and --ksr2-sweep; the exact gate runs on both lanes; a MoE-overlap probe for the shared-expert pair. Kill switch VLLM_GLM53_MK_GEMM2 (default off) ahead of the persistent launch. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…-deep-gemm-perf-667f67
…이유(상주 커널은 MoE 옆에서 진다), v2 레인 EXP-21 Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…h --stamps2 for the v2 unit timeline Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Bugbot couldn't run - usage limit reachedBugbot is counted against Cursor usage for this user or team, and this run hit a usage or spend limit. A user or team admin can review and increase usage limits in the Cursor dashboard. (requestId: serverGenReqId_6f73e42d-cc3d-4f47-9cbe-9bcd5aea83a4) |
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
There was a problem hiding this comment.
🟡 Changes recommended
It introduces a small but concrete CLI crash bug in the new trace tool and includes several future-dated doc headers that should be corrected for timeline clarity.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
Adds an opt-in MK-GEMM v2 (“non-persistent” grid) execution lane for GLM-5.3 megakernel GEMMs, plus supporting probes/bench tooling and documentation to explain/validate the trace-driven performance diagnosis. Default serving behavior remains unchanged unless VLLM_GLM53_MK_GEMM2=1 is enabled.
Changes:
- Introduces
mk_gemm2_kerneland host-side dispatch/plan helpers gated byVLLM_GLM53_MK_GEMM2/VLLM_GLM53_MK_NBUF2. - Extends probes/bench and adds trace tooling to measure per-launch position timings and MoE-overlap exposure for v1 vs v2.
- Updates runbook/measurements/docs and
tests/test_logic.pyto reflect v2 contracts and new knobs.
File summaries
| File | Description |
|---|---|
| tools/trace_launch_positions.py | New trace analysis tool to aggregate per-step/per-position kernel timings and overlaps. |
| tests/test_logic.py | Updates contract assertions to include GEMM v2 lane and related build/bench hooks. |
| RUNBOOK_KERNEL_CAMPAIGN2.md | Adds EXP-21 runbook entry describing MK-GEMM v2 rollout gates and operational cautions. |
| probes/mk_gemm_moe_overlap_probe.py | New probe to measure shared-expert GEMM exposure under routed MoE for v1 vs v2. |
| probes/megakernel_glm53_bench.py | Extends benchmark to time v2 lane, diff outputs, sweep ksr, and read v2 stamps. |
| overlay/modules/glm53_megakernel/README.md | Documents v2 lane rationale, mechanics, and validation workflow. |
| overlay/modules/glm53_megakernel/glm53_megakernel.py | Adds MK_NBUF2_DEF build flag and logs served GEMM lane plan in selftest fingerprint. |
| overlay/modules/glm53_megakernel/glm53_megakernel.cu | Implements mk_gemm2_kernel, adds v2 dispatch/plan APIs, and occupancy-based sizing. |
| MEASUREMENTS.md | Adds “30차” diagnosis write-up and expected impact section for MK-GEMM v2. |
| build/glm53/glm53_megakernel.py | Mirrors overlay build changes for the glm53 build variant. |
| build/glm53/glm53_megakernel.cu | Mirrors overlay CUDA changes for the glm53 build variant. |
| build/dsv4/glm53_megakernel.py | Mirrors overlay build changes for the dsv4 build variant. |
| build/dsv4/glm53_megakernel.cu | Mirrors overlay CUDA changes for the dsv4 build variant. |
Review details
- Files reviewed: 13/13 changed files
- Comments generated: 1
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Comment on lines
+63
to
+65
| path, pat = sys.argv[1], sys.argv[2] | ||
| out_csv = sys.argv[sys.argv.index("--csv") + 1] if "--csv" in sys.argv else None | ||
| ev = load_kernel_events(path) |
choiceoh
pushed a commit
that referenced
this pull request
Sep 4, 2026
…mk-kernel-improvements-a8b6ea -- both lanes kept, my EXP-17 renumbered EXP-22 The other session's non-persistent v2 lane (mk_gemm2_kernel, VLLM_GLM53_MK_GEMM2, 30차) and this branch's local-quant kernel (mk_gemm_lq_kernel, VLLM_GLM53_MK_LOCALQ, 29차) share the diagnosis -- a resident barrier kernel serialises with the MoE kernel on the aux stream -- and differ in prescription. Both stay, both default off: mk_run_gemm dispatches the v2 lane first, then the plan; the bench keeps the --gemm2 columns beside the local column and --gemm-sweep; probe_exact runs the v1 pair of kernels on the shared fixture and the v2 lanes after it; the ledger orders 30차 above 29차; the v2 kernel's scalar e4m3 conversion is back for its own pack. test_logic 44,646 OK; nvcc gemm 80 / lq 122 / gemm2 124 registers. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
This was referenced Sep 4, 2026
choiceoh
pushed a commit
that referenced
this pull request
Sep 4, 2026
…31차로 PR #304 는 #301(672430e)에서 갈라져 나왔고, 그 뒤 main 이 #300·#302·#303·#305· #306·#307·#308(#309)로 앞서 나가면서 충돌 상태(dirty)가 됐다. 코드 파일은 겹치지 않았고(KDA 원패스·kpool 은 #304 단독), 겹친 것은 문서·프로필·테스트 넷이다. **RUNBOOK 충돌 2건** (1) 새 EXP 절의 삽입 위치: main 이 EXP-16 뒤에 EXP-22(로컬 양자화)·EXP-21(v2 레인)을 넣었고 #304 는 같은 자리에 EXP-20(미세 융합 묶음 2)을 넣었다. 경쟁이 아니라 가산이므로 셋 다 두되 기존 배열(최신 먼저)을 따라 22 → 21 → 20 순으로 둔다. (2) "순서와 근거" 절: main 은 평평한 번호 목록을 지우고 자기가 새로 만든 상태 표로 보내는 안내문으로 바꿨고(`채택·기각·닫힘은 위 상태 표에 있다`), #304 는 그 목록에 18번(EXP-20)을 덧붙였다. **main 의 재구성을 취하되**, 사라질 뻔한 #304 의 항목을 main 의 "다음 부팅 창에서" 목록에 8번으로 옮겨 실험이 목록과 함께 증발하지 않게 했다. **원장 번호 충돌 (자동 병합이 못 보는 것)** #304 는 자기 측정을 `★29차` 로 적었는데, 그 사이 main 에 다른 29차(메가커널 로컬 양자화, #307)와 30차(비상주 v2 레인, #305)가 먼저 들어왔다. 텍스트로는 충돌하지 않아 병합 결과에 **서로 다른 29차 둘**이 남았다. 28차 항목이 세운 선례("원장 번호 27 은 PR #290 이 쓰고 있어 28 로 적는다")대로 다음 빈 번호인 **31차** 로 옮겼고, 왜 옮겼는지 절 머리에 한 줄 남겼다. 측정값 자체는 손대지 않았다. 같은 번호를 가리키던 참조 6곳을 함께 고쳤다: MEASUREMENTS 절 제목, RUNBOOK EXP-20 표의 기각 축, `glm53_kda_onepass/README.md`, `moe_gate_sm121/README.md`, `profiles/glm53.env`, `tests/test_logic.py`. main 이 자기 29차를 가리키는 참조 (`VLLM_GLM53_MK_LOCALQ` 주석)는 그대로 두었다. **자동 병합분 검증** - `profiles/glm53.env`: main 의 `VLLM_GLM53_MK_LOCALQ` 와 #304 의 세 노브 (`KDA_DUAL_GEMM`·`KDA_ONEPASS`·`KPOOL_UPDATE_DIRECT_POS`) 전부 기본 0 으로 잔존. `MODULES=` 목록은 main 의 항목 하나도 잃지 않고 `glm53_kda_onepass` 만 늘었다. - overlay ↔ build 사본 3쌍(`glm53_kda_onepass.py`, `glm5next_kda.py`, `sparse_attn_indexer_kpool.py`) 모두 동일. - `tests/test_logic.py`: 양쪽 테스트 공존, `micro-fusion bundle 2 contracts` 통과. 검증: `tests/test_logic.py` all OK (6106 checks; 이 호스트엔 torch 부재로 일부 SKIP). 충돌 마커 0, 변경된 .py 전부 파싱 OK, `run_micro_fusion_check.sh` bash -n OK. GPU 검증(프로브 VERDICT·트레이스 물리확인)은 하지 않았다 — 이 환경에 GPU 가 없다. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01JTghtTjyfT23vuhZwHzBSM
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
왜
운영자 지시: deep_gemm 197발 14.06 ms(09-01) / mk_gemm 185발 14.13 ms — mk gemm 이 deep gemm 보다 훨씬 빠르게.
09-04 18:42 무장 트레이스를 발사 위치별로 읽으니(
tools/trace_launch_positions.py, 22 스텝 중앙값) 185발 중 공유 전문가 down [4096×512] 42발이 오프라인 18 µs → 서빙 135 µs(5.7 ms/스텝, 41%)였다. MoE b12x 커널(smem 90 KB, 48블록)·mk_gemm(69.6 KB)·deep_gemm(92.7 KB) 셋 중 어느 둘도 한 SM(102.4 KB)에 같이 못 앉는다. deep_gemm 은 독립 블록이라 같은 GEMM 을 MoE 꼬리 안에서 끝냈고(09-01 스톡 트레이스 36 µs, 노출 0), 상주 mk_gemm 은 먼저 자리를 얻은 블록이 publish 배리어에서 마지막 블록(= MoE 종료 후)까지 기다린다. aux 쌍의 임계경로 노출은 층당28 µs(1.18 ms/스텝); 메인 스트림 99발은 발사당 고정비 1520 µs + 정적 타일 배분 꼬리(dense gate_up 48타일=grid: 119 µs / 바닥 62). 원장 30차 §1~2.무엇
mk_gemm2_kernel— 같은 W4A8 GEMM 을 비상주 grid(n/128 × ksr 독립 블록)로:mk_w4x4, 같은 두 LUT 즉치값·sign prmt) → smem 69.6 → 37 KB, SM 당 2블록(__launch_bounds__(256, 2); ptxas 124 reg, 스필 0).__syncthreads하나(A 더블버퍼), ksr > 1 슬라이스는 fp32 partial 배정 + 마지막 도착 폴드(결정적, zero pass 없음). ksr 규칙: 유닛 ≥ 2×상주 슬롯, 슬라이스 ≥ 4 k-블록(VLLM_GLM53_MK_KSR2로 강제).VLLM_GLM53_MK_GEMM2(기본 0 — 이 PR 은 서빙 동작을 바꾸지 않는다). 부팅 핑거프린트에 서빙 레인·in_proj 플랜 줄.--gemm2 both(v1/v2 나란히, 출력 diff — 비분할 형상은 비트 동일),--ksr2-sweep,--stamps2; exact 게이트 두 레인 × ksr 5종;probes/mk_gemm_moe_overlap_probe.py(공유 전문가 쌍의 MoE 아래 노출, 두 레인); 프로덕션 형상 4종 추가.검증
-Xptxas -v): 5 커널 모두 스필 0, mk_gemm2 124 레지스터.tests/test_logic.pyall OK (44,636 checks).~/mkab/gap_bench.sh가 자동 실행(g2.*.out) → 원장 30차 §4 에 후속 PR 로 기록. 기대: 메인 99발 6.4 → ~4.9 ms, aux 노출 28 → ~5 µs/층, 트레이스 mk_gemm 합 14.1 → ~7 ms.주의
SMLP(공유 전문가 융합, 상주+배리어)를 aux 스트림에 켜면 이 진단의 직렬화가 되살아난다 — 융합은 v2 구조 위에서.
🤖 Generated with Claude Code
Note
Medium Risk
Opt-in CUDA inference path with numerical gates (bit-identical on unsplit shapes); default remains v1, but enabling v2 changes decode GEMM scheduling and must pass bracket gates before promotion.
Overview
Adds an optional MK-GEMM v2 path (
mk_gemm2_kernel,VLLM_GLM53_MK_GEMM2=1, default off) so standalone W4 GEMMs can share SMs with routed MoE instead of the persistentmk_gemmgrid that serializes on the aux stream.Kernel/host: Non-persistent grid
(n/128)×ksr, no publish barrier or shared A staging, per-block activation quant, registermk_w4x4expansion (~37 KB smem, two blocks/SM), deterministic k-slice fold.mk_run_gemmbranches to v2 when enabled; bench hooksgemm2_plan,set_gemm2,read_ts2; build flagVLLM_GLM53_MK_NBUF2. Boot self-test logs active lane and in_proj plan.Validation & ops: Bench gains
--gemm2 both,--ksr2-sweep,--stamps2and more decode shapes; newmk_gemm_moe_overlap_probe.pyandtrace_launch_positions.py. Docs (MEASUREMENTS 30차, RUNBOOK EXP-21, README) explain trace-driven diagnosis;test_logic.pyasserts v2 contracts.Reviewed by Cursor Bugbot for commit bee951b. Bugbot is set up for automated code reviews on this repo. Configure here.