probe(megakernel): MK-KDA 11차 — 배리어 17µs 카운터화 기각 확증 + 스탬프 진단 도구

MEASUREMENTS.md

@@ -1858,6 +1858,26 @@ C>1 은 요청마다 수락률이 달라 배치 구성이 흔들려 단일 정
 **먹힌 것**: W4 확장의 로컬 메모리 바이트 배열 → 레지스터 워드 산술(−20~33%; 타일 우선 팩 + cp.async 3단, 더 깊으면 손해); MHC p2 청크 축약을 24 레인에 + 4×4 sinkhorn 16 레인(bar2 13.5 → 5.4 µs; T=8 46→36, T=32 76→65); PDL 연쇄 발사(gemm 발사당 −3~5%, `VLLM_GLM53_MK_PDL`, 기본 off); ksr 호스트 계산·폴드 나눗셈 제거; 단계 스탬프(`-DMK_PHASE_TS`, gemm/mhc, wait 누적, 상한 프로브 `MK_PROBE_SKIP`).
 **안 먹힌 것(기록)**: W 채움 끌어올림(배리어 대기로 자리만 옮김), exact wait + 깊이 4·5, last-arriver 폴드, 동적 유닛 배분(n=6416 +1~2%), x cp.async 스테이징(프롤로그가 채움 전체를 기다림), smem 예산 공유(W8 −4~7% → 별도 인스턴스), mhc p1 fn 1회 읽기·fn 재배열(토큰당 지연이 지배), m≤16 mma 가드(언롤 안 break), 스위즐 3종.
 
+## ★메가커널 11차 — 배리어 17 µs 의 정체: 카운터화는 중립, "남은 것은 커널 밖" 확증 (2026-09-02 심야, srv4, 기각·복귀)
+
+10차 결산이 커널 내에 남겨둔 마지막 항목(배리어 ~17 µs)을 실험으로 닫았다. 스탬프(diag)가 보던 배리어 대기 med 6.1/4.2/7.4/2.0 µs 를 producer-consumer 카운터로 바꾸는 세 판을 돌렸다:
+
+| 판 | 설계 | 스팬 med (µs) | 판정 |
+|---|---|---|---|
+| 기준선(HEAD) | 배리어 4개 | **172.8** (min 170.1) | — |
+| 전체 카운터 | in_proj 타일→conv·gates, 헤드별 conv+f게이트→delta, 헤드별 delta+g게이트→norm 페어(워프 단위), norm\|o_proj 배리어만 유지 | 176.3 | 역효과 |
+| 부분 카운터 | conv 를 gates 앞으로(p0 꼬리와 겹침) + delta 진입만 카운터, delta\|norm·norm\|o_proj 배리어 복원 | 173.9 | 동일 |
+
+구현은 동작했다(교착 0, 수치 동일급, 그래프 재현 정확성은 consume-카운터 프로토콜 — 마지막 소비자가 도달 카운터를 지나간 뒤 초과화). **먹히지 않은 기전 넷**: (1) 배리어 대기 med 의 대부분은 **임계 경로 밖 흡수** — 놀고 있던 블록이 다른 블록의 임계 경로(p0 마지막 타일, delta 토큰 사슬)가 끝나기를 기다리던 자리였고, 카운터화는 그 대기를 스핀으로 옮기기만 한다. 배리어의 진짜 임계 비용은 producing phase 의 도착 스프레드 ~2-3 µs 뿐. (2) p0 뒤 체인은 **진짜 데이터 의존성으로 직렬** — gates 는 in_proj 마지막 타일(48-50)을, delta 는 자기 헤드의 conv+gates 를, o_proj 는 전 헤드의 norm 을 기다린다. (3) conv 조기 시작은 p0 의 W4 스트림과 DRAM 에서 충돌해 conv 자체 +4 µs. (4) delta 조기 시작의 상태 쓰기 홍수(8 MB/레이어)가 꼬리와 겹쳐 뒤 배리어 대기 2.0→5.9.
+
+`VLLM_GLM53_MK_KSR_IN` 스윕(같은 조건): r=1 206.7 / r=3 177.7 / r=4(모델) 175.7~182 / r=6 175.3 — r=1 만 확실히 나쁘고 나머지는 ±4 µs 노이즈 밴트. 모델의 r=4 유지, 채택 없음.
+
+**판정: 기각, HEAD 로 복귀.** 10차의 "단독 커널은 각자의 하한 근처, 남은 것은 커널 밖이다"가 직접 실험으로 확증됐다. 배리어 17 µs 는 비용이 아니라 **일이 없는 블록의 자리**였다. 다음 지렛대는 결산이 가리킨 대로 서빙 무장(브래킷)과 체인 융합뿐이다.
+
+부산물: ① `probes/diag_kda_stamps.py` 랜딩 — 8차·10차는 스크래치로만 존재하던 KDA 스탬프 리더(`VLLM_GLM53_MK_PHASE_TS=1` + `read_kda_ts`). ② README 의 stale 한 "노는 16 블록이 o_proj 팩을 `prefetch.global.L2` 로 예열" 문장 정정(10차에서 기각된 코드가 살아있는 것처럼 읽혔다).
+
+**측정 환경 주의(srv4)**: 서빙 glm53-worker 가 UMA 70~80 GB 를 잡은 채 공존 — 세션 중 OOM 3 회(벤치 재시도로 회피), 같은 바이너리의 스팬이 170→177 로 드리프트. 이 환경의 미세 A/B(±4 µs) 는 한 부팅 안에서만 유효하며, 전체 벤치의 KDA µs 도 이보다 신뢰할 수 없다(srv2 대비 부풀림: KDA 302~452 vs 280, stock 1827~1921 vs 640~1760 — stock 열은 부팅 상태 의존 편차가 크다는 6차 관찰과 일관).
+
 ## ★메가커널 10차 — KDA 2라운드: 스팬 199 → 176, 벤치 320 → 280 (2026-09-02 밤, srv2)
 
 8차 뒤 스팬 199 µs 에서 phase 단위로 다시: (1) p1(gates)·p2(conv)는 둘 다 p0 의 qkv 만 읽으므로 한 phase 로 합쳐 배리어 하나 제거(배리어 2 대기 3.4→0; 스팬 199→196). (2) o_proj 의 split-K 스윕(`VLLM_GLM53_MK_KSR_OUT` 프로브 노브): 모델 선택 r=3 40.8 µs, r=1 35.5, **r=2 33.4**, r=4 43.8 → 슬라이스가 8 k-블록보다 짧으면 cp.async 파이프라인이 채워지기 전에 끝나 "라운드 수"의 이득이 벽시계로는 손해. 비용 모델에 슬라이스 ≥ 8 k-블록 상한. (3) conv: 런타임 인덱스 람다(`st[]`/`kept[]` 로컬 메모리) → 토큰 루프를 스펙 창(8)으로 언롤하고 선택식으로: 6.6→4.7. (4) gates: 워프당 내적(행 재사용, 5단 셔플) → cp.async 32행 타일 + `mma.sync m16n8k16 bf16` GEMM: **14.5→6.2**(하한 4.5). (5) **기각** — delta 동안 노는 16 블록이 o_proj 팩 4.5 MB 를 `prefetch.global.L2` 로 예열: p5 35.1→29.5 로 예열은 먹히지만 delta 가 33.8→45.9(토큰 사슬의 로드가 예열 트래픽 뒤에 줄을 섬), 스팬 176→182, 벤치 280→309. 페이싱(~100 GB/s)하면 남는 이득이 2~3 µs 라 보류.

overlay/modules/glm53_megakernel/README.md

@@ -186,18 +186,24 @@ Bench (srv2, PDL on): T=8 27.4 us, T=32 42.0 us vs stock 32.8 / 71.6
 in_proj 76 | gates 6 | conv 4 | delta 34 | norm 0.5 | o_proj 35 |
 barriers ~17 = **176 us** per layer-step (402 before the phase
 stamps went in; stock's five kernels 640+). Four grid barriers: gates and
-conv share a phase (both read only in_proj's output). The gates are a
+conv share a phase (both read only in_proj's output). The barrier waits
+are NOT a remaining lever: replacing bar1/bar2/bar3 with per-tile and
+per-head arrival counters measured neutral-to-worse and was reverted
+(MEASUREMENTS.md 11차) -- the med waits are idle blocks absorbing while
+the true-dependency chain (gates needs the LAST in_proj tiles, delta its
+head's conv+gates, o_proj every head) runs; only the producing phase's
+arrival spread (~2-3 us) is on the critical path. The 16 head-less blocks
+during delta also stay idle on purpose: warming L2 with the o_proj pack
+from them measured a net loss (10차). The gates are a
 cp.async + `mma.sync m16n8k16 bf16` GEMM over 32-row weight tiles; the
 conv is unrolled to the 8-token spec window (the host refuses a wider
-`max_query_len`); the delta rule runs two blocks per head (rows split, S
-register-resident, per-token state stores staged through smem) while the
-16 head-less blocks warm L2 with the o_proj pack (`prefetch.global.L2`);
-p4 emits the o_proj's fp8 A tiles itself so p5 starts without a prologue;
-split-K never makes a slice shorter than 8 k-blocks (o_proj: r=2, not the
-cost model's 3). `-DMK_PHASE_TS=1` + `read_kda_ts` give the per-phase,
+`max_query_len`); p4 emits the o_proj's fp8 A tiles itself so p5 starts
+without a prologue; split-K never makes a slice shorter than 8 k-blocks
+(o_proj: r=2, not the cost model's 3). `-DMK_PHASE_TS=1` +
+`read_kda_ts`, read by `probes/diag_kda_stamps.py`, give the per-phase,
 per-block stamps; the fixture's `mk_run(drain=True)` keeps its own 10 MB
 state clones from polluting the first phase; `VLLM_GLM53_MK_KSR_IN/OUT`
-force a split for probing (MEASUREMENTS.md 8차, 10차).
+force a split for probing (MEASUREMENTS.md 8차, 10차, 11차).
 
 ## Review fixes already folded in (2026-09-01)
 

probes/diag_kda_stamps.py

@@ -0,0 +1,98 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+"""MK-KDA phase-stamp diagnosis (srv4 scratch container, never serving).
+
+Builds the extension with -DMK_PHASE_TS=1 (VLLM_GLM53_MK_PHASE_TS=1 must be
+set BEFORE the extension compiles), runs the acc=3 fixture with drain=True
+and prints the per-phase end stamps, the surviving barrier's wait, and the
+span. 11차: phases 1-3 hand work to their consumers through the
+g_mk_kda_* readiness counters, so a phase's end stamp includes its own
+producers' waits and the old bar1/bar2/bar3 waits no longer exist.
+
+    python3 /repo/probes/diag_kda_stamps.py [--acc 3] [--reps 5]
+
+Stamp slots (glm53_megakernel.cu, g_mk_kda_ts[block][16]) -- the slot
+semantics belong to whatever kernel is compiled, check the MK_KDA_TS
+markers in the .cu before reading this table; the shipped (barrier)
+order: 0 kernel entry | 1 p0 in_proj done | 2 bar1 out | 3 gates done |
+  4 (conv start) | 5 conv done | 6 bar2 out | 7 delta done | 8 bar3 out |
+  9 norm done | 10 bar4 out | 11 p5 o_proj done
+"""
+from __future__ import annotations
+
+import argparse
+import statistics
+import sys
+
+sys.path.insert(0, "/usr/local/lib/python3.12/dist-packages")
+
+import torch  # noqa: E402
+
+NB = 48  # grid cap; the probe refuses anything else
+
+
+def _med(xs):
+    return statistics.median(xs)
+
+
+def main() -> int:
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--acc", type=int, default=3)
+    ap.add_argument("--reps", type=int, default=5)
+    args = ap.parse_args()
+
+    from vllm.model_executor.layers import glm53_megakernel as mk
+
+    torch.cuda.init()
+    ext = mk._build()
+    major, minor, sms, smem = ext.probe_device()
+    print(f"device cc={major}.{minor} sms={sms} smem_optin={smem}")
+    assert (major, minor, sms) == (12, 1, 48), "not a GB10"
+
+    fx = mk._KdaFixture(acc=args.acc)
+    got = fx.mk_run(drain=True)  # warm the build/JIT paths
+    torch.cuda.synchronize()
+
+    spans = []
+    for rep in range(args.reps):
+        fx.mk_run(drain=True)
+        torch.cuda.synchronize()
+        kda = ext.read_kda_ts()  # [NB*16], cleared on read
+
+        def s(slot):
+            return [kda[b * 16 + slot] for b in range(NB) if kda[b * 16 + slot]]
+
+        if len(s(0)) != NB:
+            print(f"rep {rep}: incomplete stamps -- skip")
+            continue
+        entry = min(s(0))
+        span = (max(s(11)) - entry) / 1e3
+        spans.append(span)
+        print(f"--- rep {rep}  span {span:.1f} us")
+        for label, slot in (("p0 in_proj", 1), ("bar1 out", 2),
+                            ("gates", 3), ("conv", 5), ("bar2 out", 6),
+                            ("delta", 7), ("bar3 out", 8), ("norm", 9),
+                            ("bar4 out", 10), ("p5 o_proj", 11)):
+            v = s(slot)
+            print(f"  {label:<12} med {(_med(v) - entry) / 1e3:8.1f}"
+                  f"  slowest {(max(v) - entry) / 1e3:8.1f}")
+        # barrier waits per block: exit - entry around each barrier. The
+        # medians are mostly IDLE blocks absorbing (the producing phase's
+        # arrival spread is the critical-path part) -- 11차 closed the
+        # "remove the barriers" question with these numbers, see
+        # MEASUREMENTS.md before acting on a large med here
+        for bar, sin, sout in (("bar1", 1, 2), ("bar2", 5, 6),
+                               ("bar3", 7, 8), ("bar4", 9, 10)):
+            ent, out = s(sin), s(sout)
+            waits = [(o - e) / 1e3 for e, o in zip(ent, out) if e and o]
+            if waits:
+                print(f"  {bar} wait: med {_med(waits):6.1f}  max {max(waits):6.1f}")
+
+    if spans:
+        print(f"\nspan over {len(spans)} reps: med {_med(spans):.1f} "
+              f"min {min(spans):.1f} max {max(spans):.1f} us")
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())