docs(runbook): EXP-17/18/19 전제를 무장(메가커널) 시대 수치로 정정

[choiceoh]

Summary

#301 머지 후 "그 사이에 이미 개선된 부분은 없었나" 대조에서 확인한 전제 변화 3건을 런북에 반영했다. 코드 충돌은 없었지만(이식 때 벡터화·노브·--after 전부 main 부재 확인), 병렬 메가커널 캠페인이 세 축의 가치 전제를 갈랐다:

• EXP-18(osar): perf(osar): 콜렉티브 위상별 사이클 계측 — 32ms 를 잔차가 아니라 실측으로 가른다 #99 시대 copy+reduce 11.8µs 전제 낡음(현 45.5µs/콜·고정비 ~22.7µs 지배, 장부 19차·two-shot 기각) → R1 천장 ~0.2-0.3ms 로 하향. R2(L2 위생)의 판정 근거는 EXP-13 프리페치와의 시너지로 승격 — osar churn ~26MB/스텝이 프리페치 예산 12MB(L2 24MB)와 직접 경쟁하는 것을 __stwt/__ldcs 가 제거.
• EXP-17(MHC 패스설정): 무장 디코드가 mk_mhc 로 TileLang mhc 179발 대체(185→14, 지도 ④ 무장 09-04) → 주 전장은 프리필 big_fuse + 잔여 14발로 축소.
• EXP-19(hc bf16): ONEPASS 미채택 + mk_mhc 초월(ONEPASS 가 노리던 융합을 메가커널이 더 깊게 수행) → 폴백 참조 측정으로 강등, 프로덕션 경로의 같은 질문은 mk_mhc 쪽 별도 축임을 명시.

Testing

• 문서 전용 변경 — tests/test_logic.py 6,043 checks all OK (로컬 잔해 디렉터리 제거 후, perf(osar,mhc,tools): copy/reduce 16B 벡터화+캐시정책, MHC 패스설정 노브, census 인접성 모드 (EXP-17/18/19) #301 정정 코멘트 참조)

---

Note

Low Risk
Documentation-only runbook edits; no runtime, auth, or deployment logic changes.

Overview
Updates Kernel Campaign II runbook sections EXP-17, EXP-18, and EXP-19 so experiment scope and expected gains match post–mk_mhc serving, not pre-megakernel launch counts.

EXP-17 (MHC TileLang pass A/B) now states that armed decode mostly replaced 179 TileLang MHC launches with mk_mhc (185→14); the knob mainly matters for the remaining decode launches plus prefill big_fuse, while kernel-level --passes probes stay valid but step-level ceilings must not use pre-mk math.

EXP-18 (osar vectorization + cache policy) replaces stale #99 copy+reduce (~11.8µs/call) assumptions with current osar cost (45.5µs/call, ~22.7µs fixed overhead), lowering R1 upside to ~0.2–0.3 ms and framing R2 judgment around EXP-13 AR prefetch synergy (L2 churn vs 12 MB prefetch budget), with bracketing on the same boot as EXP-13.

EXP-19 (--hcweight) is reframed as a fallback-path reference probe after ONEPASS was never adopted and mk_mhc subsumed the fusion goal; production fp32 hc weight questions belong on a separate mk_mhc axis, with legacy numbers kept as historical stock-chain notes.

^{Reviewed by Cursor Bugbot for commit 4cf3138. Bugbot is set up for automated code reviews on this repo. Configure here.}

[cursor]

Bugbot couldn't run - usage limit reached

Bugbot 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_dd59c46d-bb45-4a91-8ae8-62928d8f3e58)

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-04T15:37:50.429663Z	5c23868	PR opened

ℹ️ 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" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

[cursor]

Bugbot couldn't run - usage limit reached

Bugbot 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_07ab5813-6922-4443-87b7-0ece3e0b4b03)

[Copilot]

🟡 Changes recommended

The runbook additions contain future-dated timestamps (2026-09-05) relative to the PR metadata date (2026-09-04), which can mislead readers about when the stated premise changes took effect.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

This PR updates the Kernel Campaign II runbook to correct the stated assumptions and expected ceilings for EXP-17/18/19 to reflect the post–mk_mhc (armed decode / megakernel) reality rather than older pre-megakernel measurements.

Changes:

• EXP-17: Notes that armed decode replaced most TileLang MHC launches (185→14), narrowing where the pass-setting knob matters.
• EXP-18: Replaces stale per-collective copy+reduce assumptions with current osar costs and reframes R2 evidence around EXP-13 prefetch synergy.
• EXP-19: Reframes --hcweight as a fallback-path reference probe after ONEPASS non-adoption and mk_mhc superseding the fusion target.

File summaries

File	Description
RUNBOOK_KERNEL_CAMPAIGN2.md	Updates EXP-17/18/19 narrative/ceilings to align with mk_mhc-era measurements and scope.

Review details

• Files reviewed: 1/1 changed files
• Comments generated: 0
• Review effort level: Lite

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[chatgpt-codex-connector]

💡 Codex Review

stkernel/overlay/modules/tp_oneshot_ar/dsv4_oneshot_shim.py

Line 380 in 5c23868

and t.data_ptr() % 16 == 0

Keep alignment eligibility rank-consistent

When different ranks receive contiguous BF16 views with different storage offsets, this local pointer check can select different collective implementations: an aligned rank enters _ext.oneshot_ar, while an unaligned rank falls back to NCCL in maybe_all_reduce, leaving both sides waiting indefinitely. CUDA allocator base alignment does not guarantee alignment of arbitrary contiguous views (for example, base[:8] versus base[1:9]), so alignment must be handled inside the kernel or agreed across all ranks before choosing the fallback.

---

stkernel/probes/mhc_glm53_bench.py

Lines 355 to 358 in 5c23868

	errs = [rel_err(outs[i], saved[m][j])
	for j, i in enumerate((2, 3, 4, 5))]
	line += (f" rel_err(max vs stock-passes ref)={max(errs):.2e}"
	+ ("" if max(errs) <= 1e-4 else " ! MISMATCH"))

Fail pass combinations that violate the numerics gate

When a TMA or warp-specialization combination exceeds the documented 1e-4 error gate, this code only prints ! MISMATCH and still exits successfully. The wrapper's if ! run_probe branch in probes/run_mhc_glm53_bench.sh lines 56-60 is explicitly responsible for classifying compile or divergence failures, so the invalid combination receives no failure verdict and can be mistaken for a viable candidate; return a nonzero status after detecting any mismatch.

---

stkernel/probes/mhc_glm53_bench.py

Line 338 in 5c23868

for m in (1, 2, 4, 8, 16):

Exercise the prefill regime in the pass sweep

The new runbook identifies prefill big_fuse as a principal remaining target after most decode MHC launches were replaced (RUNBOOK_KERNEL_CAMPAIGN2.md lines 595-598), but this sweep only tests token counts 1-16, which the dispatcher routes through the small-FMA decode branch (build/glm53/tilelang.py lines 736-737). Since the pass override affects every specialization, these measurements neither validate numerics nor select a winner for the primary prefill workload, and the prescribed probe can incorrectly close or advance the axis based solely on the residual decode path.

---

stkernel/overlay/modules/tp_oneshot_ar/dsv4_oneshot_ar.cu

Line 247 in 5c23868

uint4 mine[VECITER];

Bound the per-thread vector stash independently of MAXEL

With the supported VLLM_DSV4_OSAR_MAXEL=8388608 override, VECITER becomes 86 and this array consumes 1,376 bytes (344 32-bit values) per thread, exceeding the per-thread register ceiling and forcing the stash into local/global memory. Every copied vector is then written to and reread from local memory across the peer wait, adding substantial traffic and L2 pressure to the large-message sweep that the MAXEL override exists to measure; retain only a small fixed number of vectors or reread src once the register-sized stash is exhausted.

---

stkernel/probes/mhc_glm53_bench.py

Lines 623 to 625 in 5c23868

	t_stock = bench_us(lambda: _call_onepass(onepass, tensors, fn, ref))
	t_bf = bench_us(lambda: _call_onepass(mhc_onepass_bf16w, tensors,
	fn_bf, var))

Measure hcweight with DRAM-cold weights

These timings graph-replay the same 1.57 MB fn or fn_bf allocation 100 times, so after warmup the weight fits in the documented 24 MB L2 and the probe measures cache-resident loads. The serving workload cycles through 45 layers of distinct weights, while EXP-19's stated benefit is specifically halving the DRAM weight-read floor, so this result can produce the wrong go/no-go decision; rotate enough distinct weights to exceed L2 or explicitly flush L2 between replays.

ℹ️ 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".