🧠 Accelerator

텐서 코어 아키텍처

개요

텐서 코어(Tensor Core)는 딥 러닝 및 과학 계산에서 행렬 곱셈-누적(Matrix Multiply-Accumulate, MMA) 연산을 가속화하기 위해 설계된 특수 하드웨어 유닛이다. NVIDIA가 2017년 Volta 아키텍처에서 처음 도입한 이래, 현대 AI 가속기의 핵심 구성 요소로 자리잡았다. 텐서 코어는 FP16, FP32, BF16, TF32, INT8 등 다양한 정밀도의 행렬 연산을 하나의 명령어로 처리하여, 기존 CUDA 코어 대비 수십 배 높은 처리량을 제공한다.

텐서 코어의 핵심 혁신은 "연산 접착(Fused Multiply-Add)" 패턴을 하드웨어 수준에서 구현한 것이다. 예를 들어, 4×4 행렬 곱셈에서 하나의 텐서 코어는 매 클럭 주기마다 64개의 FP16 곱셈과 64개의 FP32 누산 연산을 동시에 수행한다. 이는 전통적인 SIMD(Single Instruction, Multiple Data) 방식의 CUDA 코어보다 훨씬 높은 연산 밀도를 달성한다.

핵심 개념

텐서 코어의 기본 동작 원리

텐서 코어는 행렬 블록(Matrix Block) 단위로 연산을 수행한다. 기본적인 동작은 다음과 같다:

D = A × B + C

여기서 A, B, C, D는 모두 행렬이며, 각 행렬의 크기는 하드웨어 구현에 따라 달라진다. NVIDIA Volta의 경우 4×4 행렬 블록을 사용하며, 각 텐서 코어는 매 클럭마다 64개의 FMA(Fused Multiply-Add) 연산을 수행한다.

행렬 블록 연산 흐름:
1. 입력 행렬 A와 B를 레지스터 파일에서 로드
2. 곱셈(Multiply) 연산 수행
3. 누산(Cumulate) 연산 수행 (이전 결과와 더하기)
4. 결과 행렬 D를 레지스터 파일에 저장

지원 데이터 타입

데이터 타입	비트 수	용도	특징
FP16	16비트	딥 러닝 훈련/추론	반정밀도 부동소수점
BF16	16비트	딥 러닝 훈련	FP32와 동일한 범위, FP16 정밀도
TF32	19비트	딥 러닝 훈련	FP32 범위 + FP16 정밀도
FP32	32비트	과학 계산	단정밀도 부동소수점
FP64	64비트	고정밀 과학 계산	이중 정밀도 부동소수점
INT8	8비트	추론	8비트 정수
INT4	4비트	추론	4비트 정수

Sparse Tensor Core

Ampere 아키텍처부터 도입된 Sparse Tensor Core는 구조화된 희소성(Structured Sparsity)을 활용하여 연산 처리량을 2배로 높인다. NVIDIA의 "2:4 희소성" 패턴은 4개 요소 중 2개만 0이 아닌 값을 가지도록 제약을 둔다.

희소성 처리 흐름:
1. 밀도 높은(Dense) 가중치를 2:4 희소성 패턴으로 프루닝
2. 남은 비영 가중치를 파인튜닝
3. 하드웨어가 0 값을 건너뛰어 연산 처리량 2배 향상

프로그래밍 모델과 데이터 공급

텐서 코어의 실제 성능은 연산 유닛 자체뿐 아니라, 워프 단위로 타일을 어떻게 공급하느냐에 크게 좌우된다. NVIDIA는 WMMA와 mma.sync 계열 명령으로 32스레드 워프가 16x16x16 수준의 타일 연산을 협력 수행하도록 노출하며, 공유 메모리 타일링과 레지스터 조각(fragment) 배치가 성능의 핵심이 된다.

AMD Matrix Core도 같은 D = A x B + C 패턴을 가속하지만, 실행 단위는 보통 64레인 wavefront이며 MFMA 명령으로 노출된다. 따라서 NVIDIA는 워프-프래그먼트 매핑과 Tensor Core 피드 경로가 중요하고, AMD는 wavefront 레지스터 분배와 lane 배치가 중요하다는 차이가 있다.

비교/분석

NVIDIA 텐서 코어 세대별 비교

세대	아키텍처	텐서 코어 수 (SM당)	지원 타입	특징
1세대	Volta (2017)	8	FP16, FP32	최초 도입
2세대	Turing (2018)	8	FP16, INT8, INT4	추론 가속 강화
3세대	Ampere (2020)	4	FP16, BF16, TF32, FP64, INT8, INT4	희소성 지원, FP64 가속
4세대	Hopper (2022)	4	FP8, FP16, BF16, TF32, FP64, INT8	FP8 지원, Transformer 엔진
5세대	Blackwell (2024)	4	FP4, FP6, FP8, FP16, BF16, TF32, FP64	FP4/FP6 지원

NVIDIA vs AMD 매트릭스 코어 비교

특성	NVIDIA Tensor Core	AMD Matrix Core
최초 도입	Volta (2017)	CDNA (2019)
실행 단위	Warp (32 threads)	Wavefront (64 lanes)
대표 명령	WMMA / mma.sync	MFMA
희소성 지원	2:4 구조화 (Ampere~)	미지원 (CDNA1~2)
특수 타입	TF32, BF16, FP64	BF16, FP64
워크로드 최적화	AI 훈련/추론, HPC	AI 훈련/추론

성능 비교 (A100 기준)

연산 타입	FP32 CUDA 코어	Tensor Core	향상 배율
FP32	19.5 TFLOPS	156 TFLOPS (TF32)	8x
FP16	78 TFLOPS	312 TFLOPS	4x
BF16	-	312 TFLOPS	-
INT8	-	624 TOPS	-
FP64	9.7 TFLOPS	19.5 TFLOPS	2x

동작 원리

행렬 블록 곱셈-누적 (MMA)

텐서 코어의 핵심 연산은 행렬 블록 곱셈-누적(Matrix Multiply-Accumulate)이다. 예를 들어, 16×16 크기의 행렬 A와 B의 곱셈을 다음과 같이 처리한다:

입력 로드: 레지스터 파일에서 16×16 행렬 A와 B를 로드
타일 분할: 각 행렬을 4×4 타일로 분할
병렬 곱셈: 4×4 타일 쌍에 대해 곱셈 연산을 병렬로 수행
누산: 곱셈 결과를 누산 레지스터에 더하기
결과 저장: 최종 16×16 결과 행렬을 메모리에 저장

정밀도 처리

텐서 코어는 입력과 출력의 정밀도가 다를 수 있다:

TF32 Tensor Core: FP32 입력을 읽어 TF32(19비트)로 변환 후 연산, FP32 출력
BF16 Tensor Core: BF16 입력, FP32 누산, BF16 출력
FP16 Tensor Core: FP16 입력, FP32 누산, FP16 출력

이러한 혼합 정밀도(Mixed Precision) 처리는 메모리 대역폭을 절감하면서도 누산 정밀도를 유지한다.

워프 수준 병렬 처리

각 텐서 코어는 32개 스레드로 구성된 워프(Warp) 수준에서 동작한다. 워프 내 모든 스레드가 동일한 MMA 명령어를 실행하며, 각 스레드는 행렬의 다른 부분을 담당한다.

장단점

장점

압도적인 처리량: FP16 기준으로 단일 텐서 코어가 64 FMA/클럭을 수행하며, 이는 CUDA 코어 대비 8배 이상의 처리량을 제공한다
혼합 정밀도 지원: 다양한 정밀도(FP16, BF16, TF32, INT8 등)를 지원하여 워크로드에 맞는 최적의 정밀도를 선택할 수 있다
희소성 가속: 2:4 구조화된 희소성을 활용하면 연산 처리량을 2배로 높일 수 있다
메모리 효율: 낮은 정밀도 입력은 메모리 대역폭을 절감하고, 높은 정밀도 누산은 정확도를 유지한다
프로그래밍 용이성: cuDNN, cuBLAS 같은 라이브러리를 통해 텐서 코어를 투명하게 사용할 수 있다

단점

행렬 크기 제약: 하드웨어가 고정된 행렬 크기(예: 16×16)를 지원하므로, 작은 행렬이나 비표준 크기에는 비효율적이다
희소성 패턴 제한: 2:4 구조화된 희소성만 지원하므로, 임의의 희소 패턴에는 가속이 불가능하다
정밀도 손실: TF32, BF16 같은 낮은 정밀도는 FP32 대비 정밀도 손실이 발생할 수 있다
설계 복잡성: 텐서 코어를 활용하려면 프로그래밍 모델과 하드웨어 특성을 깊이 이해해야 한다