《GPU Mode》
L102
2025 · FP4 training v2
High priority
transcript · failed
Quartet v2
L069 Quartet v1 의 후속. “Native FP4 training can be optimal for LLMs” 의 main thesis 가 — Blackwell 의 native FP4 instruction (B200 / RTX 5090) 위에서 어떻게 production-ready 한 training stack 으로 이어지는가. QuTLASS kernel library, FP-Quant integration (Hugging Face / vLLM), MXFP4 vs NVFP4 의 trade-off, 그리고 v1 → v2 의 알고리즘적 개선들. 자막은 실패했지만 paper (arXiv 2505.14669) 와 official repo (IST-DASLab/Quartet, IST-DASLab/qutlass) 의 공개 자료로 재구성.
FP4 training
MXFP4
NVFP4
QuTLASS
Blackwell
B200
microscaling
Hadamard rotation
QAT
A
Speakers
Andrei Panferov · Erik Schultheis
IST-DASLab · Quartet / QuTLASS authors
§ 01강의가 풀려는 문제· FP4 native training
“4 bit 로 사전학습이 가능한가, 그리고 그게 정말 compute-optimal 인가”
2024 년까지의 합의 — “weight 만 4-bit 으로 (post-training quantization) 는 가능, 학습 자체를 4-bit 으로는 불안정”. Blackwell 이 native FP4 곱셈 명령을 가져오면서 이 합의가 흔들렸다. 학습 자체를 FP4 로 (MXFP4 / NVFP4) 하면서도 BF16 baseline 에 비교 가능한 정확도를 낼 수 있다 — Quartet 의 main thesis.
강의가 던지는 세 개의 질문.
- FP4 training 이 compute-optimal 인 영역이 정말로 존재하는가 — Chinchilla 류 scaling law 의 4-bit 확장.
- 학습이 무너지지 않게 하려면 알고리즘에 무엇을 추가해야 하는가 — Hadamard rotation, stochastic rounding, microscaling.
- Blackwell 위에서 BF16 대비 몇 배 빠른가 — 이론적 4× 가 실제 kernel 에서 어떻게 실현되는가.
v1 의 결과 + v2 의 motivation
v1 (NeurIPS 2025 accepted) 은 “30M Llama-style 모델 위 3B token 의 small-scale 검증” 까지. v2 는 — 1B / 7B / 8B 의 본격 규모, MXFP4:MXFP8 QAT 의 실제 production 통합, B200 위에서의 4× speedup 실증 — 으로 확장.
“FP4 가 처음에는 ‘작은 모델만’ 의 토이로 여겨졌다. v2 는 production 규모에서 break-even 을 보여주는 게 목표.”학습 노트 · 확인 필요
§ 02v1 → v2 의 변경점· algorithmic deltas
v1 이 풀지 못한 자리들이 v2 에서 어떻게 닫혔는가
v1 의 한계는 셋이었다 — (1) 검증 규모 작음, (2) 학습 후반의 정확도 drift, (3) Blackwell 외 hardware fallback 부재. v2 의 변경점이 그 셋을 직접 겨냥.
larger-scale validation
v1: 30M × 3B tokens. v2: 1B / 7B / 8B 모델 × 수백 B tokens. Llama-3.1-8B 까지의 본격 사전학습.
MXFP4:MXFP8 QAT
forward 는 MXFP4, weight 는 MXFP8 로 유지하는 mixed-precision. accuracy 회복의 새 sweet spot.
Hadamard fused with quant
v2 에서 Hadamard rotation 을 quantization kernel 에 fuse. v1 의 별도 kernel call 이 사라지고 launch overhead 감소.
Block-scale reordering
microscaling 의 scale tensor 의 layout 을 GEMM kernel 친화적으로 재배치. CUTLASS 기반.
FlashInfer backend
B200 위에서 FlashInfer 의 attention path 와 결합. inference 측 통합.
backward kernel prototype
v1 은 forward 위주. v2 는 backward 의 prototype kernel 까지 — true end-to-end FP4 training 의 시작.
vLLM / Transformers 통합
FP-Quant 이라는 wrapper 를 통해 Hugging Face Transformers 와 vLLM 에서 직접 사용 가능. production path 확보.
Abs-Max scheme
Quartet scheme 외에 Abs-Max scheme 도 추가 — 다른 quantization 정책들과의 비교.
§ 03학습 안정성 개선· Hadamard · stochastic rounding
왜 FP4 학습이 “자연스럽게” 무너지는가, 그리고 무엇을 추가해야 막히는가
naïve 한 FP4 학습은 거의 항상 NaN/Inf 또는 plateau 로 끝난다. 원인은 셋 — (1) outlier activation 이 4-bit 의 좁은 dynamic range 를 깬다, (2) round-to-nearest 가 누적 bias 를 만든다, (3) gradient 의 magnitude 가 forward 와 다르다. 각각의 답.
Hadamard rotation — outlier 분산
activation 의 분포는 보통 heavy-tailed — 몇 개의 큰 값이 평균을 끌고 가는 모양. 이걸 그대로 4-bit 으로 quantize 하면 small majority 가 0 에 collapse. 답은 random orthogonal rotation: x' = H · x, H 가 Hadamard matrix. rotation 후의 분포는 거의 Gaussian 이 되어 outlier 가 분산됨.
Hadamard 의 장점 — recursive 하게 빠르게 (FFT 처럼 O(n log n)) 계산 가능. 그리고 rotation 의 inverse 가 자기 자신 (sign flip 만). decoding 비용이 거의 없음.
v2 에서는 이 Hadamard 를 quantization kernel 에 fuse. memory bandwidth 의 추가 cost 가 사라진다.
Stochastic rounding — 누적 bias 제거
round-to-nearest 는 정확하지만 — 값이 항상 같은 방향으로 (위로 or 아래로) 이동. 학습 step 이 누적되면 systematic drift. stochastic rounding 은 두 인접 값으로 확률적으로 round, 기댓값이 정확한 값.
비용 — random number 가 필요. v2 의 kernel 은 — counter-based PRNG (Philox) 로 register 안에서 처리. memory 추가 없음.
Microscaling — block 별 scale
한 tensor 전체에 scalar scale 하나는 부족. 작은 block (보통 32 element) 마다 자기 scale 을 유지 (block-wise quantization). MX (microscaling) 또는 NV (NVIDIA) format 모두 같은 motif.
FIG · MXFP4 와 NVFP4 의 bit layout32-element block per scale
MXFP4
e2m1 + E8 scale
4-bit elem + E8 shared
NVFP4
e2m1 + E4M3 scale
4-bit elem + FP8 shared
FP8 (e4m3)
tensor-wide scale
8-bit elem
MXFP4 vs NVFP4 의 차이 — 같은 4-bit element (e2m1) 위에서 scale 의 형태만 다르다. MXFP4 는 power-of-two scale (E8M0), NVFP4 는 finer scale (E4M3 FP8). 정확도와 hardware 친화도가 trade-off — MXFP4 는 multiply 가 단순한 shift, NVFP4 는 더 정밀하지만 scale 곱셈이 필요.
두 format 의 정착
2024–2025 시점의 state — MXFP4 가 OCP (Open Compute Project) 표준, NVFP4 가 NVIDIA proprietary. AMD 는 MXFP4, NVIDIA Blackwell 은 NVFP4 native 지만 MXFP4 도 지원. Quartet v2 는 MXFP4:MXFP8 QAT 를 default 로 — open standard 친화.
§ 04정확도 vs 속도· scaling laws
“같은 compute budget” 에서 FP4 가 진짜 BF16 을 이기는가
paper 의 가장 큰 contribution — scaling law for low-precision. compute (FLOPs) 가 동일할 때 — FP4 가 BF16 보다 더 많은 token 을 학습할 수 있고, 이로 인해 같은 compute 에서 lower loss 에 도달한다는 주장.
scaling law 의 frame
표준 Chinchilla — model size N · token D · loss L. L = f(N, D). 이 식이 dtype 별로 다르다는 게 paper 의 contribution.
같은 compute C = 6 · N · D 에서 — FP4 는 BF16 의 4× throughput → 4× 더 많은 D 를 처리. 단 — quantization 의 정확도 손실이 D 의 이득을 상쇄. 두 효과의 net 이 양인가 음인가.
v2 의 결과 — 중대형 (1B–8B) 영역에서는 net 양. 즉 FP4 가 compute-optimal. 작은 모델 (100M 미만) 에서는 양자화 손실이 더 커서 BF16 이 우위.
실측의 직관
BF16 baseline
throughput 1×
1×
MXFP4:FP8 QAT
v2 default
~2.6×
Native MXFP4 fwd
B200
~3.6×
Theoretical FP4 max
B200 spec
4×
위 수치는 paper · qutlass README · NVIDIA Blackwell 광고치를 참고한 일반 추정. 실제 실험 setup 마다 다름. 확인 필요.
“정확도 갭” 의 모양도 paper 에서 직접 quantify. 1B 모델에서 — MXFP4:MXFP8 QAT vs BF16 의 perplexity gap 이 일관되게 작고, 8B 에서는 거의 noise 수준이라는 보고. (정확한 수치는 paper 본문 참조 — § 12 의 확인 항목)
§ 05hardware 가속· QuTLASS GEMM
4× 가 “이론” 이 아닌 “실제 kernel 에서” 도달하는 길
FP4 의 4× speedup 은 — Blackwell 의 native instruction 을 “어떻게 부르는가” 에 결정된다. 잘못 부르면 BF16 보다 느려지기도. QuTLASS 는 그 호출 path 를 정형화한 CUTLASS-기반 라이브러리.
LAYER 1
CUTLASS 3.x backbone
collective::CollectiveMma 의 mainloop 위에 — FP4 specific 한 cp.async.bulk + WGMMA 변형.
L101 의 hierarchy 위에 직접 build.
LAYER 2
block-scale reordering
microscaling 의 scale tensor 를 GEMM kernel 의 access pattern 에 맞게 미리 재배치. quantization kernel 의 마지막 step 에서 fused.
LAYER 3
fused Hadamard + quant kernel
activation 을 받자마자 — Hadamard rotation 후 즉시 quantize 후 SMEM/GMEM 으로. 별도 kernel launch 없음.
LAYER 4
Blackwell native FP4 mma
sm_100a / sm_120a 의 새 wgmma 변형. FP4 element 를 받아 FP32 accumulator 에 누적. 4× tensor core throughput.
LAYER 5
FlashInfer 통합 (inference)
attention path 의 PagedAttention / FlashAttention 변형이 FP4 weight 를 직접 받음. 추론 측 가속.
LAYER 6
FP-Quant wrapper
Hugging Face Transformers 의 nn.Linear 를 FP4 layer 로 swap. user-facing 한 줄 변경.
왜 CUTLASS 위인가
FP4 의 mainloop 는 — A 와 B 의 dtype, accumulator dtype, scale 의 layout 모두 다르다. CUTLASS 의 collective 추상이 — 한 type alias 변경으로 새 mainloop instantiation — 을 가능하게 함. 이게 QuTLASS 가 빠르게 진화하는 이유.
QuTLASS 의 README 는 — Llama-3.1-8B 의 inference 에서 BF16 대비 최대 4× 추론 속도, MXFP4:MXFP8 QAT 학습에서 BF16 pseudo-quant 대비 30% 빠름 — 을 보고. 4× inference 와 1.3× training 의 격차는 — backward pass 가 아직 prototype 이라는 사실의 직접적 결과.
§ 06실측· Llama 1B / 7B / 8B
v2 가 실제로 보여준 숫자
paper / repo 에서 보고된 측정의 큰 그림. 정확한 수치는 paper 본문에서 확인 필요.
Llama-style 30M × 3B (v1 baseline)
v1 의 reference. MXFP4 pretraining 이 BF16 와 거의 동등한 perplexity. 작은 규모에서의 sanity check.
Llama-3.1-1B
v2 의 main scale-up. perplexity gap 이 1% 미만. compute-optimal 이라고 주장 가능한 첫 점.
Llama-3.1-8B QAT
MXFP4:MXFP8 QAT. BF16 pseudo-quant 대비 30% 빠른 학습 + 정확도 회복. production 에 가장 가까운 실증.
B200 inference
FP4 GEMM 이 BF16 대비 최대 4×. attention 포함 end-to-end token throughput 도 큰 폭 상승.
RTX 5090 (sm_120a)
consumer 환경에서의 첫 native FP4 시연. 작은 모델에서 4× 에 근접.
downstream eval
MMLU / HellaSwag / ARC 등 다중 벤치마크 — FP4 학습 모델이 BF16 baseline 의 ±1% 안에 있다는 보고.
“이 paper 의 가장 큰 메시지는 단일 모델이 아니라 — ‘FP4 가 작동하는 영역’ 의 윤곽 을 그렸다는 점. 그 윤곽이 1B–8B 라는 production-friendly zone 을 포함한다.”학습 노트
§ 07다른 4-bit 방법과 비교· GPTQ · AWQ · NF4 · FP8
같은 “4-bit” 이라도 무엇이 4-bit 인가가 다르다
PTQ
GPTQ / AWQ
학습 후 weight 만 4-bit 으로 quantize. activation 은 FP16. inference-only. 가장 흔한 deployment path. 학습 가속 없음.
PTQ
NF4 (QLoRA)
정규분포 친화 4-bit format. fine-tuning 시 weight 를 NF4 로 freeze + LoRA. 학습은 BF16/FP16.
QAT
MXFP4:MXFP8 (Quartet v2)
forward 는 MXFP4, weight 는 MXFP8. 학습 자체가 4-bit aware. native instruction 활용.
native
Pure MXFP4 forward (Quartet v1/v2)
forward 와 backward 모두 MXFP4. v1 의 scope. v2 는 backward prototype.
FP8
Transformer Engine FP8
e4m3 / e5m2 의 8-bit. Hopper 에서 native. 두 배 빠름. FP4 는 그 2× 더.
INT4
SmoothQuant / OmniQuant 등
integer 기반 4-bit. INT4 weight + INT8 activation. dynamic range 좁음. PTQ 위주.
핵심 분기 — PTQ vs QAT vs native. PTQ 는 학습 가속이 없고 inference 만. QAT 는 학습 awareness 가 있지만 BF16 simulation. native (Quartet v2 의 자리) 는 학습 자체가 native FP4 instruction. “학습 가속 + 정확도 보장” 은 native 만 약속할 수 있음.
§ 08production 적용· FP-Quant · vLLM 통합
“paper 에서 production 으로” 의 path
algorithm 이 production 에 닿으려면 — kernel + framework + serving stack 모두에서 native 지원이 필요. v2 가 만들어낸 path 를 정리.
Quartet repo
research 코드. main_setup.sh 로 30M Llama-style pretraining 재현. notebook 으로 benchmark.
QuTLASS
production-grade FP4 kernel library. CUTLASS 위. sm_100a/sm_120a target. CUDA 12.8+.
FP-Quant
사용자 친화 wrapper. Hugging Face Transformers 의 model 에 FP4 mode 를 한 줄 옵션으로.
vLLM 통합
FP-Quant 가 만든 model 을 vLLM 이 serve. inference path 가 native FP4.
FlashInfer backend
attention 의 PagedAttention / 변형들이 FP4 weight 와 직접 호환. inference 의 attention 측.
B200 cloud 가용성
2025 년 mid 부터 B200 cloud 가용. 5090 desktop 도 native FP4. consumer-friendly.
production checklist
FP4 를 실 배포에 끼우려면 — (1) GPU 가 sm_100a 이상 (B200) 또는 sm_120a (5090). (2) CUDA 12.8+. (3) FP-Quant 로 model 변환. (4) vLLM 의 FP4 path 활성화. (5) downstream eval 로 정확도 검증. — 다섯 단계가 모두 통과되면 native 4× inference + 정확도 유지.
§ 09미래· FP4 sparse · MoE · attention native
v3 가 풀어야 할 자리
- full backward FP4 — v2 의 prototype 을 production 으로. 진정한 “native FP4 training”.
- FP4 + 2:4 sparsity — Blackwell 의 structured sparsity 와 결합. 추가 2× speedup 가능.
- MoE 의 expert quant — 큰 expert 를 FP4 로, 작은 router 는 BF16. 자연스러운 fit.
- attention 자체의 FP4 — QK · softmax · PV 의 모든 단계를 FP4. 현재는 weight 만 FP4, attention compute 는 FP8/BF16.
- FP6 / FP8 mixing — 어느 layer 는 FP4, 어느 layer 는 FP6/FP8. layer-wise sensitivity 의 자동 결정.
- federated FP4 — L099 와 결합. consumer GPU 의 FP4 학습.
“4-bit training 이 production 에 들어오면 — 같은 cluster 가 4× 모델 capacity 를 처리. compute 시장 자체의 dynamics 가 바뀐다.”학습 노트 · 확인 필요
§ 10기억할 메모· key takeaways · refs
다시 열었을 때 5분 안에 손에 잡혀야 할 것
main thesis
native FP4 training (B200 / 5090) 이 1B–8B 영역에서 compute-optimal. v2 가 그 윤곽을 production 규모에서 실증.
v2 의 4 가지 변경
(1) larger-scale 검증, (2) MXFP4:MXFP8 QAT, (3) Hadamard fused with quant, (4) backward kernel prototype.
학습 안정성 trio
Hadamard rotation (outlier) + stochastic rounding (bias) + microscaling (block-wise). 이 셋이 빠지면 FP4 학습은 거의 항상 무너짐.
MXFP4 vs NVFP4
같은 4-bit element, 다른 scale format. MXFP4 가 OCP open standard, NVFP4 가 NVIDIA proprietary. v2 default 는 MXFP4.
QuTLASS
CUTLASS 기반 FP4 GEMM library. block-scale reordering + fused Hadamard + native WGMMA. 4× inference 의 key.
production path
QuTLASS → FP-Quant → Hugging Face / vLLM. 한 줄로 배포 가능한 단계까지 v2 가 끌어옴.
PTQ vs QAT vs native
GPTQ/AWQ (PTQ) → 학습 가속 없음. QLoRA (QAT) → 일부. Quartet (native) → 학습 자체 가속.
남은 대상
attention native FP4, full backward, structured sparsity 결합, MoE 의 expert quant. v3 의 후보.
Quartet v1
L069 · 같은 라인의 prequel
손에 새기기 — 실습 시퀀스
- Quartet v1 reference 재현 —
main_setup.sh 로 30M Llama-style 학습. perplexity 가 BF16 와 비교 가능한지 직접 확인.
- Hadamard 의 효과 측정 — Hadamard 없이 학습 → diverge 또는 plateau 확인. 그 다음 켜고 비교.
- Stochastic rounding ablation — round-to-nearest 와 stochastic rounding 의 학습 곡선 비교. systematic drift 직접 본다.
- QuTLASS GEMM benchmark — Llama-3.1-8B size 의 GEMM 한 개를 BF16 vs MXFP4 로 측정. B200 / 5090 환경에서 4× 가 정말 나오는지 확인.
- FP-Quant 변환 — Hugging Face Llama-3.1-8B 를 FP-Quant 로 변환. perplexity 비교.
- vLLM serve — 변환된 모델을 vLLM 으로 serve. token throughput 측정.
- MMLU 평가 — 변환 전/후 모델의 MMLU 점수 비교. “정확도 손실” 의 실측.
§ 11다른 강의로 이어지는 길· connections
L102 가 다른 강의들과 어떻게 엮이는가
§ 12열린 질문· open questions
다음에 다시 들었을 때 직접 검증해야 할 것들
- 정확한 perplexity gap — 1B / 8B 에서의 BF16 baseline 대비 % gap. paper 본문에서 확인. (확인 필요)
- full backward kernel 의 상태 — “prototype” 이라는 표현 — 어디까지 통합 / 어디부터 forward-only인지. (확인 필요)
- v3 일정 — 다음 release 의 timeline 또는 우선순위. (확인 필요)
- MoE 통합 — DeepSeek-V3 / Mixtral 의 expert 를 FP4 로 — 실험 진행 중인지. (확인 필요)
- structured sparsity 와의 결합 실증 — 2:4 + FP4 의 net throughput. (확인 필요)
- RTX 5090 (sm_120a) 의 실측 — desktop 환경에서의 정확한 수치. (확인 필요)
- FP-Quant 의 license / 공개 범위 — Hugging Face 통합의 path. (확인 필요)
검증 메모
이 노트의 알고리즘 설명은 arXiv 2505.14669 (Quartet paper) 와 IST-DASLab/qutlass · IST-DASLab/Quartet 의 README로부터 재구성. 강의 자체의 자막은 실패. 실측 수치는 paper · README 의 일반 보고를 인용했으며 강의에서 직접 인용된 정확한 숫자는 영상 재시청으로 확인 필요.