diff --git a/MEASUREMENTS.md b/MEASUREMENTS.md index f79a9223..28b3ab07 100644 --- a/MEASUREMENTS.md +++ b/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 라 보류. diff --git a/overlay/modules/glm53_megakernel/README.md b/overlay/modules/glm53_megakernel/README.md index 9d8e2fde..2f4175b5 100644 --- a/overlay/modules/glm53_megakernel/README.md +++ b/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) diff --git a/probes/diag_kda_stamps.py b/probes/diag_kda_stamps.py new file mode 100644 index 00000000..66821bcc --- /dev/null +++ b/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())