컴퓨터공학 300 주제 시리즈의 099번째 글이다. 전체 지도는 여기.

한 줄 요약

CPU 가 소수의 복잡한 코어로 한 스레드의 지연 시간을 줄이는 데 트랜지스터를 쓴다면, GPU 는 단순한 연산 유닛 수천 개와 수만 개의 스레드로 처리량을 극대화한다. 메모리를 기다리는 스레드 묶음은 제쳐 두고 준비된 묶음을 실행하는 방식으로 지연을 “줄이지 않고 숨긴다”.

왜 필요한가

딥러닝 학습과 추론, 과학 계산, 영상 처리는 같은 연산을 엄청난 양의 데이터에 반복한다. 이런 일에서 GPU 는 CPU 보다 훨씬 많은 연산을 같은 시간에 해낸다. 그런데 모든 일이 GPU 에서 빨라지지는 않는다. 분기가 많거나, 데이터가 작거나, CPU–GPU 사이 전송이 잦은 작업은 오히려 느려질 수 있다.

GPU 가 어떤 구조인지 알면 “왜 배치 크기를 키우면 빨라지는가”, “왜 GPU 사용률 100% 인데 느린가”, “왜 if 문이 비싼가” 같은 질문에 답할 수 있다.

핵심 개념

CPU 와 GPU 의 설계 철학

CPU (지연 최적화)                      GPU (처리량 최적화)
┌───────────────────────┐            ┌───────────────────────────────┐
│ ┌──────┐ ┌──────┐     │            │ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ... │
│ │ 큰코어│ │ 큰코어│     │            │ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ... │
│ │분기예측│ │비순차 │     │            │ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ▪▪▪▪▪▪▪▪ ... │
│ └──────┘ └──────┘     │            │   작은 연산 유닛 수천 개         │
│      큰 캐시            │            │   작은 캐시, 거대한 레지스터 파일  │
└───────────────────────┘            └───────────────────────────────┘
  CPU GPU
목표 스레드 하나를 최대한 빨리 전체 처리량 최대화
코어 수 개~수백 개, 복잡함 수천 개 연산 유닛, 단순함
지연 대응 큰 캐시, 분기 예측, 비순차 실행으로 줄인다 다른 스레드로 전환해 숨긴다
메모리 큰 용량, 상대적으로 좁은 대역폭 작은 용량, 매우 넓은 대역폭(GDDR, HBM)
잘하는 일 분기 많은 일반 코드, 낮은 지연 대규모 데이터 병렬 연산

NVIDIA 용어로 본 GPU 구조

GPU 제조사마다 용어가 다르다. 이 글은 문서가 가장 풍부한 NVIDIA CUDA 용어를 쓴다.

  • SM (Streaming Multiprocessor): GPU 의 기본 블록. 연산 유닛 여러 개, 워프 스케줄러, 큰 레지스터 파일, 공유 메모리(L1 과 함께 쓰는 프로그래머 관리 온칩 메모리)를 갖는다. GPU 하나에 SM 이 수십~백여 개 있다.
  • 스레드 → 블록 → 그리드: 프로그래머는 커널 함수를 수만~수백만 개 스레드로 실행한다. 스레드는 블록으로 묶이고, 블록 하나는 SM 하나에 배정된다. 같은 블록의 스레드끼리만 공유 메모리와 동기화(__syncthreads())를 쓸 수 있다.
  • 워프(warp): 하드웨어는 스레드를 32개씩 묶어 워프 단위로 명령을 내린다. 워프의 32개 스레드는 같은 명령어를 각자의 데이터에 실행한다.

SIMT 와 분기 분산

NVIDIA 는 이 실행 방식을 SIMT(Single Instruction, Multiple Threads)라고 부른다. 프로그래머에게는 스레드마다 독립적인 코드로 보이지만, 하드웨어는 워프 단위로 같은 명령을 실행한다(#098 의 SIMD 와 닮았다).

워프 안의 스레드들이 if 에서 서로 다른 쪽으로 갈리면 하드웨어는 두 경로를 차례로 실행하고, 각 경로에서 해당하지 않는 스레드는 꺼 둔다. 이것을 분기 분산(warp divergence)이라 한다. 최악이면 효율이 경로 수만큼 나뉜다. 최신 세대는 스레드별 독립 스케줄링을 지원하지만, 갈라진 경로를 동시에 실행하는 것은 아니므로 성능 원칙은 같다. 같은 워프의 스레드는 같은 길로 가게 데이터를 배치한다.

지연 숨기기와 점유율

GPU 의 전역 메모리 접근은 수백 사이클이 걸릴 수 있다. GPU 는 이를 캐시로 줄이려 하기보다, SM 하나에 워프를 여러 개 올려 두고 메모리를 기다리는 워프 대신 준비된 워프를 매 사이클 골라 실행한다. 레지스터 파일이 거대한 이유가 이것이다. 모든 워프의 레지스터를 동시에 들고 있어서 전환 비용이 거의 없다.

SM 에 동시에 올라간 워프 수의 비율을 점유율(occupancy)이라 한다. 스레드당 레지스터나 공유 메모리를 많이 쓰면 올릴 수 있는 워프가 줄어 지연을 숨기기 어려워진다.

메모리 병합

워프의 32개 스레드가 연속된 주소를 읽으면 하드웨어가 이를 적은 수의 큰 메모리 트랜잭션으로 합친다(coalescing). 흩어진 주소를 읽으면 트랜잭션이 여러 개로 쪼개져 대역폭을 낭비한다. CPU 의 캐시 라인과 공간 지역성(#092) 이야기가 GPU 에서는 더 극단적으로 나타난다.

산술 강도와 루프라인

GPU 의 연산 능력은 메모리 대역폭보다 훨씬 빨리 늘었다. 그래서 작업이 “연산에 묶이는지, 메모리에 묶이는지” 가 핵심이다. 이를 판단하는 지표가 산술 강도(연산 수 ÷ 메모리에서 옮긴 바이트 수)다. 루프라인 모델은 산술 강도가 낮은 작업은 메모리 대역폭이, 높은 작업은 연산 성능이 상한이라고 정리한다.

  • 벡터 덧셈 c = a + b: 원소당 1번 연산에 12바이트 이동. 강도가 매우 낮아 메모리에 묶인다.
  • 행렬 곱 C = A @ B: 데이터를 재사용할수록 강도가 n 에 비례해 커진다. GPU 가 행렬 곱에 강한 이유이고, 딥러닝 연산을 가능한 한 큰 행렬 곱으로 바꾸려는 이유다. 텐서 코어 같은 행렬 전용 유닛도 여기에 맞춰져 있다.

직접 해 보기

GPU 가 없어도 구조의 핵심 두 가지는 계산으로 확인할 수 있다. python3 로 실행해 확인했다.

import random
WARP = 32

def warp_cycles(conds, cost_then=10, cost_else=10):
    """한 워프가 if/else 를 지날 때 걸리는 사이클(SIMT: 두 경로를 차례로 실행)"""
    t = any(conds)                 # 어떤 스레드라도 then 으로 가면 then 경로 실행
    e = not all(conds)             # 어떤 스레드라도 else 로 가면 else 경로 실행
    return t * cost_then + e * cost_else

random.seed(1)
N = 32 * 1000
cases = {
    "uniform (all true)":   [True] * N,
    "by warp (i//32 even)": [(i // WARP) % 2 == 0 for i in range(N)],
    "by thread (i even)":   [i % 2 == 0 for i in range(N)],
    "random":               [random.random() < 0.5 for _ in range(N)],
}
for name, c in cases.items():
    cyc = sum(warp_cycles(c[w:w + WARP]) for w in range(0, N, WARP))
    ideal = sum(10 for _ in range(0, N, WARP))
    print(f"{name:22s} cycles {cyc:6d}  efficiency {ideal / cyc:.0%}")

def intensity_vec_add(n):  return n / (3 * 4 * n)                 # c = a + b
def intensity_matmul(n):   return 2 * n**3 / (3 * 4 * n * n)      # C = A @ B, 이상적 재사용
print(f"vector add         {intensity_vec_add(10**6):.3f} FLOP/byte")
for n in (64, 1024, 8192):
    print(f"matmul n={n:<5d}    {intensity_matmul(n):.1f} FLOP/byte")

출력:

uniform (all true)     cycles  10000  efficiency 100%
by warp (i//32 even)   cycles  10000  efficiency 100%
by thread (i even)     cycles  20000  efficiency 50%
random                 cycles  20000  efficiency 50%
vector add         0.083 FLOP/byte
matmul n=64       10.7 FLOP/byte
matmul n=1024     170.7 FLOP/byte
matmul n=8192     1365.3 FLOP/byte

조건이 참인 스레드 비율은 “by warp” 와 “by thread” 가 똑같이 절반인데, 워프 단위로 갈리면 효율 100%, 스레드 단위로 섞이면 50% 다. 산술 강도는 행렬 곱이 벡터 덧셈보다 수천 배 높다. 행렬 곱 계산은 각 행렬을 메모리에서 딱 한 번만 옮긴다고 가정한 이상값이고, 실제 커널은 타일링으로 이 값에 가까워지려 한다. 사이클 비용 10 은 설명용 가정값이다.

현업에서는

  • 쿠버네티스에서 GPU 쓰기: GPU 는 장치 플러그인이 nvidia.com/gpu 같은 확장 리소스로 노드에 광고하고, 파드는 limits 에 정수 개수로 요청한다. 기본적으로 GPU 는 쪼개 쓰지 않고 통째로 할당된다. 여러 파드가 한 GPU 를 나눠 쓰려면 MIG(지원 GPU 에서 하드웨어 분할), MPS, 타임 슬라이싱 같은 별도 구성이 필요하다.
  • “GPU 사용률 100%” 의 함정: nvidia-smi 의 GPU 사용률은 샘플링 구간 중 커널이 하나라도 돌던 시간 비율에 가깝다. SM 이 얼마나 꽉 찼는지가 아니다. 실제 효율은 프로파일러로 SM 점유율과 메모리 처리량을 봐야 안다.
  • 전송 비용: CPU 메모리에서 GPU 메모리로 옮기는 PCIe 전송(#097 DMA)은 GPU 내부 대역폭보다 훨씬 느리다. 작은 연산마다 데이터를 왕복시키면 GPU 가 놀게 된다. 데이터를 GPU 에 오래 두고, 고정(pinned) 메모리와 비동기 전송으로 계산과 전송을 겹친다.
  • 배치 크기: 추론 서버가 요청을 모아 배치로 처리하는 이유는 산술 강도와 점유율을 높이기 위해서다. 대신 개별 요청의 지연은 늘어난다. 처리량과 지연의 거래(#090)가 여기서도 나온다.

확인 문제

  1. CPU 와 GPU 가 메모리 지연에 대응하는 방식의 차이는.
  2. 워프 분기 분산이란 무엇이고 어떻게 줄이는가.
  3. 벡터 덧셈이 GPU 에서 연산 성능을 다 쓰지 못하는 이유는.
  4. 쿠버네티스 파드가 GPU 0.5 개를 기본 설정으로 요청할 수 없는 이유는.

풀이

  1. CPU 는 큰 캐시·분기 예측·비순차 실행으로 지연 자체를 줄이고, GPU 는 많은 워프를 올려 두고 기다리는 워프 대신 다른 워프를 실행해 지연을 숨긴다.
  2. 같은 워프의 스레드가 분기에서 서로 다른 경로로 가 두 경로를 차례로 실행하느라 효율이 떨어지는 현상이다. 같은 워프의 스레드가 같은 경로로 가도록 데이터와 작업을 배치해 줄인다.
  3. 연산 한 번에 옮겨야 할 데이터가 많아(산술 강도가 낮아) 메모리 대역폭이 먼저 한계에 닿기 때문이다.
  4. GPU 같은 확장 리소스는 정수 단위로만 요청할 수 있고, 장치 플러그인이 GPU 를 통째로 할당하기 때문이다. 공유하려면 MIG·타임 슬라이싱 같은 별도 구성이 필요하다.

더 읽을거리 (References)