제1부 GPU 아키텍처

1.6 SM 내부 III — 점유율 수계산·limiter·튜닝

버전 주의: 분모(SM 최대 상주 워프)·레지스터 파일·공유 메모리 용량은 Compute Capability별로 다르다(시점 의존). 본문 수치는 산수 시범용 예시다.

1.6.1 이 장에서 배우는 것 / 선수 지식

1.4~1.5에서 SM 구조와 워프 스케줄링을 봤다. 이 장은 그 둘을 잇는 정량 지표 — 점유율(occupancy) 을 직접 계산하고, 무엇이 점유율을 낮추는지(limiter) 짚고, "점유율을 맹목적으로 높이는 것이 답이 아닌" 이유(점유율 vs ILP)를 다룬다. 이 장으로 SM 내부 3부작이 끝난다.

배우는 것:

• 점유율 정의·계산·해석. 레지스터·공유 메모리·블록/워프 한도가 어떻게 occupancy를 낮추는지 직접 계산.
• 점유율을 낮추는 limiter를 도구로 식별.
• 점유율 vs ILP 트레이드오프와 "충분한 점유율" 개념.
• 점유율을 과신하는 함정.

선수 지식: 1.4(레지스터 파일·레지스터 압박), 1.5(워프 스케줄링·stall). 점유율을 실제 커널 설계에 적용하는 실행 구성(그리드·블록)은 2.3, 도구로 진단하는 프로파일링은 2.14.

1.6.2 점유율 — 정의

점유율(occupancy) = SM에 실제로 상주한 워프 수 ÷ 그 SM이 지원하는 최대 상주 워프 수.

$\mathrm{occupancy} =$ (동시 상주 워프 수) / (SM 최대 상주 워프 수)

분모(SM 최대 상주 워프)는 Compute Capability로 정해진 하드웨어 한도(예: 48 또는 64 워프). 분자(실제 상주 워프)는 여러 자원 한도 중 가장 제약이 큰 것이 결정한다. 즉 occupancy를 낮추는 한도가 네 갈래다:

1. 워프/블록 하드 한도: SM당 최대 워프 수, 최대 블록 수(둘 다 CC별 상수).
2. 레지스터: 스레드당 레지스터 × 블록당 스레드 → 블록당 레지스터. SM 레지스터 파일 ÷ 블록당 레지스터 = 상주 가능 블록 수(1.4).
3. 공유 메모리: 블록당 공유 메모리 사용량. SM 공유 메모리 ÷ 블록당 공유 메모리 = 상주 가능 블록 수.
4. 블록 크기 정렬: 워프는 32 단위. 블록당 스레드가 32의 배수가 아니면 마지막 워프에 빈 레인 생김(낭비).

이 넷을 각각 "상주 가능 블록 수"로 환산해 최솟값을 취하고, ×블록당 워프 수로 상주 워프를 구한다.

1.6.3 수계산 워크스루 (수치 예시)

가정(예시값 — 실제는 대상 GPU의 CC로 확인):

• SM 최대 상주 워프 $= 64$
• SM 최대 상주 블록 $= 32$
• SM 레지스터 파일 $= 65{,}536$ (32-bit reg)
• SM 공유 메모리 $= 100\,\mathrm{KB}$ (구성 가능, 예시)

커널 A: 블록당 256 스레드, 스레드당 32 reg, 블록당 공유 메모리 8 KB.

• 블록당 워프 $= 256 / 32 = 8$ 워프
• 블록당 레지스터 $= 32 \times 256 = 8{,}192$ → 레지스터 기준 상주 블록 $= 65{,}536/8{,}192 = 8$ 블록
• 공유 메모리 기준 상주 블록 $= 100\,\mathrm{KB} / 8\,\mathrm{KB} = 12$ 블록
• 블록 수 하드 한도 $= 32$ 블록
• 워프 하드 한도로부터 블록 $= 64$ 워프 $/ 8$ 워프 $= 8$ 블록
• 상주 블록 $= \min(8, 12, 32, 8) = 8$ 블록
• 상주 워프 $= 8 \times 8 = 64$ 워프
• $\mathrm{occupancy} = 64 / 64 = 100\%$ ← 레지스터·워프 한도가 동시에 묶음

커널 B: 같은 블록(256 스레드), 스레드당 레지스터를 64로 늘림(복잡해짐), 공유 메모리 동일 8 KB.

• 블록당 레지스터 $= 64 \times 256 = 16{,}384$ → 레지스터 기준 상주 블록 $= 65{,}536/16{,}384 = 4$ 블록
• (공유 메모리·하드 한도는 커널 A와 같음: 12, 32, 8)
• 상주 블록 $= \min(4, 12, 32, 8) = 4$ 블록 ← 레지스터가 병목
• 상주 워프 $= 4 \times 8 = 32$ 워프
• $\mathrm{occupancy} = 32 / 64 = 50\%$

레지스터를 32→64로 올렸더니 점유율이 100%→50%로 반토막 났다. 이것이 레지스터 압박(1.4)의 직접적 결과다. nvcc의 __launch_bounds__나 -maxrregcount로 레지스터를 제한해 점유율을 끌어올릴 수 있으나, 레지스터를 줄이면 레지스터 스필(spill) 이 로컬 메모리(=글로벌 메모리 대기)로 새어 오히려 느려질 수 있다 — 함정 1.6.6.

정직: 위 수계산은 "레지스터 총량 ÷ 블록당 레지스터"로 단순 나눗셈했지만, 실제 하드웨어는 레지스터를 블록·워프 단위로 특정 배수에 맞춰 할당할 수 있어(정렬·granularity, 세대 의존) 실제 상주 블록 수가 수계산보다 살짝 낮게 나올 수 있다. 정확한 값은 수계산이 아니라 Occupancy API·Nsight Compute(1.6.3.1·2.14)로 확인 — 이 장의 수계산은 개략을 파악하는 근사치다.

커널 C: 공유 메모리를 키운 경우 — 블록당 32 KB, 256 스레드, 스레드당 32 reg.

• 공유 메모리 기준 상주 블록 $= 100\,\mathrm{KB} / 32\,\mathrm{KB} = 3$ 블록 (정수 내림)
• 레지스터 기준 $= 8$ 블록, 하드 한도 $=$ 32·8 블록
• 상주 블록 $= \min(8, 3, 32, 8) = 3$ 블록 ← 공유 메모리가 병목
• 상주 워프 $= 3 \times 8 = 24$ 워프
• $\mathrm{occupancy} = 24 / 64 = 37.5\%$

공유 메모리를 많이 쓰는 타일링 커널(2.7)은 이렇게 공유 메모리가 점유율을 낮춘다. 타일 크기·공유 메모리 구성(L1과의 분할, 1.7)이 occupancy 튜닝 변수다.

1.6.3.1 도구로 확인

수계산은 개략 추정용이고, 실무에선 도구로 확인한다:

• CUDA Occupancy API: cudaOccupancyMaxActiveBlocksPerMultiprocessor(...), cudaOccupancyMaxPotentialBlockSize(...) — 런타임에 이론 점유율과 권장 블록 크기를 계산.
• nvcc --ptxas-options=-v (또는 -Xptxas -v): 컴파일 시 스레드당 레지스터·정적 공유 메모리·스필 바이트를 출력.
• Nsight Compute(2.14): 이론 점유율 vs 달성(achieved) 점유율을 모두 보고, 점유율을 낮춘 한도(Limiter: registers / shared memory / block size)를 짚어준다.

  --ptxas-options=-v 예시 출력 (개략, 환경에 따라 다름)
  ptxas info : Used 64 registers, 8192 bytes smem, ... 0 bytes spill stores
                    └ 스레드당 reg      └ 블록 공유메모리   └ 스필 (0이 이상적)

1.6.4 트레이드오프 — 점유율 vs ILP, 그리고 "충분한 점유율"

지연을 은닉하는 길은 둘이다(1.2·1.5):

1. TLP(thread-level parallelism, ≈ 점유율): 워프를 많이 상주시켜 한 워프가 정지된 사이 다른 워프 발행.
2. ILP(instruction-level parallelism): 한 워프 안에서 독립 명령을 늘려(로드 여러 개 선발행, 누산기 여러 개) 의존성 구멍을 자체 메움.

레지스터를 더 써서 ILP를 키우면(언롤·다중 누산기) 점유율은 내려가지만, 워프당 발행할 독립 작업이 많아져 낮은 점유율로도 지연이 은닉될 수 있다. 고전적 사례가 잘 튜닝된 GEMM 커널 — 점유율이 25~50%라도 ILP와 데이터 재사용으로 정점에 가까운 처리량을 낸다.

세 가지 방식 — "누구에게 맞나":

견해: "무조건 100% 점유율"을 목표로 삼지 말 것. 점유율은 수단이지 목적이 아니다. "지연을 은닉할 만큼 충분한" 점유율이면 된다 — 그 이상은 레지스터·공유메모리를 줄여 ILP·재사용을 해칠 수 있다.

비엔비디아 정직성: 점유율·워프 개념은 엔비디아 SIMT 모델 고유 용어지만, "많은 동시 스레드로 메모리 지연을 은닉한다"는 발상 자체는 다른 처리량 가속기(AMD wavefront·occupancy 등)에도 공통이다. 용어와 정확한 자원 한도만 다르다.

1.6.5 중간 점검 (체크포인트)

• [ ] 레지스터·공유 메모리·블록/워프 하드 한도 네 갈래로 상주 블록을 각각 환산해 최솟값으로 점유율을 수계산한다.
• [ ] 어느 자원이 limiter인지 식별하고, 그에 맞는 대응(레지스터/공유메모리/블록 크기 조정)을 고른다.
• [ ] "점유율 50%인데 처리량 정점"이 문제가 아닐 수 있음을 설명한다(ILP·재사용).
• [ ] --ptxas-options=-v의 spill stores가 0이 아니면 레지스터 압박이 로컬 메모리로 새는 신호임을 안다.

점검 질문: 어떤 커널의 이론 점유율이 50%인데 처리량이 정점이라면 — 문제인가? (정답 방향: 아니다. 지연이 충분히 은닉됐고 ILP·재사용이 좋으면 50%로도 정점일 수 있다. limiter와 실제 처리량을 같이 봐야 판단 가능.)

1.6.6 흔한 함정

• 점유율 과신. "점유율 100% = 빠름"은 오해. 점유율은 지연 은닉의 상한 능력일 뿐, 실제 성능은 메모리 패턴·연산 밀도·ILP·캐시 적중에 좌우된다. 50%에서 정점인 커널이 흔하다. 항상 프로파일러의 limiter와 실제 처리량을 같이 봐라(2.14).
• 레지스터를 무리하게 줄여 스필 유발. -maxrregcount·__launch_bounds__로 레지스터를 줄여 점유율을 올렸는데, 줄인 레지스터가 로컬 메모리로 스필되면 그 접근이 사실상 글로벌 메모리 지연이라 더 느려진다. --ptxas-options=-v의 spill stores/loads가 0이 아니면 경고.
• 이론 점유율 vs 달성 점유율 혼동. 이론은 자원 한도가 허용하는 상한, 달성은 실제 실행 평균이다. 둘이 크게 벌어지면 tail effect(마지막 웨이브에 블록이 적어 SM이 유휴 상태가 되는 현상)나 워크로드 불균형·블록 수 부족을 의심. 그리드 크기를 SM 수의 배수에 가깝게 키우는 것이 도움.
• 블록 크기를 32의 배수로 안 잡음. 블록당 스레드가 32의 배수가 아니면 마지막 워프에 빈 레인이 생긴다. 거의 항상 32의 배수(보통 128·256).
• 공유 메모리 한도를 잊음. 레지스터만 보고 점유율을 추정하다 공유 메모리가 실제 병목인 경우를 놓침(타일링 커널, 2.7).
• 세대 차이를 일반화. "내 GPU에선 64 reg에서 50%였으니 다른 GPU도"는 위험. 레지스터 파일·최대 워프·공유 메모리가 CC별로 달라 같은 커널의 점유율이 세대마다 바뀐다 — 대상 GPU로 재계산.

1.6.7 연습문제

연습 1 (점유율 수계산). SM 최대 상주 워프 64, 레지스터 파일 65,536, 공유 메모리 100 KB. 블록당 128 스레드, 스레드당 40 reg, 블록당 공유 메모리 16 KB인 커널의 이론 점유율을 구하라. limiter는?

연습 2 (레지스터 압박 실험). -maxrregcount를 점점 낮췄더니 점유율은 오르는데 실행 시간이 어느 지점부터 오히려 늘었다. 무슨 일이 일어났고 어떤 출력을 보고 확인하나?

연습 3 (블록 크기 선택). 같은 알고리즘을 블록 128·256·512로 빌드해 점유율·실행 시간을 측정하려 한다. 어떤 도구·플래그로 이론 점유율과 권장 블록 크기를 얻고, 측정 비교 지표는 무엇이어야 하나?

연습 4 (함정 식별, 견해 포함). 동료가 "커널이 느린데 점유율이 60%밖에 안 돼서 그렇다. 100%로 높이자"고 한다. 어떻게 검증·반박하거나 동의할 것인가? 어떤 데이터가 필요한가?

연습 5 (심화, 공유 메모리 구성). L1과 공유 메모리가 같은 물리 SRAM을 분할하는 세대에서, 공유 메모리를 적게 쓰도록 구성을 바꾸면 점유율에 어떤 영향이 있나? 트레이드오프는? (1.7 선행 참조 가능)

1.6.8 요약 · 다음 장

• 점유율 = 상주 워프 / 최대 상주 워프. 레지스터·공유 메모리·블록/워프 하드 한도 네 갈래를 상주 블록으로 환산해 최솟값을 취하면 수계산된다. 도구는 Occupancy API·-ptxas -v·Nsight Compute.
• 점유율을 낮추는 limiter를 식별해 그에 맞는 자원을 조정한다(레지스터 ≠ 항상 답).
• 점유율은 수단이지 목적이 아니다. 100%가 곧 빠름이 아니며, 레지스터를 무리하게 줄이면 스필로 역효과. 점유율 vs ILP를 균형을 맞추고, 항상 limiter와 실제 처리량을 함께 본다.

SM 내부 3부작(1.4~1.6)이 끝났다. 다음 장: 1.7 온칩 메모리 I — 레지스터·공유 메모리·뱅크. SM이 쓰는 온칩 메모리의 물리 구조와 뱅크 충돌을 본다. 점유율을 낮추던 공유 메모리가 왜 그렇게 빠르고 어떻게 충돌하는지를 여기서 해부한다.