챕터별 상세
GPU 하드웨어 아키텍처 복습
NVIDIA GPU(A100, H100, B200)의 핵심 구성 요소인 SM(Streaming Multiprocessor)과 메모리 계층 구조를 비교한다. 레지스터, L1 캐시, L2 캐시, HBM 순으로 용량은 커지지만 대역폭은 낮아지는 특성을 가진다. 특히 B200은 H100 대비 HBM 크기가 80GB에서 192GB로, 대역폭은 3.35 TB/s에서 8 TB/s로 대폭 향상되었다. 이러한 하드웨어 제약 사항을 이해하는 것이 커널 최적화의 시작점이다.
HBM(High Bandwidth Memory)은 GPU 외부의 큰 메모리이며, SM 내부의 레지스터와 공유 메모리는 매우 빠르지만 용량이 작다.
GPU 프로그래밍 모델: Thread와 Block
GPU 연산은 Thread, Thread Block(CTA), Grid의 계층 구조로 조직된다. 개별 Thread는 데이터를 처리하고, Block은 공유 메모리를 통해 통신하는 Thread들의 집합이며, Grid는 전체 Block의 모음이다. Element-wise 연산은 Thread 단위로 충분하지만, Softmax나 Matrix Multiplication 같은 연산은 Thread 간 통신이 필요하여 Block 단위의 설계가 필수적이다. Triton은 이러한 Thread Block 단위의 사고를 기반으로 커널을 작성하게 한다.
CUDA 프로그래밍의 기본 단위인 Thread와 이들의 묶음인 Block이 어떻게 하드웨어 SM에 할당되는지가 성능의 핵심이다.
성능 저하의 원인: Control Divergence와 Occupancy
Warp 내의 Thread들이 서로 다른 명령어를 실행해야 할 때 발생하는 Control Divergence 문제를 다룬다. Warp는 32개의 Thread가 동시에 같은 명령을 실행(Lockstep)해야 하므로, if-else 분기가 발생하면 순차적으로 실행되어 효율이 급감한다. 또한 Occupancy는 SM이 가질 수 있는 최대 Warp 대비 실제 실행 중인 Warp의 비율을 의미하며, 레지스터 사용량이 많을수록 Occupancy가 낮아져 성능이 제한될 수 있다. 하지만 Occupancy가 낮더라도 각 Thread가 더 많은 일을 수행하는 Thread Coarsening 기법을 통해 전체 성능을 높일 수도 있다.
Warp는 GPU에서 명령어를 실행하는 최소 단위이며, 모든 Thread가 같은 길을 가야 가장 빠르다.
메모리 최적화: Bank Conflict와 Coalescing
공유 메모리 접근 시 발생하는 Bank Conflict와 HBM 접근 시의 Memory Coalescing을 설명한다. 공유 메모리는 32개의 Bank로 나뉘어 있으며, 여러 Thread가 동시에 같은 Bank의 다른 주소에 접근하면 요청이 직렬화되어 성능이 떨어진다. HBM 접근 시에는 32개 Thread의 요청을 하나의 128-byte 트랜잭션으로 묶는 Coalescing이 일어나야 대역폭을 최대로 활용할 수 있다. Matrix Multiplication 연산 시 이러한 메모리 접근 패턴을 최적화하기 위해 Swizzling 같은 기법이 사용된다.
메모리 접근을 '정렬'하여 한 번에 많은 데이터를 가져오는 것이 GPU 성능 최적화의 핵심이다.
벤치마킹과 프로파일링 실습
PyTorch를 이용해 실제 연산 시간을 측정하는 올바른 방법을 제시한다. GPU 연산은 비동기적이므로 `torch.cuda.synchronize()`를 호출하여 연산 완료를 기다려야 정확한 측정이 가능하다. 또한 Warmup 단계를 거쳐 초기 컴파일 오버헤드를 제거하고, 여러 번 실행하여 평균값을 구해야 한다. PyTorch Profiler를 사용하면 실제 어떤 CUDA 커널이 호출되는지, HBM 읽기/쓰기에 시간이 얼마나 소요되는지 상세히 분석할 수 있다.
GPU 연산은 CPU와 비동기로 작동하므로 단순히 CPU 시간을 측정하면 안 된다.
Kernel Fusion의 효과 분석
GeLU 활성화 함수를 예로 들어 Naive 구현과 Kernel Fusion의 성능 차이를 비교한다. Naive 구현은 각 연산마다 HBM에서 데이터를 읽고 다시 쓰는 과정을 반복하여 메모리 대역폭 병목이 발생한다. 반면 Triton으로 작성된 퓨전 커널은 데이터를 한 번만 읽어 SM 내부에서 모든 연산을 처리하고 결과만 다시 쓴다. 실험 결과 퓨전된 커널이 Naive 구현보다 약 5배 이상 빠른 성능을 보여주며, 이는 연산량보다 메모리 접근이 병목인 'Memory-bound' 연산에서 특히 중요하다.
Kernel Fusion은 여러 연산을 하나의 커널로 합쳐 중간 결과를 메모리에 쓰지 않고 재사용하는 기법이다.
Triton 커널 작성법: GeLU 구현
Triton을 사용하여 GeLU 커널을 직접 작성하는 과정을 보여준다. `@triton.jit` 데코레이터를 사용하며, 각 Block이 처리할 데이터의 오프셋을 계산하고 `tl.load`로 데이터를 가져온다. 이후 표준 Python 문법과 유사한 Triton 연산자를 사용해 수식을 계산하고 `tl.store`로 결과를 저장한다. Triton은 복잡한 CUDA C++ 코드 없이도 효율적인 PTX 코드를 생성해주며, Thread Coarsening과 같은 최적화를 자동으로 수행한다.
Triton은 OpenAI에서 만든 GPU 프로그래밍 언어로, Python과 유사한 문법으로 고성능 커널을 작성할 수 있게 해준다.
Tiling 기법과 Matrix Multiplication
데이터가 SM의 공유 메모리보다 클 때 사용하는 Tiling 기법을 설명한다. 전체 행렬을 작은 Tile 단위로 쪼개어 순차적으로 공유 메모리에 로드하고 연산하는 방식이다. Matrix Multiplication의 경우 A 행렬의 행 Tile과 B 행렬의 열 Tile을 가져와 부분합을 계산하고 누적한다. 이 방식은 HBM 접근 횟수를 획기적으로 줄여 연산 밀도(Arithmetic Intensity)를 높이며, 현대적인 GPU 연산의 핵심 알고리즘이다.
Tiling은 큰 문제를 작은 조각으로 나누어 빠른 로컬 메모리에서 해결하는 전략이다.
AI 분석 전체 내용 보기
AI 요약 · 북마크 · 개인 피드 설정 — 무료
출처 · 인용 안내
원문 발행 2026. 04. 29.수집 2026. 04. 29.출처 타입 YOUTUBE
인용 시 "요약 출처: AI Trends (aitrends.kr)"를 표기하고, 사실 확인은 원문 보기 기준으로 진행해 주세요. 자세한 기준은 운영 정책을 참고해 주세요.
