perf(osar): 스핀이 기다리는 대상과 경쟁하지 않게 — 백오프 + 도착한 플래그 재독 제거 - #100
Merged
Conversation
가정: 미귀속 32ms 는 콜렉티브 지연이다. #99 의 위상 계측이 이 가정을 곧 확정 하거나 기각하지만, 가정이 맞다면 손댈 곳은 커널 개수가 아니라 대기 방식이다. 두 스핀 루프가 읽는 플래그는 전부 피어가 쓴다 -- ack_seq 는 프록시가, rxf 는 피어 NIC 가 RDMA 로. 그 플래그를 최고 속도로 읽으면 그 쓰기가 가시화되는 바로 그 경로에 트래픽을 계속 얹는 셈이다. 기다리는 대상과 경쟁한다. - 백오프: SPIN_HOT(8) 회는 그대로 뜨겁게 돌고(짧은 대기는 대가를 안 낸다), 이후 __nanosleep 을 128ns 에서 두 배씩 4096ns 상한까지. 상한이 이 변경이 더할 수 있는 지연을 묶는다 -- 콜렉티브당 최악 4us 이고, 의심하는 대기는 그보다 두 자릿수 크다. - 링 가드에서 tx_seq 를 루프 밖으로. 이 런치는 아직 게시 전이고 같은 스트림의 직전 런치는 이미 회수됐으므로 스핀 중 움직일 수 없다. 재독할 이유가 있는 것은 프록시가 올리는 ack_seq 뿐이라 루프 트래픽이 절반이 된다. - 피어 대기에서 이미 도착한 플래그를 다시 읽지 않는다. 옛 형태는 매 패스마다 세 개를 전부 재평가했다 -- 답이 바뀔 수 없는 것까지. 도착 여부를 기억하고 아직 없는 첫 번째만 읽는다. 단축평가 모양은 그대로다. 정확성은 안 건드린다. 대기 조건도, 펜스 위치도, 카운터 불변식도 그대로이고 바뀐 것은 "얼마나 자주, 무엇을 읽는가" 뿐이다. 측정 순서 주의: #99(계측만) 부팅으로 위상 분해를 먼저 받고 그 다음 이걸 얹어야 한다. 둘을 한 부팅에 섞으면 wait 이 줄어든 것이 이 변경 덕인지 원래 그랬는지 못 가른다. tests/test_logic.py 669 검사 통과. 컴파일 미검증 -- 부팅 전 컴파일 검사 필수.
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_933eca7c-ff4b-4d2f-8664-694b7ac6fe57) |
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. |
![cursor[bot]](https://avatars.githubusercontent.com/in/1210556?s=60&v=4){kind=small}
There was a problem hiding this comment.
Left a non-blocking comment; did not approve. Cursor Bugbot was present but completed as skipped (usage/spend limit), so human review is needed. No reviewers were assigned (no independent reviewers besides the author).
Sent by Cursor Approval Agent: Pull Request Router and Approver
choiceoh
added a commit
that referenced
this pull request
Aug 31, 2026
C=1 디코드 66.0ms 를 두 축(스펙 on/off, 커널 개수)에서 갈랐다. 오늘 이전까지 이 스텝에 붙은 설명은 전부 잔차 귀속이었고 네 번 뒤집혔다. 원장에 실측만 남긴다. - 스텝 예산: 전문가 가중치 읽기 41.0ms(62%, 대역폭 91% — 짤 것 없음) + 나머지 25.0ms - 드래프터 = 0: SPEC=0 이 31.3ms, k=7 이 66.1ms. 차 34.7ms 를 전문가 증가분 0.706x(58.1-8)=35.4ms 가 거의 전부 설명한다. D = -0.7ms. 부수로 0.706ms/전문가 기울기가 완전히 다른 축에서 독립 검증됐다 - 커널 가격 실측 5.4us (앞선 상한 15.4us 의 1/3). 잔여 전체를 없애도 12.8%, 우리 소유는 1.1% — -42 커널짜리 PR 은 CV 아래라 단독 측정 불가 - osar 위상: wait 이 AR 시간의 88%. #100 이 그것을 44% 줄였으나 종단 step/s 는 +0.1% — 종단이 이긴다. AR 은 이 스텝의 주범이 아니다 - 기각된 가설 넷과 그 반증 - 드래프터 3종 위치별 수용률 비교. GLM-5.3 DSpark 는 전 위치에서 DFlash2 보다 낮고(0.95 vs 1.53) 서빙 불가 — 대안 아님 - 판정 채널은 C=1 step/s 하나. 예외는 수락률 자체를 움직이는 변경 - 계측 규율: 파생량 전에 환산 계수부터, 종단과 어긋나면 종단이 이긴다, 프로파일은 시간이 아니라 개수, 인구조사 분모는 추정 말고 세기 tests/test_logic.py 671 검사 통과. Co-authored-by: choiceoh <astra7471@gmail.com>
Merged
4 tasks
choiceoh
added a commit
that referenced
this pull request
Sep 4, 2026
…접성 모드 (EXP-17/18/19) (#301) - 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 이 고치지 않는다.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
가정: 미귀속 32 ms 는 콜렉티브 지연이다. #99 의 위상 계측이 이 가정을 곧 확정하거나 기각하지만, 가정이 맞다면 손댈 곳은 커널 개수가 아니라 대기 방식이다 — 개수 축은 실측 5.4 µs/커널 · 우리 소유 148 개 = 1.1% 로 천장이 이미 보인다.
두 스핀 루프가 읽는 플래그는 전부 피어가 쓴다 —
ack_seq는 프록시가,rxf는 피어 NIC 가 RDMA 로. 그 플래그를 최고 속도로 읽으면 그 쓰기가 가시화되는 바로 그 경로에 트래픽을 계속 얹는 셈이다. 기다리는 대상과 경쟁한다.Changes
SPIN_HOT(8) 회는 뜨겁게, 이후__nanosleep128 ns → ×2 → 4096 ns 상한tx_seq호이스팅ack_seq뿐 → 루프 트래픽 절반정확성
건드리지 않는다. 대기 조건도, 펜스 위치도, 카운터 불변식도 그대로이고 바뀐 것은 얼마나 자주, 무엇을 읽는가 뿐이다.
tx_seq + 1 > ack_seq + RING. 호이스팅은 값이 상수임을 이용할 뿐 조건을 약화하지 않는다>= nxt. 단축평가 모양도 그대로 — 아직 없는 첫 번째 하나만 읽는다__nanosleep은__CUDA_ARCH__ >= 700가드 안에 있다#99(계측만) 부팅으로 위상 분해를 먼저 받고 그 다음 이걸 얹어야 한다. 둘을 한 부팅에 섞으면
wait이 줄어든 것이 이 변경 덕인지 원래 그랬는지 못 가른다 — 이 브링업이 오늘 이미 한 번 겪은 실수다(#93·#96·#98 이 한 부팅에 섞여 귀속이 사라졌다).판정:
[osar] phase us/collective기준선wait감소 확인Verification
tests/test_logic.py669 검사 통과🤖 Generated with Claude Code
Note
Medium Risk
Performance-only changes to synchronization spins on a cross-node AllReduce hot path; logic is intended equivalent but timing and long-run behavior should be validated on hardware (#99 baseline first).
Overview
Tight spin loops in
k_oneshotwaited on RDMA-visible flags (ack_seq, peerrxf) by hammering the same memory path peers write through. This PR addsosar_backoff: a few hot polls, then__nanosleepwith doubling delay up to 4096 ns (Volta+), capping worst-case added wait per collective.The ring-space guard still waits on
tx_seq + 1 > ack_seq + RING, buttx_seqis read once outside the loop (it cannot advance before this launch publishes) so onlyack_seqis re-polled. The peer-arrival wait still requires allNPEERflags>= nxt, but tracks which peers already landed and stops re-reading settled flags, backing off only while the first missing peer is still behind.Wait conditions, fences, and
done_ctrpublishing logic are unchanged—only poll frequency and which fields are re-read.Reviewed by Cursor Bugbot for commit 76b741b. Bugbot is set up for automated code reviews on this repo. Configure here.