GPU 아키텍처 이해: CUDA 코어, Tensor Core와 병렬 연산 최적화

GPU가 처리량에 집중하는 이유

GPU(Graphics Processing Unit)는 본래 3D 그래픽 렌더링을 위해 설계됐지만, 많은 데이터를 동시에 처리하는 능력을 바탕으로 범용 컴퓨팅(GPGPU)의 핵심 하드웨어가 됐다. 인공지능, 과학 계산, 데이터 분석처럼 같은 연산을 대량 데이터에 반복 적용하는 작업에서 특히 강점을 보인다.

CPU는 복잡한 제어 흐름과 순차 처리에 맞춰져 있다. 반면 GPU는 단순한 연산을 수천 개 데이터에 병렬 적용하도록 설계된 프로세서다. 이 차이는 다음과 같은 설계 선택으로 이어진다.

처리량(Throughput) 중심 설계
수천 개의 경량 코어
제한된 제어 로직, 확장된 ALU
높은 메모리 대역폭
지연 숨김(Latency Hiding) 기법

특성	CPU	GPU
코어 수	수 개 ~ 수십 개	수천 개
코어 복잡도	높음	낮음
클록 속도	높음 (3-5 GHz)	낮음 (1-2 GHz)
캐시 크기	큼	작음
제어 로직	복잡	단순
최적화 대상	지연(Latency)	처리량(Throughput)
병렬성 수준	낮음	높음

⤢✕

CUDA 코어는 SM 안에서 병렬 실행된다

CUDA(Compute Unified Device Architecture) 코어는 NVIDIA GPU의 기본 연산 유닛이다. 부동소수점과 정수 연산을 수행하는 스칼라 프로세서가 대량으로 동작하면서 병렬 처리를 구성한다.

CUDA 코어에는 FP32 유닛, INT32 유닛, 삼각함수와 지수함수 등을 처리하는 특수 함수 유닛(SFU), 메모리 접근을 담당하는 로드/스토어 유닛이 포함된다.

이 코어들은 스트리밍 멀티프로세서(Streaming Multiprocessor, SM) 단위로 묶인다. SM은 독립적인 스케줄러와 공유 리소스를 가진 GPU의 기본 실행 블록이다. Ampere 아키텍처 기준 SM은 다음 리소스로 설명할 수 있다.

CUDA 코어: 64개 FP32
Tensor 코어: 4개
RT 코어: 1개 (레이 트레이싱용)
L1 캐시/공유 메모리: 128KB (구성 가능)
레지스터 파일: 256KB
워프 스케줄러: 4개

⤢✕

워프의 실행 방식이 분기 비용을 만든다

워프는 32개의 스레드를 묶는 실행 단위이며, 하나의 워프에 속한 스레드는 동일한 명령어를 동시에 실행한다. 이 실행 모델을 SIMT(Single Instruction, Multiple Thread)라고 한다.

워프 내 32개 스레드는 동일한 PC(Program Counter)를 공유한다. 조건 분기가 생겨 스레드가 서로 다른 경로를 선택하면, 각 경로는 순차적으로 실행되고 해당 경로를 실행하지 않는 스레드는 마스킹된다. 이 분기 발산(Branch Divergence)은 성능 저하 요인이며, 실행 후에는 컨버전스(Convergence)로 재동기화된다.

⤢✕

커널 실행은 그리드(Grid), 블록(Block), 스레드(Thread) 계층으로 나뉜다. 그리드는 모든 블록을 포함하는 최상위 단위이고, 블록은 SM에 할당되는 스레드 그룹이다. 같은 블록의 스레드는 공유 메모리를 통해 통신할 수 있으며, 스레드는 고유한 ID를 가진 최소 실행 단위다.

계층	구성	통신 수단	동기화
그리드	블록의 집합	전역 메모리	커널 종료 시
블록	스레드의 집합	공유 메모리	__syncthreads()
워프	32 스레드	워프 내 셔플	암시적
스레드	최소 단위	로컬 메모리	N/A

Tensor Core가 행렬 연산을 맡는 방식

Tensor Core는 NVIDIA Volta 아키텍처(2017)에서 도입된 행렬 연산 전용 하드웨어 유닛이다. 딥러닝 훈련과 추론의 핵심 연산인 행렬 곱셈-덧셈(Matrix Multiply-Accumulate, MMA)을 가속한다.

Tensor Core는 4x4 행렬 연산을 단일 사이클에 수행하고 혼합 정밀도(Mixed Precision)를 지원한다. CUDA 코어 대비 최대 8배 처리량을 제공하며, 딥러닝과 HPC 워크로드에 맞춰 설계됐다.

세대	입력 타입	출력 타입	연산
Volta	FP16	FP16/FP32	MMA
Turing	FP16, INT8, INT4	FP16/FP32, INT32	MMA
Ampere	FP16, BF16, TF32, FP64, INT8, INT4	다양	MMA
Hopper	FP8, FP16, BF16, TF32, FP64, INT8	다양	MMA, Transformer Engine

혼합 정밀도 훈련에서는 입력과 가중치를 FP16으로 저장하고 계산해 메모리와 대역폭을 절약한다. 누적(Accumulation)은 FP32로 수행해 정밀도를 유지한다.

⤢✕

Warp Matrix Multiply-Accumulate API(WMMA API)는 Tensor Core에 접근하는 방법을 제공한다.

// WMMA 행렬 선언
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

// 행렬 로드
wmma::load_matrix_sync(a_frag, a_ptr, 16);
wmma::load_matrix_sync(b_frag, b_ptr, 16);
wmma::fill_fragment(c_frag, 0.0f);

// MMA 연산
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

// 결과 저장
wmma::store_matrix_sync(c_ptr, c_frag, 16, wmma::mem_row_major);

Hopper 아키텍처에서 도입된 Transformer Engine은 FP8 데이터 타입으로 트랜스포머 모델의 훈련과 추론을 가속한다. FP8은 FP16 대비 2배 처리량을 제공하며, 동적 스케일링으로 정밀도를 유지하고 대규모 언어 모델(LLM) 훈련에 최적화됐다.

GPGPU 성능은 데이터 배치와 자원 사용량에서 갈린다

GPGPU(General-Purpose computing on GPU)는 그래픽 외 범용 계산에 GPU를 사용하는 방식이다. 과학 계산, 금융 모델링, 이미지 처리, 머신러닝 등에서 CPU 대비 수십 배의 성능 향상을 달성한다.

GPU에 적합한 작업은 데이터 병렬성이 높고, 연산 밀도(Compute Intensity)가 높으며, 메모리 접근이 규칙적이고 분기가 적은 문제다.

메모리 계층도 성능에 직접 영향을 준다.

메모리 유형	범위	지연 시간	크기	용도
레지스터	스레드	~1 사이클	한정적	지역 변수
공유 메모리	블록	~20 사이클	~100KB	스레드 간 통신
L1 캐시	SM	~30 사이클	~100KB	자동 캐싱
L2 캐시	GPU 전체	~200 사이클	~6MB	전역 캐싱
전역 메모리	GPU 전체	~400 사이클	~80GB	대용량 데이터

⤢✕

워프에 속한 스레드가 연속된 메모리 주소에 접근하면 결합 접근(Coalesced Access)이 이뤄지고 단일 메모리 트랜잭션으로 처리된다. 반대로 접근 주소가 흩어지면 비효율이 커진다.

// 결합 접근 (최적)
data[threadIdx.x] = input[threadIdx.x];

// 비결합 접근 (비효율)
data[threadIdx.x] = input[threadIdx.x * stride];

공유 메모리는 32개의 뱅크로 구성된다. 동일 뱅크에 동시에 접근하면 직렬화가 발생하므로 데이터 배치도 고려해야 한다.

점유율(Occupancy)은 SM에서 활성 워프 수와 최대 가능 워프 수의 비율이다. 높은 점유율은 지연 숨김에 유리하지만, 항상 최고 성능을 보장하지는 않는다. 블록당 스레드 수, 스레드당 레지스터 사용량, 블록당 공유 메모리 사용량이 점유율에 영향을 준다. 레지스터 스필링을 줄이고 공유 메모리를 효율적으로 사용하면서 블록 크기를 실험적으로 조정하는 방식이 필요하다.

CUDA 스트림은 독립적인 작업 큐다. 서로 다른 스트림의 연산은 병렬로 실행될 수 있으며, 커널 실행과 메모리 전송을 겹치고 다중 커널을 동시에 실행하거나 CPU-GPU 파이프라이닝을 구성하는 데 사용된다.

⤢✕

세대 변화와 GPU 간 연결

NVIDIA GPU 아키텍처는 CUDA 지원부터 Tensor Core와 Transformer Engine까지 확장돼 왔다.

세대	출시	주요 특징
Tesla	2006	최초 CUDA 지원
Fermi	2010	L1/L2 캐시, ECC
Kepler	2012	동적 병렬화, Hyper-Q
Maxwell	2014	전력 효율성 향상
Pascal	2016	NVLink, HBM2
Volta	2017	Tensor Core 도입
Turing	2018	RT Core, INT8/INT4
Ampere	2020	3세대 Tensor Core, TF32
Hopper	2022	Transformer Engine, FP8
Blackwell	2024	5세대 Tensor Core

NVLink는 GPU 간 고속 상호연결 기술로 PCIe 대비 최대 10배 이상의 대역폭을 제공한다. NVSwitch는 다수 GPU를 완전 연결(Full-Mesh) 토폴로지로 연결하는 스위치다. Unified Memory는 CPU와 GPU 메모리를 단일 주소 공간으로 통합해 프로그래밍을 단순화한다.

CUDA 코어의 데이터 병렬 실행, Tensor Core의 행렬 연산 가속, 그리고 메모리 계층과 점유율 관리가 GPU 활용의 중심을 이룬다. GPU가 강점을 보이는 워크로드인지 먼저 판단하고, 워프 발산과 메모리 접근 패턴, 비동기 실행 구성을 함께 점검해야 한다.