제1부 GPU 아키텍처

1.11 텐서코어 I — MMA·워프 수준 행렬곱

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

딥러닝·HPC 실행 시간 대부분은 행렬곱(GEMM)과 그 변형(컨볼루션·어텐션)에 쓰인다. 텐서코어는 이 연산만을 위해 SM 안에 박은 전용 행렬 가속기다. 일반 CUDA 코어가 스칼라 FMA를 하나씩 처리하는 동안, 텐서코어는 작은 행렬 블록 전체의 곱·누산을 한 명령으로 끝낸다. 이 장은 그 메커니즘(MMA·프래그먼트) 과 호출 방법을, 세대 진화·트랜스포머 엔진은 1.12로 잇는다.

배우는 것:

• 텐서코어가 가속하는 MMA(Matrix Multiply-Accumulate) $D = A\cdot B + C$를 한 명령으로 처리하는 의미.
• FMA 대비 처리량이 왜 도약하는지(명령당 산술량·곱셈기 밀도)를 타일 MMA로 설명.
• 워프 수준 MMA와 프래그먼트(fragment) 개념.
• 텐서코어 GEMM을 WMMA·cuBLASLt로 호출하는 코드 골격과 "실전엔 라이브러리"라는 판단.

선수 지식: 1.4~1.5(SM 내부 실행 유닛), 1.7(피연산자 출처: 레지스터·공유 메모리), 1.10(루프라인 — 텐서코어가 코어를 너무 빠르게 해 메모리 바운드로 넘어가는 이유). 정밀도 포맷의 비트 구성·거동은 1.13~1.14 전담 — 이 장은 "무엇을 먹는가" 수준만.

사양 주의(S-5): 세대별 형상·정밀도·처리량 수치는 시점 의존, 확인 필요. 정확한 수는 화이트페이퍼·CUDA 프로그래밍 가이드로.

1.11.2 왜 행렬 전용 유닛인가 — FMA 한 번 vs 행렬 한 번

CUDA 코어의 기본은 FMA(d = a*b + c): 스칼라 곱 1 + 덧셈 1을 한 명령에. 행렬곱 $C = A\cdot B$(M×K · K×N)는 출력 원소당 K번 FMA, 전체 $M\cdot N\cdot K$번. 텐서코어는 작은 행렬 블록 전체를 한 명령으로 곱·누산한다:

핵심은 명령당 산술량이다. 한 MMA가 m·n·k 곱·누산을 수행하므로 같은 명령 발행으로 수십~수백 배 산술을 끌어낸다. 발행·디코드·스케줄 오버헤드가 산술량에 비해 무시할 만큼 작아지고, 실리콘은 곱셈기 어레이를 빽빽이(systolic-array) 채워 면적·전력 효율을 극대화한다.

텐서코어는 독립 칩이 아니라 SM 내부 실행 유닛이다(1.4). 같은 SM의 CUDA 코어·레지스터 파일·공유 메모리·스케줄러와 자원을 공유한다. 이 위치에서 두 사실이 따라온다:

1. 텐서코어를 기아 상태로 두지 않으려면 공유 메모리·레지스터로 피연산자를 충분히 빨리 날라야 한다 → 타일링(2.7)·비동기 복사가 중요.
2. 텐서코어가 너무 빨라서 데이터를 못 대면 워크로드가 메모리 바운드로 넘어간다(1.10 루프라인). 텐서코어의 실질적 효과는 "연산 천장을 높이는 것"이고, 그 천장을 쓰려면 데이터 공급이 따라와야 한다.

1.11.3 타일 MMA와 처리량 도약 원리

텐서코어는 한 번에 작은 타일을 처리하고, 큰 GEMM은 이 타일 연산을 K 방향으로 누적·반복해 만든다.

결정적인 두 가지:

1. 입력은 저정밀(FP16/BF16/FP8…), 누산은 고정밀(보통 FP32). A·B는 좁은 포맷으로 읽어 메모리·대역폭을 아끼고, 부분곱의 합은 FP32 어큐뮬레이터에 쌓아 누적 오차(특히 K가 클 때)를 막는다. 이것이 혼합정밀도의 하드웨어적 본질(상세 1.14).
2. 데이터 재사용. 한 번 적재한 A·B 타일은 여러 출력 원소에 재사용된다. 타일 변을 키울수록 산술/바이트 비(I)가 올라가 메모리 천장에서 멀어진다(루프라인, 1.10).

처리량 도약은 곱셈기 개수 × 명령 효율의 결과다:

   CUDA 코어 FP32:  SM당 FP32 FMA 유닛 × 클럭 × 2
   텐서코어:        SM당 텐서코어 수 × (명령당 m·n·k) × 클럭 × 2
       ──▶ 같은 면적에 곱셈기를 훨씬 빽빽이 → FP16/BF16 텐서 처리량이
           FP32 CUDA 코어 대비 수 배~수십 배(세대·정밀도 의존)

정밀도를 더 좁히면(FP8·FP4) 같은 실리콘에서 더 많은 곱셈기를 가동해 처리량이 또 뛴다 — 세대가 갈수록 더 좁은 포맷을 추가하는 이유(1.12). 구체 배수·TFLOPS는 시점 의존, 확인 필요.

1.11.3.1 워프 수준 MMA와 프래그먼트

텐서코어 명령은 개별 스레드가 아니라 워프(32 스레드) 전체가 협력해 발행한다. 입력·출력 타일의 원소는 워프의 32개 스레드 레지스터에 분산 저장되며, 이 "한 스레드가 든 행렬 조각"을 프래그먼트(fragment) 라 한다.

프래그먼트 안에서 어느 스레드가 어느 원소를 쥐는지(레이아웃)는 하드웨어가 정의하며 불투명하다 — 프로그래머가 직접 인덱싱하지 않는다. 그래서 WMMA·MMA API는 "로드→MMA→스토어"를 하나로 제공하고 레이아웃을 감춘다. 이 불투명성이 정렬·형상 제약(1.12)의 근원이다. 세대가 갈수록 저수준 인터페이스가 등장했다: WMMA(워프 수준 C++) → 인라인 PTX mma.sync → Hopper의 워프그룹 wgmma(여러 워프 묶음, 비동기). 라이브러리가 이 저수준 명령을 세대별로 골라 쓴다.

1.11.4 워크스루 — 텐서코어 GEMM 호출

목표: 텐서코어로 GEMM을 실행하는 두 경로(저수준 WMMA, 고수준 cuBLASLt)를 골격으로 맛본다. 빌드·드라이버 전제는 0.3·A.1.

1.11.4.1 경로 A — WMMA API (저수준, 워프 수준 MMA)

WMMA(mma.h, nvcuda::wmma)는 워프 수준 MMA를 C++로 노출한다. 16×16×16 타일 하나를 처리하는 커널 골격:

// wmma_skel.cu  (골격 — 형상·정렬 제약 준수 필요)
#include <mma.h>
using namespace nvcuda;
__global__ void wmma_gemm(const half* A, const half* B, float* C, int M, int N, int K) {
    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> acc_frag;
    wmma::fill_fragment(acc_frag, 0.0f);           // 누산기 0
    for (int k = 0; k < K; k += 16) {              // K 방향 누산(1.11.3)
        wmma::load_matrix_sync(a_frag, A + /*타일 오프셋*/ 0, K); // 메모리→프래그먼트
        wmma::load_matrix_sync(b_frag, B + /*타일 오프셋*/ 0, K);
        wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag);      // 텐서코어 1명령
    }
    wmma::store_matrix_sync(C + /*타일 오프셋*/ 0, acc_frag, N, wmma::mem_row_major);
}

컴파일(개략): nvcc -arch=sm_80 wmma_skel.cu -o wmma_skel(SM 7.0+). 실제 동작 커널은 타일 좌표·경계 처리·공유 메모리 스테이징을 더 써야 하며, 여기선 호출 골격만. 전부 직접 작성하는 것은 학습용이고 실전에선 거의 라이브러리에 맡긴다(경로 B, 1.12).

1.11.4.2 경로 B — cuBLASLt (고수준, 실전 권장)

라이브러리 경로는 형상·정렬·세대별 최적 커널 선택을 알아서 한다. 상세 API·에필로그 융합은 3.1.

// 핸들/디스크립터 설정 일부 생략
//   입력 타입: CUDA_R_16F (FP16),  연산 타입: CUBLAS_COMPUTE_32F → 텐서코어 FP32 누산
//   cublasLtMatmulAlgoGetHeuristic(...) 으로 세대 최적 알고리즘 질의
cublasLtMatmul(ltHandle, matmulDesc, &alpha, A, Adesc, B, Bdesc,
               &beta, C, Cdesc, D, Ddesc, &algo, workspace, workspaceSize, stream);
//   compute/입력 타입을 TF32·BF16·FP8 등으로 바꾸면 해당 텐서코어 경로를 고른다.

핵심: 텐서코어 on/off를 직접 켜는 스위치가 아니라, 입력/연산 타입·형상·정렬이 조건을 만족하면 라이브러리가 텐서코어 커널을 고른다. FP32 순정으로 두면 CUDA 코어 경로, FP16/BF16/TF32로 바꾸면 텐서코어 경로 — 이 차이가 곧 "텐서코어 on/off" 실험이다.

주의: 작은 행렬·메모리 바운드 형상에서는 텐서코어를 켜도 도약이 안 보인다(1.10). 큰·정사각에 가깝고 K가 충분히 큰 GEMM에서 도약이 가장 잘 드러난다. 측정 시 워밍업·반복 평균·첫 호출 알고리즘 탐색 제외(2.14).

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

• [ ] FMA와 MMA의 차이를 "명령당 산술량"으로 한 문장에 설명한다.
• [ ] 타일 MMA에서 입력은 저정밀, 누산은 FP32인 이유(누적 오차·대역폭)를 제시한다.
• [ ] 프래그먼트가 "워프 스레드들에 분산된 행렬 조각, 레이아웃 불투명"임을 설명한다.
• [ ] cuBLASLt에서 입력/연산 타입을 바꾸는 것이 곧 텐서코어 경로 선택임을 안다.

증상: 텐서코어로 바꿨는데 안 빨라짐 → 형상이 작거나 메모리 바운드(1.10), 또는 정렬·형상 제약 위반으로 폴백(프로파일러로 실제 커널 확인, 2.14).

1.11.6 흔한 함정

• 정렬·형상 제약 무시. 차원(M·N·K)이 타일 배수여야 하고 포인터·leading dimension 정렬 요구가 있다. 어기면 라이브러리가 조용히 CUDA 코어 폴백. 차원을 타일 배수로 패딩하라.
• "텐서코어 켰다"고 믿지만 안 켜짐. 타입·형상·정렬이 조건을 못 맞추면 안 쓴다. 프로파일러로 실제 커널·SM 유닛 활용을 확인(2.14).
• 누산을 좁혀 정확도 붕괴. 입력을 FP16/FP8로 낮춰도 누산은 FP32 유지. K 큰 GEMM에서 누적 오차가 폭발한다.
• 피크 미달을 당연시. 발표 피크는 이상적 형상·완벽한 데이터 공급 가정. 실측은 루프라인 상의 달성 가능 상한과 비교(1.10).

1.11.7 연습문제

연습 1 (FMA vs MMA 산수). 한 MMA가 16×16×16 타일을 처리할 때 곱·덧셈(FLOP) 수는? 같은 산술을 스칼라 FMA로 하면 명령 몇 개?

연습 2 (누산 정밀도). A·B를 FP16, 누산도 FP16으로 하면 K=4096 GEMM에서 무엇이 문제고 왜 FP32 누산인가?

연습 3 (안 빨라진다). cuBLASLt FP16 GEMM이 FP32 대비 거의 안 빨라졌다. 점검 순서 3가지+.

1.11.8 요약 · 다음 장

• 텐서코어는 SM 내부 행렬 전용 유닛으로 $D = A\cdot B + C$(MMA)를 한 명령으로 타일 단위 곱·누산한다. 명령당 산술량이 커 발행 오버헤드가 희석되고 곱셈기 밀도로 처리량이 도약한다.
• 패턴은 일관: 좁은 입력 + FP32 누산. 입력·출력은 워프 스레드에 분산된 프래그먼트로 다뤄진다(레이아웃 불투명).
• 실전에선 라이브러리(cuBLASLt·cuDNN·CUTLASS)에 맡기고, 프로파일러로 "텐서코어 경로를 탔는지"를 검증한다.

다음 장: 1.12 텐서코어 II — 세대 진화·트랜스포머 엔진. Volta→Blackwell의 세대별 정밀도 추가, FP8 학습을 자동화하는 트랜스포머 엔진, 그리고 "직접 짤까 라이브러리에 맡길까"의 결정을 본다.