제1부 GPU 아키텍처

1.12 텐서코어 II — 세대 진화·트랜스포머 엔진

1.12.1 이 장에서 배우는 것 / 선수 지식

1.11에서 텐서코어의 메커니즘(MMA·프래그먼트)과 호출법을 봤다. 이 장은 세대별 진화(Volta→Blackwell), FP8 학습을 자동화하는 트랜스포머 엔진, 그리고 "직접 짤까 라이브러리에 맡길까"의 결정을 다룬다.

배우는 것:

세대별 진화와 각 세대가 추가한 정밀도·기능을 개략 정리한다.
트랜스포머 엔진의 동적 정밀도 스케일링이 무엇을 하는지 설명한다.
텐서코어를 쓰는 추상화 수준(WMMA / cuBLASLt / CUTLASS / 도메인 라이브러리)을 워크로드별로 고른다.

선수 지식: 1.11(MMA·프래그먼트·호출). 정밀도 포맷의 비트 구성·수치 거동은 1.13~1.14, 추론 양자화는 5.4.

사양 주의(S-5): 세대별 형상·정밀도·처리량은 시점 의존. 미발표·로드맵(Blackwell 세부)은 발표 기준, 변동 가능. 정확한 수는 화이트페이퍼로.

1.12.2 세대별 진화 (개략 — 수치 확인 필요)

텐서코어는 세대마다 (1) 지원 정밀도 추가, (2) 형상·처리량 확대, (3) 비동기·소프트웨어 기능을 더해 왔다. 큰 줄기만:

세대(아키텍처)	대표 추가/특징	정밀도(개략)
Volta (1세대)	텐서코어 최초 도입. FP16 입력·FP32 누산 혼합정밀도	FP16
Turing (2세대)	정수 추론 가속 확대	FP16, INT8, INT4
Ampere (3세대)	TF32(FP32 대체 학습 경로)·BF16, 구조적 희소성(2:4), 비동기 복사로 급식 개선	TF32, BF16, FP16, INT8
Hopper (4세대)	트랜스포머 엔진 + FP8, 워프그룹 `wgmma`(비동기·대형 타일), TMA로 데이터 공급	FP8, BF16, FP16, TF32, INT8
Blackwell (차세대)	2세대 트랜스포머 엔진, FP4(및 더 좁은 포맷) 초저정밀 추론 — 발표 기준, 변동 가능	FP4, FP8, …

정직성(S-5): 각 세대가 정확히 어떤 타입·형상·희소성 모드를 지원하는지, FP4/FP6 신규 포맷의 정의·처리량은 해당 아키텍처 공식 문서로 확인. 미발표 항목은 발표 기준이며 양산 가용성과 다를 수 있다. 일관 흐름: "세대가 갈수록 더 좁은 정밀도를 추가하고, 데이터 공급(비동기 복사·TMA)과 소프트웨어 자동화(트랜스포머 엔진)를 강화한다."

지원 정밀도가 입력으로 무엇을 먹는지(상세 비트 구성은 1.13):

FP16/BF16: 학습·추론 주력. BF16은 지수 범위가 FP32와 같아 오버플로에 강함.
TF32: Ampere 도입. 입력은 FP32처럼 받되 내부에서 가수를 줄여 텐서코어로 처리 — 코드 변경 거의 없이 FP32 학습을 가속.
INT8/INT4: 정수 추론(양자화). 추론 처리량·메모리 절감.
FP8(E4M3·E5M2): Hopper 도입. LLM 학습·추론 저정밀.
FP4: Blackwell 세대 초저정밀 추론(발표 기준).

세대별 진화의 시스템 측면(Hopper의 TMA·클러스터, Blackwell의 FP4·NVLink)은 4.2~4.5(데이터센터 세대 심화)에서 다룬다.

1.12.3 트랜스포머 엔진 — 동적 정밀도 스케일링

FP8처럼 매우 좁은 포맷은 처리량은 좋지만 동적 범위가 좁다 — 값이 표현 범위를 벗어나면 오버플로(∞)나 언더플로(0)로 망가진다. 트랜스포머처럼 텐서별 스케일이 천차만별인 모델에서, 한 텐서엔 맞는 스케일이 다른 텐서엔 전부 0이 되기도 한다.

트랜스포머 엔진(Transformer Engine, TE) 은 Hopper와 함께 등장한 라이브러리+하드웨어 협력 기능으로 이를 자동 스케일링으로 푼다:

핵심 둘:

혼합정밀 자동화 — 어느 연산을 FP8로 내리고 어느 부분(누산·민감 연산)을 BF16/FP32로 둘지 TE가 결정·관리.
동적 스케일링 — 텐서별 amax를 추적해 스케일을 자동 조정함으로써, 프로그래머가 수동으로 양자화 스케일을 튜닝하지 않아도 FP8 학습이 수렴.

실무에선 PyTorch 등에서 TE의 Linear·LayerNorm 레이어로 모델을 구성하거나 프레임워크 통합으로 쓰며, FP8 경로 디테일은 라이브러리가 감춘다. Blackwell의 2세대 TE는 더 좁은 포맷(FP4 등)으로 확장(발표 기준) — 위에서 설명한 텐서별(per-tensor) 스케일링은 FP8에는 맞지만, FP4는 가수가 1비트뿐이라 동적 범위가 훨씬 좁아 텐서 하나에 스케일 하나로는 부족한 경우가 많다. 그래서 FP4 계열은 보통 더 잘게 분할한 블록마다 별도 스케일을 공유하는 방식(마이크로스케일링, OCP MX 포맷 계열 — 1.13.3의 "OCP MX (Blackwell)")을 쓴다 — per-tensor 하나가 아니라 텐서 안에 작은 블록마다 스케일이 여러 개라는 정도로 이해하면 충분하다. 정확한 블록 크기·스케일 포맷은 표준·구현 의존(확인 필요). 정밀도 포맷의 수치 거동과 양자화 일반론은 1.13~1.14, 추론 양자화는 5.4, 학습 통합은 6.2에서.

1.12.4 트레이드오프 — 직접 짤까, 라이브러리에 맡길까

텐서코어를 쓰는 방법은 추상화 수준으로 나뉜다. 거의 모든 실전의 답은 라이브러리지만 선택지를 펼친다:

① WMMA / 인라인 mma.sync·wgmma 직접 작성

누구에게: 라이브러리가 못 다루는 특수 융합을 직접 작성하거나, 텐서코어를 학습·연구하는 경우.
대가: 정렬·형상·레이아웃·세대별 명령 차이를 수동으로 처리. 세대 바뀌면 재튜닝. 생산성 낮음.

② cuBLAS·cuBLASLt (GEMM) — 3.1

누구에게: 표준 GEMM·배치 GEMM·에필로그 융합(bias·activation)을 텐서코어로. 대부분의 행렬곱.
대가: 라이브러리가 지원하는 패턴에 한정.

③ CUTLASS / CuTe (템플릿 라이브러리) — 3.4~3.5

누구에게: 라이브러리 GEMM으론 부족하지만 ①보다 생산적인 커스텀 텐서코어 커널(특수 융합·새 정밀도·새 형상). cuDNN·어텐션 커널의 토대.
대가: 학습 곡선 가파름. 그래도 ①보다 세대 이식성·재사용성 우수.

상위 도메인 라이브러리: cuDNN(3.2~3.3), 트랜스포머 엔진(이 장), TensorRT(5.2~5.6)가 내부에서 텐서코어를 호출한다.

견해: 텐서코어를 직접 잘 쓰는 것은 라이브러리 팀의 전업이다. 응용 개발자는 "내 GEMM·컨볼루션·어텐션이 텐서코어 경로를 타는지"를 프로파일러로 검증하고, 안 타면 정밀도·형상·정렬을 라이브러리가 좋아하게 맞추는 데 시간을 쓰는 것이 비용 대비 효과가 가장 크다. WMMA 직접 작성은 학습·연구·극한 융합에 한정하라.

비엔비디아 대안(정직): 다른 벤더도 행렬 가속 유닛을 둔다 — AMD Matrix Core(rocBLAS·Composable Kernel), 구글 TPU systolic MXU, 인텔 AMX, 여러 NPU. 발상(전용 행렬 유닛 + 저정밀 입력·고정밀 누산)은 공통이고, 차이는 소프트웨어 생태계 성숙도·정밀도 지원·통합 도구에 있다. 엔비디아 우위는 칩이 아니라 cuBLAS·cuDNN·CUTLASS·TE·TensorRT 라이브러리 스택의 폭과 성숙도에 크게 기댄다(0.1).

1.12.5 중간 점검 (체크포인트)

[ ] 세대 흐름을 "Volta=FP16 → Turing=INT → Ampere=TF32·BF16·2:4 → Hopper=FP8·TE·wgmma → Blackwell=FP4(발표 기준)"로 요약한다.
[ ] 트랜스포머 엔진이 텐서별 amax 추적·동적 스케일링으로 FP8 학습을 자동화함을 설명한다.
[ ] 표준 GEMM 융합엔 cuBLASLt, 커스텀엔 CUTLASS, 직접 작성은 학습·극한용임을 고른다.

1.12.6 흔한 함정

FP8/FP4를 스케일링 없이 사용. 동적 범위가 좁아 스케일을 안 맞추면 오버/언더플로로 학습 발산. 트랜스포머 엔진·양자화 스케일(1.14·5.4)을 거쳐야 한다.
세대 차이를 일반화. "Volta에서 되던 WMMA 형상·명령이 Hopper에서 최적"이라 가정하지 말 것. 세대마다 권장 명령(mma.sync→wgmma)·타일·정밀도가 다르다. 라이브러리가 이를 흡수.
TE를 "켜면 무비용"으로 여김. FP8 학습은 큰 모델·클러스터에서 검증 후 도입할 처리량 카드다. 품질 지표를 반드시 측정.
추상화 과잉 선택. 표준 GEMM에 WMMA를 직접 작성하는 것은 생산성·이식성 낭비. 라이브러리부터.

1.12.7 연습문제

연습 1 (세대 매칭). 다음이 처음 등장한(개략) 세대를 짝지어라: TF32, FP8·트랜스포머 엔진, 구조적 희소성(2:4), FP4(발표 기준).

정답

TF32 → Ampere, 2:4 희소성 → Ampere, FP8·TE → Hopper, FP4 → Blackwell(발표 기준). FP16 텐서코어 자체는 Volta. (시점 의존, 확인 필요.)

연습 2 (추상화 선택, 견해 포함). 표준 배치 GEMM에 bias+GELU 융합을 텐서코어로 하고 싶다. WMMA·cuBLASLt·CUTLASS 중 무엇을 고르나?

정답

1순위 cuBLASLt — 에필로그(bias·activation) 융합 지원으로 표준 패턴이면 가장 적은 노력으로 텐서코어 경로(3.1). 라이브러리가 못 다루는 특수 융합·형상이면 CUTLASS(3.4). WMMA 직접은 학습·연구·극한 커스텀 아니면 비추천.

연습 3 (TE 역할). FP8로 LLM을 학습할 때 TE가 자동으로 해주는 일 두 가지와, 수동으로 하기 어려운 이유를 설명하라.

정답

(1) 텐서별 동적 스케일링 — amax를 추적해 FP8 범위에 값을 맞춰 오버/언더플로 방지. (2) 혼합정밀 관리 — 어느 연산을 FP8로, 어느 부분을 BF16/FP32로 둘지 결정. 수동으로 하면 레이어·스텝마다 변하는 텐서 스케일을 일일이 추적·튜닝해야 해 비현실적이고, 잘못 잡으면 학습이 발산한다.

1.12.8 요약 · 다음 장

세대 흐름(확인 필요): Volta(FP16) → Turing(INT) → Ampere(TF32·BF16·2:4) → Hopper(FP8·TE·wgmma) → Blackwell(FP4, 발표 기준). 일관 흐름은 "더 좁은 정밀도 + 데이터 공급·자동화 강화".
트랜스포머 엔진은 텐서별 amax 기반 동적 스케일링으로 FP8 학습을 자동화한다.
실전 결론: 텐서코어는 라이브러리(cuBLASLt·cuDNN·CUTLASS·TE·TensorRT)에 맡기고, 프로파일러로 경로를 검증하라.

다음 장: 1.13 수치 정밀도 I — 부동소수 포맷. 텐서코어가 "무엇을 먹는가"를 넘어, FP32·TF32·FP16·BF16·FP8·FP4·INT8의 비트 구성·동적 범위·수치 거동을 비트 단위로 파고든다.