choiceoh · choiceoh · Sep 2, 2026 · Sep 2, 2026
diff --git a/MEASUREMENTS.md b/MEASUREMENTS.md
@@ -1858,6 +1858,28 @@ 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종.
 
+## ★메가커널 9차 — MHC 재구성: T=8 30.7→27.4, T=32 59.4→42.0 µs (2026-09-02, srv2)
+
+MHC 스탬프(main, T=8 / T=32): p1 13.4 / 37.8, bar1 3.1 / 3.0, p2 0.1(최대 4.0~4.4: 토큰당 워프 하나라 블록 0 에 몰림) + bar2 5.3 / 5.3, p3 0.8 / 2.6, bar3 1.4 / 1.5, p4 1.1 / 1.9 → 스팬 24.9 / 52.3, 이벤트 32.8 / 61.2.
+
+| 실험 | 결과 (T=8 / T=32, 이벤트 µs) | 배운 것 |
+|---|---|---|
+| p1 의 청크별 fn 96 KB smem 상주, 48 블록 | 43.0 / 75.7 (악화) | 96-로드 채움이 지연 지배, 토큰마다 smem 96 읽기, 동시성 1/3 |
+| 꼬리(p2·p3·p4)를 토큰의 last-arriver 가 처리, 배리어 3개 제거 | 34.6 / 86.0 (악화) | 쌍 16t+15 를 맡는 블록이 t, t+9, t+18, t+27 의 꼬리를 직렬로 |
+| 꼬리 티켓 큐(유휴 블록이 도착 카운터 대기 후 처리, 마지막 퇴장 블록이 재무장) | 34.8 / 59.4 | T=32 동률, T=8 은 꼬리 ~13 µs 가 그대로 더해짐 |
+| 꼬리 분리 스탬프 | p2 10.7, p34 10.0 (T=8) | 3 블록/SM 이면 컴파일러가 레지스터를 85 로 묶어 꼬리가 스필 → L1 이 거의 없는 SM 에서 로컬 메모리 왕복 |
+| p34 로드를 p2 밑에서 선발행 | 32.8 / 67.6; 96 블록(2/SM) | 레지스터가 늘어 2 블록/SM 이 되자 꼬리는 p2 4.7 / p34 0.7 로 정상, 대신 p1 동시성 손실(p1 end 11.7→17.5) |
+| **fn 을 레지스터(96개)에, 블록당 토큰 3개, 큐 꼬리** | **30.7 / 51.2** (stock 32.8 / 71.7) | T=8 의 p1 11.7 µs 는 128쌍 동시 × fn 96 KB = 12 MB 의 L2→SM 트래픽(L2 속도 한계); 레지스터 상주로 1/3. 48 유닛만 있어 96 블록의 절반이 놀았다 |
+
+| 꼬리를 레지스터 절약형으로(레인별 셔플, p34 선로드는 워프 1~7만) | 34.8 / 45.1 | 꼬리 p2 12→8.8 |
+| 적응형 그룹(grid/16) + 다음 토큰 x/residual 선발행 | 36.9 / — | T=8 p1 end 15.0→14.8 무변화: 토큰당 사슬 ~5 µs 가 병목 |
+| **mix 계수 선로드 + 8레인 그룹 축약(32 로드 + 3 셔플)** | **27.4 / 41.8** | 레지스터가 늘어 48 블록(1/SM); 토큰당 사슬 ~3.3 µs |
+| hc 상수 선로드, sinkhorn 근사 내장함수, 단일 레인 4×4 sinkhorn | 27.4 / 42.0 (동일) | 프로브 창(4.5 µs)은 줄었지만 벤치 불변 — `%globaltimer` 프로브 자체가 워프 0 의 사슬을 창마다 ~0.5 µs 늘린다. 실제 꼬리는 ~5 µs |
+
+**최종(c69f6d8)**: MHC **T=8 27.4 µs, T=32 42.0 µs**(stock 32.8 / 71.6 → −16% / −41%; 이번 라운드 시작 30.7 / 59.4). 수치 1.25e-7 / 3.86e-7(게이트 1e-3). 구조: 격자 배리어 0개, (청크, 토큰 그룹) 블록 + fn 레지스터 상주, 전치 smem 축약, 꼬리 티켓 큐(마지막 퇴장 블록이 재무장), p3+p4 레지스터 융합.
+
+**교훈**: (1) 축약·발행 사슬이 있는 커널은 SM당 동시 블록 수가 성능이다 — 큰 smem/레지스터로 동시성을 줄이면 같은 코드도 느려진다; (2) "마지막 도착자가 꼬리를 처리"는 매핑에 따라 한 블록에 직렬로 쌓인다 — 티켓 큐가 안전; (3) 3 블록/SM 목표는 컴파일러의 레지스터 상한(85)을 뜻하므로 꼬리 같은 희소 경로의 배열이 스필된다; (4) 프로브는 측정 대상을 바꾼다 — 벤치(이벤트)로 확정할 것.
+
 ## ★★메가커널 8차 — MK-KDA 단계 귀속과 재작성: 커널 스팬 402 → 199 µs (2026-09-02, srv2)
 
 KDA 커널에 단계 스탬프(`-DMK_PHASE_TS`, `read_kda_ts`, 스크래치 `diag_kda.py`)를 넣고 6 phase 를 귀속시켰다. 시작점(acc=3, 스팬 402 µs): in_proj 104, **gates 68**, conv 4(+배리어 9.7), **delta 160**(헤드당 블록 하나, 나머지 32블록은 배리어 4 에서 158 대기), norm 2.6, o_proj 41.
@@ -1975,59 +1997,3 @@ W4 루프의 누적 스탬프(srv2): n=6416 루프 127 µs = mma_fold 41 + expan
 **먹힌 것(순서대로, W4 x2 n=6416)**: A 타일 밀집 128 B + XOR 스위즐(4-way 충돌 제거) 145→141; 확장의 바이트-레인 SIMD(LUT 닫힌 식, `__vcmpeq4/__vadd4/__byte_perm`, 비트 동일) →128; mma k 스텝 완전 언롤 →122; stage_a L2 로드 선발행 →118; m ≤ 16 에서 두 번째 m 타일 제거(`mma_fold<MT>`, 제네릭 람다) →112. W8 도 같은 절감으로 6416 +10%·4096 동률·2048 +4%·1024 +8% 까지.
 **지속 스트림 한계**: srv2 마이크로벤치에서 cp.async·TMA 1D·TMA 2D(텐서맵 128 B 스위즐)·레이아웃(타일 우선/행 우선 128~2 KB 런)·깊이 3~8 모두 ~135~150 GB/s. 프로세스의 **첫 커널 배치만 225 GB/s**(부스트) — 마이크로벤치는 워밍업 뒤 값만 믿을 것. W8 이 deepgemm 에 지는 남은 10% 는 스트림이 아니라 고정비.
 **W4 채택은 운영자 결정**(서빙 수치 변경, by-design 1.24e-01): README 의 품질 브래킷(9/9 + 한국어 0/16 + pos-1 수용률 2%p) 뒤에만.
-
-## ★★준비 커널 통합 (`glm53_prep_fused`) — 스텝의 12% 가 GPU 유휴, 그 자리에 발사 하나 (2026-09-02, srv4 오프라인)
-
-**발견 (9월 1일 트레이스 재분석, rank 3, 229 스텝, 스트림 합집합)**: 프로파일된 72.1 ms
-스텝에서 GPU 가 비는 시간 8.9 ms(12%)가 한 곳에 몰려 있다 — 드래프터 그래프와
-타깃 그래프 사이의 eager 입력 준비 5.7 ms(aten 1,028 호출·memcpy 45·1~3 us 커널
-~100개, 커널 간 공백 50~430 us = 호스트가 발사를 못 따라감), 타깃 그래프 제출
-`cudaGraphLaunch` 1.43 ms(1,640 노드, 50스텝 중앙, p10 1.34 p90 1.59), 드래프터 앞
-DtoH 대기 0.6, 스텝 전환 0.8. 프로파일러가 호스트를 부풀리므로 프로파일러 없는
-앵커는 기존 원장의 "디코드 중 GPU 93%" (같은 정의) → 실제 유휴 4~5 ms/66 ms. dflash
-가 async scheduling 을 끈다는 기록(위 표)과 일치한다. 준비 구간이 그 대부분이다.
-**천장 = 준비 구간 자체, 스텝의 4~6%**; 읽는 바이트 0.
-
-같은 트레이스에서 그래프 안 elementwise 글루 605개의 정체도 잡았다(STEP_KERNEL_MAP
-보충 분해 2): 풀어텐션 인덱서 275(우리 kpool 파일), KDA 의 split-뷰 복사 136(MK-KDA
-가 흡수), MoE shared 덧셈 42, 나머지 ~150 — 커널당 CUPTI 2.0 us + 간격 0.2 us,
-전부 합쳐 2.7 ms(부풀린 값)라 호스트 유휴가 그래프 안 글루 전체보다 크다.
-
-**구현**: `overlay/modules/glm53_prep_fused` — 균일 spec-verify 스텝(전 요청 드래프트
-7개, FULL 그래프, 요청 패딩 없음)에서 `prepare_inputs` + `prepare_attn` + KV 그룹 7개
-빌더가 쓰는 **모든 persistent 버퍼**를 pinned H2D 1회 + Triton 발사 1회 + deep_gemm
-스케줄 1회로 쓴다. 메타데이터 dict 는 형상별 캐시(FULL 리플레이는 버퍼만 읽고,
-dict 는 드래프터에 넘어가 무시된다). 설치는 러너/빌더 파일 17개의 preimage 를
-고정(드리프트 → DISARM), plan 은 캡처 직후 live 러너에서 만들고 커널을 미리
-컴파일한다. 러너 파일은 덮지 않는다(메서드 패치).
-
-**오프라인 게이트 (srv4 새 컨테이너 `glm53:sm121-fi618`, 마운트 오버레이 30개,
-`probes/run_prep_fused_check.sh`)**: stock 빌딩블록(input_batch Triton 3종·
-BlockTables gather/slot·GDN FULL 복사·MLA req_id·indexer uniform-decode +
-compressed slot + deep_gemm 스케줄) vs fused 커널, 프로덕션 기하(그룹 7, MLA 폭
-2052, indexer 256토큰 페이지/factor 4/storage 64, kpool 4, mamba 폭 8), 무작위
-배치 C=1~4 **60/60 회 46개 텐서 bit-exact**.
-
-| C=1 스텝당 | stock 빌딩블록 | fused |
-|---|---|---|
-| 발사 | 커널 ~30 + memcpy ~8 | 커널 3 + memcpy 1 |
-| 벽시계(호스트+GPU) | 2,556 us | 201 us |
-| GPU 이벤트 | (호스트 대기 포함 2,562) | 200 = 커널 88.5 + deep_gemm 스케줄 64.4 + 복사 |
-
-stock 열은 빌딩블록만이라 러너의 aten ~1,000 호출은 빠져 있다 — 실제 절감은 이보다
-크고, 상한은 트레이스의 준비 구간(프로파일러 없이 3~5 ms). fused 의 GPU 200 us 는
-임계경로에 새로 올라오는 비용이다(그중 64 us 는 stock 도 내는 deep_gemm 스케줄).
-커널 88 us 는 요청당 직렬 루프(그룹 7 gather + 토큰 8 × 폭 2052 행 복사)라 더 줄일
-수 있으나, 절감 규모(ms) 대비 작아 섀도 부팅 뒤로 미룬다.
-
-**미측정 (서빙)**: 섀도 부팅(`VLLM_GLM53_PREP_FUSED=shadow`, drift=0 확인) → EXP-7
-브래킷(C=1 step/s). 물리 기전 확인 = 켠 부팅 트레이스에서 준비 구간의 memcpy·
-`at::native::*` 가 사라지고 `_glm53_prep_fused_kernel` 하나가 남는 것. 기본값 0.
-
-**부수 발견(미조치)**: (1) kpool tail 의 원형 슬롯 매핑은 러너가 빌더에 `positions`
-를 안 넘겨 이 이미지에서 잠들어 있다 — `glm53_tail_slot_persistent` 의 고정 버퍼도
-그래서 무효이고, tail 그룹은 generic 매핑(문서 자체가 "pos >= kpool 이면 블록 0 으로
-붕괴"라 적은 것)을 쓴다. C>=2 수치 축이라 별도 브래킷. (2) 인덱서 fp32 head-gate
-`torch.mm` 이 cuBLAS gemmSN 2블록 커널로 층당 86 us × 11(CUPTI) — 우리
-`glm53_prefill_fastpath.py:402` 소유. (3) 드래프터 커널 합 ~3 ms(CUPTI; fc 투영
-814 us eager bf16 168 MB 포함) — 원장 D≈0 과 긴장, 직접 측정 전 판단 보류.
diff --git a/RUNBOOK_KERNEL_CAMPAIGN2.md b/RUNBOOK_KERNEL_CAMPAIGN2.md
@@ -248,59 +248,6 @@ EXTRA_ENV="VLLM_GLM53_MEGAKERNEL=1 VLLM_GLM53_MK_GEMM=1" \
 - 첫 암 부팅은 확장 컴파일(~1분, `/root/.mk_build`)만큼 느려진다.
 - MK-KDA는 3단계 섀도 로그가 깨끗하기 전에 4단계에 올리지 않는다.
 
-## EXP-7 — 준비 커널 통합 (`glm53_prep_fused`, 2026-09-02 추가)
-
-디코드 스텝의 **호스트 쪽** 입력 준비를 Triton 발사 하나로 접는 모듈. 트레이스
-(2026-09-01, rank 3, 229 스텝)에서 스텝의 GPU 유휴가 8.9 ms/72 ms(12%) 이고,
-그중 5.7 ms 가 드래프터 그래프와 타깃 그래프 사이의 eager 준비 구간이다 —
-`prepare_inputs` + `prepare_attn` + KV 그룹 7개의 메타데이터 빌더가 스텝마다
-aten 호출 ~1,000개, memcpy ~45개, 1~3 us 짜리 커널 ~100개를 내고, dflash 는
-스케줄러가 동기라 그 시간 내내 GPU 가 빈다. 프로파일러 없는 앵커는 원장의
-"디코드 중 GPU 93%" (같은 정의) → 실제 4~5 ms/스텝. 천장: 그 구간 자체,
-스텝의 4~6%. 읽는 바이트는 0.
-
-| 스텝의 준비 구간 | stock | fused |
-|---|---|---|
-| H2D memcpy | ~45 | 1 (idx_mapping, pinned) |
-| 커널 발사 | ~100 | 1 (+ deep_gemm 스케줄 메타 1 + 복사 1) |
-| aten 호출 | ~1,000 | ~15 |
-
-적용 조건(하나라도 아니면 stock): 전 요청이 spec-verify 이고 드래프트가 꽉 참
-(`decode_query_len - 1`), FULL cudagraph 디스패치, 요청 패딩 없음, 프리필 없음,
-적응형 검증·DCP·PCP·PP·LoRA 없음. 설치 시 러너/빌더 파일 17개의 preimage 를
-고정하고(드리프트 → DISARM), 첫 적격 스텝에서 live 러너로 plan 을 만들며
-기하가 다르면 그 부팅은 stock.
-
-사다리(순서대로, 건너뛰기 없음):
-
-```bash
-# 1. 순수 논리
-python3 tests/test_logic.py                      # test_glm53_prep_fused_contracts
-
-# 2. srv4 새 컨테이너: stock 빌딩블록 vs fused 커널, 무작위 배치 60회 bit-exact
-IMAGE=glm53:sm121-fi618 bash probes/run_prep_fused_check.sh --trials 60
-
-# 3. 섀도 부팅: fused 먼저, 그 위에 stock, 버퍼 전부 diff — stock 이 진실
-EXTRA_ENV="VLLM_GLM53_PREP_FUSED=shadow" bash launchers/start-glm53-nvfp4-tp4.sh
-#    bench-tp4 1회 동안 [prep-fused] shadow ... drift=0 확인
-
-# 4. 브래킷: base -> cand -> base, C=1 step/s
-EXTRA_ENV="VLLM_GLM53_PREP_FUSED=1" bash launchers/start-glm53-nvfp4-tp4.sh
-```
-
-주의:
-
-- 부팅 로그의 `[prep-fused] installed mode=...` 와 `[prep-fused] plan built:` 가
-  판정 근거다. `preimage drift -> DISARM` 이나 `plan build failed` 가 보이면 그
-  부팅은 stock 이고, 그 브래킷 셀은 무효다.
-- 물리 기전 확인: 켠 부팅의 트레이스에서 준비 구간의 `Memcpy HtoD/DtoD` 와
-  `at::native::*` 커널이 사라지고 `_glm53_prep_fused_kernel` 하나가 남아야 한다.
-- 수치는 stock 과 bit-exact 가 계약이라 품질 게이트는 형식상 통과해야 하지만,
-  섀도 drift=0 없이는 arm 하지 않는다.
-- 이 모듈은 kpool tail 의 원형 슬롯 매핑을 **건드리지 않는다**(러너가 빌더에
-  `positions` 를 안 넘겨 그 매핑은 이 이미지에서 잠들어 있고, fused 는 현행
-  generic 매핑을 그대로 재현). 그 활성화는 C>=2 수치 변경이라 별도 브래킷.
-
 ---
 
 ## 순서와 근거
@@ -311,7 +258,6 @@ EXTRA_ENV="VLLM_GLM53_PREP_FUSED=1" bash launchers/start-glm53-nvfp4-tp4.sh
 4. **EXP-4 (bproj)** — 천장 0.9%, 최우선순위 아님. 다음 창으로 미뤄도 됨.
 5. **EXP-5 (프리필)** — 캡처 1회가 관문. KDA 스윕은 그 다음.
 6. **EXP-6 (메가커널)** — 프로브와 섀도가 먼저 수치·계약을 닫고 브래킷.
-7. **EXP-7 (준비 커널 통합)** — bit-exact 프로브 → 섀도 부팅 → 브래킷. 호스트 유휴 4~5 ms 가 표적이라 GPU 커널 축과 독립.
 
 ## 금지 (기존 판정 유지 — 재조사하지 않는다)
 

diff --git a/STEP_KERNEL_MAP.md b/STEP_KERNEL_MAP.md
@@ -215,54 +215,6 @@ self `copy_`는 ATen 쇼트서킷으로 커널 미발생. conv→recurrent 융
 상태 읽기가 지배라 0.1~0.2% — 채택 바 아래로 확정. 교훈: 인구조사의 그룹 합계는
 **소스를 읽기 전까지는 개선 여지가 아니라 의문점**이다.
 
-## 보충 분해 2 (2026-09-02 — 9월 1일 트레이스 재분석, 커널 밖)
-
-앞의 지도는 GPU 커널만 셌다. 같은 트레이스(229 스텝)를 **스트림 합집합**으로
-보면 스텝의 12%가 GPU 유휴이고, 그 위치가 하나다.
-
-| 구간 (스텝 72.1 ms, 프로파일러 하) | 유휴 | 무엇 |
-|---|---|---|
-| 1.6 ms 지점 | 0.6 ms | DtoH 회수 뒤 드래프터 그래프 발사 대기 |
-| 6.6~12.3 ms | **5.7 ms** | eager 입력 준비: aten 호출 1,028개, memcpy 45개, 1~3 us 커널 ~100개, 커널 사이 50~430 us 공백 (호스트가 발사를 못 따라감) |
-| 12.5~14.0 ms | **1.43 ms** | `cudaGraphLaunch` — 1,640 노드 타깃 그래프 제출 호스트 소요 (50스텝 중앙, p10 1.34 p90 1.59); 노드당 ~0.87 us |
-| 스텝 끝 | 0.8 ms | 샘플러 → 다음 스텝 전환 |
-| 합 | 8.9 ms | |
-
-프로파일러가 호스트를 부풀리므로 절대값은 과장이다. 프로파일러 없는 앵커는
-원장의 "디코드 중 GPU 93%" (같은 정의) → 실제 유휴 4~5 ms/스텝. 원장의
-"호스트 병목" 기각은 32 ms 잔차의 설명으로서 맞았고, 5 ms 짜리 항목으로는
-살아 있다. dflash 가 async scheduling 을 자동으로 끈다는 기록(#1455)과 준비
-구간 동안 GPU 가 완전히 비는 트레이스가 일치한다.
-
-준비 구간의 코드: `v1/worker/gpu/input_batch.py`·`block_table.py`(Triton 5개),
-`model_states/mamba_hybrid.py`, KV 그룹 7개(MLA+indexer 1, kpool tail 1,
-mamba 4, 드래프터 1)의 빌더 — GDN 빌더 4개가 각각 `to`·`sub`·`arange`·`index`×2·
-`copy_`×5 를 내고(트레이스의 4회 반복 패턴), indexer 빌더가 `floor_divide`×3·
-`diff`·fill·`_prepare_uniform_decode`·deep_gemm 스케줄을 낸다. 전부 이미지 파일이고
-오버레이돼 있지 않다. 접수 방식은 `overlay/modules/glm53_prep_fused`(EXP-7):
-파일을 덮지 않고 러너 메서드를 패치하며 preimage 를 고정한다.
-
-**그래프 안 글루 605개의 정체**(같은 트레이스, 위치 귀속):
-
-| 위치 | 개/스텝 | 무엇 | 처리 |
-|---|---|---|---|
-| 풀어텐션 11층 인덱서 | 275 | 꼬리 풀 강제 8개, seq_len 유도, fill, 확장·변환 주변 | 우리 `sparse_attn_indexer_kpool.py` — `VLLM_GLM53_KPOOL_FUSED_TOPK`(기본 off, 미측정)에 접기 |
-| KDA 34층 | 136 | conv 뒤 `reshape` 가 split 뷰를 q/k/v 세 번 복사 | MK-KDA 가 흡수 |
-| MoE 42층 | 42 | shared expert 덧셈 | osar copy-in 에 접기 |
-| 나머지 | ~150 | 드래프터 inductor 글루, 꼬리 native RMSNorm, eager memcpy | 작음 |
-
-커널당 점유는 CUPTI 로 길이 중앙 2.0 us + 그래프 안 간격 0.2 us — 원장의
-5.4 us 상한과 부합하고, 605개 전부가 2.7 ms(부풀린 값)다.
-
-**읽다가 나온 것**: (1) 인덱서의 fp32 head-gate `torch.mm(hidden.float(), _wp_fp32)`
-가 cuBLAS gemmSN 2블록 커널로 층당 86 us, 11층 0.95 ms/스텝(CUPTI) — 우리
-`glm53_prefill_fastpath.py:402` 소유, split-K 로 수 us 감. (2) 드래프터 fc 투영
-814 us(eager bf16, 5층 hidden cat, 168 MB 읽기, `ReplicatedLinear` 라 fp8 dense
-패턴 밖) 포함 드래프터 커널 합 ~3 ms(CUPTI) — 원장 D≈0 과 긴장, 직접 측정 전
-판단 보류. (3) `KpoolTailMetadataBuilder` 의 원형 tail 슬롯 매핑은 러너가
-`positions` 를 안 넘겨 이 이미지에서 잠들어 있다(generic 매핑 사용) —
-`glm53_tail_slot_persistent` 의 고정 버퍼도 그래서 무효; C>=2 수치 축.
-
 ## 재현
 
 ```bash