choiceoh · choiceoh · Sep 4, 2026 · Sep 4, 2026 · Sep 4, 2026 · Sep 4, 2026
diff --git a/MEASUREMENTS.md b/MEASUREMENTS.md
@@ -2033,6 +2033,78 @@ top-k 슬롯이 2048 고정이라 장문맥에서도 디코드 이득은 ~1% 천
   355 토큰). 아키텍처가 같아도 커널 선택·리덕션 순서가 부팅마다 달라 근접 동률 argmax 가 뒤집힌다.
   따라서 **그리디 diff 로는 서빙 수치 변경을 판정할 수 없다** — 게이트는 품질 9/9·한국어 0/16·수용률
   프로파일이고, §6 에 남긴 "off 부팅과 diff" 후속은 폐기한다.
+## ★★메가커널 27차 — 서빙 PDL 은 꺼져 있었다, MoE 커널은 우리 레인과 같은 속도, L2-warm 은 W4 GEMM 에 7% (2026-09-04, srv2)
+
+전략 문서(2026-09-04 아티팩트 "메가커널 대대적 개선 방안")의 후보 중 운영자가 고른 셋 — 서빙
+PDL, AR 대기 중 다음 커널 L2 프리페치, MK_SEG_MOE go/no-go — 를 `probes/run_mk_probe.sh`(b12x
+파일까지 마운트하는 일반 러너)로 srv2 빈 창(다른 세션의 DRAFTW4 브래킷이 끝난 뒤)에서 쟀다.
+
+**1. 서빙 PDL (EXP-12).** 발견: `mk_pdl_enabled()` 는 env `VLLM_GLM53_MK_PDL` 을 읽는데 glm53.env
+에도 `ab-glm53.sh` 의 cand 팔에도 없었고 벤치 프로브만 켰다 → 09-03 무장 부팅(21차 트레이스)의
+173발은 전부 PDL 없이 돌았다. `probes/mk_pdl_graph_check.py`(그래프 캡처 안 체인): gemm→gemm→gemm 과
+mhc→gemm 의 리플레이가 eager 와 비트 동일, 재리플레이 동일 — PDL on/off 모두. 24발 그래프
+(n=k=4096, 8.4 MB 팩 24개 회전, DRAM-cold): **off 58.0 → on 53.6 µs/발사(−7.6%)**. 173발/스텝이면
+−0.76 ms(1.1%). 프로필 기본 1, cand 팔에 명시. 수치 불변이라 EXP-6 브래킷의 cand 팔에 얹는다.
+
+**2. MK_SEG_MOE go/no-go (EXP-14).** 서빙이 만드는 b12x 래퍼(같은 기하·같은 디스패치 오버레이;
+C=1 은 64 pairs 라 static 백엔드)를 디코드 형상(8토큰·top-8)에서 고유 전문가 U 별로, 가중치
+DRAM-cold(8 세트 순환, 세트당 1 GB), 그래프 리플레이로:
+
+| U | µs/호출 | MB | GB/s |
+|---|---|---|---|
+| 8 / 16 / 24 | 148 / 311 / 451 | 28 / 57 / 85 | 191 / 182 / 188 |
+| 32 / **40** / 48 | 573 / **720** / 859 | 113 / 142 / 170 | 198 / **197** / 198 |
+| 56 / 64 | 985 / 1135 | 198 / 227 | 201 / 200 |
+| MK W4 레인 n=6416 (12팩 그래프, PDL) | 76.7 | 15.0 | **196** |
+| torch sum (읽기만, 180 MB) | 736 | 180 | 245 |
+
+U=40 의 720 µs 는 21차의 층당 738 µs 와 맞는다 — 서빙의 고유 전문가 수가 ~40 이라는 것과
+실효 ~197 GB/s 가 둘 다 확인됐다. **판정: b12x = 레인의 103% → 90% 규칙으로 축을 닫는다.**
+남은 20%(196~200 vs 읽기 245)는 두 커널에 공통인 "발사 안 구조" 몫이다: 레인은 루프만 244
+GB/s 이고 발사당 ~15 µs 의 고정비(4차)가 있으며, MoE 커널은 층당 한 발사라 고정비가 아니라
+타일 구조다. persistent 스트림이 240 에 닿는다는 증거는 아직 없다(레인의 발사당 196~222). 그러니
+MK_SEG_MOE 의 상한 −4~5 ms 는 그 가정 위에 있고, 착수하려면 먼저 2~3일짜리 "persistent 전문가
+타일 스트림" 마이크로커널로 240 도달 여부를 재야 한다. 지금은 보류.
+
+**3. AR 프리페치의 소비자 이득 (EXP-13 게이트 3).** n=6416 W4 GEMM 한 발: DRAM-cold 86.0 µs →
+팩 15 MB 를 L2 에 미리 읽어 둔 뒤 79.9 µs(**−7%**). W4 GEMM 은 L2-warm 이어도 80 µs 다 — DRAM 이
+아니라 발행(cp.async/LSU, 2차)에 묶여 있어서, 프리페치가 사는 것은 DRAM 지연분뿐이다. 회당
+~10 MB 를 데워도 −4 µs/AR, 스텝당 100회면 **−0.4~0.6 ms(0.6~0.9%)**. 전략 문서의 −1.5~2.5 ms 는
+"L2 는 DRAM 의 4배" 가정이었고 그 가정이 틀렸다. 구현은 랜딩했다(기본 off, 힌트 없으면 커널
+byte-identical): 4랭크 disttest(`t_wait` 힌트/무힌트)와 브래킷은 부팅이고, 단독으로는 CV 1.7%
+아래라 EXP-6+12 위에 얹어서만 판정한다.
+
+**4. AR 커널 빌드 (EXP-13 게이트 1).** `probes/osar_build_check.py`: `prefetch.global.L2` 가 sm_121a
+ptxas 를 통과하고 `oneshot_ar_hint`·`phase_counters` 가 묶였다. PASS.
+
+**5. 드래프터 꼬리 인구조사 (09-03 트레이스, 깨끗한 디코드 스텝 5·6).** 샘플러 앵커에서 다음 스텝의
+준비 커널까지 136 커널, 커널 합 6.28 ms(스팬 6.1~7.1): bf16 cutlass GEMM 33 × 122 = 4.0 ms(EXP-10 이
+서빙되면 ~1.25), 타깃 헤드 fp8 814, `k_oneshot` 11 × 72 = 788(forward 안의 45 보다 느리다 — 꼬리
+구간의 랭크 편차), `kernel_mha` 5 × 29 = 145, AllGather 2 × 54 = 108, 글루 ~50개 0.33 ms. 드래프터
+메가커널의 상한 = 글루 0.33 + 발사 간극 ~0.35 + MK 발사 고정비 30 × ~10 µs ≈ **−0.8~1.0 ms
+(1.2~1.5%)**; AR 0.79 ms 는 융합해도 남는다.
+
+**6. 전략 문서의 상한 대 프로브 (정정).** 운영자 물음 "실질 이득이 적다면 왜 처음엔 크다고
+봤나" 에 대한 답:
+
+| 항목 | 전략 상한 | 프로브 뒤 | 틀린 가정 |
+|---|---|---|---|
+| AR 프리페치 | −1.5~2.5 ms | −0.4~0.6 | 소비자가 DRAM 바운드라는 가정. 2차의 "발행 병목" 을 L2-warm 에 적용하지 않았다 |
+| MK_SEG_MOE | −4~8 ms | 닫힘 | MoE 190 을 레인의 루프 전용 244 와 비교. 같은 형태(그래프 안 발사당)면 레인 196 = MoE 197 |
+| 샤딩 샘플링 | −0.4 ms | −0.1 | 25차의 AllGather 409~567 µs 는 프리필 문맥 스텝의 값. 깨끗한 디코드는 2 × 54 µs → **후보에서 내림** |
+| 드래프터 메가커널 | −1~1.5 ms | −0.8~1.0 | 대체로 맞음(위 5) |
+| 서빙 PDL | −0.9~1.7 ms | −0.76 | 범위 안 |
+
+공통 원인: 상한을 목표와 같은 측정 형태(그래프 안 발사당, 깨끗한 디코드 스텝, 확인된 병목)로
+만들지 않았다. 남은 큰 항목 — k 축소(SPEC=0 대 k=7 실측에 앵커), W4 팔(발사당 실측), EXP-7(호스트
+유휴 실측) — 은 그 형태의 측정 위에 있다.
+
+**교훈**: (1) 노브에 독자가 있어도 서빙이 그 값을 나른다는 보장은 없다 — 프로브가 켜는 env 는
+프로필에도 있어야 하고, 무장 부팅의 fingerprint 에 찍혀야 한다. (2) "대역폭 바닥" 판정은 같은
+바이트를 우리 커널이 얼마나 빨리 흘리는지와 견줘야 한다 — 역산값끼리 맞는 것은 검증이 아니다.
+(3) L2-warm 이 이득이 되려면 소비자가 DRAM 바운드여야 한다 — W4 GEMM 은 아니었다.
+(4) 상한은 목표와 같은 측정 형태로 — 루프 전용 스탬프·프로파일러 꼬리·미확인 병목 가정으로 만든
+상한은 프로브 하나에 절반 이하로 준다.
 
 ## ★★★26차 — 부팅의 메모리 절벽: fp8-dense 패스가 로드 중에 여러 번 돌았고, W4 팩 빌더가 예약 메모리를 흘렸다 (2026-09-04, 4노드 계측)
 

diff --git a/RUNBOOK_KERNEL_CAMPAIGN2.md b/RUNBOOK_KERNEL_CAMPAIGN2.md
@@ -483,6 +483,97 @@ MLA(rope 64 · topk 512 · 압축기 · 슬라이딩 윈도)·GEMM(dense 가 블
   프로브의 `--stock dispatch` 가 부팅이 실제로 타는 팔을 재고 `hit` 열로 그걸
   말한다.
 
+## EXP-12 — 서빙 PDL (`VLLM_GLM53_MK_PDL=1`, 2026-09-04 추가)
+
+메가커널 발사는 PDL(programmatic dependent launch)로 튜닝돼 있다 — 다음 MK 커널이
+앞 커널이 비운 SM 에서 시작해 꼬리 동안 첫 W 타일을 당긴다(연속 2발 발사당
+−17~19%, 2차). 그런데 드라이버는 env 를 읽고 **프로브만 그것을 켰다**: 프로필에도
+`ab-glm53.sh` 의 cand 팔에도 없었으므로 지금까지의 무장 부팅은 전부 PDL 없이
+돌았다(173발/스텝). EXP-6 종단 무효과의 용의자 1번.
+
+- 수치 불변: 모든 MK 커널이 앞 커널 출력을 읽기 전에 `griddepcontrol.wait` 하고,
+  wait 앞의 채움은 가중치뿐(mk_gemm_phase 의 hoist). 그래프 캡처 안 체인 형태는
+  `probes/mk_pdl_graph_check.py` 가 판정한다(gemm→gemm→gemm, mhc→gemm 리플레이
+  비트 동일 + PDL on/off 발사당 µs).
+- 배선: `profiles/glm53.env` 기본 1, `ab-glm53.sh` cand 팔에 명시. 세그먼트가
+  하나도 무장되지 않으면 무효(inert).
+- 상한: 173발 × 5~10 µs = −0.9~1.7 ms/스텝. 단독 부팅 금지 — EXP-6 브래킷의 cand
+  팔에 얹는다(수치 축 하나 규칙: PDL 은 속도만 바꾼다).
+
+```bash
+bash probes/run_mk_probe.sh probes/mk_pdl_graph_check.py                     # PDL=1
+VLLM_GLM53_MK_PDL=0 bash probes/run_mk_probe.sh probes/mk_pdl_graph_check.py # 대조
+```
+
+**결과(2026-09-04, srv2, 27차)**: 체인 리플레이 비트 동일(on/off 모두), 24발 그래프 off 58.0 →
+on 53.6 µs/발사(−7.6%). 프로필 기본 1. 남은 것 = EXP-6 브래킷의 cand 팔.
+
+## EXP-13 — AR 대기 중 다음 커널 가중치 L2 프리페치 (`VLLM_GLM53_AR_PREFETCH`, 2026-09-04 추가)
+
+`k_oneshot` 의 대기(회당 38.7/45.5 µs, 스텝당 ~100회)는 모든 랭크에서 DRAM 이
+노는 시간이다. 커널이 `HintArgs`(최대 8개 (포인터, 바이트))를 받아 워프 1~7 이
+`prefetch.global.L2` 로 걷고 스레드 0 은 그대로 플래그를 돈다. 힌트는 **학습**:
+MK 드라이버의 발사(`_gemm_call`·`_mhc_call`·`_kda_launch`)가 읽을 가중치를
+`note_consumer` 로 알리고, shim 이 "타깃 forward 의 몇 번째 콜렉티브 뒤인가" 로
+파일해 두었다가 캡처 시 발사에 굳힌다. forward 경계는 컴파일 영역 위의
+`Glm5NextForConditionalGeneration.forward`. MK-GEMM 이 무장돼야 배울 것이 있다.
+
+- 게이트(순서대로): (1) `probes/osar_build_check.py` — 새 명령의 ptxas 통과(실패면
+  부팅이 NCCL 로 조용히 떨어진다), (2) `probes/oneshot_ar_disttest.py` 4랭크 —
+  12 MB 힌트/무힌트의 `t_wait` 와 maxerr 0(힌트가 NIC 쓰기를 밀어 대기를 늘리면
+  손해), (3) `moe_decode_stream_probe.py` 의 `gemm cold` vs `gemm L2-warm` 행 —
+  소비자 쪽 이득의 단위, (4) 플릿 브래킷: EXP-6+12 위에 `VLLM_GLM53_AR_PREFETCH=1`
+  (caller env). 수치 불변 → step/s 만. 부팅 로그에 `[osar] prefetch hints learned:
+  N collectives, X MB` 가 없으면 무장이 아니다.
+- 상한: 임계 랭크의 대기 ≈ 전송 ~20 µs = 4.6 MB → −1.5~2.5 ms/스텝(전략 문서).
+- 예산 노브: 1 = 12 MB/콜렉티브, N = N MB(1..20; L2 24 MB).
+- **결과(2026-09-04, 27차)**: 게이트 1 PASS(빌드), 게이트 3 = n=6416 W4 GEMM cold 86.0 → L2-warm
+  79.9 µs(−7%): 소비자가 DRAM 바운드가 아니라 상한은 **−0.4~0.6 ms/스텝**으로 내려갔다(전략 문서의
+  −1.5~2.5 는 철회). 단독 판정 불가 — EXP-6+12 위에 얹어서만. 게이트 2(4랭크 disttest)는 부팅.
+
+## EXP-14 — MK_SEG_MOE go/no-go (2026-09-04 추가, 프로브만)
+
+21차의 "MoE 는 대역폭 바닥" 은 190 GB/s 를 **추정한 고유 전문가 수 ~40** 에서
+역산한 값이라 순환이다. `probes/moe_decode_stream_probe.py` 가 서빙이 만드는
+b12x 래퍼(같은 기하, 같은 디스패치 오버레이; C=1 은 64 pairs 라 static 백엔드)를
+디코드 형상(8토큰·top-8)에서 고유 전문가 U=8..64 별로, 가중치 DRAM-cold(8 세트
+순환), 그래프 리플레이로 재고 같은 바이트를 MK W4 레인이 스트리밍하는 속도(팩
+12개 연속, PDL)와 견준다.
+
+- 판정 규칙: b12x 가 레인의 90% 이상이면 축을 닫는다(원장에 기록). 아래면
+  (레인 − b12x) 비율 × 31 ms 가 세그먼트의 상한이고, 설계는 전략 문서 4장(48블록
+  persistent, FC1 (전문가, n타일) 유닛 → 전문가별 완료 카운터로 열리는 FC2 동적
+  큐, 공유 전문가 = 41번째 전문가, b12x nvfp4 레이아웃 제자리 읽기, A4→A8).
+- 실제 서빙 U 는 다음 부팅에서 로그 한 줄로 확정한다(프로브는 U 별 곡선만 준다).
+- **결과(2026-09-04, 27차)**: b12x static U=40 = 197 GB/s, 레인 n=6416 = 196 GB/s → **103%, 닫힘**.
+  읽기 전용 참조(torch sum) 245 GB/s 와의 20% 는 두 커널 공통의 발사 안 구조 몫이라, 재개 조건은
+  "persistent 전문가 타일 스트림 마이크로커널이 240 에 닿는다" 는 2~3일 프로브의 양성이다. 보류.
+
+## EXP-15 — 드래프터 fc 를 타깃 헤드·샘플러 아래로 (`glm53_dflash_early_fc`, 2026-09-04 추가)
+
+디코드 꼬리는 타깃 헤드 → 로짓 AllGather → 거부 샘플러 → 드래프터 순인데, 드래프터의 첫
+GEMM `fc`(aux 은닉 [토큰, 5×4096] → 4096; 레인에서 301 µs)는 타깃 forward 가 끝나는 순간
+입력이 다 있다. stock 은 `propose()` 안에서 샘플러 뒤에 계산한다. 그 사이 구간(AllGather 는
+패브릭, 샘플러는 소형 커널)은 DRAM 이 놀아 fc 가 공짜로 흐른다. 생산자 = `GPUModelRunner.
+execute_model` 래퍼(forward 뒤 side stream 에서 cat + fc, 영속 버퍼 + 이벤트), 소비자 =
+드래프터 오버레이의 `combine_hidden_states`(같은 토큰 수의 대기 결과만, 이벤트 대기 뒤). 소비가
+`precompute_and_store_context_kv` 와 드래프터 그래프보다 앞이라 MK 발사끼리 겹치지 않는다.
+
+- 수치 동일(같은 커널·같은 입력). 노브 `VLLM_GLM53_DFLASH_EARLY_FC=1`, 기본 0, 생산자 실패 시
+  부팅 동안 자동 해제.
+- 상한 ~0.3 ms/스텝(27차 인구조사: fc 만 그 시점에 입력이 준비된 꼬리 GEMM). EXP-10 위에서만
+  의미(fc 가 레인에 있어야 함). 단독 판정 불가 — EXP-10 브래킷의 cand 에 얹는다.
+
+## EXP-16 — 드래프터 메가커널 (제안, 착수 전 승인; 2026-09-04)
+
+27차 인구조사(깨끗한 디코드 스텝): 꼬리 136 커널 중 GEMM 33(EXP-10 뒤 MK 발사 ~31), AR 11,
+`kernel_mha` 5, 글루 ~50개 0.33 ms, 발사 간극 ~0.35 ms. 융합으로 없앨 수 있는 것은 글루 + 간극 +
+MK 발사 고정비(30 × ~10 µs) ≈ **−0.8~1.0 ms(1.2~1.5%)**; AR 0.79 ms 와 mha 는 남는다. 브래킷
+해상도(CV 1.7%) 아래라 단독으로는 판정할 수 없고, 공사는 MK-KDA 급(2주). 설계는 층당 두 발사:
+[norm → conv.prepare → qkv GEMM] 과 [conv.finish → post-norm → mlp_conv.prepare → gate_up →
+act → mlp_conv.finish → down], 어텐션(`kernel_mha`)은 stock 유지. **운영자 승인 뒤 착수** — 정정된
+상한을 본 뒤의 결정이어야 한다.
+
 ## 브래킷 자동화 — `bench/bracket.py` (도구, 판정 아님)
 
 `leg`(살아있는 서버에 rep 기록) + `judge`(기록 판정) 2중 명령. 원장 규율을 코드로
@@ -506,6 +597,9 @@ MLA(rope 64 · topk 512 · 압축기 · 슬라이딩 윈도)·GEMM(dense 가 블
 9. **EXP-9 (head-gate split-K)** — 단독 부팅 금지, EXP-7 부팅에 얹는다.
 10. **EXP-7 이 붙은 뒤 드래프터 D 를 다시 잰다** — 9월 1일 트레이스에서 드래프터 ~4.3 ms 는 다음 스텝의 호스트 준비 유휴 뒤에 숨어 있었다(그래서 D≈0). 은신처가 사라지면 임계경로에 올라온다(천장 ~6%; fc GEMM 809 us 는 K=20480 직렬 스케줄이라 split-K 후보). #104 를 지금 재론하는 것이 아니라 조건이 바뀐 뒤의 재측정이다. 오프라인으로 닫을 >1% 레버는 더 없다(STEP_KERNEL_MAP 보충 분해 3).
 11. **EXP-10 (드래프터 GEMM → MK W4)** — **닫힘, 기본값(2026-09-04, 28차)**: 서빙된 브래킷 C=1 step/s 15.95 → 16.235(+1.8%), 수용률·품질·한국어·프리필 게이트 통과. 첫 브래킷의 0 은 컴파일 캐시가 옛 bf16 그래프를 서빙한 탓 — 노브는 이제 캐시 키이고 부팅 로그가 `drafter lane serving: 30 of 31` 로 증명한다. MK-MLA 서빙 사망(스크래치 재할당)도 같은 항목에서 수정.
+12. **EXP-12 (서빙 PDL)** — 프로브(그래프 체인) 뒤 EXP-6 브래킷의 cand 팔에 얹는다. 단독 부팅 없음.
+13. **EXP-13 (AR 프리페치)** — 컴파일 → 4랭크 disttest → 브래킷(EXP-6+12 위). 수치 불변.
+14. **EXP-14 (MK_SEG_MOE go/no-go)** — 프로브 하나가 착수 여부를 정한다. 90% 규칙.
 
 12. **EXP-11 (dsv4 에 MK_SEG_MHC)** — 2단계는 **부팅 없음**: srv2 에서 서빙 컨테이너가
     비었을 때(`docker ps`) `bash probes/run_megakernel_bench.sh --profile dsv4 --iters 20`.

diff --git a/bench/bracket.py b/bench/bracket.py
@@ -39,6 +39,7 @@
         "VLLM_GLM53_PREP_FUSED", "VLLM_GLM53_ASYNC_DFLASH",
         "VLLM_GLM53_MHC_SMALLM", "VLLM_DFLASH_PREP_WARMUP",
         "VLLM_GLM53_MK_PDL", "VLLM_GLM53_MK_KSR_OUT",
+        "VLLM_GLM53_AR_PREFETCH", "VLLM_GLM53_DFLASH_EARLY_FC",
     ) if k in os.environ]
     + sorted(k for k in os.environ
              if k.startswith(("VLLM_GLM53_MK_", "VLLM_GLM53_KPOOL")))