10.2 양자화(Quantization) - FP8 학습과 추론의 시대

10.2 양자화(Quantization) - FP8 학습과 추론의 시대

2025-12-22, G30DR

1. 서론: 트랜스포머 싱귤래리티와 정밀도의 종말

2025년 현재, 인공지능 학계와 산업계는 ’트랜스포머 싱귤래리티(Transformer Singularity)’라 불리는 특이점에 도달했다. 모델의 파라미터 수는 무어의 법칙을 비웃기라도 하듯 기하급수적으로 증가하여, 이제 단일 모델이 수조(Trillion) 단위의 파라미터를 갖는 것이 놀랍지 않은 시대가 되었다.1 이러한 모델의 거대화는 필연적으로 하드웨어, 특히 GPU 메모리와 메모리 대역폭(Memory Bandwidth)의 한계를 시험하게 만들었다. 과거 딥러닝의 표준이었던 FP32(단정밀도)는 이미 수년 전 역사 속으로 사라졌고, 지난 3년간 사실상의 표준(De facto Standard)으로 군림했던 BF16(Brain Floating Point 16)조차 이제는 병목의 주원인으로 지목되고 있다.

바야흐로 FP8(8-bit Floating Point)의 시대가 도래했다. 이는 단순히 데이터를 더 작게 압축하여 메모리를 절약하려는 시도가 아니다. 이는 엔비디아(NVIDIA)의 H100 Hopper 아키텍처에서 시작되어 B200 Blackwell 아키텍처에서 완성된, 하드웨어와 소프트웨어, 그리고 알고리즘의 공진화(Co-evolution)가 만들어낸 새로운 컴퓨팅 패러다임이다.3 과거 INT8(8비트 정수) 양자화가 주로 추론(Inference) 단계에서 제한적으로 사용되었던 것과 달리, FP8은 LLM(Large Language Model)의 사전 학습(Pre-training)과 미세 조정(Fine-tuning), 그리고 추론을 모두 아우르는 범용 포맷으로 자리 잡았다.

특히 2024년 말 공개된 DeepSeek-V3와 같은 모델은 FP8 혼합 정밀도(Mixed Precision) 학습을 통해 수조 토큰 규모의 대규모 학습이 가능함을 입증하며, FP8이 더 이상 실험실의 이론적 유희가 아닌 상용 모델 개발의 핵심 도구임을 증명했다.2 본 장에서는 OCP(Open Compute Project) 표준에 기반한 FP8의 기술적 세부 사항, 이를 지원하는 최신 GPU 아키텍처의 진화, 그리고 실제 학습과 추론 파이프라인에서 발생하는 수치적 불안정성(Numerical Instability)을 극복하기 위한 최신 기법들을 심도 있게 분석한다.

2. 수의 물리학: FP8의 이중 구조와 OCP 표준

FP8이 기존의 정수 기반 양자화(INT8)와 가장 극명하게 차별화되는 지점은 데이터의 분포 특성에 따라 포맷을 이원화했다는 점이다. 8비트라는 극도로 제한된 정보 공간(256개의 상태) 안에서 ’표현 범위(Dynamic Range)’와 ’정밀도(Precision)’라는 상충되는 두 가치를 모두 만족시키는 것은 불가능하다. 따라서 OCP 표준은 FP8을 E4M3와 E5M2라는 두 가지 상호 보완적인 형식으로 정의하여 이 딜레마를 해결하고자 했다.7

2.1 포맷 정의 및 비트 할당

FP8의 두 가지 변형은 각각의 목적에 맞게 비트를 배분한다. 부호 비트(Sign bit)는 1비트로 고정되지만, 지수(Exponent)와 가수(Mantissa)의 비율에 따라 그 성격이 완전히 달라진다.

포맷 (Format)부호 (Sign)지수 (Exponent)가수 (Mantissa)편향 (Bias)최대 표현값 (Max Value)정밀도 (Machine Epsilon)주요 용도
FP8 E4M314374482^{-3} = 0.125가중치(Weight), 활성화(Activation)
FP8 E5M21521557,3442^{-2} = 0.25구배(Gradient), 옵티마이저 상태
BF16 (참고)1871273.39 \times 10^{38}2^{-7} \approx 0.0078일반 학습 (FP32 대체)
FP16 (참고)15101565,5042^{-10} \approx 0.00097구형 혼합 정밀도 학습

표 10.2.1 FP8 포맷(E4M3, E5M2)과 기존 포맷(BF16, FP16)의 구조적 비교 7

2.2 E5M2: 광활한 다이내믹 레인지의 확보

E5M2 포맷은 IEEE 754 표준의 FP16(반정밀도) 형식을 그대로 축소한 형태와 유사하다. 지수 비트가 5개라는 것은 FP16과 동일한 지수 범위를 갖는다는 것을 의미하며, 이는 FP8이 표현할 수 있는 수의 크기 범위(Magnitude)가 상당히 넓다는 것을 시사한다.

  • 구조적 특징: 1개의 부호 비트, 5개의 지수 비트, 2개의 가수 비트로 구성된다.
  • 다이내믹 레인지: 최대 약 \pm 57,344까지 표현 가능하다. 이는 무한대(Infinity)와 NaN(Not a Number)을 명시적으로 표현하는 IEEE 754 표준을 준수하기 때문에, 수치적 오버플로우(Overflow)가 발생했을 때 이를 감지하고 처리하는 기존의 로직을 그대로 사용할 수 있다는 장점이 있다.10
  • 활용 전략: 역전파(Backpropagation) 과정에서 계산되는 구배(Gradient)는 레이어를 거치며 그 크기가 기하급수적으로 변할 수 있다. 특히 손실 함수(Loss Function)의 초반부에서는 구배의 크기가 매우 크거나 작을 수 있어 넓은 다이내믹 레인지가 필수적이다. 또한, Adam이나 RMSProp과 같은 옵티마이저의 2차 모멘텀(Second-order momentum)은 값의 범위가 넓기 때문에 E5M2가 적합하다.8
  • 한계: 가수가 단 2비트에 불과하여 유효숫자의 정밀도가 극도로 낮다. 이는 1.0과 다음 표현 가능한 수 사이의 간격이 0.25나 됨을 의미한다. 이러한 낮은 해상도는 미세한 가중치 업데이트가 필요한 학습 후반기에 수렴을 방해하는 요인이 될 수 있다.8

2.3 E4M3: 정밀도를 위한 타협과 혁신

E4M3 포맷은 지수 비트를 하나 희생하여 4개로 줄이는 대신, 가수 비트를 3개로 늘려 정밀도를 확보했다. 이는 FP8 도입 초기 가장 논란이 되었던 부분이기도 했으나, 현재는 LLM 학습의 핵심 포맷으로 자리 잡았다.

  • 구조적 특징: 1개의 부호 비트, 4개의 지수 비트, 3개의 가수 비트로 구성된다.
  • 비표준적 설계: E4M3는 정해진 비트 공간을 최대한 활용하기 위해 IEEE 754 표준을 위반하는 과감한 설계를 채택했다. 구체적으로, 무한대(Inf) 표현을 제거하고, NaN 표현을 단 하나의 비트 패턴(부호와 상관없이 0x7F 등)으로 제한했다. 이렇게 확보한 비트 패턴을 일반 수치를 표현하는 데 할당함으로써, 표현 가능한 최대값을 448까지 확장하고 수의 밀도를 높였다.8
  • 활용 전략: 신경망의 순전파(Forward pass)에서 사용되는 가중치(Weight)와 활성화(Activation) 값은 정규화(Normalization) 레이어 등을 통해 비교적 좁은 범위 내에 분포하도록 제어된다. 따라서 \pm 448의 범위는 대부분의 경우 충분하며, 오히려 추론의 정확도를 결정짓는 것은 수치 간의 미세한 차이를 구분하는 정밀도다. 따라서 E4M3는 가중치와 활성화 텐서의 저장 포맷으로 표준화되었다.13
  • 기계 엡실론: E4M3의 기계 엡실론은 2^{-3}=0.125로, E5M2에 비해 2배 더 높은 정밀도를 제공한다. 이는 모델이 학습하거나 추론할 때 정보 손실을 최소화하여 성능 저하(Degradation)를 막는 방파제 역할을 한다.8

2.4 포맷 선택의 딜레마와 최신 경향

초기 엔비디아의 트랜스포머 엔진(Transformer Engine) 가이드라인은 순전파에는 E4M3를, 역전파의 구배에는 E5M2를 사용하는 하이브리드 방식을 권장했다.13 그러나 2025년 시점의 최신 연구와 DeepSeek-V3 등의 사례는 이러한 관행에 도전하고 있다.

DeepSeek-V3는 모든 텐서(가중치, 활성화, 구배)에 E4M3 포맷을 적용하는 파격을 선보였다.6 이는 블록 단위 스케일링(Block-wise Scaling) 기술이 발전함에 따라, 좁은 다이내믹 레인지 문제를 국소적인 스케일링으로 해결할 수 있게 되었기 때문이다. 즉, 전체 텐서가 \pm 448을 넘더라도, 작은 블록 단위로 나누어 스케일링하면 각 블록 내부의 값들은 E4M3 범위 내에 안착시킬 수 있다. 이는 FP8 포맷의 선택이 고정된 규칙이 아니라, 모델의 아키텍처와 스케일링 전략에 따라 유동적으로 결정되어야 함을 시사한다.

3. 하드웨어 가속의 진화: Hopper를 넘어 Blackwell로

FP8의 이론적 이점은 하드웨어의 물리적 지원 없이는 실현될 수 없다. 엔비디아는 2022년 H100(Hopper)을 통해 FP8 텐서 코어를 처음 선보였고, 2025년 주류가 된 B200(Blackwell)에서는 이를 완성형에 가깝게 진화시켰다. 이 두 아키텍처의 차이를 이해하는 것은 효율적인 FP8 커널을 작성하는 데 필수적이다.

3.1 H100 Hopper와 트랜스포머 엔진의 한계

H100 GPU는 4세대 텐서 코어를 탑재하여 FP16 대비 2배의 이론적 처리량(Throughput)인 약 4 PFLOPS(SXM 모델 기준)를 제공한다.3 H100의 등장은 ’트랜스포머 엔진’이라는 개념을 대중화시켰는데, 이는 소프트웨어 계층에서 텐서의 값을 모니터링하다가 FP8로 표현 가능한 범위 내에 들어오면 자동으로 정밀도를 낮추어 연산을 수행하는 메커니즘이다.14

그러나 H100의 FP8 구현에는 몇 가지 근본적인 한계가 있었다.

  1. 텐서 단위 양자화의 강제: H100의 텐서 코어는 연산 효율성을 위해 행렬 전체에 대해 하나의 스케일링 인자(Scalar Scale Factor)만을 적용하는 구조에 최적화되어 있었다. 이는 이상치(Outlier)가 하나만 존재해도 전체 행렬의 정밀도가 희생되는 결과를 낳았다.12
  2. 누적기(Accumulator) 정밀도 문제: 행렬 곱셈(C = A \times B + C) 과정에서 A와 B는 FP8이지만, 그 곱의 합인 C는 정밀도를 잃지 않기 위해 FP16 또는 FP32로 누적되어야 한다. H100은 FP32 누적기를 지원하지만, 이 과정에서의 데이터 이동 오버헤드가 병목으로 작용했다.18
  3. 디퀀타이제이션 비용: 블록 단위 양자화를 소프트웨어적으로 구현하려면 CUDA 코어(SM)를 사용하여 별도의 디퀀타이제이션 연산을 수행해야 했는데, 이는 텐서 코어의 속도 이점을 상쇄할 만큼 비용이 컸다.

3.2 Blackwell B200과 MXFP8의 하드웨어 네이티브 지원

2025년의 표준인 Blackwell 아키텍처는 H100의 한계를 극복하고 FP8 성능을 비약적으로 향상시켰다. B200은 FP8 텐서 코어 성능이 약 9 PFLOPS에 달해 H100 대비 2.25배 이상의 성능 향상을 이뤘으며, FP4 정밀도까지 지원하여 최대 18 PFLOPS의 연산 능력을 보여준다.3

가장 혁신적인 변화는 마이크로스케일링(Microscaling, MX) 포맷의 하드웨어 네이티브 지원이다. Blackwell 텐서 코어는 OCP MXFP8 사양을 실리콘 레벨에서 구현했다.

  • 하드웨어 구조: 32개의 인접한 원소(Element)를 하나의 블록으로 묶고, 이 블록마다 별도의 스케일링 인자(E8M0 포맷)를 적용한다.21
  • tcgen05 명령어: Blackwell은 tcgen05와 같은 5세대 텐서 코어 명령어를 도입했다. 이 명령어는 텐서 메모리(Tensor Memory)에서 데이터를 로드하면서 실시간으로 블록 스케일링을 적용하여 행렬 곱셈을 수행한다.22 즉, 기존에 CUDA 코어에서 수행하던 복잡한 디퀀타이제이션 로직이 텐서 코어 파이프라인 내부로 흡수된 것이다.
  • 성능 영향: 이를 통해 개발자들은 소프트웨어 오버헤드 없이 미세 입도(Fine-grained) 양자화를 적용할 수 있게 되었으며, 이는 이상치가 많은 트랜스포머 모델의 어텐션 레이어에서도 정밀도 저하 없이 FP8을 사용할 수 있게 하는 결정적인 기술적 도약이다.5

3.3 누적기 정밀도와 승격(Promotion) 전략

FP8 연산이라 하더라도, 내부적인 누적(Accumulation)은 반드시 고정밀도로 수행되어야 한다. 수천 개의 FP8 값을 더할 때 발생하는 정밀도 손실(Swamping)은 모델의 학습을 발산시키는 주원인이다.

Blackwell의 텐서 코어는 내부적으로 더 높은 정밀도(일반적으로 14비트 이상의 중간 포맷 또는 FP32)로 부분합(Partial Sum)을 계산한다. DeepSeek-V3와 같은 최신 모델 학습 커널은 청크 단위 승격(Chunk-based Promotion) 전략을 사용한다. 즉, 텐서 코어 내부에서 일정 개수(예: 128개)의 곱셈-덧셈 연산(FMA)을 수행한 후, 그 중간 결과를 FP32 레지스터로 이동(Promotion)시켜 누적한다.6 이 방식은 텐서 코어의 빠른 연산 속도와 FP32의 수치적 안정성 사이의 균형을 맞추는 현대적 GPU 프로그래밍의 핵심 패턴이다.

4. FP8 학습(Training)의 수학과 방법론

FP8을 추론에 사용하는 것은 상대적으로 수월하다. 이미 학습된 가중치는 고정되어 있고, 활성화 값의 범위도 보정(Calibration)을 통해 예측 가능하기 때문이다. 그러나 FP8로 모델을 처음부터 학습(Pre-training)하는 것은 차원이 다른 문제다. 학습 초기에는 가중치와 구배가 매우 불안정하게 변동하며, 역전파 과정에서의 수치적 오차가 증폭되어 학습 실패로 이어질 수 있다.

4.1 스케일링 전략: 지연된 스케일링(Delayed Scaling) vs. 즉시 스케일링

FP8의 제한된 다이내믹 레인지 안에 실수(Real number) 데이터를 매핑하기 위해서는 적절한 스케일링 인자(S)를 찾아 X_{fp8} = \text{Clamp}(X_{fp32} / S) 형태로 변환해야 한다. 여기서 S를 어떻게 결정하느냐가 학습의 성패를 가른다.

  1. 즉시 스케일링(Just-in-Time Scaling): 매 연산마다 텐서의 절댓값 최댓값(Amax)을 계산하여 스케일링한다. 가장 이상적이지만, 메모리 읽기/쓰기를 두 번(Amax 계산 1회, 양자화 1회) 해야 하므로 메모리 대역폭 낭비가 심하여 실시간 학습에는 부적합하다.17
  2. 지연된 스케일링(Delayed Scaling): NVIDIA Transformer Engine이 채택한 방식으로, 과거 N번의 반복(Iteration) 동안 관측된 Amax 값들의 최대치를 기반으로 현재의 스케일링 인자를 결정한다.14
  • 메커니즘: S_{t} = \max(\text{Amax}_{t-1}, \text{Amax}_{t-2}, \dots, \text{Amax}_{t-N}).
  • 장점: 추가적인 메모리 접근 없이 높은 처리량을 유지할 수 있다.
  • 단점: 학습 초기의 급격한 값 변화(Spike)에는 대응하기 어렵다. 과거 데이터에 기반하기 때문에 현재 텐서가 갑자기 커지면 오버플로우가 발생할 수 있다.

4.2 DeepSeek-V3의 혁신: 미세 입도(Fine-grained) 온라인 양자화

DeepSeek-V3는 기존의 지연된 스케일링이 가진 한계를 극복하기 위해 타일 및 블록 단위의 온라인(Online) 양자화를 도입했다.6

  • 블록 구성: 활성화(Activation) 텐서는 1 \times 128 크기의 타일 단위로, 가중치(Weight)는 128 \times 128 크기의 블록 단위로 스케일링 인자를 각각 계산한다.
  • 온라인 계산: 지연된 스케일링 대신, 현재 블록의 Amax를 즉시 계산하여 적용한다. “계산 비용이 비싸다“는 통념과 달리, DeepSeek팀은 커널 최적화(Kernel Fusion)를 통해 Amax 계산과 양자화를 하나의 커널 안에서 수행함으로써 오버헤드를 무시할 수 있는 수준으로 낮췄다.
  • E4M3 단일화: 이러한 정교한 로컬 스케일링 덕분에, DeepSeek-V3는 구배를 포함한 모든 텐서에 E4M3 포맷을 적용할 수 있었다. 국소적인 범위에서는 E4M3의 좁은 레인지로도 충분히 값을 표현할 수 있기 때문이다. 이는 혼합 포맷(E4M3/E5M2) 관리의 복잡성을 제거하고 메모리 효율을 극대화했다.15

4.3 확률적 라운딩(Stochastic Rounding)의 필수성

구배(Gradient)를 FP8로 양자화할 때 가장 치명적인 문제는 ’사라지는 업데이트(Vanishing Updates)’다. 구배 값은 0에 가까운 매우 작은 값들이 많은데, 이를 가장 가까운 FP8 값으로 반올림(Round-to-nearest)하면 대부분 0이 되어버리거나(Underflow), 특정 값으로 편향(Bias)되어 학습이 정체된다.

이를 해결하기 위해 확률적 라운딩(Stochastic Rounding) 기법이 필수적으로 적용된다.14

  • 원리: 어떤 실수 x가 두 표현 가능한 FP8 값 a와 b 사이에 위치할 때 (a \le x \le b), 다음과 같은 확률로 반올림을 수행한다.
    P(\hat{x} = b) = \frac{x - a}{b - a}, \quad P(\hat{x} = a) = 1 - P(\hat{x} = b)

  • 효과: 이렇게 하면 양자화된 값 \hat{x}의 기댓값 E[\hat{x}]이 원래 값 x와 같아진다 (E[\hat{x}] = x). 즉, **비편향 추정량(Unbiased Estimator)**이 되어, 수만 번의 반복 학습 동안 누적되는 구배의 방향성을 통계적으로 보존할 수 있다. 이는 저정밀도 학습의 수렴성을 수학적으로 보장하는 핵심 장치다.

5. 추론(Inference) 가속과 KV 캐시 혁명

학습된 모델을 서비스하는 추론 단계에서 FP8의 효용성은 메모리 대역폭 절약과 KV 캐시(Key-Value Cache) 용량 증대에서 극대화된다. 추론 비용의 대부분은 연산(Compute)이 아닌 메모리 이동(Memory Movement)에서 발생하기 때문이다.

5.1 vLLM과 FP8 KV 캐시

LLM 추론, 특히 긴 문맥(Long Context)을 처리할 때 GPU 메모리의 대부분은 KV 캐시가 점유한다. KV 캐시란 이전에 계산된 토큰들의 Key와 Value 텐서를 저장해두는 공간으로, 문맥 길이가 길어질수록 선형적으로 증가한다. KV 캐시를 FP16(16비트)에서 FP8(8비트)로 전환하면 메모리 사용량을 즉각적으로 50% 줄일 수 있다.24

  • 성능 이득: 메모리 절약은 곧 처리량(Throughput) 증대로 이어진다. 동일한 GPU 메모리(예: H100 80GB)에서 배치 크기(Batch Size)를 2배 이상 늘리거나, 2배 더 긴 문맥을 처리할 수 있다. vLLM 벤치마크에 따르면, Llama 3 70B 모델에서 FP8 KV 캐시 적용 시 최대 3배의 처리량 향상이 관측되었다.25
  • 포맷 선택: vLLM 문서는 KV 캐시용으로 E4M3를 추천한다.24 E5M2는 다이내믹 레인지는 넓지만 정밀도가 낮아, 어텐션 메커니즘에서 중요한 정보(Key/Value 간의 미세한 연관성)를 잃을 위험이 크기 때문이다. 다만, E4M3의 좁은 범위를 보완하기 위해 텐서별 스케일링 인자(Per-tensor scale factor)를 반드시 함께 저장해야 한다.

5.2 FP8 vs INT8: 왜 트랜스포머는 FP8인가?

과거 CNN(Convolutional Neural Networks) 시절에는 INT8이 추론의 표준이었다. 그러나 트랜스포머 기반 LLM에서는 FP8이 확실한 우위를 점한다. 그 이유는 데이터의 분포 특성(Distribution Logic)에 있다.

  1. 이상치(Outlier) 문제: 트랜스포머 모델의 활성화 값, 특히 어텐션 레이어의 출력값은 꼬리가 긴(Heavy-tailed) 분포를 가지며, 극단적인 이상치가 빈번하게 등장한다.12
  2. 선형 vs 비선형 양자화:
  • INT8: 모든 구간의 간격이 동일한 선형 양자화다. 이상치를 표현하기 위해 스케일링 인자를 키우면, 0 근처의 작은 값들이 모두 하나의 값(0)으로 뭉개지는 현상이 발생한다. 이를 해결하려면 복잡한 이상치 분리(Outlier Suppression) 알고리즘이 필요하다.
  • FP8: 지수 비트를 사용하는 비선형 양자화다. 0 근처에서는 매우 조밀하게, 큰 값 영역에서는 듬성듬성하게 값을 표현할 수 있다. 이는 LLM의 데이터 분포 특성과 기하학적으로 완벽하게 일치한다.
  1. 정확도 비교: 연구 결과에 따르면, FP8-E4 포맷은 별도의 재학습(Quantization Aware Training, QAT) 없이 학습 후 양자화(Post-Training Quantization, PTQ)만으로도 INT8보다 월등히 높은 추론 정확도를 보여준다. 특히 BERT나 ViT와 같은 모델에서 FP8은 원본 FP16 성능을 거의 그대로 유지하는 반면, INT8은 현저한 성능 저하를 보인다.26

5.3 Llama 3와 FP8 추론의 실제

2024년 출시된 Meta의 Llama 3.1 405B 모델은 FP8 추론의 실효성을 증명하는 대표적인 사례다. 이 모델은 BF16으로 학습되었으나, 공식적으로 FP8 양자화 버전을 배포하고 있다. 벤치마크 결과, FP8 버전은 BF16 버전과 비교하여 품질 차이가 거의 없으면서도 단일 노드(8x H100)에서 실행 가능한 메모리 효율성을 제공한다.5 이는 FP8이 단순히 “가능한” 옵션이 아니라, 초거대 모델 배포를 위한 “필수적인” 표준이 되었음을 의미한다.

6. 마이크로스케일링(MX)과 블록 양자화의 표준화

2025년 시점에서 가장 중요한 기술적 흐름은 텐서 단위(Per-tensor) 양자화에서 블록 단위(Block-wise) 양자화로의 완전한 전환이다. OCP는 이를 Microscaling (MX) 포맷으로 표준화하여 하드웨어와 소프트웨어의 파편화를 막고 있다.9

6.1 MXFP8의 작동 원리 및 구조

MXFP8은 텐서를 32개 원소 크기의 미세 블록(Micro-block)으로 나눈다. 이 블록 크기(32)는 GPU 워프(Warp) 크기와 메모리 트랜잭션 단위(32바이트 또는 128바이트)를 고려한 하드웨어 친화적인 설계다.

  • 공유 지수(Shared Exponent): 각 블록은 하나의 공통된 스케일링 인자를 가진다. 이 스케일링 인자는 E8M0 포맷으로 정의된다. 즉, 가수 없이 8비트 지수만으로 이루어진 값이며, 이는 2^{E - 127} 형태의 2의 거듭제곱(Power of 2) 값만을 표현한다.14
  • 하드웨어 효율성: 스케일링 인자를 2의 거듭제곱으로 제한한 것은 신의 한 수다. 이렇게 하면 양자화/디퀀타이제이션 과정에서 복잡한 곱셈 연산 대신 단순한 비트 시프트(Bit Shift) 연산만을 사용하면 되므로 하드웨어 구현 비용이 획기적으로 낮아진다.
  • 데이터 구조: 32개의 FP8 데이터(32바이트)마다 1개의 E8M0 스케일(1바이트)이 붙는 구조다. 따라서 메모리 오버헤드는 약 3.1% (1/32)에 불과하다.

6.2 지역성(Locality)의 승리

MXFP8의 핵심 철학은 데이터의 지역성(Locality) 활용이다. 텐서 전체를 보면 값의 범위가 10^{-5}에서 10^4까지 매우 넓지만, 인접한 32개 값만 보면 그 범위가 상대적으로 매우 좁다. 따라서 각 블록에 맞는 최적의 로컬 스케일을 적용하면, 전체 텐서의 다이내믹 레인지를 희생하지 않으면서도 각 블록 내부에서는 높은 정밀도를 유지할 수 있다.5 이는 이상치가 발생하더라도 그 영향이 해당 블록 32개 원소에만 국한되므로 전체 연산을 오염시키지 않는다는 강력한 이점을 제공한다.

7. 소프트웨어 생태계: CUTLASS와 커널 프로그래밍

하드웨어가 MXFP8을 지원하더라도, 이를 활용할 수 있는 소프트웨어 커널이 없다면 무용지물이다. 2025년 현재, 엔비디아의 CUTLASS 라이브러리와 Triton 언어는 FP8 프로그래밍의 핵심 도구다.

7.1 CUTLASS 3.x와 Hopper/Blackwell 지원

CUTLASS는 CUDA C++ 템플릿 라이브러리로, 고성능 GEMM 커널 작성의 사실상 표준이다. 3.x 버전부터 도입된 CuTe 레이아웃 엔진은 복잡한 텐서 메모리 접근 패턴을 추상화하여, 개발자가 Blackwell의 tcgen05 명령어와 같은 저수준 기능을 쉽게 호출할 수 있게 해준다.29

  • TMA (Tensor Memory Accelerator): FP8의 성능을 극대화하려면 데이터를 글로벌 메모리에서 공유 메모리(Shared Memory)로, 다시 레지스터로 이동시키는 파이프라인이 멈추지 않아야 한다. CUTLASS는 Hopper와 Blackwell의 TMA를 활용하여, 비동기(Asynchronous)적으로 데이터를 복사하면서 동시에 텐서 코어가 연산을 수행하는 Overlap 전략을 자동으로 구현한다.29
  • WGMMA (Warpgroup Matrix Multiply-Accumulate): Hopper 아키텍처에서는 워프 그룹(128 스레드) 단위로 협업하여 행렬 곱을 수행하는 wgmma 명령어가 핵심이었다. CUTLASS는 이를 통해 FP8 텐서 코어의 성능을 100% 가까이 이끌어낸다.

7.2 사용자 정의 커널의 시대

DeepSeek-V3와 같은 선도적인 그룹은 표준 라이브러리에 의존하지 않고 직접 최적화된 FP8 커널을 작성하기도 한다. 특히 DualPipe와 같은 파이프라인 병렬화 전략이나, **통신과 연산의 중첩(Communication-Computation Overlap)**을 위해 FP8 데이터를 압축하여 전송하는 맞춤형 커널들이 개발되고 있다.2 이는 FP8이 단순한 데이터 타입 변경을 넘어, GPU 커널 레벨의 최적화 경쟁을 촉발시켰음을 보여준다.

8. 결론: FP4를 향한 가교

FP8은 이제 선택이 아닌 필수다. H100의 등장으로 가능성이 타진되었고, DeepSeek-V3와 같은 모델을 통해 대규모 학습에서의 안정성이 입증되었으며, Blackwell 아키텍처와 MXFP8 표준을 통해 완성 단계에 이르렀다. E4M3와 E5M2라는 두 가지 무기, 그리고 마이크로스케일링이라는 정밀 타격 기술은 수천억 파라미터 모델을 효율적으로 학습하고 추론할 수 있게 하는 ’민주화’의 열쇠가 되고 있다.

하지만 10.2절에서 살펴본 FP8의 시대는 종착역이 아니다. Blackwell 아키텍처는 이미 **FP4(4-bit Floating Point)**를 지원하고 있으며, 연구자들은 더 낮은 비트 수에서 정보를 보존하는 방법을 탐구하고 있다.3 FP8에서 축적된 블록 스케일링, 확률적 라운딩, 그리고 혼합 정밀도 학습의 노하우와 수학적 기반은 향후 도래할 FP4 시대의 밑거름이 될 것이다. 트랜스포머 싱귤래리티는 정밀도를 낮춤으로써 지능의 총량을 높이는 역설적인 방향으로 가속화되고 있다. 이제 우리는 비트(Bit)의 다이어트가 지능의 폭발을 이끄는 새로운 컴퓨팅의 새벽을 목격하고 있다.

9. 참고 자료

  1. DeepSeek-V3 Explained: Optimizing Efficiency and Scale - ADaSci, https://adasci.org/deepseek-v3-explained-optimizing-efficiency-and-scale/
  2. deepseek-ai/DeepSeek-V3 - Hugging Face, https://huggingface.co/deepseek-ai/DeepSeek-V3
  3. Blackwell vs Hopper: A Deep Dive GPU Architecture Comparison | IntuitionLabs, https://intuitionlabs.ai/articles/blackwell-vs-hopper-gpu-architecture-comparison
  4. Comparing NVIDIA’s B200 and H100: What’s the difference? - Civo.com, https://www.civo.com/blog/comparing-nvidia-b200-and-h100
  5. Faster Training Throughput in FP8 Precision with NVIDIA NeMo | NVIDIA Technical Blog, https://developer.nvidia.com/blog/faster-training-throughput-in-fp8-precision-with-nvidia-nemo/
  6. Four unique takeaways from Deepseek v3 - AWS Builder Center, https://builder.aws.com/content/2rJj1WkztSfYwVfsIibhWxeqMf1/four-unique-takeaways-from-deepseek-v3
  7. Low precision floating point types — HIP 6.4.43483 Documentation, https://rocm.docs.amd.com/projects/HIP/en/docs-6.4.1/reference/low_fp_types.html
  8. FP8 Precision in Deep Learning - Emergent Mind, https://www.emergentmind.com/topics/fp8-precision
  9. OCP Microscaling Formats (MX) Specification Version 1.0 - Open Compute Project, https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf
  10. Understanding FP8 and Mixed Precision Training | by Noel Akkidas - Medium, https://medium.com/@noelakkidas/understanding-fp8-and-mixed-precision-training-0a76c7f2c3ac
  11. FP8 Formats for Deep Learning, https://arxiv.org/pdf/2209.05433
  12. Floating-Point 8: An Introduction to Efficient, Lower-Precision AI Training - NVIDIA Developer, https://developer.nvidia.com/blog/floating-point-8-an-introduction-to-efficient-lower-precision-ai-training/
  13. Mixed Precision Training — NVIDIA NeMo Framework User Guide, https://docs.nvidia.com/nemo-framework/user-guide/24.09/nemotoolkit/features/mixed_precision.html
  14. Using FP8 and FP4 with Transformer Engine - NVIDIA Documentation, https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/examples/fp8_primer.html
  15. DeepSeek-V3 Technical Report - arXiv, https://arxiv.org/html/2412.19437v1
  16. H100 GPU - NVIDIA, https://www.nvidia.com/en-us/data-center/h100/
  17. Using FP8 with Transformer Engine - NVIDIA Documentation, https://docs.nvidia.com/deeplearning/transformer-engine-releases/release-1.2.1/user-guide/examples/fp8_primer.html
  18. Efficient Matrix Multiplication via Low-Rank Approximation with FP8 Acceleration - arXiv, https://arxiv.org/html/2511.18674v1
  19. Fused FP8 4-Way Dot Product With Scaling and FP32 Accumulation - IEEE Xplore, https://ieeexplore.ieee.org/document/10579354/
  20. Comparing Blackwell vs Hopper | B200 & B100 vs H200 & H100 | Exxact Blog, https://www.exxactcorp.com/blog/hpc/comparing-nvidia-tensor-core-gpus
  21. Per-Tensor and Per-Block Scaling Strategies for Effective FP8 Training - NVIDIA Developer, https://developer.nvidia.com/blog/per-tensor-and-per-block-scaling-strategies-for-effective-fp8-training/
  22. 1.5x faster MoE training with custom MXFP8 kernels - Cursor, https://cursor.com/blog/kernels
  23. DeepSeek Technical Analysis — (5) FP8 Training | by Jinpeng Zhang - Medium, https://dataturbo.medium.com/deepseek-technical-analysis-5-fp8-training-ff34768727b8
  24. Quantized KV Cache - vLLM, https://docs.vllm.ai/en/latest/features/quantization/quantized_kvcache/
  25. vLLM brings FP8 inference to the open source community - Red Hat Developer, https://developers.redhat.com/articles/2024/07/15/vllm-brings-fp8-inference-open-source-community
  26. FP8 versus INT8 for efficient deep learning inference, https://arxiv.org/pdf/2303.17951
  27. LMSYS finds minimal differences between bf16 and fp8 Llama-3.1-405b in Chatbot Arena : r/LocalLLaMA - Reddit, https://www.reddit.com/r/LocalLLaMA/comments/1fil2an/lmsys_finds_minimal_differences_between_bf16_and/
  28. INT v.s. FP: A Comprehensive Study of Fine-Grained Low-bit Quantization Formats - arXiv, https://arxiv.org/html/2510.25602
  29. CUTLASS 3.8.0 · NVIDIA cutlass · Discussion #2125 - GitHub, https://github.com/NVIDIA/cutlass/discussions/2125
  30. Overview — NVIDIA CUTLASS Documentation, https://docs.nvidia.com/cutlass/latest/overview.html
  31. DeepSeek v3 and R1 Model Architecture: Why it’s powerful and economical - Fireworks AI, https://fireworks.ai/blog/deepseek-model-architecture