perf(osar): 페이로드를 안 가진 블록은 두 스핀을 건너뛴다 — 디코드에서 피어 폴러 256→16 - #93
Conversation
#90 이 그리드를 ARGRID 로 고정한 것은 카운터 불변식(런치마다 정확히 ARGRID 증가) 때문이고, 디코드에서 그 대가가 드러난다. 콜렉티브 하나의 n 은 hidden 4096 이고 blockDim 은 256 이라, blockIdx 16..255 인 240 개 블록은 복사도 리듀스도 하지 않는다. 그런데도 링 가드 스핀, __threadfence_system, 피어 대기 스핀을 전부 낸다. 피어 대기가 비싸다. RDMA 도착을 기다리는 창 내내 256 개 블록이 같은 volatile 플래그 3 개를 폴링한다 -- 관측하려는 그 캐시라인에 트래픽을 쏘는 셈이다. 원소를 하나도 안 가진 블록에는 두 스핀 다 필요 없다. 가드는 그 블록이 쓰지 않는 슬롯을 보호하고, 피어 플래그는 그 블록이 읽지 않는 데이터를 여닫는다. 복사 뒤 fence 도 순서 지을 쓰기가 없으니 공허하다. atomicAdd 만 무조건 유지한다 -- 런치마다 ARGRID 를 더한다는 불변식이 마지막-블록 판정을 성립시킨다. - owns = (blockIdx.x == 0) || (blockIdx.x * blockDim.x < n) - 링 가드 / 복사 후 fence / 피어 대기 + 그 뒤 fence 를 owns 로 게이팅 - owns 는 블록 단위로 균일하므로 __syncthreads 발산 없음 - 블록 0 은 무조건 owns -- n == 0 콜렉티브도 가드를 타게 해서, 아무도 확인 안 한 슬롯에 게시하는 일이 없다 n 이 그리드를 채우면 동작은 이전과 동일하다. 커널 수는 1 개 그대로라 #91 의 부팅 지문(kernels=1)도 변하지 않는다. 효과는 부팅으로 재야 한다. 디코드에서 피어 플래그 폴러가 256 -> 16 으로 준다.
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_0100fbd2-c065-4da7-a2a0-d65d439b16e2) |
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.
Did not approve: Cursor Bugbot was present but skipped (usage/spend limit), so the required automated-review signal did not complete successfully. Human review is needed; a reviewer was requested where assignment succeeded. No Slack summary was posted (no Slack tool).
Sent by Cursor Approval Agent: Pull Request Router and Approver
스텝 고정비 42.0ms 중 dense 가중치 9.9ms 를 뺀 32.1ms 가 아직 미귀속이다. 지금까지 이 값에 붙은 설명은 전부 잔차 귀속이었고 세 번 뒤집혔다: 커널 개수 가설(상한 15.4us/커널)은 실측 5.4us/커널 로 1/3 이 됐고, 호스트 병목 가설은 GPU 사용률 93% 로 기각됐다. 지금 서 있는 가설(콜렉티브 대기)도 사용률 93% + 전력 15W 평평이라는 간접 증거뿐이다. 직접 재는 편이 빠르다. - Ctrl 의 pad[8] 을 위상 누적기 5 개 + pad[3] 로 바꾼다. 총 64B 그대로라 tx/rx 오프셋이 안 변하고 피어의 rx_base/rxf_base 가 유효하다. 단조 증가이며 리셋하지 않는다 -- done_ctr 과 같은 이유로 cudagraph 재생이 리셋을 못 본다. - k_oneshot 이 clock64() 를 5 지점에서 읽어 guard / copy / wait / reduce 로 가른다. 기록은 블록 0 스레드 0 만 한다. #93 이후 블록 0 은 항상 owns 라 모든 위상을 걷는 유일하게 대표성 있는 표본이다. 콜렉티브당 한 스레드의 store 5 회. - 호스트 리더는 프록시 스레드에 둔다. 시임에 두면 안 된다 -- 풀 디코드 cudagraph 아래서 파이썬 진입점은 캡처 때 한 번 돌고 다시는 안 돌지만 커널은 재생마다 누적한다. 프록시는 부팅 내내 콜렉티브마다 도는 유일한 호스트 코드다. 랭크 0 만, 10 초마다, 직전 보고 이후의 델타로 찍는다(느린 워밍이 정상 상태를 뭉개지 않게). 절대 us 는 SM 클럭 1592MHz 가정만큼만 정확하다(nvidia-smi clocks.sm 이 부하 전 구간에서 평평했다). 다만 네 위상의 비는 그 가정에 전혀 의존하지 않고, 우리가 알고 싶은 것이 바로 그 비다. 성능 변경이 아니라 계측기다. 컴파일 검사 없이 부팅하지 말 것 -- 실패한 부팅이 13 분이다. Co-authored-by: choiceoh <astra7471@gmail.com>
Summary
#90 이 그리드를
ARGRID로 고정한 것은 카운터 불변식(런치마다 정확히 ARGRID 증가) 때문이고, 디코드에서 그 대가가 드러난다. 콜렉티브 하나의n은 hidden 4096,blockDim은 256 이므로blockIdx16..255 인 240 개 블록은 복사도 리듀스도 하지 않는다. 그런데도 링 가드 스핀 ·__threadfence_system· 피어 대기 스핀을 전부 낸다.피어 대기가 비싸다. RDMA 도착을 기다리는 창 내내 256 개 블록이 같은 volatile 플래그 3 개를 폴링한다 — 관측하려는 그 캐시라인에 트래픽을 쏘는 셈이다. #90 본문이 남긴 미측정 항목("원자가 시스템 메모리로 떨어지는 비용은 대상기에서 실측 필요 — 빈 블록당 원자 1회 × 256")의 후속이다.
Changes
const bool owns = (blockIdx.x == 0) || (blockIdx.x * blockDim.x < n);owns로 게이팅atomicAdd는 무조건 유지 — 런치마다 ARGRID 를 더한다는 불변식이 마지막-블록 판정을 성립시킨다owns—n == 0콜렉티브도 가드를 타게 해서, 아무도 확인 안 한 슬롯에 게시하지 않는다안전성 논증
owns는blockIdx만의 함수라 블록 단위로 균일 →__syncthreads발산 없음게시 블록이 빈 블록일 수 있으나 무해하다: 마지막 증가자이므로 모든 복사 fence 가 이미 지났고,
slot은 같은nxt에서 계산된다.n이 그리드를 채우면(프리필) 동작은 이전과 완전히 동일하다.Verification
__global__1 개 유지 → fix(osar): 부팅 로그에 소스 지문(md5+커널 수) — 어느 판본이 도는지 로그로 가린다 #91 의 부팅 지문kernels=1불변[osar] source md5=… kernels=1+self-test PASS+real=Truebench-dec.py 6, rep 별 수락률 보존)로 잰 팔 C(1커널·k=7)🤖 Generated with Claude Code
Note
Medium Risk
Changes ordering and synchronization in a cross-node RDMA AllReduce kernel; correctness relies on the documented invariants around
ownsand unconditionalatomicAdd.Overview
Decode-sized AllReduces (
n≈ hidden 4096, fixed 256-block grid) leave most blocks with no copy/reduce work, but they still spun on the ring guard and peer RDMA flags and issued__threadfence_system— adding cache-line traffic during the expensive peer wait.This PR introduces
owns(block 0always, else blocks withblockIdx.x * blockDim.x < n) and gates the ring guard, post-copy fence, peerrxfwait, and pre-reduce fence behind it.atomicAddondone_ctrstays unconditional so each launch still contributes exactly ARGRID increments and last-blocktx_seqpublish remains sound; block 0 stillownswhenn == 0so the ring guard runs.When
nfills the grid (prefill), behavior matches the previous kernel.Reviewed by Cursor Bugbot for commit 67f8597. Bugbot is set up for automated code reviews on this repo. Configure here.