How Deep Learning Computes

Warp와 Block Shape

cudawarpblockperformance

CUDA에서 thread는 코드상 최소 실행 단위처럼 보인다. 하지만 NVIDIA GPU 하드웨어는 thread를 보통 32개씩 묶어 실행한다.

1 warp = 32 threads

block 안의 warp

blockDim.x = 128이면 block 하나에는 warp가 4개 있다.

warp 0: threadIdx.x 0  ~ 31
warp 1: threadIdx.x 32 ~ 63
warp 2: threadIdx.x 64 ~ 95
warp 3: threadIdx.x 96 ~ 127

warp 안에서의 위치는 lane이라고 부른다.

int warp_id = threadIdx.x / 32;
int lane_id = threadIdx.x % 32;

왜 32의 배수로 잡는가

block thread 수가 40이면:

warp 0: 32 threads 사용
warp 1: 8 threads만 사용

두 번째 warp의 나머지 lane은 놀 수 있다. 그래서 초반에는 block 총 thread 수를 보통 32의 배수로 잡는다.

128 threads/block
256 threads/block

너무 크게 잡으면 안 되는 이유

block 하나의 thread 수에는 하드웨어 제한이 있다. 일반적으로 최대 1024 threads/block이다.

또 너무 큰 block은 SM에 동시에 올라갈 수 있는 block 수를 줄일 수 있다. 그러면 메모리 지연을 숨기는 능력이 떨어지고 occupancy가 낮아질 수 있다.

초반 기본값은 다음처럼 잡으면 무난하다.

1D elementwise:
  256

2D matrix:
  16 x 16 = 256
  32 x 8  = 256

3D tensor:
  8 x 8 x 4 = 256

인터랙티브 설계

  • block 안 thread들을 32개씩 색이 다른 warp로 묶어 보여준다.
  • block 총 thread 수를 40, 64, 128, 256으로 바꿔 warp가 어떻게 채워지는지 보여준다.
  • 2D block shape (16,16), (32,8), (8,32)가 모두 256 threads라는 점을 보여준다.

확인

  • blockDim.x = 256이면 warp는 몇 개인가?
  • blockDim.x = 40이면 왜 비효율적일 수 있는가?
  • 2D block에서 (16,16)과 (32,8)은 총 thread 수가 어떻게 같은가?