perf(osar,mhc,tools): copy/reduce 16B 벡터화+캐시정책, MHC 패스설정 노브, census 인접성 모드 (EXP-17/18/19) - #301
Conversation
…접성 모드 (EXP-17/18/19) - osar k_oneshot: copy/reduce 를 16B 렌으로(R1) + __stwt/__ldcs/__ldg 캐시정책과 bfloat162 팩 변환(R2). 요소별 연산 순서·rn 불변이라 스칼라 원조와 비트 동일, 펜스·done_ctr·프리페치 기계장치 불별. 미정렬 입력은 shim 이 NCCL 로 라우팅. L2 위생(~26MB/스텝 churn 제거)은 #100 중립 판정이 예측 못 하는 부류 — 브래킷 판정. 이미지 환경 컴파일+링크 검증(md5=b0275622). - glm53_mhc_tilelang: VLLM_GLM53_MHC_PASSES(tma|ws|none, 기본 off) — 이미지가 TMA lowering·warp specialization 을 끈 채 놔둔 것의 A/B 노브. unset/무효값 = 스톡과 동일. 부팅 로그에 engine-confirmed 라인. - probes/mhc_glm53_bench: --passes(4콤보 × 독립 컨테이너, 참조 저장/대조 rel≤1e-4)·--hcweight(bf16 가중치 수치 실험, 전위 오류와 양자화 오차 분리) 모드. - census.py: --after REGEX 인접성 모드(스트리밍 판독기 위에). 08-31 트레이스로 osar 에필로그 축 기각 확정(원장 기록) 및 vllm-mhc PDL 봉인 실측 기록. - 런북 EXP-17/18/19 추가, 순서·금지 갱신. - 참고: tests 의 고아 모듈 하이진 실패(glm53_v2_hard_constraint_guard)는 main 순정에서도 재현되는 선행 결함으로 이 PR 이 고치지 않는다.
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
Bugbot couldn't run - usage limit reachedBugbot is counted against Cursor usage for this user or team, and this run hit a usage or spend limit. A user or team admin can review and increase usage limits in the Cursor dashboard. (requestId: serverGenReqId_75ce131c-0f46-4cd7-bf54-021a49ee37da) |
There was a problem hiding this comment.
🔵 Needs a closer look
It changes a decode hot-path CUDA allreduce kernel and introduces compiler pass-config knobs whose correctness/perf impact depends on hardware validation beyond what can be confirmed from the diff alone.
Pull request overview
This PR ports three performance/analysis efforts onto the current main baseline: (1) a 16-byte vectorized + cache-policy rewrite of the osar one-shot allreduce copy/reduce hot path, (2) an opt-in VLLM_GLM53_MHC_PASSES knob to A/B TileLang pass-config choices across all MHC kernels, and (3) a new census.py --after “adjacency” mode to empirically validate “what runs right after X” claims in traces.
Changes:
- Vectorize
k_oneshotcopy/reduce to 16B (8×bf16) lanes and apply cache hints (__ldg/__stwt/__ldcs) with explicit 16B alignment routing/guards. - Add import-time
VLLM_GLM53_MHC_PASSESparsing + wiring, plus probe harness support (--passes,--hcweight) and a contract test. - Add
census.py --after(+--depth) adjacency counting and update runbook/measurements/docs for EXP-17/18/19.
File summaries
| File | Description |
|---|---|
| tests/test_logic.py | Adds test_mhc_passes_knob and wires it into the full test run. |
| RUNBOOK_KERNEL_CAMPAIGN2.md | Documents EXP-17/18/19 procedures and ordering/forbidden axes updates. |
| probes/run_mhc_glm53_bench.sh | Adds --passes sweep orchestration across isolated containers. |
| probes/mhc_glm53_bench.py | Implements --passes reference gating and --hcweight bf16-weight probe. |
| overlay/modules/tp_oneshot_ar/README.md | Documents the 16B vector lane + cache policy rationale and constraints. |
| overlay/modules/tp_oneshot_ar/dsv4_oneshot_shim.py | Routes misaligned tensors to NCCL via a 16B alignment eligibility check. |
| overlay/modules/tp_oneshot_ar/dsv4_oneshot_ar.cu | Implements 16B vectorized copy/reduce and adds a 16B alignment TORCH_CHECK. |
| overlay/modules/glm53_mhc_tilelang/tilelang_kernels.py | Adds VLLM_GLM53_MHC_PASSES parsing + pass_configs mutation before first @tilelang.jit. |
| overlay/modules/glm53_mhc_tilelang/README.md | Documents the new pass-config knob and probe workflow. |
| MEASUREMENTS.md | Records adjacency-based rejection of an osar-epilogue fusion idea and PDL sealing note. |
| census.py | Adds --after adjacency mode and basic arg parsing for --depth. |
| build/glm53/tilelang_kernels.py | Mirrors the overlay TileLang pass-config knob into the composed build copy. |
| build/glm53/dsv4_oneshot_shim.py | Mirrors 16B alignment eligibility check into composed build copy. |
| build/glm53/dsv4_oneshot_ar.cu | Mirrors 16B vectorized copy/reduce + alignment guard into composed build copy. |
| build/dsv4/dsv4_oneshot_shim.py | Mirrors 16B alignment eligibility check into composed build copy. |
| build/dsv4/dsv4_oneshot_ar.cu | Mirrors 16B vectorized copy/reduce + alignment guard into composed build copy. |
Review details
- Files reviewed: 16/16 changed files
- Comments generated: 3
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
| if _a == "--after" and _i + 1 < len(_args): | ||
| after = _args[_i + 1]; _i += 2 | ||
| elif _a == "--depth" and _i + 1 < len(_args): | ||
| depth = int(_args[_i + 1]); _i += 2 | ||
| elif _path is None: |
| frozen = getattr(kernels, "_DENEB_MHC_PASSES", None) | ||
| if frozen is None: | ||
| raise RuntimeError( | ||
| "VLLM_GLM53_MHC_PASSES was not set at import; the wrapper must " | ||
| "freeze a combo per process (stock combo saves the reference)" | ||
| ) |
| src = open(_overlay_source("overlay/tilelang_kernels.py"), | ||
| encoding="utf-8").read() |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: ec999f2549
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| line += (f" rel_err(max vs stock-passes ref)={max(errs):.2e}" | ||
| + ("" if max(errs) <= 1e-4 else " ! MISMATCH")) |
There was a problem hiding this comment.
When a TMA/WS combination produces rel_err > 1e-4, this branch only adds ! MISMATCH to the output and still returns successfully. The --passes wrapper recognizes compile/divergence failures only when run_probe exits nonzero, so it never prints the advertised failed verdict for numerical divergence and the sweep can treat a broken combination as viable. Raise an error or otherwise return a failing status after detecting any mismatch.
Useful? React with 👍 / 👎.
| op() | ||
| torch.cuda.synchronize() | ||
| line += f" onepass={bench_us(op):7.1f}us" |
There was a problem hiding this comment.
When a compiler-pass combination miscompiles only mhc_onepass_tilelang, the pair comparison remains clean because the result of op() is discarded and only its runtime is recorded. The wrapper explicitly enables ONEPASS for every combination, so EXP-17 can select a fast but numerically invalid one-pass result even though the default-pass one-pass check says nothing about TMA/WS-generated code. Compare the returned one-pass tensors against the stock reference before reporting the combination as a candidate.
Useful? React with 👍 / 👎.
| const uint4 *src4 = reinterpret_cast<const uint4 *>(src); | ||
| uint4 *tx4 = reinterpret_cast<uint4 *>(c->tx[slot]); | ||
| const int nv = n >> 3; | ||
| uint4 mine[VECITER]; |
There was a problem hiding this comment.
When the supported VLLM_DSV4_OSAR_MAXEL override is raised to its 8,388,608-element ceiling, VECITER becomes 86 and this array alone requires 344 32-bit register values per thread, exceeding CUDA's 255-register per-thread limit before accounting for the rest of the kernel. The compiler must therefore spill the stash to local memory, making all 12,288 threads round-trip the cached source lanes through device memory across the peer wait and defeating the optimization precisely for the larger-message sweep the override exists to run. Keep the register stash bounded and reload or otherwise handle additional grid-stride iterations without a MAXEL-sized per-thread array.
Useful? React with 👍 / 👎.
| # shape checks above: every rank runs the same producer code, and the | ||
| # caching allocator's blocks are 512B-aligned, so alignment is a | ||
| # property of the (identical) view chain, not of allocator state. | ||
| and t.data_ptr() % 16 == 0 |
There was a problem hiding this comment.
If ranks provide the same legal contiguous BF16 collective shape through views with different storage offsets, this pointer test can pass on one rank and fail on another; allocator base alignment does not make rank-dependent view offsets identical. The aligned rank then enters OSAR and waits for peer RDMA traffic while the unaligned rank returns None and enters NCCL through cuda_communicator.py, splitting the collective and deadlocking exactly as the shim's post-commit policy is meant to prevent. Handle unaligned input inside the OSAR path, such as with a scalar prefix/path, or perform an all-rank eligibility decision instead of falling back rank-locally.
Useful? React with 👍 / 👎.
|
정정: "main 순정에서도 재현되는 선행 결함"은 틀린 주장이었다. #301 머지 후 확인 결과, 고아 모듈 하이진 실패( 부수 정정: 리뷰 중 "스위트가 10408행에서 중단돼 뒤 112개 테스트가 미실행으로 머지됐다"고 판단했던 것도 같은 착오 — 중단은 로컬 잔해 때문이었고, 잔해 제거 후 미실행분 포함 전량 통과로 #301 의 회귀 없음을 재확认했다. |
지도(#300) 다음 칸: 문서가 말하는 구조·기본값·흡수량을 지금 상태와 맞췄다. 코드 변경 0. - README: 구성 표가 "overlay/ 단일 manifest + 16파일" 이었다 -- 루트 overlay/manifest.tsv 는 없어졌고 모듈(42) × 프로필(dsv4 18 / glm53 25 / qwen38 1)을 compose-overlays.sh 가 build/<profile>/manifest.tsv 로 렌더한다. 배포 레이아웃에 glm53 경로(~/overlays/glm53, 4노드 동일)와 A/B base 가 MEGAKERNEL=0 을 명시하는 이유, 재배포가 브래킷 대상이라는 규율 추가. tools/ 절 신설 -- census.py(소유권·인접성), trace_step_composition, trace_step_tail, trace_common 과 각 도구의 함정, 플릿 CPU 규율. - profiles/README: 모듈 × 프로필 표가 실제 프로필과 어긋나 있었다 -- 쓰이는 모듈 10개 누락, 삭제된 glm53_async_dflash 잔존. 기계적으로 일치시켰고 (누락 0 확인), glm53 상태(bring-up, blocked -> 커널 캠페인 · 매일 부팅)와 기본 노브 행을 넣었다. "이식 가능한 모듈 넷" 은 매니페스트 전 행이 absent 인 모듈 열 개로 정정하되, 형식(계약)과 내용(모델 무관)은 다른 명제라고 단서를 달았다. - MEASUREMENTS: 두 스택(dsv4 프로덕션 / glm53 캠페인)이 한 원장에 섞여 있다는 안내와 기준선 제목의 스택 명시. glm53 기준선 블록 신설 -- 28차 §8 부팅의 요약이다. 쓰다가 오귀속을 잡았다: 그 20:36 부팅에는 MK_PDL 이 없었다 (PR #290 은 20:57 머지) -> "현행 프로필 = 이 기준선 + PDL 한 칸, 그 칸의 종단 수치는 아직 없다" 로 적었고 다른 두 문서의 출처도 분리했다. - RUNBOOK EXP-6 · 메가커널 README: 세그먼트 흡수량이 예측치(MHC 179->45)로만 있었다. 실측 열 추가 -- MHC 179->89(+잔여 7), quant+GEMM 376->187, KDA 미무장, MLA 22->11. 천장 4.9 ms 는 낙관으로 판정: 커널 -304 발인데 시간이 준 자리는 MHC -0.5 와 양자화 -1.5 뿐이다. - KERNEL_CAMPAIGN: 1차 캠페인의 두 결론("overlay-ownable 타깃 없음", "디코드는 논외")이 뒤집힌 경위와, 여전히 유효한 부분(MoE 47% 는 대역폭 바닥)을 후속 블록으로. - RUNTIME_GUARD_AUDIT: manifest 경로 각주(감사 결론은 이미지에 대한 것이라 오버레이가 늘어도 무효화되지 않는다). #301 반영: census --after 인접성 모드를 README 도구 표에, MHC_PASSES 노브를 프로필 표에, TileLang mhc PDL 봉인과 우리 MK PDL 이 다른 축이라는 포인터를 지도에 넣었다. 병합된 census.py 로 지도의 수치(1,582 / 402 / 25.4%)를 재확인. 테스트: 44,627 checks OK. Co-authored-by: choiceoh <choiceoh@srv4.tail7fec17.ts.net> Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
#301 머지 후 대조 확인 — 145커밋 사이에 생긴 변화가 세 축의 전제를 갈랐다: - EXP-18(osar): #99 시대 copy+reduce 11.8µs 전제는 낡음(현 45.5µs/콜, 고정비 ~22.7µs 지배) → R1 천장 ~0.2-0.3ms 로 하향. R2 의 판정 근거는 L2 위생×EXP-13 프리페치 시너지(osar churn 26MB/스텝이 프리페치 예산 12MB 와 직접 경쟁하는 것을 stwt/ldcs 로 제거)로 승격. - EXP-17(MHC 패스설정): 무장 디코드가 mk_mhc 로 TileLang mhc 179발을 대체(185→14) — 노브의 주 전장은 프리필 big_fuse + 잔여 14발로 축소. - EXP-19(hc bf16): ONEPASS 미채택 + mk_mhc 초월 — 폴백 참조 측정으로 강등, 프로덉션 경로의 같은 질문은 mk_mhc 쪽 별도 축임을 명시.
#301 머지 후 대조 확인 — 145커밋 사이에 생긴 변화가 세 축의 전제를 갈랐다: - EXP-18(osar): #99 시대 copy+reduce 11.8µs 전제는 낡음(현 45.5µs/콜, 고정비 ~22.7µs 지배) → R1 천장 ~0.2-0.3ms 로 하향. R2 의 판정 근거는 L2 위생×EXP-13 프리페치 시너지(osar churn 26MB/스텝이 프리페치 예산 12MB 와 직접 경쟁하는 것을 stwt/ldcs 로 제거)로 승격. - EXP-17(MHC 패스설정): 무장 디코드가 mk_mhc 로 TileLang mhc 179발을 대체(185→14) — 노브의 주 전장은 프리필 big_fuse + 잔여 14발로 축소. - EXP-19(hc bf16): ONEPASS 미채택 + mk_mhc 초월 — 폴백 참조 측정으로 강등, 프로덉션 경로의 같은 질문은 mk_mhc 쪽 별도 축임을 명시.
…31차로 PR #304 는 #301(672430e)에서 갈라져 나왔고, 그 뒤 main 이 #300·#302·#303·#305· #306·#307·#308(#309)로 앞서 나가면서 충돌 상태(dirty)가 됐다. 코드 파일은 겹치지 않았고(KDA 원패스·kpool 은 #304 단독), 겹친 것은 문서·프로필·테스트 넷이다. **RUNBOOK 충돌 2건** (1) 새 EXP 절의 삽입 위치: main 이 EXP-16 뒤에 EXP-22(로컬 양자화)·EXP-21(v2 레인)을 넣었고 #304 는 같은 자리에 EXP-20(미세 융합 묶음 2)을 넣었다. 경쟁이 아니라 가산이므로 셋 다 두되 기존 배열(최신 먼저)을 따라 22 → 21 → 20 순으로 둔다. (2) "순서와 근거" 절: main 은 평평한 번호 목록을 지우고 자기가 새로 만든 상태 표로 보내는 안내문으로 바꿨고(`채택·기각·닫힘은 위 상태 표에 있다`), #304 는 그 목록에 18번(EXP-20)을 덧붙였다. **main 의 재구성을 취하되**, 사라질 뻔한 #304 의 항목을 main 의 "다음 부팅 창에서" 목록에 8번으로 옮겨 실험이 목록과 함께 증발하지 않게 했다. **원장 번호 충돌 (자동 병합이 못 보는 것)** #304 는 자기 측정을 `★29차` 로 적었는데, 그 사이 main 에 다른 29차(메가커널 로컬 양자화, #307)와 30차(비상주 v2 레인, #305)가 먼저 들어왔다. 텍스트로는 충돌하지 않아 병합 결과에 **서로 다른 29차 둘**이 남았다. 28차 항목이 세운 선례("원장 번호 27 은 PR #290 이 쓰고 있어 28 로 적는다")대로 다음 빈 번호인 **31차** 로 옮겼고, 왜 옮겼는지 절 머리에 한 줄 남겼다. 측정값 자체는 손대지 않았다. 같은 번호를 가리키던 참조 6곳을 함께 고쳤다: MEASUREMENTS 절 제목, RUNBOOK EXP-20 표의 기각 축, `glm53_kda_onepass/README.md`, `moe_gate_sm121/README.md`, `profiles/glm53.env`, `tests/test_logic.py`. main 이 자기 29차를 가리키는 참조 (`VLLM_GLM53_MK_LOCALQ` 주석)는 그대로 두었다. **자동 병합분 검증** - `profiles/glm53.env`: main 의 `VLLM_GLM53_MK_LOCALQ` 와 #304 의 세 노브 (`KDA_DUAL_GEMM`·`KDA_ONEPASS`·`KPOOL_UPDATE_DIRECT_POS`) 전부 기본 0 으로 잔존. `MODULES=` 목록은 main 의 항목 하나도 잃지 않고 `glm53_kda_onepass` 만 늘었다. - overlay ↔ build 사본 3쌍(`glm53_kda_onepass.py`, `glm5next_kda.py`, `sparse_attn_indexer_kpool.py`) 모두 동일. - `tests/test_logic.py`: 양쪽 테스트 공존, `micro-fusion bundle 2 contracts` 통과. 검증: `tests/test_logic.py` all OK (6106 checks; 이 호스트엔 torch 부재로 일부 SKIP). 충돌 마커 0, 변경된 .py 전부 파싱 OK, `run_micro_fusion_check.sh` bash -n OK. GPU 검증(프로브 VERDICT·트레이스 물리확인)은 하지 않았다 — 이 환경에 GPU 가 없다. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01JTghtTjyfT23vuhZwHzBSM
Summary
세션 3종 작업을 현재 main(=#300) 위로 이식한 것 — 로컬 브랜치가 main 에
145+ 커밋 뒤처진 베이스(osar 프리페치·메가커널 캠페인 이전)였고, 낡은 베이스
그대로 머지하면 선행 작업을 되돌리므로(#230 전례) main 기반 재이식 후 전량
재검증했다.
Changes
k_oneshotR1+R2: copy/reduce 16B(8×bf16) 벡터 렌 +__stwt(tx쓰기-스루)·
__ldcs(rx 스트리밍)·__ldg캐시정책 +__bfloat1622float2팩 변환. 요소별 연산 순서·rn 불변 = 스칼라 원조와 비트 동일. 펜스·done_ctr·
48블록·EXP-13 프리페치 기계장치 불변. 미정렬 입력은 shim
_eligible이 NCCL라우팅. L2 위생(~100 콜렉티브 × ~256KB = L2 크기급 churn 제거)이 R2 본체 —
perf(osar): 스핀이 기다리는 대상과 경쟁하지 않게 — 백오프 + 도착한 플래그 재독 제거 #100 중립 판정이 예측 못 하는 부류, C=1 브래킷이 판정(EXP-18).
VLLM_GLM53_MHC_PASSES(기본 off): 이미지가 모든 mhc 커널에서 끈TMA lowering·warp specialization의 A/B 노브. unset/무효값 = 스톡 바이트
동일, 부팅 로그 engine-confirmed 라인(EXP-17).
--passes(4콤보 × 독립 컨테이너, 참조 저장/rel≤1e-4대조, 조합 실패도 판정으로 기록)·
--hcweight(bf16 가중치, 전위 오류게이트와 양자화 오차 보고 분리)(EXP-19).
census.py --after: 인접성 모드(스트리밍 판독기 위에 이식, 08-31트레이스에서 동일 수치 재생산으로 검증). 이 도구로 osar 에필로그 축 기각과
vllm-mhc PDL 봉인을 확정 — 원장 2건 기록(스코핑 명시 포함).
Testing
tests/test_logic.py신규test_mhc_passes_knob(파서+배선) 포함대상 테스트 전부 통과 — 단 main 순정에서도 실패하는 선행 결함 1건
(고아 모듈
glm53_v2_hard_constraint_guard하이진)은 이 PR 이 고치지않는다(클린 체크아웃에서 재현 확인).
glm53:v13-b12x, sm_121a + libverbs,torch.utils.cpp_extension) 컴파일+링크 통과 —md5=b0275622,kernels=1. 부팅 로그 지문과 일치.
census.py --after k_oneshot실트레이스 스모크: 328,164 이벤트,102.0/스텝, +1 어드저시전 mhc_fused 82.4% — 이식 전 수치와 동일.
실행 대기.
Note
Medium Risk
Changes the custom one-shot allreduce kernel on the decode hot path, though outputs are claimed bit-exact and misaligned inputs fall back to NCCL; MHC pass overrides are opt-in and default off.
Overview
osar
k_oneshot(EXP-18) rewrites the copy/reduce phases to 16-byte bf16 vectors: copy values are stashed in registers for reduce (no secondsrcread), with__ldg/__stwt/__ldcson src/tx/rx and packed bf16↔fp32 converts while keeping element order and rounding unchanged (bit-identical to the scalar path). 16-byte misaligned tensors are excluded in the Python shim_eligibleand routed to NCCL; the CUDA entry point adds a matching alignment guard. Same changes are mirrored across dsv4/glm53 build and overlay.cu/shim copies.Adds import-time
VLLM_GLM53_MHC_PASSES(tma/ws/ combos /none) to optionally re-enable TileLang TMA lowering and warp specialization for all MHC kernels, with boot logging andtest_mhc_passes_knob.probes/run_mhc_glm53_bench.sh --passesruns four isolated-container combos with a saved stock reference and rel≤1e-4 gate;--hcweightprobes bf16 HC weights on onepass (transcription vs quantization errors separated).census.py --after REGEX [--depth N]counts the next kernels on the same CUDA stream after a match—used to document rejecting osar epilogue fusion and sealing vllm-mhc PDL on GB10. MEASUREMENTS and RUNBOOK add EXP-17/18/19 and updated ordering/forbidden axes.Reviewed by Cursor Bugbot for commit ec999f2. Bugbot is set up for automated code reviews on this repo. Configure here.