제1부 GPU 아키텍처

1.2 SIMT·워프·지연 은닉

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

1.1에서 "GPU는 지연을 은닉한다"는 철학을 봤다. 이 장은 어떻게 은닉하는지의 메커니즘 — SIMT 실행 모델·워프·다이버전스·지연 은닉을 산수로 익힌다. CUDA 성능 사고의 출발점이며, 이후 점유율(1.6)·실행 구성(2.3)·프로파일링(2.14)이 모두 이 모형 위에 구축된다.

이 장을 마치면 다음을 할 수 있어야 한다.

• SIMT 실행 모델을 워프(warp) 단위 lockstep으로 설명하고, 다이버전스(divergence)가 무엇이며 왜 성능을 저하시키는지 안다.
• SIMT와 SIMD(CPU 벡터)의 프로그래밍 모델 차이를 구분한다.
• "많은 워프로 메모리 지연을 은닉한다"는 메커니즘을 산수(리틀의 법칙에 따른 정성적 추정)로 설명한다.
• 산술 강도로 커널이 메모리 바운드인지 연산 바운드인지 정성적으로 판단한다(루프라인 개요, 본격은 1.10).

선수 지식: 1.1(처리량 컴퓨팅 철학). 워프 = 32 스레드 묶음이라는 점부터 시작한다.

정직성(S-5): 지연 사이클·워프 수 등은 설명용 개략값이다. 핵심은 "지연 ↔ 필요 워프 수"의 정성적 방향성.

1.2.2 SIMT 실행 모델 — 워프 단위 lockstep

GPU가 ALU를 그렇게 밀집시킬 수 있는 비결이 SIMT(Single Instruction, Multiple Threads)다. 핵심 개념은 워프(warp) — 엔비디아 GPU에서 32개 스레드의 묶음이다(오랫동안 32로 고정, 시점 의존).

규칙: 한 워프의 32 스레드는 같은 명령을 같은 사이클에 실행한다(lockstep). 각 스레드는 자기 레지스터·자기 데이터를 갖지만, "다음에 무슨 명령을 실행할지"는 워프 단위로 하나만 결정된다.

이것이 "하나의 명령 fetch/decode를 32 레인이 공유"하므로 제어 로직을 32분의 1로 아껴 그 면적을 ALU에 할애할 수 있는 이유다(1.1.3). 워프는 발행의 단위이자, 뒤에서 보듯 메모리 접근의 단위이기도 하다.

1.2.2.1 다이버전스 — lockstep이 깨질 때

문제는 분기(if/else)다. 한 워프 안의 스레드들이 조건에 따라 서로 다른 경로로 가면, 하드웨어는 32 레인에 하나의 명령만 줄 수 있으므로 두 경로를 순차로 실행하고 각 경로에서 해당되지 않는 레인을 마스크(비활성)한다. 이를 워프 다이버전스(warp divergence)라 한다.

if (threadIdx.x % 2 == 0) A(); else B(); — 같은 워프 안에서 짝/홀로 갈리면 하드웨어가 두 경로를 직렬로 실행하며 각 단계에서 해당 없는 레인을 마스크한다(■ = 활성, · = 마스크·유휴). 두 경로 시간이 합산되고 각 단계에서 절반 레인은 유휴다.

다이버전스가 없으면(32 스레드가 모두 같은 경로) 분기는 무비용에 가깝다. 다이버전스가 심하면 최악 32-way로 갈라져 그 구간의 유효 처리량이 1/32까지 떨어질 수 있다. 그래서 커널을 작성할 때 분기를 워프 경계에 정렬하거나(같은 워프는 같은 경로) 분기 없는 산술로 바꾸는 것이 중요하다. (최신 세대의 독립 스레드 스케줄링은 재수렴을 더 유연하게 다루지만 비용 자체가 사라지진 않는다 — 세대 의존.)

또 하나: 워프 내 32 스레드가 글로벌 메모리에서 이웃한 주소를 함께 읽으면 하나의 넓은 트랜잭션으로 묶인다(병합, coalescing). 흩어진 주소를 읽으면 트랜잭션이 분할되어 대역폭을 낭비한다. 병합은 1.8에서 깊이 다룬다 — 여기서는 "워프가 메모리 접근의 단위이기도 하다"만 기억한다.

1.2.3 SIMT를 SIMD로 오해하지 마라

SIMT와 SIMD(AVX 같은 벡터 명령)는 둘 다 "하나의 명령으로 여러 데이터"를 처리하지만, 프로그래밍 모델이 다르다. 이것을 섞으면 GPU를 보는 관점이 계속 어긋난다.

요점: SIMT는 "스칼라 스레드를 작성하되, 하드웨어가 32개씩 lockstep으로 묶어 SIMD처럼 실행한다"는 모델이다. 그래서 코드는 직렬처럼 읽히지만 성능은 워프·다이버전스·병합 같은 SIMD적 현실에 지배된다. 이 이중성을 이해하는 것이 CUDA 성능 사고의 출발점이다.

1.2.4 지연 은닉 — 처리량의 핵심 원리

이제 핵심. GPU는 메모리 지연을 줄이지 않는다(작은 캐시). 대신 은닉한다. 방법: SM(Streaming Multiprocessor)에 준비된 워프를 다수 상주시켜 두고, 한 워프가 메모리를 기다리는 동안 스케줄러가 다른 준비된 워프로 즉시 교체해 ALU를 유휴 상태로 두지 않는다.

CPU의 컨텍스트 스위치는 비용이 크다(레지스터 저장·복원). GPU는 다르다 — 모든 워프의 레지스터가 거대한 레지스터 파일에 상주하므로 워프 전환은 사실상 0 사이클(어느 워프를 디스패치할지 고르기만 함). 그래서 워프를 많이 상주시킬수록 지연을 잘 은닉한다(레지스터 파일의 물리 구조는 1.4).

여기서 점유율(occupancy) 개념이 나온다: SM이 실제로 상주시킨 워프 수 ÷ SM이 상주시킬 수 있는 최대 워프 수. 점유율이 높을수록 지연을 은닉할 워프가 많다(다만 점유율이 높다고 항상 빠른 것은 아니다 — 1.6에서 정밀하게). 무엇이 점유율을 제한하는가(레지스터·공유 메모리 사용량)는 1.6의 주제다.

1.2.4.1 워크스루 — 지연을 은닉하는 산수

지연 은닉을 수치로 확인해 보자. 모두 설명용 개략값이다(실제 세대마다 다름, S-5).

가정:

• 메모리 접근 지연: 400 사이클 (글로벌 메모리, 개략).
• 한 워프가 메모리 요청 사이에 평균 8 사이클의 독립 산술을 실행.
• 워프 발행은 사이클당 1 워프라 가정(단순화).

질문: ALU를 100% 바쁘게 하려면 몇 개의 워프가 SM에 동시에 있어야 하나?

리틀의 법칙(Little's Law) 요지: 한 워프가 8 사이클 일하고 400 사이클 기다린다면, 그 워프 하나가 "일하는 비율"은 8/(8+400) ≈ 2%다. ALU를 꽉 채우려면 빈 시간을 다른 워프로 메워야 한다.

• 필요 워프 수 $\approx$ (계산 + 대기) / 계산
  • $= (8 + 400) / 8$
  • $= 408 / 8$
  • $\approx 51$ 워프

즉 약 51개 워프(≈ 51 × 32 = 1,632 스레드)가 한 SM에 동시에 떠 있어야 400 사이클 지연을 완전히 은닉한다. 만약 워프가 10개뿐이면? 10 × 8 = 80 사이클 분량의 일을 한 뒤 다 같이 메모리를 기다리게 되어, 매 라운드 400 − 80 = 320 사이클 동안 ALU가 유휴 상태다 → 가용 처리량의 약 80%를 낭비한다.

교훈 세 가지:

1. 지연이 길수록(또는 산술 강도가 낮을수록) 더 많은 워프가 필요하다.
2. 워프(스레드)를 충분히 생성하지 못하면 지연이 노출되어 GPU가 제 성능을 못 낸다 — "작은 문제"가 GPU에서 느린 이유.
3. 그래서 GPU 커널 설계의 제1 목표는 충분한 병렬성(많은 워프)을 노출하는 것이다.

(실전에서는 워프당 동시에 여러 메모리 요청을 띄울 수 있고[memory-level parallelism], 명령 수준 병렬성[ILP]도 있어 필요 워프 수가 위 단순 산수보다 적을 수 있다. 점유율 vs ILP의 트레이드오프는 1.6. 그래도 "지연 ↔ 필요 워프 수"의 정성적 방향성은 그대로다.)

1.2.5 산술 강도 — 메모리 바운드 vs 연산 바운드 (개요)

위 산수에서 "산술 8 사이클 vs 메모리 400 사이클"의 비가 결정적이었다. 이 비를 일반화한 것이 산술 강도(arithmetic intensity, AI)다.

산술 강도 $=$ 수행한 연산 수(FLOPs) / 이동한 바이트 수(Bytes) [FLOP/Byte]

요지: 메모리에서 데이터 1바이트를 가져올 때마다 그걸로 연산을 몇 번 하느냐. 강도가 높으면 데이터를 충분히 재사용하는 것이고, 낮으면 메모리만 들락거리며 ALU는 유휴 상태다.

• 강도가 낮으면 메모리 바운드 — 대역폭이 천장이다. 더 빠른 ALU·텐서코어는 도움이 안 되고, 데이터 재사용을 늘려(타일링·융합) 강도를 키워야 한다.
• 강도가 높으면 연산 바운드 — ALU/텐서코어가 천장이다. 정밀도를 낮추거나(1.13) 텐서코어를 쓰는 것이(1.11) 답이다.

이 표가 시사하는 바: 딥러닝의 많은 연산(요소별·정규화·작은 행렬)은 의외로 메모리 바운드다. 그래서 융합(fusion, 여러 연산을 한 커널로 묶어 중간 결과를 메모리에 안 내림 — 3.3·3.6)이 큰 효과를 낸다. 반대로 큰 GEMM은 연산 바운드라 텐서코어·저정밀도가 결정적이다. 이 산술 강도와 두 천장을 한 그림에 겹쳐 놓는 루프라인 모델은 1.10에서 본격적으로, 측정으로 판별하는 법은 2.14(프로파일링)에서 다룬다.

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

여기까지 이해했다면 다음을 막힘없이 답할 수 있어야 한다.

• [ ] 워프가 무엇이고(32 스레드 lockstep), 다이버전스가 왜 성능을 저하시키는지 그림으로 설명한다.
• [ ] SIMT와 SIMD의 프로그래밍 모델 차이를 제시한다("스칼라로 작성하되 32개가 묶인다").
• [ ] 지연 400 사이클·산술 8 사이클일 때 필요한 워프 수를 암산으로 추정한다(≈ 51).
• [ ] 벡터 덧셈 vs 큰 GEMM을 산술 강도로 메모리/연산 바운드로 분류한다.

이해가 어렵다면: 워프 산수는 1.2.4.1을, SIMT가 SIMD처럼 보이면 1.2.3 표를 다시.

1.2.7 흔한 함정

• SIMT를 SIMD로 오해(가장 흔함). "벡터 폭을 직접 다뤄야 한다"고 생각하면 CUDA를 잘못 짠다 — 스칼라 스레드로 작성하되 워프·다이버전스·병합을 의식해야 한다(1.2.3).
• 워프를 충분히 생성하지 않음. 문제가 작아 스레드가 적으면 지연이 노출돼 GPU가 CPU보다 느릴 수도 있다.
• 다이버전스를 무시한 if/else 남발. 워프 안에서 경로가 갈리면 직렬화된다. 분기를 워프 경계에 맞추거나 산술로 푼다.
• 모든 커널을 연산 바운드로 착각해 텐서코어·저정밀도에만 집착. 메모리 바운드 커널엔 효과가 없다. 먼저 산술 강도로 판별하라(1.10·2.14).
• lockstep을 "스레드가 항상 동기"라고 오해. 같은 워프는 명령 진행이 묶이지만, 서로 다른 워프는 비동기다. 워프 간·블록 간 동기화는 명시적(2.8).
• 지연 은닉을 "캐시가 크다"로 착각. 은닉의 주력은 캐시가 아니라 많은 워프 + 거대 레지스터 파일 + 높은 대역폭이다.

1.2.8 연습문제

전문가 독자용 — 직접 산수하고 사고를 점검한다. 정답은 접어 둔다.

연습 1 (지연 은닉 산수). 메모리 지연 600 사이클, 워프가 메모리 요청 사이 평균 6 사이클의 독립 산술을 한다. 지연을 완전히 은닉하려면 SM에 최소 몇 워프가 필요한가? 32 스레드/워프라면 총 스레드는?

연습 2 (다이버전스). 한 워프(32 스레드)에서 if (threadIdx.x < 4) heavy(); else light();이고 heavy()가 light()의 10배 시간이다. 이 구간 시간은 다이버전스 없는 이상적 경우 대비 대략 어떻게 되나? 조건이 threadIdx.x < 32(전원 동일)라면?

연습 3 (산술 강도). FP32(4바이트) 기준, 벡터 덧셈 c[i]=a[i]+b[i]의 산술 강도(FLOP/Byte)를 구하고 메모리·연산 바운드 중 무엇인지 판단하라(읽기 2회 + 쓰기 1회).

연습 4 (융합의 효과). y = relu(gelu(x))를 (a) 커널 두 개로 따로, (b) 하나로 융합해 실행한다. 산술 강도 관점에서 왜 (b)가 빠른가?

1.2.9 요약 · 다음 장

• SIMT: 32 스레드 워프가 lockstep으로 같은 명령을 실행한다. 스칼라처럼 작성하되 워프·다이버전스·병합이 성능을 지배한다 — SIMD와 혼동하지 마라.
• 지연 은닉이 처리량의 비밀: 거대 레지스터 파일 위의 많은 워프를 스케줄러가 0-사이클로 교체해 ALU를 채운다. 충분한 워프(병렬성)를 노출하는 것이 커널 설계 제1 목표.
• 산술 강도가 커널의 메모리/연산 바운드를 정한다(루프라인 개요 → 1.10). 딥러닝은 GEMM(연산 바운드)과 요소별(메모리 바운드)이 섞여 있다.

다음 장: 1.3 칩 상위 구조 — GPC·TPC·SM 계층. 이 추상 모델이 실제 실리콘에서 어떤 물리적 계층으로 구현되는지, "코어 수"가 왜 성능을 단순 의미하지 않는지를 데이터시트로 읽는다.