Chapter 4

GPU 아키텍처와 워프 스케줄링

스레드를 발사한 뒤, 칩 안에서 실제로 무슨 일이 벌어지는가

4.1 현대 GPU 아키텍처: SM이라는 작은 공장

지금까지 우리는 GPU를 "수천 개 코어가 든 마법 상자"로 다뤘다. 이 장에서 그 상자를 열어 본다. 현대 NVIDIA GPU는 스트리밍 멀티프로세서(Streaming Multiprocessor, SM)의 배열로 구성되어 있다. SM 한 개가 작은 공장이고, GPU는 그런 공장을 수십 개 갖춘 산업단지인 셈이다.

구체적인 숫자를 보자. 데이터센터급 칩 A100(컴퓨트 캐퍼빌리티 8.0)은 108개의 SM을, H100(9.0)은 132개의 SM을 갖는다. 한 SM 안에는 다음이 들어 있다.

SM 바깥에는 모든 SM이 공유하는 L2 캐시(A100 40 MB, H100 50 MB)와, 칩 밖의 HBM DRAM(A100 40~80 GB, H100 80 GB+, 대역폭 1.5~3 TB/s)이 있다. 메모리 계층은 위에서 아래로 점점 크고 점점 느려진다 ─ 레지스터(수 클럭) → 공유 메모리/L1(20 클럭 안팎) → L2(200 클럭 안팎) → DRAM(400~800 클럭). 이 차이를 직관적으로 받아들이는 것이 5장 공유 메모리 최적화의 출발점이다.

GPU ├─ SM 0 ├─ SM 1 ... ├─ SM 107 ┐ │ ┌────┐ │ ┌────┐ │ ┌────┐ │ on-chip │ │ R/L1/Shared │ │ R/L1/Shared │ │ │ │ Tensor │ │ │ Tensor │ │ │ │ Cores │ │ │ Cores │ │ │ └────┘ │ └────┘ │ └────┘ ┘ └────┬───────┬─── ... ────────┘ └──── L2 cache (40~50 MB) ────┐ on-chip ┘ HBM DRAM (40~80 GB) ─ off-chip

그림 4.1 — A100급 GPU의 계층 구조 개략

4.2 블록 스케줄링: SM에 누가 들어갈까

커널을 발사하면 그리드는 블록 단위로 SM들에 분배된다. 한 블록은 통째로 한 SM에 들어가고, 일단 들어가면 그 블록은 그 SM 안에서 끝까지 산다. 다른 SM으로 옮겨 가지 않는다. 이걸 블록 단위 원자성이라 부른다.

한 SM에는 여러 블록이 동시에 거주할 수 있다. 자원만 허락하면 말이다. A100은 SM당 최대 32 블록, 최대 2048 스레드까지 거주시킬 수 있다(예를 들어 256 스레드 블록 8개, 1024 스레드 블록 2개 등). 이 거주 가능 비율을 점유율(occupancy)이라 부른다 ─ 4.7절의 주제다.

블록을 SM에 매핑하는 일은 GPU 하드웨어 스케줄러가 동적으로 한다. 프로그래머가 "블록 17은 SM 5에" 같은 지정을 할 수 없다(그리고 그게 의도된 설계다). 그래서 같은 코드가 SM이 14개인 노트북 GPU에서도 132개인 H100에서도 그냥 돈다 ─ 큰 칩에서는 더 많은 블록이 동시에 살게 될 뿐이다. 이 성질이 다음 절의 핵심 약속, 투명한 확장성으로 이어진다.

4.3 동기화와 투명한 확장성

같은 블록 안의 스레드끼리는 협력할 수 있다. 협력의 두 도구가 (1) 공유 메모리, (2) __syncthreads()다. 다음 5장에서 본격적으로 쓸 패턴을 미리 보면 이렇다.

__global__ void exampleSync(float* X) {
    __shared__ float tile[256];
    int tid = threadIdx.x;

    tile[tid] = X[blockIdx.x * 256 + tid];   // ① 글로벌 → 공유
    __syncthreads();                          // ② 모두 ①을 끝낼 때까지 대기

    // 이 시점부터는 같은 블록 누구나 tile 전체를 안전하게 읽을 수 있다
    float v = tile[(tid + 1) % 256];
    // ...
}

__syncthreads()는 같은 블록 안의 모든 스레드가 이 줄에 도달할 때까지 막는 장벽(barrier)이다. 비용은 보통 수 클럭 정도지만, 워프 일부가 멀리 가 있으면 그만큼 다 같이 기다리니 워프 다이버전스(4.5절)와 결합하면 위험하다.

그런데 블록 사이에는 동기화가 없다. 한 커널 호출 안에서 블록 0과 블록 1이 서로 데이터를 주고받으며 기다리게 할 방법이 (기본적으로) 없다는 뜻이다. 이게 처음엔 제약처럼 보이지만, 사실은 GPU의 가장 큰 약속을 떠받치는 설계다.

블록끼리 독립이라는 건 GPU 하드웨어가 블록을 어떤 순서로 어떤 SM에 던져 넣어도 결과가 같다는 뜻이다. 그래서 같은 .cu 코드가 SM 14개짜리 GPU에선 14개 SM에 순서대로 흘려 들어가고, 132개짜리 GPU에선 132개 SM이 한꺼번에 받아 처리한다. 프로그래머가 칩 종류별로 코드를 다시 짜지 않아도 된다. 이걸 투명한 확장성(transparent scalability)이라 부른다. CPU의 멀티스레딩에서는 흔히 못 누리는 호사다.

감 잡기

"블록 간 동기화가 필요하면 어떻게 하는가?" 답은 단순하다 ─ 커널을 끝내고 다시 발사한다. 한 커널의 모든 블록이 끝난 시점이 자연스러운 글로벌 동기화 지점이 된다. 큰 알고리즘은 흔히 여러 개의 작은 커널로 쪼개진다.

4.4 워프와 SIMD: 32명 한 조

한 블록에 256 스레드가 있다고 하자. SM은 이 256명을 한 사람씩 따로 일 시키지 않는다. 32명씩 묶어 한 단위로 다룬다. 이 32 스레드 묶음을 워프(warp)라 부른다. 워프는 NVIDIA GPU 하드웨어의 기본 실행 단위다 ─ 256 스레드 블록은 사실 8 워프다.

워프 안의 32 스레드는 매 사이클마다 같은 명령을 동시에 실행한다. 데이터만 다르다. 이 모델을 SIMT(Single Instruction, Multiple Thread)라 부른다 ─ 본질적으로 SIMD인데 프로그래밍 모델이 한 스레드 시점으로 보이게 추상화돼 있는 것이다. 어떻게 보면 32명이 한 줄로 서서, 지휘자(스케줄러)가 "더하기!" 하면 모두 자기 데이터에 더하기를 하고, "메모리 읽기!" 하면 모두 자기 주소에서 읽는다.

이 사실은 두 가지 실용적 결과를 낳는다.

첫째, blockDim.x는 32의 배수가 좋다. 그렇지 않으면 마지막 워프가 일부만 살아 있는 상태로 발사된다. 가령 blockDim.x = 100이면 한 블록은 4 워프(128 스레드)를 잡아먹지만 28명이 빈자리로 시작한다. 그 28명도 매 명령에 같이 따라다니지만 결과는 마스크되어 버려진다 ─ 자원 낭비다. 그래서 블록 크기는 거의 항상 32, 64, 128, 256, 512, 1024 중에서 고른다.

둘째, 워프 안 32개의 메모리 접근은 한 트랜잭션으로 묶을 수 있다. 32 스레드가 연속한 128바이트(=32 × 4바이트 float)를 읽으면 한 번의 DRAM 트랜잭션으로 끝난다. 만약 워프 안에서 주소가 흩어져 있으면 트랜잭션이 여러 번으로 쪼개진다. 이게 메모리 합치기(coalescing) 문제이고, 6장의 주제다.

4.5 컨트롤 다이버전스: if 한 줄의 비용

워프가 32명이 같은 명령을 따라간다는 사실에서 자연스러운 질문이 나온다. "if/else로 갈라지면 어떻게 되나?" 답은 직관적이지만 처음 듣기엔 충격적이다. 워프는 양쪽 분기를 모두 직렬로 실행한다. 일부 스레드는 if 쪽으로 갈 때 나머지는 마스크되어 NOP을 따라가고, 그다음 else 쪽을 실행할 때 역할이 바뀐다. 두 분기를 다 돈 뒤 다시 합쳐진다. 이걸 컨트롤 다이버전스(control divergence)라 한다.

비용은 단순하다. 분기 두 갈래 중 어느 한 워프 안에 둘 다 활성 스레드가 있으면, 둘 다 도는 만큼의 시간이 든다. 워프 안에서 16명이 if, 16명이 else면 두 갈래의 일을 다 한 비용을 그 워프가 짊어진다. 32명 모두 같은 쪽이면 다이버전스 비용은 0이다 ─ 이쪽이 핵심이다.

그래서 다음과 같은 코드는 의외로 무해할 수 있다.

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
    out[i] = doWork(in[i]);
}

n이 1024이고 blockDim.x가 256이면 블록은 4개. 마지막 블록에서 한 워프가 다이버전스를 겪을 가능성도 있지만, 대부분의 워프는 32명 모두 i < n을 만족하거나 모두 만족하지 않는다. 워프 경계에 분기가 정렬되어 있으면 다이버전스 비용은 거의 없다.

반면 다음 같은 코드는 괴롭다.

if (threadIdx.x % 2 == 0) {
    a();
} else {
    b();
}

워프 안에서 짝수와 홀수 스레드가 정확히 반반 갈린다 ─ 32 스레드 모두에서 다이버전스가 발생한다. 두 분기가 둘 다 무거우면 워프의 처리량이 사실상 절반으로 떨어진다.

함정

다이버전스는 분기뿐 아니라 루프 길이에서도 생긴다. for (int k = 0; k < lengths[tid]; ++k)처럼 스레드마다 다른 횟수를 도는 루프는, 한 워프가 가장 긴 lengths[tid]만큼 끝까지 따라간다. 짧은 스레드들은 다 끝났는데도 마스크된 채 같이 도는 셈이다. 이런 패턴이 보이면 의심해야 한다.

4.6 워프 스케줄링과 레이턴시 톨러런스

여기서 GPU 설계의 가장 영리한 부분이 등장한다. SM 한 개에 워프 64개(예: 2048 스레드)가 거주하고 있는데 워프 스케줄러는 4개라고 하자. 매 클럭 스케줄러는 자기가 담당하는 워프 무리 중에서 "지금 실행 가능한" 워프 한 개를 골라 명령을 발사한다.

"실행 가능"이 핵심이다. 워프가 글로벌 메모리 읽기를 시작하면 결과가 도착할 때까지 400~800 클럭이 걸린다. 그동안 그 워프는 잠재워둔다. 스케줄러는 다른 실행 가능한 워프를 즉시 골라 발사한다. 메모리에서 결과가 도착해 잠자던 워프가 깨어나면 다시 풀에 들어간다. 이 제로 오버헤드 컨텍스트 스위치가 GPU의 비밀 무기다.

왜 오버헤드가 0인가? CPU의 컨텍스트 스위치는 레지스터를 메모리에 저장하고 다른 스레드의 레지스터를 다시 로드해야 한다 ─ 수백 클럭. GPU에서는 SM의 거대한 레지스터 파일(64K 레지스터)이 거주 중인 모든 워프의 상태를 동시에 들고 있다. 스레드 A의 레지스터와 스레드 B의 레지스터가 물리적으로 다른 자리에 살아 있어서 스위치 자체가 자유다.

그래서 GPU는 메모리 레이턴시를 "더 빠른 메모리"로 풀지 않는다(그러기엔 너무 비싸다). 대신 충분히 많은 워프로 가린다. 한 SM에 워프가 충분히 많으면 어떤 워프가 메모리를 기다릴 때 다른 워프가 일을 한다. 이걸 레이턴시 톨러런스(latency tolerance)라 한다. CPU의 비순차 실행(out-of-order)이 한 스레드 안에서 명령 단위로 같은 일을 하는 거라면, GPU는 워프 단위로 그것을 한다 ─ 훨씬 거칠지만 훨씬 큰 규모로.

4.7 자원 분할과 점유율

한 SM에 워프를 많이 거주시키고 싶지만 자원이 한정돼 있다. 한 SM의 자원은 다음 셋이 핵심이다.

한 블록의 스레드 수, 한 스레드가 쓰는 레지스터 수, 블록당 공유 메모리 크기 ─ 이 셋이 거주 가능한 블록 수를 동시에 제약한다. 가장 빡빡한 제약이 천장이 된다.

예를 들어 블록당 256 스레드, 스레드당 32 레지스터를 쓴다고 하자. 한 블록의 레지스터 사용량은 256 × 32 = 8,192. SM의 65,536으로 나누면 8 블록까지 가능. 한 블록이 16 KB 공유 메모리를 쓰면 164/16 = 10.x이므로 공유 메모리는 10 블록까지 허용. 거주 스레드 한도 2048/256 = 8 블록. 가장 빡빡한 게 8이니 점유율은 8 × 256 = 2048 / 2048 = 100%.

이번엔 스레드당 레지스터를 64로 늘리면 어떻게 될까? 256 × 64 = 16,384. 65,536/16,384 = 4 블록만 거주. 거주 스레드는 4 × 256 = 1024 / 2048 = 50% 점유율로 떨어진다. 레지스터 사용량이 점유율을 결정하는 흔한 케이스다. 그래서 --ptxas-options=-v로 출력되는 레지스터 수와 공유 메모리 크기를 항상 보는 습관이 중요하다.

점유율 ≠ 성능

"점유율 100%가 항상 빠른가?"라는 질문에 답은 아니다이다. 점유율은 레이턴시를 가릴 워프가 충분한지를 나타내는 신호일 뿐이다. 레지스터를 늘려서라도 한 스레드가 더 많은 일을 캐싱해 가지고 있으면 메모리 트래픽 자체가 줄어 더 빨라질 수 있다. 일반적으로 50~70% 점유율이면 충분하다는 게 경험적 합의다 ─ 점유율을 짜내려고 알고리즘을 망치지 말 것.

4.8 디바이스 속성 쿼리

방금 등장한 숫자(SM 수, 레지스터 한도, 공유 메모리 크기 등)는 칩마다 다르다. 런타임에서 자기 GPU의 사양을 그대로 들춰 볼 수 있다.

#include <cuda_runtime.h>
#include <stdio.h>

int main(void) {
    int devCount;
    cudaGetDeviceCount(&devCount);
    for (int d = 0; d < devCount; ++d) {
        cudaDeviceProp p;
        cudaGetDeviceProperties(&p, d);
        printf("Device %d: %s (CC %d.%d)\n", d, p.name, p.major, p.minor);
        printf("  SM 개수            : %d\n", p.multiProcessorCount);
        printf("  SM당 최대 스레드   : %d\n", p.maxThreadsPerMultiProcessor);
        printf("  블록당 최대 스레드 : %d\n", p.maxThreadsPerBlock);
        printf("  워프 크기          : %d\n", p.warpSize);
        printf("  SM당 레지스터 수   : %d\n", p.regsPerMultiprocessor);
        printf("  블록당 공유메모리  : %zu bytes\n", p.sharedMemPerBlock);
        printf("  글로벌 메모리      : %.1f GB\n",
               p.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
        printf("  메모리 클럭        : %d MHz\n", p.memoryClockRate / 1000);
        printf("  메모리 버스 폭     : %d bit\n", p.memoryBusWidth);
    }
    return 0;
}

여기서 출력되는 값들이 점유율 계산기와 메모리 대역폭 추정의 입력이 된다. 메모리 이론 대역폭은 memoryClockRate × memoryBusWidth × 2 (DDR) / 8로 GB/s가 나온다. 자기 GPU의 천장을 알아야 자기 코드의 효율을 평가할 수 있다.

NVIDIA는 이 계산을 자동화하는 occupancy calculator를 제공한다. 과거에 엑셀 시트였고, 지금은 런타임 API에 내장돼 있다 ─ cudaOccupancyMaxActiveBlocksPerMultiprocessor()가 그것이다. 커널 함수와 블록 크기, 동적 공유 메모리 크기를 주면 SM당 거주 가능한 블록 수를 돌려준다. Nsight Compute는 한 발 더 나가서 측정된 점유율, 워프 활성도, 메모리 대역폭 활용률까지 한 화면에 보여준다. 처음부터 이 도구들과 친해지자.

4.9 정리

이 장에서 그동안 추상적이던 GPU 내부가 구체적인 풍경이 됐다. SM이라는 작은 공장, 그 안의 워프 스케줄러, 32명 한 조의 워프, 분기에서의 다이버전스, 메모리를 기다리는 워프와 일하는 워프를 자연스럽게 교대시키는 zero-overhead 스케줄링, 그리고 그 모든 것을 가능하게 하는 거주 워프 풀 ─ 점유율. 다음 장부터는 이 풍경 위에서 진짜 최적화를 시작한다. 5장에서 공유 메모리 타일링으로 행렬 곱의 산술 강도를 끌어올리고, 6장에서 메모리 합치기로 DRAM 트래픽을 절반으로 줄인다. 이 두 장에서 단순 naive 커널이 5~10배 빨라지는 모습을 직접 보게 될 것이다.

이 챕터에서 챙길 것