제1부 GPU 아키텍처
1.5 SM 내부 II — 워프 스케줄링·지연 은닉·dual-issue
버전 주의: 스케줄러 개수·디스패치 폭·발행 정책은 세대마다 다르다(시점 의존, 확인 필요). 본문은 동작 원리에 집중한다.
1.5.1 이 장에서 배우는 것 / 선수 지식
1.4에서 SM의 물리 구조(처리 블록·레지스터 파일)를 봤다. 이 장은 그 위에서 워프 스케줄러가 매 사이클 무엇을 하는가 — 준비된 워프를 골라 발행하고, 정지(stall)된 워프를 건너뛰며, dual-issue로 두 파이프를 채우는 메커니즘을 본다. 이것이 1.2의 "지연 은닉"이 실제로 일어나는 곳이다.
배우는 것:
- 워프 스케줄러가 매 사이클 "준비됨/막힘"을 어떻게 가르고 발행하는가.
- 대표 stall 사유(Long/Short Scoreboard·Wait·Barrier·MIO·Not Selected)와 그 진단 의미.
- dual-issue(슈퍼스칼라 발행)의 조건과 이득.
- LD/ST 유닛·SFU의 역할과 병목.
선수 지식: 1.2(지연 은닉의 큰 그림), 1.4(처리 블록·레지스터 파일). 점유율 정량 계산은 1.6, 프로파일러로 stall을 보는 법은 2.14.
1.5.2 스케줄러는 매 사이클 무엇을 하나
한 처리 블록의 워프 스케줄러는 그 블록에 할당된 워프들(여러 커널 블록의 워프가 섞일 수 있다)을 관리한다. 매 사이클:
- 상주 워프들의 상태를 본다. 각 워프는 "다음 명령이 발행 가능한가"로 나뉜다 — 준비됨(ready) 또는 막힘(stalled).
- 준비된 워프가 있으면 그중 하나(또는 dual-issue면 한 워프에서 2개, 또는 두 워프에서 각 1개 — 세대 정책마다 다름)를 골라 명령을 발행한다.
- 준비된 워프가 하나도 없으면 그 사이클은 발행 슬롯이 빈다(stall) — 처리량 손실.
워프가 "정지하는" 이유(stall reason)는 프로파일러(2.14)에서 핵심 진단 신호다. 대표적으로:
| stall 사유 | 의미 |
|---|---|
| Long Scoreboard | 메모리(글로벌/로컬) 응답 대기 — 가장 흔한 큰 구멍 |
| Short Scoreboard | 공유 메모리 등 짧은 의존성 대기 |
| Wait | 산술 명령 결과 의존성(파이프 지연) 대기 |
| Barrier | __syncthreads() 등 블록 동기화에서 형제 워프 대기 |
| MIO Throttle | LD/ST·SFU 등 특정 파이프 포화 |
| Not Selected | 준비됐으나 이번 사이클에 다른 워프가 선택됨 (=경쟁, 건강한 신호) |
지연 은닉이 잘 되는 SM은 "Long Scoreboard로 정지된 워프가 많아도, 그 사이 발행할 준비된 워프가 늘 있는" 상태다. "Not Selected"가 많다는 것은 준비된 워프가 넘쳐 발행 슬롯을 두고 경쟁한다는 뜻 — 점유율이 충분하다는 건강한 신호다. 반대로 발행 슬롯이 자주 비면(준비 워프 부족) 점유율·ILP·메모리 패턴을 의심한다.
1.5.2.1 지연 은닉을 그림으로
이것이 "왜 점유율이 중요한가"의 본질이다. 메모리 지연이 길수록(수백 사이클), 그 동안 발행을 이어가려면 더 많은 상주 워프가 필요하다(1.2의 리틀의 법칙에 따른 정성적 추정). 다만 점유율만이 지연을 은닉하는 수단은 아니다 — 한 워프 안에서 독립적인 명령을 늘리는 ILP(예: 메모리 로드를 여러 개 미리 기동하기)로도 같은 구멍을 메울 수 있다. 이 둘의 트레이드오프가 1.6의 핵심이다.
1.5.2.2 dual-issue / 슈퍼스칼라 발행
세대에 따라 처리 블록의 스케줄러는 한 사이클에 명령을 둘 이상 발행할 수 있다(dual-issue). 두 명령이 서로 다른 실행 파이프(예: 하나는 FP32 산술, 하나는 LD/ST)로 가고 의존성이 없을 때 동시에 내보낸다. 같은 FP32 파이프로 두 개를 넣을 수는 없다(파이프 폭 제한). 그래서 dual-issue의 이득은 명령 믹스가 다양할 때(산술+메모리+SFU가 섞일 때) 커진다.
세대 주의: 이 한 스케줄러의 dual-issue(한 워프에서 독립 2명령/사이클)는 Kepler·Maxwell·Pascal의 특징이었다. Volta 이후(Turing·Ampere·Hopper·Ada·Blackwell)는 각 스케줄러가 워프당 1명령/사이클을 발행하고, SM의 4개 스케줄러가 합쳐 최대 4명령/사이클을 낸다(SM 수준 슈퍼스칼라). 모던 GPU에선 한 스케줄러의 dual-issue를 가정하지 말 것.
같은 워프에서 독립인 두 명령이거나 두 워프에서 각각 — 세대 정책마다 다르다. 서로 다른 파이프로 가므로 두 파이프가 동시에 가동되어 발행 처리량이 오른다.
1.5.3 LD/ST 유닛과 SFU
- LD/ST(Load/Store) 유닛: 워프의 메모리 명령에 대해 32 스레드의 주소를 모아 메모리 시스템에 요청을 발행한다. 32 스레드의 접근이 연속·정렬돼 있으면 적은 수의 트랜잭션으로 묶이고(coalescing), 흩어져 있으면 트랜잭션이 분할돼 효율이 떨어진다. LD/ST 유닛 수가 워프 메모리 명령의 발행률을 제한하므로, 메모리 바운드 커널에서 병목이 될 수 있다(stall 사유
MIO Throttle·Long Scoreboard). 병합·뱅크·캐시 동작의 상세는 1.7~1.8, 대역폭·메모리 바운드 분석은 1.9~1.10. - SFU(Special Function Unit):
sinf·cosf·expf·rsqrtf·역수·로그 등 초월함수의 하드웨어 근사를 제공한다. SFU 처리량은 보통 FP32 코어보다 훨씬 낮다(블록당 SFU 레인이 적음). 따라서 초월함수를 워프당 많이 쓰면 SFU가 병목이 된다.--use_fast_math나__sinf같은 내장 함수는 SFU 근사를 직접 써서 빠르지만 정밀도를 양보한다 — 정밀도 요구에 따라 선택.
1.5.4 중간 점검 (체크포인트)
- [ ] "워프 스케줄러가 매 사이클 준비된 워프를 골라 발행하고, 준비된 워프가 없으면 슬롯이 빈다"를 설명하고, 지연 은닉이 왜 상주 워프 수에 의존하는지 그림으로 보인다.
- [ ] 대표 stall 사유를 구분하고, "Not Selected가 많다 = 건강한 신호"를 설명한다.
- [ ] dual-issue가 명령 믹스가 다양할 때 이득이 큰 이유를 제시한다(같은 파이프엔 두 개 못 넣음).
- [ ] LD/ST 병합 여부·SFU 처리량이 병목이 될 수 있음을 안다.
점검: 프로파일러에서 발행 슬롯이 자주 비고 stall이 Long Scoreboard에 몰려 있다면 → 무엇을 의심하나? (방향: 점유율 부족 또는 메모리 바운드. 준비 워프를 늘리거나 메모리 패턴을 개선.)
1.5.5 흔한 함정
- stall이 보이면 무조건 점유율을 높이려 함. stall 사유부터 봐야 한다. Long Scoreboard 지배 + 슬롯 빔이면 점유율/메모리가, "Not Selected" 지배면 이미 충분(다른 병목).
- dual-issue를 "항상 2배"로 기대. 같은 파이프엔 둘을 못 넣고, 의존성이 있으면 동시 발행 불가. 이득은 명령 믹스·독립성에 달림.
- SFU 남발. 초월함수를 워프당 많이 쓰면 SFU가 병목. 정밀도가 허용되면 fast-math 내장 함수로, 아니면 호출 수를 줄인다.
- 비병합 메모리 접근. 워프 32 스레드가 흩어진 주소를 읽으면 트랜잭션이 분할되어 LD/ST·대역폭을 낭비(1.8).
1.5.6 연습문제
연습 1 (stall 해석). 프로파일러가 한 커널에서 "Not Selected"가 지배적이라 보고한다. 점유율을 더 높이면 빨라질까?
정답
아닐 가능성이 크다. "Not Selected"가 많다는 것은 준비된 워프가 이미 넘쳐 발행 슬롯을 두고 경쟁한다는 뜻 — 점유율은 충분하다. 병목은 다른 곳(연산 처리량·특정 파이프·메모리 대역폭)일 수 있다. 점유율 대신 그쪽을 본다(1.6·1.10).
연습 2 (dual-issue). 한 커널이 FP32 fma만 연달아 실행한다. dual-issue로 발행 처리량이 2배가 될까?
정답
아니다. 두 명령이 같은 FP32 파이프로 가면 동시 발행이 안 된다. dual-issue 이득은 서로 다른 파이프(FP32 + LD/ST + SFU 등)로 갈 독립 명령이 섞일 때 나온다. 순수 FP32 연쇄는 FP32 파이프 폭이 천장.
연습 3 (SFU 병목). 활성화 함수에 expf를 대량 호출하는 커널이 느리다. 어디를 의심하고 어떻게 푸나?
정답
SFU 병목을 의심(SFU 처리량이 FP32보다 낮음). 프로파일러에서 MIO/SFU 관련 throttle 확인. 정밀도가 허용되면 __expf/fast-math로 SFU 근사를 쓰거나, 호출 횟수를 줄이는 알고리즘(룩업·근사)으로 바꾼다.
1.5.7 요약 · 다음 장
- 워프 스케줄러는 매 사이클 준비된 워프를 골라 발행한다. 준비 워프가 없으면 슬롯이 비어 처리량을 잃는다 — 지연 은닉은 충분한 상주 워프(또는 ILP)에 달려 있다.
- stall 사유가 진단의 핵심: Long Scoreboard(메모리)·Barrier·MIO·"Not Selected"(건강). 슬롯이 비는지부터 본다.
- dual-issue는 서로 다른 파이프로 갈 독립 명령이 섞일 때 이득. LD/ST(병합)·SFU(낮은 처리량)는 병목 후보.
다음 장: 1.6 SM 내부 III — 점유율 수계산·limiter·튜닝. 레지스터·공유 메모리·블록/워프 한도를 상주 블록으로 환산해 점유율을 직접 계산하고, 점유율 vs ILP 트레이드오프와 흔한 함정을 다룬다.