estudy 한국어/English

GPU 스레드 실행 모델

3,181 단어·7 분·원문(.md)

리눅스 os 계층의 백엔드 엔지니어가 다루는 POSIX 스레드 (pthread)는 수메가바이트의 독립된 스택 메모리를 가지며 커널이 컨텍스트 스위칭을 제어한다. 반면, 단일 gpu 커널 실행시 생성되는 수백만개의 gpu 스레드는 어떻게 os 커널의 개입 없이 하드웨어 레벨에서 초고속으로 스케줄링되고 동시에 연산을 얻을 수 있을까?

백엔드 시스템의 스레드 모델과 GPU의 대규모 병렬 스레드 모델은 단어만 같을 뿐, 하드웨어 자원을 제어하고 명령어를 실행하는 메커니즘이 완전히 다르다.

  • 백엔드 os 스레드(MIMD) 기반: 다중 명령어, 다중 데이터 (Multiple Instruction, Mulitple Data) 구조를 따른다. 각 스레드는 프로그램 카운터와 독립된 가상 메모리 스택 영역을 소유하며, CPU 코어들과 독립적으로 전혀 다른 코드 라인을 실행할 수 있다.
  • GPU 스레드 (SIMT 기반): 단일 명령어, 다중 스레드 (Signle Instruction, Multiple threads) 구조를 따른다. 수백만 개의 스레드가 생성되지만, 하드웨어적으로 물리적인 명령어 실행 스케줄링 단위인 워프 구조에 묶여 격자 형태로 동작한다.
  • Warp: 워프는 NVIDIA GPU에서 명령어를 동시에 실행하는 하드웨어의 최소 기본 단위다. 항상 32개의 스레드가 하나의 묶음으로 구성되며, 이 32개의 스레드는 1클럭마다 완전히 동일한 소스코드 라인(명령어)를 데이터만 다르게 하여 실행한다.
  • 스레드 블록 (Thread Block) 및 SM: 프로그래머가 논리적으로 묶어놓은 스레드 집합 (최대 1024개)이다. 이 블록은 gpu 내부의 물리적인 연산 코어 덩어리인 SM (Stremaing Multiprocessor)에 통째로 할당되며, 블록 내부의 스레드들은 초고속 고유 SRAM 영역인 공유 메모리를 통해 상호 통신하고 동기화할 수 있다.

물리적/구조적 병목 현상 정의 #

하드웨어 스레드 스케줄러 계층에서 백엔드와 GPU는 각각 제어 메커니즘의 한계로 인해 독특한 병목을 겪게된다.

  • OS 컨텍스트 스위칭 오버헤드 (TLB/캐시 플러시): CPU 백엔드에서 활성 스레드수가 물리 코어 수를 초과하면, OS 커널은 타이머 인터럽트를 발생시켜 스레드를 교체한다. 이때 레지스터 세트를 커널 스택에 백업하고 가상 메모리 매핑 정보 Page Table을 교체하며, CPU 캐시(L1/L2)가 오염되어 막대한 시스템 레이턴시(수 마이크로초)가 발생한다.
  • 워프 디바이전스: GPU 스레드가 소스 코드 내에서 조건문 if else를 만나 32개 스레드의 실행 경로가 갈라질 때 발생하는 하드웨어 마비 현상이다. SIMT 아키텍처 특성상 워프는 2개의 다른 명령어를 한 클럭에 실행할 수 없으므로 if 조건에 만족하는 16개의 스레드가 실행되는 동안 나머지 16개 스레드의 실행 유닛은 강제로 락에 걸려 대기하게 되며 이로 인해 하드웨어 처리량이 정확히 반토막 난다.
  • 레지스터 압박에 따른 occupancy(점유율)저하: SM 내부의 물리 레지스터 파일 크기는 고정되어 있다. 만약 개별 GPU 스레드가 너무 많은 변수를 선언하여 스레드당 레지스터 할당량이 늘어나면 SM은 물리적 한계로 인해 스레드 블록을 많이 상주시키지 못한다. 이로 인해 대기 워프가 부족해져 메모리 지연을 숨기지 못하는 하드웨어 유후 상태가 발생한다.

아키텍처 혁신 및 해결 방식 #

GPU는 명령어 제어를 극단적으로 단순화하고 하드웨어 배선을 최적화하며 백엔드 스레드의 스케줄링 한계를 해결했다.

  • 하드웨어 워프 스케줄러 Zero-Overhead Scheduling: GPU 스레드 스케줄링은 os 커널 소프트웨어가 관여하지 않는다. SM 내부의 전용 하드웨어 칩인 워프 스케줄러가 클럭마다 점수판(scoreboard)를 검사하여 메모리 IO 대기가 끝나 실행 가능한 상태가 된 워프를 선택해 즉시 연산 유닛에 할당한다. 모든 워프의 스레드 레지스터 상태가 물리 칩 위에 상시 보존되어 있으므로 컨텍스트 스위칭 비용이 완전히 0이다.
  • 실행 마스킹 (Execution Masking) 메커니즘: 워프 다이버전스가 발생하면 하드웨어는 내부의 특별한 레지스터 마크인 조건부 마스크 비트를 활성화한다. 32비트 마스크 맵을 통해 이번 클럭에서 연산 결과를 저장할 스레드 (1)와 연산을 무시하고 클럭만 보낼 스레드 0을 전기적으로 제어하여 단일 명령어 파이프라인 안에서 분기 로직을 강제로 소화해낸다.

작동 원리 #

호스트 cpu에서 gpu 커널을 호출했을때 하드웨어 그리드와 워프 파이프라인에서 데이터와 명령어가 움직이는 저수준 물리 경로를 추적한다.

  1. 커널 런칭 및 GigaThread 스케줄러 가동: CPU가 CUDA 드라이버를 통해 kernel<<<Grid, Black>>>()을 호출하면 gpu 내부의 글로벌 하드웨어 스케줄러인 GigaThread 엔진이 이 요청을 수신한다.
  2. 스레드 블록의 물리 SM 분배: GigaThread 엔진은 커널이 요구한 스레드 블록들을 사용 가능한 전국의 SM 들로 물리적으로 분배한다. 이때 한 SM에 몇 개의 블록이 들어갈지는 스레드당 레지스터 요구량과 공유 메모리 크기를 계산하여 런타임에 결정된다.
  3. 스레드 워프 쪼개기 (Warp 소화): SM에 안착한 스레드 블록은 물리적으로 32개의 스레드씩 쪼개져 워프 스케줄러의 대기 슬롯에 등록된다. 이때 스레드 id 0번부터 31번까지가 0번 워프, 32번부터 63번까지가 1번 워프로 기계적으로 고정 배정된다.
  4. Instruction Fetch 및 SIMT 데이터 패스 리치: 워프 스케줄러가 대기중인 워프 하나를 선택하면, 해당 워프의 프로그램 카운터가 가리키는 명령어 코드가 L1 Instruction Cache에 인입된다. 디코더를 통과한 단일 명령 신호는 SM 내부의 전력 공급 버스를 타고 32개의 ALU (또는 Tensor Core) 컨트롤 러안에 동시에 연결 된다.
  5. 레지스터 파일 동시 엑세스: 32개의 스레드는 레지스터 파일 SRAM의 서로 다른 주소 섹션에 동시 접근하여 피연산자 데이터를 꺼내고 할당된 32개의 물리 연산 파이프라인을 통과시킨 뒤 결괏값을 다시 자신의 하드웨어 레지스터 공간에 Write-back

시스템 로우레벨 메트릭 및 프로파일링 로그 분석 #

GPU 하드웨어 스레드 실행 상태와 병목 현상을 진단하기 위해 프로파일러 및 커널 내부 카운터가 출력하는 로우 레벨의 메트릭 의미를 분석한다.

NVIDIA Nsight Compute ncu 핵심 하드웨어 카운터 메트릭 #

  • sm__wraps_active.avg.pct_of_peak (Active Warps): 물리적으로 SM에 로드되어 자원을 점유하고 있는 워프의 비율이다 이 값이 낮다면 호스트 CPU 측의 커널 발급 지연이나 극심한 레지스터 부족으로 GPU에 스레드가 채워지지 못하고 있음을 뜻한다.
  • smsp_warp_issue_stalled_barrier.per_warp (Barrier Stall): 스레드 블록 내에서 스레드 간 동기화 명령어 __syncthreads()를 만나 다른 워프가 도달할 때까지 파이프라인이 멈춰 선 횟수다. 알고리즘의 워크로드 분배가 불균형함을 시사한다.
  • smsp__thread_inst_executed_per_inst_executed.ratio (Divergence Metric): 워프 다이버전스를 잡아내는 핵심 지표고 이론적 최대치는 32이다 명령어 1번 수행시 32개의 스레드가 모두 살아있음이고, 만약 이 메트릭 로그가 16.5같이 낮게 찍히면 조건문 분기로 인해 매 클럭마다 절반에 가까운 코어가 연산에 참여하지 못하고 전력 낭비중이라는 뜻이다.

CUDA 컴파일러 nvcc의 리소스 할당 로우레벨 로그 --ptxas-options=-v 옵션 활성화 시 출력 #

ptxas info    : Compiling entry function '_Z11train_kernelPfS_' for 'sm_80'
ptxas info    : Used 64 registers, 1024 bytes smem, 320 bytes cmem[0]

컴파일러가 에이전트 커널을 분석한 결과 스레드당 64개의 레지스터를 고정 점유함을 명시한다.

Ampere 아키텍처 sm_80의 sm당 최대 레지스터 한계는 65,536개이므로, 이 커널은 단일 SM에 최대 65536/64 = 1024개의 스레드 즉 32개 워프만 상주시킬 수 있어 하드웨어 이론적 최대치 2048 스레드의 50퍼 수준으로 하드웨어가 제한되는 병목 상태를 로그를 통해 진단할 수 있다.

Example #

AI 프레임워크나 런타임 환경에서 워프 실행 모델 최저고하를 위해서는

스레드 블록 차원 설정 공식과 컴파일 설정을 알면 좋은데 파이토치의 커스텀 CUDA확장 카드를 작성하거나 Triton 컴파일러 설정을 조율할 대, 블록당 스레드수는 하드웨어 워프크기의 배수로 맞추는것이 절대 규칙이다.

# Triton 컴파일러 또는 커스텀 가속기 레이어 설정 예시
# 블록 내부 스레드 크기가 32의 배수가 아닌 250 등 애매한 수치로 정의되면, 
# 마지막 워프는 26개의 스레드가 물리적으로 비어있는 채(Execution Masked) 실행되어 자원이 영구 낭비됨
BLOCK_SIZE_X = 128  # 32 * 4 (정확히 4개의 워프로 스케줄러 슬롯 구성)
BLOCK_SIZE_Y = 2    # 총 256 스레드 블록 구성

Warp-level Primitive 활용을 통한 조건문 제거 (소프트웨어 최적화): if-else 분기문을 통한 동기화를 피하기 위해, 레지스터 내부에서 32개 워프 스레드가 공유 메모리 없이 서로의 데이터를 직접 교환하는 하드웨어 셔플 명령어(__shfl_sync)를 사용하여 분기 다이버전스를 코드 레벨에서 회피.

// 안티 패턴: 분기로 인한 워프 다이버전스 유발
if (threadIdx.x == 0) {
    shared_mem[0] = local_val;
}
__syncthreads();

// 최적화 패턴: 하드웨어 워프 데이터 버스를 직접 태워 분기문과 동기화 오버헤드를 모두 제거
// 32개 스레드가 동시에 0번 스레드의 레지스터 값을 자신의 레지스터로 복사 (1클럭 소요)
float warp_shared_val = __shfl_sync(0xFFFFFFFF, local_val, 0);
MLOps/gpu_thread.md