① GPU 아키텍처

이 장의 목표는 딱 하나입니다 — "GPU는 왜 빠른가, 그 안에서 무슨 일이 벌어지는가"의 직관을 잡는 것. 이 그림이 머릿속에 그려지면, 다음 장의 CUDA 코드가 왜 그렇게 생겼는지, 스케줄링이 왜 그렇게 도는지가 술술 이해됩니다. 용어가 많이 나오지만 외우려 하지 말고, "아 이런 게 있구나" 정도로 읽어주세요.

1.1 CPU vs GPU — 왜 둘은 다르게 생겼나

CPU와 GPU는 둘 다 "계산하는 칩"이지만, 완전히 다른 목표로 만들어졌습니다. 비유로 시작할게요.

• CPU = 박사 몇 명. 어렵고 복잡한 문제를 하나씩 빠르게 풉니다. 똑똑하지만 사람 수가 적어요(코어 수~수십 개).
• GPU = 초등학생 수천 명. 한 명 한 명은 단순 계산만 하지만, 수천 명이 같은 문제를 동시에 풀면 어마어마한 양을 처리합니다.

그래서 "어려운 문제 하나"는 CPU가, "쉬운 문제 수천 개"는 GPU가 잘합니다. 좀 더 정확히 말하면:

• CPU = 지연시간(latency) 최적화 — 작업 하나를 최대한 빨리 끝내도록 설계. 큰 캐시, 분기 예측, 비순차 실행 같은 "똑똑한 제어 장치"에 칩 면적을 많이 씁니다. 한 명이 빠름
• GPU = 처리량(throughput) 최적화 — 같은 연산을 수천~수만 데이터에 동시에. 똑똑한 제어 장치는 줄이고 계산기(ALU)를 최대한 많이 박아넣습니다. 여럿이 많이

	CPU	GPU
코어 수	수~수십 개	수천 개 (CUDA 코어)
코어 하나의 성능	매우 빠름·복잡	단순·느림
캐시(빠른 임시 메모리)	크고 계층적	상대적으로 작음
"느린 메모리" 대처법	큰 캐시로 줄임	많은 일꾼으로 가림
잘하는 일	순차적·분기 많은 로직	대규모 데이터 병렬
비유	박사 몇 명	초등학생 수천 명

👤 CPU 개발자를 위한 — 개념 1:1 대응표

CPU 개발만 해오셨다면, GPU 용어가 익숙한 단어인데 의미가 미묘하게 달라 헷갈리기 쉽습니다. 먼저 이 표로 "내가 알던 그것과 무엇이 같고 다른지"를 잡아두세요. (각 항목은 해당 장에서 자세히 다룹니다.)

CPU에서 알던 것	GPU의 대응	핵심 차이 (이게 의문의 원인)
코어 (완전한 코어)	SM	SM 하나가 "코어"에 가까움. "CUDA 코어 수천 개"의 코어는 코어가 아니라 ALU 레인임
스레드 (pthread)	thread (CUDA)	GPU 스레드는 초경량. OS가 안 만들고, 수천 개가 생기며, 32개씩 warp로 묶여 같은 명령을 함께 실행
SIMD 레인 (AVX 등)	warp의 32 레인	GPU는 SIMD를 "스레드"로 추상화 → 스칼라 코드처럼 짜도 병렬화됨
malloc / memcpy	cudaMalloc / cudaMemcpy	GPU는 주소 공간이 따로. CPU 포인터를 GPU에서 못 씀, 복사 필요
함수 호출 (동기)	커널 <<<>>> (비동기)	호출 = "큐에 주문 넣기". CPU는 결과를 안 기다리고 다음 줄로
스레드 풀 / 작업 큐	스트림(stream)	스트림은 "일꾼"이 아니라 순서 큐. 다른 스트림끼리 병렬
pthread_barrier	__syncthreads()	블록 안에서만 됨. 블록(=다른 SM)끼리는 커널 도중 동기화 불가
L1/L2 캐시 (자동)	shared memory (수동)	shared는 자동 캐시가 아니라 프로그래머가 직접 채우는 스크래치패드
OS 스케줄러 (선점적)	HW warp 스케줄러 (비선점적)	OS가 개입 안 함. warp 전환은 0-비용이지만, 블록은 끝까지 실행(선점 약함)
context switch (비쌈)	warp 전환 (≈0 비용)	모든 warp의 상태가 레지스터에 동시에 상주 → 저장/복원이 없음

1.2 SM 내부 구조 — GPU 안의 "작은 공장"

GPU 칩 하나를 열어보면, 안에 SM(Streaming Multiprocessor)이라는 단위가 수십~수백 개 박혀 있습니다. SM은 우리말로 풀면 "스트리밍 멀티프로세서" — 쉽게 말하면 그 자체로 작동하는 작은 프로세서(공장 한 동)입니다. GPU는 이 작은 공장들의 집합이에요.

GPU (큰 공장 단지)
 ├─ SM 0 ─┐  각 SM(공장 한 동)이 일감(블록)을 받아 독립적으로 처리
 ├─ SM 1  │  누가 어느 공장에 일감을 나눠줄지는
 ├─ ...   │  GPU 안의 'GigaThread 엔진'이 정함
 └─ SM N ─┘

중요한 점: 우리가 GPU에 일을 시키면, 그 일은 "블록(block)" 단위로 쪼개져 SM에 하나씩 배정됩니다 (블록은 ②장에서 자세히). 즉 SM이 실제로 일하는 곳이에요.

SM 안에는 무엇이 들어있나

SM 하나를 더 확대하면, 다음 부품들이 들어 있습니다 (예시 수치는 Ampere/Hopper 세대 기준이고, 아키텍처마다 조금씩 다릅니다):

부품	쉽게 말하면	역할
CUDA 코어 (~128개/SM)	계산기(일꾼)	덧셈·곱셈 같은 일반 산술
Tensor 코어 (~4개/SM)	행렬 전문 계산기	행렬 곱-누적(딥러닝에 특화)
Warp 스케줄러 (~4개/SM)	작업 반장	매 순간 어떤 일꾼 그룹을 일 시킬지 결정
레지스터 파일 (~256KB/SM)	일꾼 개인 메모장	각 스레드의 지역 변수 저장
Shared memory / L1 (~최대 228KB/SM)	공동 작업대	블록 안 스레드들이 함께 쓰는 빠른 메모리
LD/ST, SFU	운반·특수계산 담당	메모리 읽고 쓰기, sin·sqrt 등

세대(아키텍처)와 Compute Capability

위 수치들이 "Ampere 기준"이라고 했는데, NVIDIA GPU는 몇 년 주기로 세대가 바뀌고 그때마다 SM 구성·자원 한도가 조금씩 달라집니다. 세대를 숫자로 표시한 게 compute capability(CC)예요 — 코드를 컴파일할 때 "어느 세대용인가"를 지정하는 기준이 됩니다(②장 컴파일에서 다시 등장).

세대	CC	대표 GPU	기억할 변화
Volta	7.0	V100	Tensor 코어 등장, warp 스레드의 독립 PC
Ampere	8.0 / 8.6	A100 / RTX 30	MIG, 비동기 복사(cp.async)
Hopper / Ada	9.0 / 8.9	H100 / RTX 40	Thread Block Cluster, TMA
Blackwell	10.x / 12.x	B200 / RTX 50	FP4 등 저정밀 Tensor 코어 강화

지금 단계에선 "세대마다 수치가 다르니, 내 GPU의 정확한 값은 cudaGetDeviceProperties나 nvidia-smi로 확인한다" 정도면 충분합니다.

1.3 SIMT · Warp — GPU가 일하는 "단위"

GPU의 일꾼(스레드)들은 한 명씩 따로 노는 게 아닙니다. 여기서 GPU 특유의 두 개념이 나옵니다.

SIMT — "한 명령을 여러 명이 동시에"

SIMT는 Single Instruction, Multiple Threads(하나의 명령, 여러 스레드)의 약자입니다. 쉽게 말하면 "반장이 '3번 줄 더해!'라고 외치면, 여러 일꾼이 각자 자기 데이터로 동시에 더하는" 방식이에요. 프로그래머는 일꾼 한 명의 입장에서 평범하게 코드를 쓰면, 하드웨어가 알아서 수천 명에게 같은 명령을 뿌려줍니다.

(참고: CPU에도 비슷한 SIMD라는 게 있는데, GPU는 이걸 "스레드"라는 더 다루기 쉬운 단위로 추상화한 게 SIMT입니다. 차이는 지금 몰라도 됩니다.)

Warp — 32명씩 묶인 한 조

스레드는 정확히 32개씩 묶여서 한 조로 움직입니다. 이 32명짜리 조를 warp(워프)라고 부릅니다(NVIDIA GPU 기준). 반장(warp 스케줄러)은 개별 스레드가 아니라 이 warp 단위로 명령을 내립니다. 즉 warp가 GPU 스케줄링의 최소 단위예요.

# GPU 일감의 계층 구조 (큰 것 → 작은 것)
Grid    (커널 전체 = 일감 전부)
 └─ Block   (한 SM/공장에 배정되는 일감 한 덩어리)
      └─ Warp   (32개 스레드 = 한 조)
           └─ Thread (일꾼 한 명)

1.4 메모리 계층 — GPU 성능의 진짜 병목

놀랍게도, 많은 GPU 프로그램이 느린 이유는 계산이 아니라 메모리 기다림 때문입니다(이걸 "memory-bound"라고 해요). 일꾼은 빠른데, 일감(데이터)이 창고에서 늦게 도착하면 다들 손 놓고 기다리는 거죠. 그래서 메모리의 종류·속도·공유 범위를 아는 게 최적화의 절반입니다.

GPU 메모리는 "가깝고 빠르지만 작은 것"부터 "멀고 느리지만 큰 것"까지 계층으로 되어 있습니다 (가까울수록 빠름):

메모리	누가 같이 쓰나(공유 범위)	대략의 지연	비유
레지스터	스레드 혼자	~1 사이클	손에 든 메모장
Shared Memory	같은 블록끼리 (SM 내부)	~20–30 사이클	책상 위 공용 작업대
L1 캐시	한 SM	~30 사이클	사무실 책장
L2 캐시	모든 SM	~200 사이클	층 공용 자료실
Global Memory (VRAM)	모든 스레드	~400–800 사이클 (느림!)	건물 밖 대형 창고
Constant Memory	모든 스레드 (읽기 전용)	캐시 적중 시 매우 빠름	벽에 붙은 공지문
Local Memory	스레드 혼자 (이름과 달리 global에 있음!)	global과 동일 (느림)	메모장이 넘쳐 창고에 맡긴 짐

(지연 수치는 세대마다 다른 대략값입니다 — "위아래 층 사이가 수십 배씩 차이난다"는 감만 잡으세요.) 마지막 두 줄을 짚고 갈게요:

• Constant memory — 커널 내내 안 바뀌는 상수(계수, 설정값 등)를 두는 64KB의 읽기 전용 공간(__constant__). 전용 캐시가 있어서, warp의 32명이 모두 같은 주소를 읽으면 방송(broadcast)처럼 한 번에 처리됩니다. "모두가 같은 값을 읽는" 패턴에 최적.
• Local memory — 함정 주의. 이름은 "로컬"인데 실제 위치는 느린 global memory입니다. 스레드가 레지스터를 너무 많이 필요로 하면, 컴파일러가 넘친 변수를 여기로 쫓아내요(register spill). 커널이 이유 없이 느리면 spill부터 의심 — nvcc --ptxas-options=-v로 확인할 수 있습니다.

핵심은 "global memory(대형 창고)는 느리다"입니다. 그래서 고성능 커널의 두 가지 황금 기법이 나옵니다:

• 메모리 병합(coalescing) — 32명이 창고의 연속된 칸을 한꺼번에 가져오면 한 번에 운반 가능 (효율↑).
• 재사용(shared memory) — 창고에서 한 번 가져온 데이터를 공용 작업대(shared)에 올려두고 여러 번 우려먹기.

이 둘은 ⑥ 최적화에서 코드와 함께 자세히 다룹니다. 지금은 "먼 메모리는 느리니, 가까운 데로 가져와 재사용한다"만 가져가세요.

1.5 Occupancy — SM을 일꾼으로 얼마나 채우나

Occupancy(점유율)는 한마디로 "한 SM(공장)에 일꾼(warp)을 얼마나 꽉 채웠나"입니다. SM이 동시에 올릴 수 있는 warp 수는 앞에서 말한 "정해진 예산" 중 가장 빡빡한 것에 막힙니다.

한도(예산)	예시 값 (Ampere)	의미
최대 warp / SM	64 warp (= 2048 스레드)	한 SM에 올릴 수 있는 일꾼 상한
최대 블록 / SM	~32개	한 SM에 올릴 수 있는 일감 덩어리 수
레지스터 / SM	65,536개	일꾼이 메모장을 많이 쓰면 → 적게 올라감
shared memory / SM	~164KB	블록이 작업대를 많이 쓰면 → 적게 올라감

예시로 계산해볼게요. 어떤 커널이 일꾼 한 명당 레지스터(메모장)를 64개 쓴다고 합시다.

일꾼 2048명 × 메모장 64개 = 131,072개 필요
그런데 SM의 메모장은 65,536개뿐 → 절반인 1024명(=32 warp)만 올릴 수 있음
→ occupancy = 32 / 64 = 50%

메모장을 적게 쓰는 커널로 바꾸면 더 많은 일꾼이 올라가 occupancy가 올라갑니다.