GPU 스레드 계층 구조와 Warp 스케줄링 | Notes

왜 이걸 배우는가?

AXPY 장에서 이상한 현상을 봤다. 계산을 2배로 늘렸는데 GPU 실행 시간은 거의 그대로였다. 메모리 접근이 병목이기 때문이다.

그렇다면 질문이 생긴다. 메모리에서 데이터를 100 cycle 동안 기다릴 때 GPU 코어들은 뭘 하는가? 그냥 놀고 있으면 낭비다.

또 다른 질문도 있다. CUDA 코드를 짜면서 Block 크기를 256으로 설정한다. 왜 256인가? 128이나 512는 안 되나? 그냥 다른 사람 코드 복붙만 하면 이유를 모른다.

이번 장에서는 GPU 내부 작동 원리를 배운다. 세 가지를 이해하게 된다.

GPU가 메모리 대기 시간을 어떻게 숨기는가. Warp Scheduling으로 수천 개 스레드를 돌려가며 실행한다. 한 그룹이 기다리면 즉시 다른 그룹으로 전환한다.
왜 같은 코드가 다양한 GPU에서 돌아가는가. 저가 스마트폰(4개 SM)부터 고급 워크스테이션(200개 SM)까지 같은 CUDA 프로그램이 실행된다. Transparent Scalability의 비밀이다.
Block 크기를 어떻게 정하는가. 너무 작으면 메모리 대기를 못 숨기고, 너무 크면 레지스터가 부족하다. 최적 지점을 찾는 방법을 배운다.

Transparent Scalability

다양한 GPU 문제

CUDA device는 매우 다양하다. SM(Streaming Multiprocessor) 개수가 천차만별이다.

저가 스마트폰, 태블릿: 1~4개 SM
일반 PC GPU: 6~80개 SM
High-end 워크스테이션: 80~200개 SM
데이터센터 서버: 수백 개 SM

같은 CUDA 프로그램을 모든 장치에서 실행해야 한다. 하지만 하드웨어가 이렇게 다르면 어떻게 가능한가?

Thread Block 설계

CUDA의 해결책은 Thread Block 개념이다.

Grid → Block → Thread

Grid는 전체 커널 실행이다. Block들의 집합으로 구성된다.

Block은 독립적인 작업 단위다. 각 Block은 하나의 SM에서 처리된다. Block들 사이에는 실행 순서가 정해져 있지 않다.

Thread는 개별 실행 단위다. Block 내부에서만 synchronization이 가능하다.

이 설계 덕분에 scalability가 생긴다.

저가 GPU (4개 SM):

4개 Block을 동시 실행
나머지 Block들은 대기
완료되면 다음 4개 처리

고가 GPU (80개 SM):

80개 Block을 동시 실행
훨씬 빠르게 완료

코드는 똑같다. 속도만 다를 뿐이다. 이것이 Transparent Scalability다.

GPU 하드웨어 구조

물리적 계층

GPU는 계층적으로 구성되어 있다.

GPU → GPC → SM → Core

![[GPU 스레드 계층 구조와 Warp 스케줄링-20251020201905127.webp|671]]

GPC(Graphics Processing Cluster)는 여러 SM을 묶어서 관리하는 상위 단위다. 최신 GPU는 여러 GPC를 가진다.

SM(Streaming Multiprocessor)이 핵심이다. 실제로 CUDA 커널이 실행되는 하드웨어 유닛이다. Thread Block이 SM에 할당되어 실행된다.

SM 내부 유닛

각 SM 내부에는 여러 처리 유닛이 있다.

SP (Streaming Processor, FP32 Core)

단정밀도 부동소수점 연산을 처리한다. 각 SP는 단일 CUDA 스레드를 위한 Scalar ALU 역할을 한다. 한 SM에 32개 이상 존재한다.

DP (Double Precision, FP64 Core)

배정밀도 부동소수점 연산을 담당한다. 과학 계산에 사용된다. SP보다 개수가 적다.

SFU (Special Function Unit)

sin, cos, square root 같은 복잡한 수학 함수를 처리한다. 각 SFU는 한 클록당 스레드마다 하나의 특수 명령어를 실행한다.

Tex (Texture Processor)

그래픽 목적으로 사용된다. Texture memory 접근을 최적화한다.

Control Unit (CU)

SM 전체를 제어한다. SM당 1개만 존재한다. 이것이 중요하다. 모든 스레드가 같은 명령어를 실행하는 이유다.

CUDA 코드 매핑

CUDA 커널을 실행하면 다음과 같이 하드웨어에 매핑된다.

Grid → Block → Warp → Thread

Grid는 전체 커널 실행이다. <<<dimGrid, dimBlock>>>으로 지정한다.

Block은 SM에 할당되는 기본 단위다. 하나의 Block은 하나의 SM에서 처리된다. Block 하나는 1개부터 1024개까지의 스레드를 가질 수 있다.

Warp는 32개 스레드의 묶음이다. 하드웨어가 실제로 동시 실행하는 단위다. 이것은 CUDA 프로그래밍 모델의 일부가 아니라 구현 세부사항이다.

Thread는 개별 실행 단위다. 각 스레드는 고유한 threadIdx를 가진다.

Block을 32로 나누면 Warp 개수가 나온다.

\text{Warps per block} = \lceil \frac{\text{Block size}}{32} \rceil

예를 들어 256 threads/block이라면 8개 Warp가 생성된다.

\frac{256}{32} = 8 \text{ warps}

만약 3개 Block이 한 SM에 할당되면 총 24개 Warp가 생긴다.

3 \times 8 = 24 \text{ warps}

중요한 점은 어느 시점에든 1개 Warp만 실행된다는 것이다. SM에 하나의 제어 유닛만 있기 때문이다.

Warp Scheduling의 작동 원리

Time Sharing으로 대기 숨기기

한 SM은 24개 Warp를 동시에 관리할 수 있다. 하지만 물리적으로는 1개 Warp만 실행한다.

메모리 접근이 발생하면 어떻게 될까?

Warp A가 메모리 읽기 명령 실행
데이터가 올 때까지 약 100 cycle 대기 필요
즉시 Warp B로 전환 (zero overhead)
Warp B가 계산 명령 실행
Warp B도 메모리 대기 → Warp C로 전환
...순환하다가 Warp A의 데이터 도착
Warp A로 다시 전환해서 계속 실행

이 과정이 Time Sharing이다. 프로세서가 쉬지 않고 계속 일한다.

CPU의 multi-tasking과 비슷하지만 훨씬 빠르다. CPU는 context switch에 수백 cycle이 걸린다. GPU는 0 cycle이다.

SPMD 실행 모델

한 Warp의 32개 스레드는 같은 명령어를 동시에 실행한다. SPMD(Single-Program, Multiple-Data) 모델이다.

SPMD model이란?
Single-Program, Multiple-Data의 약자다. 모든 스레드가 같은 프로그램을 실행하지만, 각 스레드는 서로 다른 데이터를 처리한다. GPU의 기본 실행 모델이다.

SM에 제어 로직이 하나만 있으므로, 한 Warp의 모든 스레드는 동일한 명령어를 실행한다. 각 스레드는 threadIdx에 따라 다른 데이터에 접근한다.

예를 들어 벡터 덧셈 커널에서 Warp 0의 32개 스레드는 모두 동시에 다음 명령어를 실행한다.

z[i] = x[i] + y[i];

z[i] = x[i] + y[i];

Thread 0은 i=0, Thread 1은 i=1, ... Thread 31은 i=31에 접근한다. 명령어는 같지만 데이터가 다르다.

Zero-Overhead Warp Switching

CPU에서 context switch는 비용이 크다. 레지스터를 메모리에 저장하고 복원해야 한다.

GPU는 다르다. SM당 65,536개 이상의 레지스터를 가진다. 각 Warp의 레지스터를 모두 물리적으로 할당해둔다.

Warp A가 레지스터 0~~2047을 사용한다면, Warp B는 레지스터 2048~~4095를 사용한다. 모든 Warp가 자기만의 레지스터 공간을 가진다.

Warp 전환 시 레지스터 저장/복원이 없다. 그냥 다른 레지스터 세트를 사용하면 된다. 전환 비용이 사실상 0이다.

이것이 GPU가 수천 개 스레드를 효율적으로 관리하는 비밀이다.

Score-boarding

Warp Scheduler는 어떤 Warp를 실행할지 어떻게 결정하나?

Score-boarding 기법을 사용한다.

Score-boarding이란?
명령어 간의 의존성을 추적하는 기법이다. 각 레지스터의 상태를 추적해서 실행 가능한 명령어를 동적으로 찾는다. Warp scheduler는 이를 통해 ready 상태의 warp를 선택한다.

Instruction Buffer에 있는 모든 instructions의 모든 register operands가 score-boarded된다.

각 Warp는 다음 상태 중 하나다.

Ready: 모든 operand가 준비됨, 실행 가능
Waiting: 메모리 읽기 대기 중

Memory read/write instruction이 있으면 해당 warp는 waiting으로 전환된다. Instruction의 operand가 모두 ready면 해당 warp가 ready 상태가 된다.

Warp Scheduler는 Ready 상태의 Warp 중 우선순위가 높은 것을 선택한다. 메모리 명령어가 완료되면 해당 Warp가 Ready로 전환된다.

필요한 Warp 개수 계산

프로세서가 쉬지 않으려면 몇 개 Warp가 필요할까? 구체적인 예제로 계산해보자.

가정:

레지스터 연산: 1 cycle
메모리 접근: 100 cycle
평균적으로 4개 명령어당 메모리 접근 1회

일반적인 CUDA 프로그램을 보면 4개 instructions마다 memory access가 발생한다. 한 Warp의 실행 패턴은 다음과 같다.

plain

연산 - 연산 - 연산 - 연산 - 메모리(100 cycle 대기)

연산 - 연산 - 연산 - 연산 - 메모리(100 cycle 대기)

4 cycle 동안 일하고 100 cycle 동안 쉰다. 프로세서가 쉬지 않으려면 몇 개 Warp가 필요할까?

\text{필요한 Warp 개수} = \frac{100}{4} = 25 \text{ warps}

25개 Warp가 있으면 한 Warp가 메모리 대기하는 동안 다른 24개가 돌아가며 실행된다. 프로세서가 쉬지 않는다.

이것이 GPU가 수천 개 스레드를 필요로 하는 이유다. 메모리 대기 시간을 숨기기 위해서다.

AXPY에서 계산 추가가 공짜였던 이유도 이제 명확하다. 메모리 접근 시간은 그대로지만, 그 사이에 다른 Warp들이 계산을 더 많이 한다. 메모리 대기와 계산이 overlap된다.

실습으로 이해하기

Warp ID와 Lane ID

Warp 내부에서 각 스레드의 위치를 나타내는 두 가지 ID가 있다.

Warp ID

SM 내에서 각 Warp의 고유 번호다. SM마다 독립적이므로 globally unique하지 않다. 다른 SM에 동일 warp id가 존재할 수 있다.

Lane ID

Warp 내에서 각 스레드의 인덱스다. 0부터 31까지 값을 가진다. 한 Warp는 32개 thread로 구성되므로 lane도 32개다.

CUDA API로는 직접 접근할 수 없다. GPU assembly instruction으로만 가져올 수 있다.

warp-lane.cu 코드

Warp ID와 Lane ID를 확인하는 device 함수를 작성한다.

__device__ unsigned warp_id(void) {
    // this is not equal to threadIdx.x / 32
    unsigned ret;
    asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
    return ret;
}

__device__ unsigned lane_id(void) {
    unsigned ret;
    asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret));
    return ret;
}

__device__ unsigned warp_id(void) {
    // this is not equal to threadIdx.x / 32
    unsigned ret;
    asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
    return ret;
}

__device__ unsigned lane_id(void) {
    unsigned ret;
    asm volatile ("mov.u32 %0, %laneid;" : "=r"(ret));
    return ret;
}

PTX(Parallel Thread Execution) 어셈블리를 inline으로 사용한다. %warpid와 %laneid는 GPU의 특수 레지스터다.

warp_id()는 threadIdx.x / 32와 다르다는 점을 주의하자. 하드웨어가 부여하는 실제 warp ID다.

커널 함수는 warp 0번만 정보를 출력한다.

__global__ void kernel_warp_lane(void) {
    unsigned warpid = warp_id();
    unsigned laneid = lane_id();
    if (warpid == 0) {
        printf("lane=%2u threadIdx.x=%2d threadIdx.y=%2d blockIdx.x=%2d blockIdx.y=%2d\\n",
               laneid, threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y);
    }
}

__global__ void kernel_warp_lane(void) {
    unsigned warpid = warp_id();
    unsigned laneid = lane_id();
    if (warpid == 0) {
        printf("lane=%2u threadIdx.x=%2d threadIdx.y=%2d blockIdx.x=%2d blockIdx.y=%2d\\n",
               laneid, threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y);
    }
}

실행 설정은 다음과 같다.

int main(void) {
    dim3 dimBlock(16, 16, 1);  // 256 threads/block
    dim3 dimGrid(2, 2, 1);     // 4 blocks
    kernel_warp_lane <<<dimGrid, dimBlock>>>();
    cudaDeviceSynchronize();
    CUDA_CHECK_ERROR();
}

int main(void) {
    dim3 dimBlock(16, 16, 1);  // 256 threads/block
    dim3 dimGrid(2, 2, 1);     // 4 blocks
    kernel_warp_lane <<<dimGrid, dimBlock>>>();
    cudaDeviceSynchronize();
    CUDA_CHECK_ERROR();
}

Block당 256개 스레드이므로 8개 Warp가 생성된다. 총 4개 Block이므로 32개 Warp가 있다.

Block과 Warp 매핑 규칙

256개 스레드를 가진 Block은 다음과 같이 Warp로 나뉜다.

\text{Warps per block} = \frac{256}{32} = 8

Warp 0: threadIdx 0~31
Warp 1: threadIdx 32~63
Warp 2: threadIdx 64~95
...
Warp 7: threadIdx 224~255

2D Block의 경우 threadIdx는 다음과 같이 계산된다.

\text{threadIdx (linear)} = \text{threadIdx.y} \times \text{blockDim.x} + \text{threadIdx.x}

예를 들어 (x=5, y=1) 위치의 스레드는 1 × 16 + 5 = 21번 스레드다. Warp 0에 속한다.

실행 결과를 보면 Block들이 임의의 순서로 실행된다. blockIdx가 순서대로 나타나지 않는다.

이것이 Transparent Scalability의 증명이다. Block들은 서로 독립적으로 실행된다. SM 개수가 다른 GPU에서도 같은 코드가 동작한다.

Lane ID와 Warp ID 부여 규칙은 시스템 의존적이다. 알 필요가 없다. CUDA는 Block 단위로 추상화되어 있다. Block 간 실행 순서도 정해지지 않는다. 하드웨어가 알아서 처리한다.

자원 제약과 최적화

SM의 자원 제한

SM은 물리적 자원에 제약이 있다. 아키텍처마다 다르지만 일반적인 예시는 다음과 같다.

최대 Block 수: SM당 32개
최대 Thread 수: SM당 2,048개
최대 Warp 수: SM당 64개
레지스터: SM당 65,536개
Shared Memory: SM당 48KB 또는 96KB

이 제약 안에서 Block 크기를 선택해야 한다.

Block 크기 선택 문제

같은 총 스레드 개수를 다른 방식으로 구성할 수 있다. 이것을 Block granularity 문제라고 한다.

2,048개 스레드를 처리하는 경우를 보자.

Threads/Block	Blocks/SM	Warps/Block	Total Warps
128	16	4	64
256	8	8	64
512	4	16	64
1024	2	32	64

Warp 개수는 같지만 성능이 다르다.

Block이 너무 작으면 (128 threads):

Warp 4개로 부족할 수 있다
메모리 대기 시간을 완전히 숨기지 못함
앞에서 계산했듯이 25 warps가 이상적이므로 부족

Block이 너무 크면 (1024 threads):

레지스터 사용량이 많으면 2개 Block만 올라감
SM당 활성 Thread가 줄어듦
SM 활용도 감소

최적 Block 크기: 일반적으로 256~512가 좋다. 하지만 커널마다 다르므로 실험해야 한다.

256을 많이 쓰는 이유가 이제 명확하다. 8 warps면 대부분의 메모리 대기를 숨길 수 있고, 레지스터도 충분히 할당 가능하다.

Occupancy

Occupancy는 SM의 최대 Thread 수 대비 실제 활성 Thread 비율이다.

\text{Occupancy} = \frac{\text{Active Threads}}{\text{Maximum Threads}}

최대 2,048개인데 1,024개가 활성이면 occupancy는 50%다.

Occupancy가 높을수록 메모리 대기를 숨기기 쉽다. 더 많은 Warp가 있으므로 time-sharing이 효과적이다.

하지만 100%가 항상 최고는 아니다. 레지스터와 shared memory 사용량에 따라 trade-off가 있다. 때로는 occupancy를 낮춰서 레지스터를 더 쓰는 게 빠를 수 있다.

CUDA Occupancy Calculator나 nvcc --ptxas-options=-v 옵션으로 확인할 수 있다.

2단계 병렬성

GPU는 두 단계로 병렬 처리한다.

1단계: Block-level 병렬성

Grid의 Block들이 여러 SM에서 동시에 실행된다. Kernel program은 grid를 생성하고, Grid는 thread blocks로 나뉘어 SMs에서 parallel execution된다.

100개 Block과 80개 SM이 있다면 80개가 동시 실행되고, 완료되면 나머지 20개가 실행된다.

이것이 Transparent Scalability를 가능하게 한다. SM이 많으면 더 많은 Block을 동시 처리하므로 빨라진다.

2단계: Warp-level 병렬성

각 SM 내에서 여러 Warp가 time-sharing으로 실행된다. Score-boarding 방식으로 parallel execution된다.

Warp 또는 Block이 종료되면 다음 Warp 또는 Block을 pick up해서 계속 실행한다.

이 2단계 병렬성 덕분에 GPU는 수천 개 스레드를 효율적으로 관리한다. Block-level로 scalability를 확보하고, Warp-level로 latency를 숨긴다.

다음 단계로

Warp Scheduling의 원리를 배웠다. GPU가 메모리 대기 시간을 어떻게 숨기는지 이해했다. Block 크기를 왜 256으로 설정하는지도 이제 안다.

다음 장에서는 메모리 시스템을 더 깊이 파고든다. Global Memory, Shared Memory, Coalescing 같은 개념을 배우면 왜 메모리 접근 패턴이 성능에 결정적인지 보인다.

그 이후에는 행렬 곱셈으로 진입한다. 계산 강도가 높고 메모리 재사용이 많은 연산에서 GPU의 진가가 드러난다. Warp Scheduling과 메모리 최적화를 모두 활용하는 실전 예제다.