feat(glm53): #227 split-K head-gate 를 main 에 반영 (v2 브랜치로 머지된 것)

MEASUREMENTS.md

@@ -2121,3 +2121,31 @@ will be disabled"). 업스트림 main 도 동일(2026-09-02). 그런데 `DSparkS
 
 **미측정**: EXP-8(부팅 1회 + 브래킷). 수용률은 움직이면 안 된다(실행 시점만 바뀜).
 V2+async 는 `max_concurrent_batches`=2 라 KV in-flight 예약이 두 배 — KV 라인 확인.
+
+## ★인덱서 fp32 head-gate: cuBLAS 2블록 47 us → split-K 10 us (`glm53_indexer_gate_splitk`, 2026-09-02, srv4 오프라인)
+
+**대상**: `attention.py` `Indexer.forward` 의 `torch.mm(hidden_states.float(), self._wp_fp32)`
+— [M,4096]×[4096,16] fp32, 층당 1회 × 11층. fp32 인 이유는 코드 주석대로 bf16 게이트
+(~1e-2)가 근소 차 풀 순위를 뒤집기 때문. cuBLAS 는 이 형상에 `gemmSN` 2블록 커널을
+고른다(48 SM 중 2개).
+
+**실측**(GB10, CUDA 그래프 리플레이, `probes/indexer_gate_check.py`; 9월 1일 트레이스의
+86 us 는 CUPTI + 공유 경합):
+
+| M | stock `torch.mm` | split-K(8) | 경로 |
+|---|---|---|---|
+| 1 | 15.5 us | 8.5 us | split-K |
+| 8 (C=1) | 50.0 us | 9.9 us | split-K |
+| 16 (C=2) | 50.1 us | 11.7 us | split-K |
+| 32 (C=4) | 15.9 us | 17.4 us | `torch.mm` 유지 |
+
+M=32 부터 cuBLAS 가 다른 커널을 고르며 빨라지므로 M<=16 만 라우팅. 대안 비교(M=8):
+`F.linear` NT 레이아웃 32 us, mul+sum 20.7 us, split-K(16) 6.2 us — 8 분할이 원자 경합과
+프로그램 수의 균형점이라 채택.
+
+**수치**: 양쪽 fp32 누적, 합산 순서만 다름. 300회/2,480행 max|diff| 2.4e-6 절대,
+6.7e-7/행 최대, top-1 뒤집힘 0, top-4 집합 변화 0. bit-exact 아님 → 품질 브래킷 대상.
+
+**천장**: 11 × 40 us ≈ 0.44 ms/스텝 = C=1 스텝의 ~0.65% (<1%, 단독 부팅 불가). EXP-7/8
+부팅에 얹어 잰다. `VLLM_GLM53_FUSED_K_GATE=1` 팔의 융합 인덱서도 같은 헬퍼를 타므로
+두 팔이 어긋나지 않는다. 기본 0 = stock 과 같은 `torch.mm` 호출. RUNBOOK EXP-9.

RUNBOOK_KERNEL_CAMPAIGN2.md

@@ -345,6 +345,33 @@ VLLM_GLM53_ASYNC_DFLASH=1 bash launchers/start-glm53-nvfp4-tp4.sh   # cand (프
   dflash 의 async off 모두 손상, 원인은 LibertAIDAI 가중치). 이 실험은 그 판정을
   재론하지 않는다.
 
+## EXP-9 — 인덱서 fp32 head-gate 를 split-K 로 (`glm53_indexer_gate_splitk`, 2026-09-02 추가)
+
+`Indexer.forward` 의 `weights = torch.mm(hidden_states.float(), self._wp_fp32)`
+([M,4096]×[4096,16] fp32, 층당 1회 × 11) 을 cuBLAS 가 2블록 `gemmSN` 커널로 답한다:
+유휴 GB10 에서 47 us, 9월 1일 트레이스(CUPTI)에서 86 us. 48 SM 에 블록 2개가 문제의
+전부라 (행, K-슬라이스 512) 마다 프로그램 하나를 띄우고 fp32 atomic 으로 모으는
+split-K Triton 커널이 같은 곱을 10 us 에 낸다. **M<=16(디코드) 만** 이 경로,
+나머지(프리필, C>=3 verify)는 stock `torch.mm` 그대로.
+
+**수치**: 양쪽 다 fp32 누적, 합산 순서만 다르다 — bit-exact 가 아니다. 오프라인
+300회/2,480행: max|diff| 2.4e-6, 행 최대 대비 6.7e-7, top-1 뒤집힘 0, top-4 집합 변화 0
+(`probes/indexer_gate_check.py`). bf16 게이트가 순위를 뒤집는 오차(1e-2)보다 네 자릿수
+아래지만 품질 브래킷은 필요하다.
+
+**천장**: 11층 × ~40 us ≈ 0.44 ms/스텝, C=1 66 ms 스텝의 **~0.65%** — 원장 규칙(1% 미만은
+단독 부팅 가치 없음)에 걸린다. EXP-7/EXP-8 과 독립이라 그 부팅에 얹어 같이 재는
+용도. 단독 부팅 금지.
+
+```bash
+# 프로필 선언 키: caller env. 트레이스에서 gemmSN 11개가 _gate_splitk_kernel 11개로
+# 바뀐 것이 켜진 증거 (부팅 로그 줄은 없다 — 그래프 안에서 층마다 호출).
+VLLM_GLM53_INDEXER_GATE_SPLITK=1 bash launchers/start-glm53-nvfp4-tp4.sh
+```
+
+- 게이트: 품질 9/9, 한국어 0/16, C=1 step/s 브래킷 base→cand→base (얹은 부팅의 것과 공유).
+- 롤백 = env 한 줄. 기본 0 = stock 과 동일한 `torch.mm` 호출.
+
 ---
 
 ## 순서와 근거

STEP_KERNEL_MAP.md

@@ -256,7 +256,8 @@ mamba 4, 드래프터 1)의 빌더 — GDN 빌더 4개가 각각 `to`·`sub`·`a
 
 **읽다가 나온 것**: (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 투영
+`glm53_prefill_fastpath.py:402` 소유, split-K 로 수 us 감 → `glm53_indexer_gate_splitk`
+(EXP-9, opt-in, 오프라인 50 → 10 us, ~0.65%/스텝). (2) 드래프터 fc 투영
 814 us(eager bf16, 5층 hidden cat, 168 MB 읽기, `ReplicatedLinear` 라 fp8 dense
 패턴 밖) 포함 드래프터 커널 합 ~3 ms(CUPTI) — 원장 D≈0 과 긴장, 직접 측정 전
 판단 보류. (3) `KpoolTailMetadataBuilder` 의 원형 tail 슬롯 매핑은 러너가

build/glm53/glm53_indexer_gate.py

@@ -0,0 +1,82 @@
+# SPDX-License-Identifier: Apache-2.0
+"""deneb fork: the sparse indexer's fp32 head-gate projection as a split-K
+Triton kernel (VLLM_GLM53_INDEXER_GATE_SPLITK).
+
+`Indexer.forward` computes `weights = torch.mm(hidden_states.float(),
+self._wp_fp32)` -- an [M, 4096] x [4096, 16] fp32 product per full-attention
+layer, kept in fp32 on purpose (bf16 head-gates flip near-tie pool rankings).
+cuBLAS answers that shape with a two-block `gemmSN` kernel: 47 us on an idle
+GB10 (86 us under CUPTI in the 2026-09-01 serving trace), eleven times per
+decode step, for 256 KB of weights. A split-K kernel that hands each
+(row, K-slice) to one program and reduces with fp32 atomics runs the same
+product in 7 us. Both accumulate in fp32; only the summation order differs
+(measured max |diff| 3e-5 on values of magnitude ~64, i.e. ~5e-7 relative,
+0 top-1 rank flips over the offline trials) -- not bit-exact, so this stays
+an opt-in behind a numerics bracket.
+
+At M > 16 (C >= 3 verify batches, prefill) cuBLAS is already fast, so the
+kernel is used only for M <= 16; larger M keeps torch.mm.
+"""
+from __future__ import annotations
+
+import os
+
+import torch
+
+from vllm.triton_utils import tl, triton
+
+ENV = "VLLM_GLM53_INDEXER_GATE_SPLITK"
+MAX_M = 16
+_SPLIT = 8
+_BLOCK_K = 128
+
+
+def gate_splitk_enabled() -> bool:
+    """Exact opt-in: only the string "1" arms; anything else is stock."""
+    return os.environ.get(ENV, "").strip() == "1"
+
+
+@triton.jit
+def _gate_splitk_kernel(x_ptr, w_ptr, out_ptr, K, N, sxm, swk, som,
+                        SPLIT: tl.constexpr, BLOCK_K: tl.constexpr, BN: tl.constexpr):
+    m = tl.program_id(0)
+    s = tl.program_id(1)
+    kper = K // SPLIT
+    k0 = s * kper
+    offs_n = tl.arange(0, BN)
+    nmask = offs_n < N
+    acc = tl.zeros([BN], dtype=tl.float32)
+    for k in range(0, kper, BLOCK_K):
+        offs_k = k0 + k + tl.arange(0, BLOCK_K)
+        xv = tl.load(x_ptr + m * sxm + offs_k)
+        wv = tl.load(w_ptr + offs_k[:, None] * swk + offs_n[None, :], mask=nmask[None, :], other=0.0)
+        acc += tl.sum(xv[:, None] * wv, axis=0)
+    tl.atomic_add(out_ptr + m * som + offs_n, acc, mask=nmask)
+
+
+def head_gate_splitk(x: torch.Tensor, w: torch.Tensor) -> torch.Tensor:
+    """fp32 [M, K] @ [K, N] for N <= 16, K a multiple of SPLIT*BLOCK_K.
+
+    x may be bf16 (cast to fp32 here, as the stock `.float()` does); w is the
+    stock `_wp_fp32` ([K, N], fp32, contiguous)."""
+    xf = x.float()
+    M, K = xf.shape
+    N = w.shape[1]
+    out = torch.zeros(M, N, device=x.device, dtype=torch.float32)
+    _gate_splitk_kernel[(M, _SPLIT)](
+        xf, w, out, K, N, xf.stride(0), w.stride(0), out.stride(0),
+        SPLIT=_SPLIT, BLOCK_K=_BLOCK_K, BN=16, num_warps=4)
+    return out
+
+
+def splitk_applicable(x: torch.Tensor, w: torch.Tensor) -> bool:
+    return (x.shape[0] <= MAX_M and w.dtype == torch.float32 and w.is_contiguous()
+            and w.shape[1] <= 16 and w.shape[0] % (_SPLIT * _BLOCK_K) == 0)
+
+
+def head_gate(x: torch.Tensor, w: torch.Tensor) -> torch.Tensor:
+    """The stock `torch.mm(x.float(), w)` unless the knob is on and the shape
+    is the small-M decode one."""
+    if gate_splitk_enabled() and splitk_applicable(x, w):
+        return head_gate_splitk(x, w)
+    return torch.mm(x.float(), w)

build/glm53/glm53_prefill_fastpath.py

@@ -34,6 +34,23 @@
 )
 from .ops.kpool_compress import fwht128_quant_fp8
 
+# deneb fork (glm53_indexer_gate_splitk): the split-K helper when that module
+# is mounted, else the stock fp32 torch.mm. Resolved once, on first call, so
+# the fused indexer forward stays loadable without the sibling file.
+_HEAD_GATE = None
+
+
+def _glm53_head_gate(x, w):
+    global _HEAD_GATE
+    if _HEAD_GATE is None:
+        try:
+            from vllm.models.glm5next.nvidia.glm53_indexer_gate import head_gate as fn
+        except ImportError:
+            def fn(x, w):
+                return torch.mm(x.float(), w)
+        _HEAD_GATE = fn
+    return _HEAD_GATE(x, w)
+
 logger = init_logger(__name__)
 
 _GLM53_SM121_MLA_PREFILL_ENV = "VLLM_GLM53_SM121_MLA_PREFILL"
@@ -399,7 +416,7 @@ def _glm53_fused_indexer_forward(
         self._wp_fp32 = (
             k_weight.data[self.head_dim :, :].t().contiguous().float()
         )
-    weights = torch.mm(hidden_states.float(), self._wp_fp32)
+    weights = _glm53_head_gate(hidden_states, self._wp_fp32)  # deneb fork (glm53_indexer_gate_splitk)
 
     k = _fused_indexer_k_norm(
         k,