GPU Architecture & Execution Model

GPU 는 계층적 병렬 구조 + 메모리 계층 + 파이프라인 기반 실행 모델로 돌아감

커널 최적화의 모든 의사결정은 이 구조에서 파생

 

1. Thread / Warp / BLock / Grid

SIMT 의 기본 단위 : Thread & Warp

• GPU 의 연산 기본 단위는 thread, 실제 실행 단위 warp
• 논리적 단위와 물리적 실행 단위라고 보면 됌
• 같은 warp 내부의  thread 들은 동일한 명령을 동시에 실행 SIMT ( Single Instruction, Multiple Threads)

Warp Divergence

• 조건 분기에서 서로 다른 경로를 선택하면 warp 는 분리되어 순차 실행
• 이는 GPU 에서 즉각적인 성능 하락 요인
• 커널 최적화에서 branch-free 또는 warp-uniform branch 를 중요하게 여기는 이유

Block & Grid

• Block : 여러 warp 들의 집합
• Block 단위로 shared memory 를 소유
• Grid : Block 들의 전체 집합

왜 Block 단위가 중요한가

• shared memory 는 block 단위로 제공 - tiling 최적화의 핵심 전제
• warp 스케줄링도 block 수준에서 이뤄짐

 

Grid
 ├─ Block (0)
 │    ├─ Warp 0 (32 threads)
 │    └─ Warp 1 (32 threads)
 ├─ Block (1)
 │    ├─ Warp 0
 │    └─ Warp 1
 ...

GPU 는 warp - block - grid 3단계 구조를 기반으로 병렬성을 조직한다.

 

2. Memory Hierarchy & Data Movement

GPU 성능의 절반 이상은 계산 속도가 아닌, 데이터를 어디서 어떻게 가져오는냐로 결정, 

다음으로 주요 메모리 계층

Global Memory

• 모든 블록에서 접근 가능
• 가장 느림 
• 대용량
• 최적화 핵심 : coalescing
  • warp 내 threads 가 연속된 주소를 읽을 때 - 한 번의 transaction
  • 랜덤 / 스트라이드 패턴 - 다중 transaction - 성능 하락

 

 

Shared Memory ( SMEM )

• Block 내부 공유 메모리
• L1 보다 빠르고, register 보단 느림
• 직접 주소 계산 가능
• tiling / double-buffering / cp.async 등의 기

Bank Conflict

• Shared memory  32 개의 bank 로 구성
• warp 내 threads 가 동일 bank 의 서로 다르느 주소를 동시에 읽으면 conflict
• broadcast 만 예외

 

L1 / L2 Cach2

• L2 : 전체 GPU 공유
• L1 : SM 전용
• CUDA kernel 에서는 명령어 / 옵션으로 L1 bypass, read-onnly cache 활용도 가능
• GEMM 같은 고정 패턴의 경우, L2 hit 유도 + shared memory reuse 가 훨씬 중요

 

Register 

• 가장 빠른 공간
• 가장 제한적
• 블록 크기와 레지스터 수가 occupancy 결정에 직접 영향

 

3. SM Execution & Scheduling

GPU 의 진짜 성능은 아래의 키워드 이해를 통해 설명

SM ( Streaming Multiprocessor)

각 SM 은 독립적인 연산 엔진

각 SM 내부에는

• warp Schdulers
• load / store units
• FP32 / FP16 / MMA pipelines
• register file
• shared memory
• instruction issue logic

여러 warp 를 interleave 실행하여 latency 를 숨긴다.

GPU 는 느린 메모리 찾는 동안 idel 하는 것이 아닌, 다른 warp 를 즉시 실행

 

Occupancy

• 한 SM 에 동시에 올라갈 수 있는 active warps 개수
• 레지스터 개수, shared memory 사용량, block 크기에 의해 결정
• 높을수록 좋은 건 아님, compute bound 에선 낮아도 최적의 성능

 

Issue & Instructino Scheduling

• warp scheduler 는 매 클럭 실행 가능한 warp를 선택
• 메모리 대기 중 warp 는 skip
• 이를 latency hiding 이라 부름
• warp 수가 너무  적거나, 메모리 패턴이 나쁘면 hiding 불가, stall 증가

 

4. Tensor Cores & MMA Pipeline

최신 GPU 에서 GEMM 중심 compute 는 Tensor Core 기반으로 돌아감

Tensor Core 의 본질

• warp 단위 연산
• 정해진 고정 크기의 연산을 실시간 실행
• mixed precision

 

MMA Instruction Pipeline

• MMA.sync 는 warp 전개가 하나의 행렬 조각을 매 클럭 누적
• 타일 구조 : 16 x 8, 16 x 16 등
• register file 에 조각들을 분배하여 저장

 

Tensor Core 핵심 포인트

• 메모리의 tile 단위 공급
• register presure 높음
• cp.async + double-buffering 필수