GPU Architecture Primer

GPU kernel 최적화를 이해하려면 먼저 GPU가 계층적으로 생겼다는 점을 잡아야 한다.

GPU
  many SMs
    many warps
      32 threads per warp

Compute hierarchy

GPU의 큰 연산 단위는 SM, Streaming Multiprocessor다.

GPU = many SMs
SM  = block이 배정되어 실행되는 장소

CUDA에서 block을 launch하면 hardware scheduler가 block들을 SM에 배치한다. block 하나는 한 SM에서 실행되고, 한 SM은 resource가 허용하는 만큼 여러 block을 동시에 실행할 수 있다.

block 안 thread들은 32개 단위의 warp로 묶여 실행된다.

1 warp = 32 threads

그래서 CUDA 성능을 볼 때는 thread 하나만 보지 말고, 같은 warp 안 thread들이 어떤 주소를 읽고 어떤 분기를 타는지 봐야 한다.

Memory hierarchy

GPU memory도 계층적이다.

register
  thread private, fastest, tiny

shared memory / L1
  block 안 thread들이 공유, fast, small

L2 cache
  all SMs shared

global memory / HBM
  largest, slowest relative to on-chip memory

HBM은 크고 대역폭도 높지만, SM 가까이에 있는 register/shared memory보다 훨씬 멀다. 그래서 빠른 kernel은 보통 이런 전략을 쓴다.

HBM에서 한 번 가져온다
shared memory/register 근처에서 최대한 재사용한다
중간 결과를 HBM에 자주 쓰지 않는다

Kernel launch

CUDA kernel은 CPU host code가 GPU device에 launch한다.

kernel<<<gridDim, blockDim>>>(...);

host code는 보통 다음 일을 한다.

1. device memory allocation
2. host -> device copy
3. kernel launch
4. device -> host copy
5. device memory free

kernel 자체는 GPU에서 많은 thread가 동시에 실행하는 device code다.

왜 이 카드가 필요한가

뒤에서 나오는 최적화는 모두 이 구조를 이용한다.

memory coalescing:
  warp 안 thread들이 연속 주소를 읽게 한다

tiling:
  HBM에서 가져온 tile을 shared memory에서 재사용한다

occupancy:
  SM에 active warp를 충분히 올려 latency를 숨긴다

fused kernel:
  중간 결과를 HBM에 쓰고 다시 읽는 일을 줄인다

확인

block은 어디에 배정되어 실행되는가?
warp는 몇 개 thread로 구성되는가?
shared memory와 global memory의 차이는 무엇인가?
HBM traffic을 줄이는 것이 왜 중요한가?