텐서 코어
더 많은 작업
- Tensor Core; 텐서 코어
- 엔비디아 GPU의 스트리밍 멀티프로세서(SM) 안에 있는 행렬 곱-누산 전용 연산 유닛. 낮은 정밀도 입력으로 작은 행렬 곱 D = A×B + C를 한 번에 계산한다
텐서 코어(Tensor Core)는 엔비디아가 2017년 볼타(Volta) 아키텍처의 Tesla V100에서 처음 넣은 행렬 연산 전용 회로다. 일반 CUDA 코어가 스레드마다 스칼라 곱셈-덧셈(FMA)을 하나씩 처리하는 데 비해, 텐서 코어는 작은 행렬 블록의 곱셈과 누산을 한 명령으로 처리한다. 딥 러닝의 학습과 추론 시간 대부분은 행렬 곱(선형층, 합성곱, 어텐션)이 차지하므로, 텐서 코어의 처리량과 지원 숫자 형식이 사실상 GPU 세대별 AI 성능을 결정한다.[1]
세대가 바뀔 때마다 지원 형식이 FP16에서 INT8·INT4, TF32·BF16·FP64, FP8, FP6·FP4로 넓어졌고 2:4 구조적 희소성, 비동기 실행, 전용 메모리 같은 기능이 붙었다. 아키텍처 전반은 엔비디아 GPU 아키텍처, 프로그래밍 모델은 CUDA 문서를 참고하고, 이 문서는 텐서 코어 자체의 동작과 세대별 차이, 쓰는 방법을 다룬다.
볼타 백서에 따르면 텐서 코어 하나는 4×4 행렬에 대해 를 계산한다. 는 FP16이고 는 FP16이나 FP32다. FP16 곱은 정밀도 손실 없이 계산된 뒤 FP32 덧셈으로 다른 곱들과 누산된다. 4×4×4 행렬 곱 하나는 곱셈-덧셈 64개이며, V100 텐서 코어는 클럭당 FMA 64개를 처리한다. SM마다 8개(처리 블록마다 2개), GPU 전체에 640개가 있어 최대 125 Tensor TFLOPS를 냈다.[1]
프로그래머는 텐서 코어 하나를 직접 다루지 않는다. 워프(32개 스레드)가 함께 16×16 같은 더 큰 타일을 나눠 들고 한 명령을 실행하면 하드웨어가 이를 여러 텐서 코어 연산으로 쪼갠다. 큰 행렬 곱(GEMM)은 다시 이 타일들의 반복으로 만든다. 그래서 행렬 크기가 타일 크기(보통 8 또는 16의 배수)에 맞지 않으면 덧채우기(padding)가 생겨 효율이 떨어진다.
입력은 낮은 정밀도, 누산은 높은 정밀도로 하는 이 구조가 혼합 정밀도 학습의 하드웨어 바탕이다. 입력 비트가 줄면 같은 회로 면적과 메모리 대역폭으로 더 많은 원소를 처리할 수 있어, 세대마다 더 짧은 형식이 추가될 때 처리량이 크게 뛰었다. 예를 들어 호퍼의 FP8은 FP16·BF16보다 처리량이 2배이고,[2] 암페어의 TF32는 같은 GPU의 CUDA 코어 FP32보다 8배 빠르다(156 대 19.5 TFLOPS).[3]
| 세대 | 아키텍처(발표) | 대표 GPU | 새로 지원한 형식·기능 |
|---|---|---|---|
| 1세대 | 볼타(2017) | Tesla V100 | FP16 입력, FP16/FP32 누산. CUDA 9의 WMMA API[1] |
| 2세대 | 튜링(2018) | RTX 20 시리즈, T4 | 추론용 INT8, INT4 모드 추가. 지포스 제품군에 처음 탑재, DLSS 같은 그래픽 기능에도 사용[4] |
| 3세대 | 암페어(2020) | A100, RTX 30 시리즈 | TF32, BF16, FP64(IEEE 준수), INT8·INT4·이진. 2:4 미세 구조적 희소성. 텐서 코어당 클럭당 FP16 FMA 256개[3] |
| 4세대 | 호퍼(2022) | H100 | FP8(E4M3, E5M2), 트랜스포머 엔진, 비동기 워프그룹 MMA(wgmma), TMA. 지원 형식은 FP8·FP16·BF16·TF32·FP64·INT8[2] |
| 4세대 | 에이다 러브레이스(2022) | RTX 40 시리즈, L4 | FP8 트랜스포머 엔진을 지포스에 도입[5] |
| 5세대 | 블랙웰(2024) | B200, GB200 | FP4·FP6, 마이크로스케일링 형식(MXFP8·MXFP6·MXFP4), NVFP4, 2세대 트랜스포머 엔진, tcgen05 명령과 텐서 메모리[6][7] |
| 5세대(강화) | 루빈(예정) | 베라 루빈 NVL72 | NVFP4·FP8 최적화. 엔비디아는 루빈 GPU의 NVFP4 추론 성능을 50 PFLOPS로 예고(예비 사양)[6] |
| 형식 | V100 | A100(밀집 / 희소) | H100 SXM5(밀집 / 희소, 발표 당시 예비치) |
|---|---|---|---|
| FP16 텐서 | 125 TFLOPS | 312 / 624 TFLOPS | 1000 / 2000 TFLOPS |
| BF16 텐서 | 지원 안 함 | 312 / 624 TFLOPS | 1000 / 2000 TFLOPS |
| TF32 텐서 | 지원 안 함 | 156 / 312 TFLOPS | 500 / 1000 TFLOPS |
| FP64 텐서 | 지원 안 함 | 19.5 TFLOPS | 60 TFLOPS |
| FP8 텐서 | 지원 안 함 | 지원 안 함 | 2000 / 4000 TFLOPS |
| INT8 텐서 | 지원 안 함 | 624 / 1248 TOPS | 2000 / 4000 TOPS |
| 참고: FP32(CUDA 코어) | 15.7 TFLOPS | 19.5 TFLOPS | 60 TFLOPS |
수치는 볼타·암페어 백서와 호퍼 기술 블로그의 표에서 가져왔다.[1][3][2] 호퍼 블로그는 같은 형식끼리 비교할 때 H100 SM의 텐서 코어 MMA 처리량이 A100 SM의 2배이고, FP8을 쓰면 A100의 16비트 대비 4배라고 설명했다.[2] 희소 수치는 아래 2:4 희소성을 썼을 때의 실효값이다.
| 형식 | 부호·지수·가수 비트 | 도입 세대 | 특징 |
|---|---|---|---|
| FP16 | 1·5·10 | 볼타 | 범위가 좁아 학습 때 손실 스케일링이 필요 |
| BF16 | 1·8·7 | 암페어 | FP32와 같은 지수 범위, 가수 정밀도는 낮음. FP16과 같은 처리량 |
| TF32 | 1·8·10 | 암페어 | FP32 범위 + FP16 정밀도. FP32 데이터를 넣으면 텐서 코어 안에서 TF32로 계산. A100에서 FP32 대비 10배(희소 20배) |
| FP8 E4M3 / E5M2 | 1·4·3 / 1·5·2 | 호퍼 | E4M3는 정밀도, E5M2는 범위 우선. FP16 대비 저장 절반, 처리량 2배 |
| FP6, FP4(E2M1) | FP4는 1·2·1 | 블랙웰 | FP4가 표현하는 값은 0, 0.5, 1, 1.5, 2, 3, 4, 6과 그 음수 |
| MXFP4 | FP4 + 32개 블록마다 E8M0(2의 거듭제곱) 스케일 | 블랙웰 | 업계 공동 마이크로스케일링 형식 |
| NVFP4 | FP4 + 16개 블록마다 E4M3 스케일 + 텐서당 FP32 스케일 | 블랙웰 | 값당 약 4.5비트. FP16 대비 메모리 약 3.5분의 1, FP8 대비 약 1.8분의 1 |
| INT8, INT4 | 정수 | 튜링 | 추론용 양자화. 하위 바이트(INT4, 1비트) WMMA 일부는 sm_90에서 제거 예정으로 공지 |
TF32·BF16 비트 구성과 TF32 성능은 암페어 백서,[3] FP8은 호퍼 블로그,[2] FP4·MXFP4·NVFP4는 엔비디아 NVFP4 소개 글에 따랐다.[7] 하위 바이트 연산은 CUDA 프로그래밍 가이드가 "미리 보기" 기능으로 분류하며 u4, s4, XOR 이진 연산 MMA 변형은 sm_90에서 제거될 예정이라고 밝혀 두었다.[8]
저정밀 형식은 값의 범위가 좁아 텐서나 블록마다 스케일 인자를 따로 둔다. 호퍼의 트랜스포머 엔진은 층마다 FP8과 16비트 중 무엇을 쓸지 고르고 다시 형변환과 스케일링을 자동으로 처리하는 소프트웨어·하드웨어 조합이다.[2] 블랙웰은 스케일 단위를 텐서에서 작은 블록(16~32개 원소)으로 줄여 4비트에서도 정확도를 지키려 한다.[7] → 신경망 양자화
암페어부터 텐서 코어는 가중치 행렬에서 연속한 4개 값 중 2개가 0인 패턴(2:4)을 하드웨어로 이용한다. 0이 아닌 값과 위치 정보(메타데이터)만 저장하므로 데이터 크기와 대역폭이 약 절반으로 줄고, 희소 텐서 코어는 0을 건너뛰어 연산 처리량이 2배가 된다. 엔비디아가 제시한 방법은 밀집 가중치로 학습한 뒤 2:4 패턴으로 가지치기하고, 남은 가중치를 추가로 미세 조정하는 것이며, 여러 비전·언어 신경망에서 추론 정확도 손실이 거의 없었다고 백서는 밝혔다.[3] PTX에서는 mma.sp, 호퍼의 wgmma.mma_async.sp, 블랙웰의 tcgen05.mma.sp가 이를 쓴다.[9] 실제 속도 이득은 행렬 크기와 커널 구현에 따라 2배보다 작을 수 있다.
| 계층 | 단위 | 대상 아키텍처 | 내용 |
|---|---|---|---|
WMMA(CUDA C++ nvcuda::wmma) |
워프 | 볼타 이후(compute capability 7.0+) | 조각(fragment) 적재·곱셈·저장 API. 레지스터 배치는 감춰져 있음 |
mma.sync(PTX) |
워프 | 볼타 이후 | m16n8k16 등 모양을 지정하는 명령. 레지스터 배치가 문서화되어 있어 직접 최적화 가능 |
wgmma.mma_async(PTX) |
워프그룹(연속한 워프 4개, 128스레드) | 호퍼(sm_90a) | 비동기 실행, 피연산자를 공유 메모리에서 행렬 기술자(descriptor)로 직접 읽음 |
tcgen05.mma(PTX) |
스레드 하나가 발행 | 블랙웰(sm_100a 등) | 누산기를 레지스터 대신 SM의 전용 텐서 메모리(CTA당 128레인×512열, 칸당 32비트)에 둠 |
| 라이브러리 | 호스트 API | 전 세대 | cuBLAS, cuBLASLt, cuDNN, CUTLASS, 트랜스포머 엔진 |
| 프레임워크 | 파이썬 | 전 세대 | 파이토치 AMP, TF32 설정, torch.compile, Triton |
WMMA와 PTX 명령의 세부는 CUDA 가이드와 PTX ISA 문서를 따랐다.[8][9] 호퍼 이후 최고 성능 커널(예: 플래시 어텐션 3)은 WMMA가 아니라 wgmma·TMA를 직접 쓰는 CUTLASS 템플릿이나 손으로 짠 커널로 만든다. 세대마다 명령 체계가 크게 바뀌므로 대부분의 개발자는 라이브러리와 프레임워크 계층만 쓴다.
CUDA 가이드의 WMMA 표에 나온 조합이다(m-n-k).[8]
- FP16 × FP16 → FP32 또는 FP16 누산: 16×16×16, 32×8×16, 8×32×16
- 부호 있는·없는 8비트 정수 × 같은 형식 → INT32 누산: 16×16×16, 32×8×16, 8×32×16
- BF16 × BF16 → FP32 누산: FP16과 같은 세 모양
- TF32 × TF32 → FP32 누산: 16×16×8
- FP64 × FP64 → FP64 누산: 8×8×4
- (실험) u4·s4 → INT32: 8×8×32, 1비트 → INT32: 8×8×128
워프 하나가 16×16×16 FP16 행렬 곱을 FP32로 누산하는, CUDA 가이드의 예제와 같은 구조다.
#include <mma.h>
using namespace nvcuda;
// 블록당 워프 1개(32스레드)로 실행: wmma_16x16<<<1, 32>>>(a, b, c);
__global__ void wmma_16x16(const half* a, const half* b, float* c) {
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::fill_fragment(c_frag, 0.0f);
wmma::load_matrix_sync(a_frag, a, 16); // 마지막 인자: leading dimension
wmma::load_matrix_sync(b_frag, b, 16);
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag); // C = A*B + C
wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}
compute capability 7.0 이상에서 nvcc -arch=sm_70(또는 실제 GPU에 맞는 값)으로 컴파일한다. 워프의 32스레드가 모두 같은 분기에서 실행해야 하며, 조각 안의 원소가 어느 스레드 레지스터에 있는지는 아키텍처마다 달라 가정하면 안 된다.[8]
import torch
# FP32 행렬 곱을 TF32 텐서 코어로 계산 허용 (암페어 이후)
torch.backends.cuda.matmul.allow_tf32 = True # 파이토치 1.12부터 기본값 False
torch.backends.cudnn.allow_tf32 = True # 합성곱은 기본값 True
# 같은 효과: torch.set_float32_matmul_precision("high")
model = torch.nn.Linear(4096, 4096).cuda()
x = torch.randn(64, 4096, device="cuda")
with torch.autocast("cuda", dtype=torch.bfloat16): # BF16 텐서 코어
y = model(x)
TF32 허용 플래그 allow_tf32는 파이토치 1.7~1.11에서는 기본으로 켜져 있었고 1.12부터 행렬 곱은 기본으로 꺼졌다. 최신 판은 torch.backends.cuda.matmul.fp32_precision 같은 새 설정을 권하고 기존 플래그는 폐지 예정이다.[10] 텐서 코어 효율을 높이려면 행렬 차원(은닉 크기, 어휘 크기, 배치)을 8이나 16, 가능하면 64의 배수로 맞추는 것이 좋다.
- CUDA 코어: 범용 스칼라·벡터 연산. 활성화 함수, 정규화, 소프트맥스처럼 행렬 곱이 아닌 연산은 여기서 처리되므로, 텐서 코어가 빨라질수록 이런 연산이 새 병목이 된다. 플래시 어텐션 4 논문이 블랙웰에서 지수 연산 유닛과 공유 메모리 대역폭이 텐서 코어만큼 늘지 않은 점을 병목으로 지목한 것이 그 예다.
- 구글 TPU의 행렬 유닛: 시스톨릭 배열로 된 큰 행렬 곱 유닛을 칩의 중심에 둔다. 텐서 코어는 범용 GPU의 SM 안에 작은 행렬 유닛을 여럿 흩어 둔 방식이다. → 구글 TPU, 신경망 처리 장치
- 지원 형식은 세대마다 다르므로, 같은 모델이라도 GPU 세대에 따라 쓸 수 있는 정밀도(예: FP8은 호퍼·에이다 이후, FP4는 블랙웰 이후)와 성능이 달라진다.
- ↑ 1.0 1.1 1.2 1.3 NVIDIA, 〈NVIDIA Tesla V100 GPU Architecture〉 백서 WP-08608-001_v1.1, 2017. images.nvidia.com
- ↑ 2.0 2.1 2.2 2.3 2.4 2.5 NVIDIA Technical Blog, 〈NVIDIA Hopper Architecture In-Depth〉, 2022. developer.nvidia.com
- ↑ 3.0 3.1 3.2 3.3 3.4 NVIDIA, 〈NVIDIA A100 Tensor Core GPU Architecture〉 백서, 2020. images.nvidia.com
- ↑ NVIDIA, 〈NVIDIA Turing GPU Architecture〉 백서, 2018. images.nvidia.com
- ↑ NVIDIA, 〈Ada Lovelace Architecture〉. nvidia.com
- ↑ 6.0 6.1 NVIDIA, 〈Tensor Cores〉. nvidia.com
- ↑ 7.0 7.1 7.2 NVIDIA Technical Blog, 〈Introducing NVFP4 for Efficient and Accurate Low-Precision Inference〉, 2025. developer.nvidia.com
- ↑ 8.0 8.1 8.2 8.3 NVIDIA, 〈CUDA Programming Guide〉 5.4.11 Warp Matrix Functions. docs.nvidia.com
- ↑ 9.0 9.1 NVIDIA, 〈Parallel Thread Execution ISA〉 9.7.16~9.7.18. docs.nvidia.com
- ↑ PyTorch, 〈CUDA semantics〉 TensorFloat-32 절. docs.pytorch.org