CPU는 적은 수의 복잡한 코어로 서로 다른 작업(스레드)을
독립적으로 빠르게 처리하도록 설계됐다. 반면 GPU는 수천 개의 단순한 코어로
같은 연산을 수많은 데이터에 동시에 반복하도록 설계됐다 — 그래서 이름이 SIMT
(Single Instruction, Multiple Threads)다.
CPU (MIMD에 가까움)
코어 수: 보통 4~64개, 코어당 트랜지스터 다수
비순차 실행, 분기 예측, 깊은 파이프라인
각 코어가 완전히 다른 명령어 실행 가능
지연시간(latency) 최소화에 최적화
예: 웹 서버 요청 처리, OS 스케줄링
GPU (SIMT)
코어(CUDA 코어) 수: 수천~수만 개, 코어당 매우 단순
분기 예측 없음, 얕은 파이프라인, 대신 병렬성으로 승부
워프(Warp) 32개 스레드가 한 몸처럼 같은 명령어 실행
처리량(throughput) 극대화에 최적화
예: 행렬 곱셈, 딥러닝, 3D 렌더링, 시뮬레이션
워프(Warp)란?
NVIDIA GPU는 스레드 32개를 워프(Warp) 한 묶음으로 스케줄링한다
(AMD는 64개짜리 웨이브프론트). 워프 안의 32개 스레드는 매 사이클마다
정확히 같은 명령어를 서로 다른 데이터에 대해 동시에 실행한다 —
마치 32명이 동일한 동작을 동시에 하는 군무와 같다. 문제는 if 문처럼
스레드마다 다른 경로로 갈라져야 할 때 생긴다. 이를
워프 분기 발산(Warp Divergence)이라 한다.
인터랙티브 데모 — 워프 분기 발산 시뮬레이터
아래 32개 상자는 워프 하나에 속한 스레드다. 각 스레드는 자신만의 데이터 값을 갖고 있다.
분기 조건을 고르고 "커널 실행"을 누르면, GPU가 실제로 각 경로를 어떻게 순차 처리하는지
사이클 단위로 재생된다.
사이클: 0실행된 고유 경로: 0평균 워프 활용률: -
커널 실행 버튼을 눌러 시뮬레이션을 시작하세요.
분기 발산은 코드에서 이렇게 생긴다
// CUDA 커널: 워프 안에서 threadIdx.x 값에 따라 분기__global__voidprocess(int* data) {
int tid = threadIdx.x;
if (data[tid] % 2 == 0) {
// 짝수인 스레드만 이 경로 실행 (나머지는 대기/마스킹)
data[tid] = data[tid] * 2;
} else {
// 홀수인 스레드만 이 경로 실행
data[tid] = data[tid] + 1;
}
// → 워프 하나에 짝수·홀수가 섞여 있으면 두 경로 모두 순차 실행됨
}
핵심 요약
SIMT 실행 모델
워프 내 32개 스레드가 매 사이클 동일 명령어를 실행. 하드웨어가 스레드별 레지스터는
따로 유지하되 명령어 인출·해석은 공유한다.
분기 발산 비용
워프 안에서 경로가 N갈래로 갈라지면, GPU는 각 경로를 순차적으로 실행하며
해당 경로에 속하지 않는 스레드는 결과를 버리는 예측 실행(마스킹) 상태로 대기한다.
실전 최적화 힌트
워프 내 스레드들이 같은 분기를 타도록 데이터를 정렬하거나, 분기 대신 산술식으로
바꾸면(branchless) 발산을 줄여 처리량을 높일 수 있다.