《GPU Mode》
L032
2024 · OCT · 19
High priority
transcript · available
Unsloth — Faster, Lower-Memory LLM Fine-Tuning
Fine-tuning 은 사전학습보다 작지만 — 사용자/예산의 자리에서는 여전히 무겁다. Daniel Han 이 동생과 같이 만든 Unsloth 는 QLoRA + 손으로 짠 Triton 커널 + 정확한 gradient 식 + 메모리 패턴 최적화 로 같은 GPU 에서 fine-tuning 을 2 배 빠르게, 메모리는 절반으로 만든 라이브러리. 그 안에서 매 layer 마다 “정확한 미적분 + 작은 Triton 커널” 의 디자인 결정이 어떻게 나왔는가의 학습 노트.
Unsloth
QLoRA
fine-tuning
Triton kernels
gradient calculus
cross-entropy
RoPE
RMSNorm
Daniel Han
D
Speaker
Daniel Han
Unsloth 공동창업 · 전 NVIDIA · 수학 출신
§ 01강의가 풀려는 문제· fine-tuning 비용
“7B 모델을 노트북 한 대에서 fine-tune 하고 싶다” — 실제로 가능한가의 답
Unsloth 의 출발점은 명확하다. fine-tuning 이 학생/소규모 팀의 손에 있어야 한다. 그런데 vanilla 길을 가면 — 7B 모델 fine-tune 에 H100 두 대가 필요하다. Unsloth 가 이 자리를 한 RTX 4090 위에서 끝낸다.
강의가 답하려는 두 줄 —
- fine-tuning 의 메모리/속도가 어디서 새는가 — vanilla Hugging Face Trainer 로 LoRA 학습 step 을 까보면 무엇이 dominant 인지.
- 그 자리를 어떻게 더 압축할 수 있는가 — 매 layer 마다 미적분으로 backward 를 다시 유도하고, Triton 커널을 손으로 짜고, 메모리 access 패턴을 다시 그린다.
Daniel 의 입장은 명시적으로 “정확한 수학으로부터 시작한다”. autograd 가 만드는 backward 가 충분히 정확하지만 충분히 빠르지는 않다. 그래서 — RMSNorm, RoPE, LayerNorm, cross_entropy 등 — 모든 layer 의 gradient 를 손으로 다시 유도하고, 한 Triton 커널로 다시 짠다.
강의의 frame
“Unsloth 는 단순히 ‘QLoRA 의 wrapping’ 이 아닙니다. 모든 layer 의 backward 를 다시 유도하고, 매 자리의 Triton 커널을 손으로 짭니다 — 그 결과 같은 결과를 더 적은 메모리, 더 빠른 시간 안에 냅니다.” QLoRA 의 hierarchy 위에 또 한 층의 손에 짠 최적화가 얹혀 있는 게 본질.
“우리는 모든 layer 의 gradient 를 종이에 직접 유도했습니다. 매 derivative 를 손으로 풀고 — 그 식이 PyTorch 의 autograd 가 만드는 것보다 어디서 더 효율적인지 봅니다. 매 자리에 작은 win 이 있어요.”Daniel Han · 11:24
§ 02fine-tuning 의 비용 구조· 왜 무거운가
“full fine-tune 의 메모리는 weight 의 4 배” — Adam 의 자리
vanilla full fine-tune 은 가장 비싼 자리. weight 의 4–6 배 메모리. Daniel 이 강의에서 짚은 분포 —
FIG · full fine-tune 의 메모리 — Llama-2 7BBF16 + Adam · seq 2048
Llama-2 7B full fine-tune 이 H100 80GB 한 대에 안 들어가는 이유. Adam state 가 weight 의 4 배. 단일 GPU 로는 불가능.
해결의 사다리 — 위로 갈수록 메모리가 줄어든다.
- Adafactor / 8-bit Adam — opt state 를 줄임. 메모리 -25% 정도.
- LoRA — weight 의 일부만 학습. trainable param 이 0.1% 로. opt state 가 거의 0.
- QLoRA — 추가로 base weight 를 4-bit 로 저장. weight 메모리 -75%.
- Unsloth = QLoRA + 위에 한 층 — Triton 커널 + gradient 재유도 + 메모리 패턴.
이 사다리를 손에 박아두면 Unsloth 의 위치가 명확해진다. Unsloth 는 새 알고리즘이 아니다 — QLoRA 의 구현을 손으로 다시 짠 결과물. 그 “손으로” 가 강의의 본론.
vanilla full FTweight + grad + Adam fp32 + activation~100 GB
+ 8-bit Adamopt state INT8~75 GB
+ LoRAtrainable = adapter only~30 GB
+ QLoRAbase weight INT4~14 GB
+ Unsloth+ Triton 커널 + 메모리 패턴~7 GB
§ 03LoRA + QLoRA· low-rank adapter
fine-tuning 의 “학습 가능한 부분” 을 0.1% 로 — 메모리 폭발의 근본 해결
Unsloth 의 base 가 되는 두 알고리즘. 강의에서 Daniel 이 빠르게 정리한 부분.
LoRA
weight 변화없음
adapter 추가A·B 두 행렬
rankr=8 / 16 / 32
trainable param~0.1% of full
opt state~0.1% of full
식W' = W + α·(A·B)
QLoRA
base weight4-bit (NF4)
adapterBF16 (LoRA 그대로)
forward 시4-bit dequant + matmul
backward 시adapter 만 학습
weight 메모리~25% of LoRA
식y = dequant(W4)·x + α·A·B·x
Unsloth
baseQLoRA 위
+ Triton 커널RoPE, RMSNorm, swiglu, CE
+ gradient 재유도autograd 우회
+ memory layoutlogit upcast inside CE
속도 win~2× (vs vanilla QLoRA)
메모리 win~50% (vs vanilla QLoRA)
QLoRA 의 작은 디테일 — NF4
QLoRA 의 4-bit 포맷은 단순 INT4 가 아닌 NF4 (NormalFloat 4-bit). weight 의 분포가 정규분포라는 가정으로 — quantile 기반의 비대칭 codebook. 같은 4 비트로 정확도가 더 높다. NF4 의 dequant 가 Unsloth 의 hot path 중 하나.
§ 04Unsloth 의 추가 최적화· Triton 커널 자리
“같은 QLoRA 인데 왜 2× 빠른가” — Triton 커널이 뽑는 자리들
Unsloth 의 본론. QLoRA 는 알고리즘 자체로 모든 win 을 가져왔다. Unsloth 는 그 위에 — 매 layer 의 forward/backward 를 손으로 짠 Triton 커널로 다시. RMSNorm, RoPE, cross_entropy, swiglu — Liger Kernel 과 겹치는 자리지만 fine-tuning 환경에 특화.
RMSNorm forward + backward
PyTorch 의 RMSNorm 이 4–7 개 elementwise 커널로 분해. Unsloth 가 한 Triton 커널로 fuse. backward 식을 손으로 유도해서 forward 의 RSTD 를 직접 재사용.
RoPE — 직접 유도한 backward
RoPE 의 backward 가 사실 — input 의 절반-회전 패턴을 그대로. 별도 saved tensor 없음. 한 Triton 커널 안에 forward + backward 가 동거.
cross_entropy — logit upcast 자리
PyTorch 는 LM head logit 을 fp32 로 upcast 후 softmax. Unsloth 는 upcast 를 cross_entropy 커널 안에서. 메모리 절감의 큰 자리.
swiglu MLP
silu(gate) * up. backward 가 명시적인 closed form 식으로 정리 가능. 한 커널 안에서 forward, backward 모두.
QLoRA dequant + matmul fusion
NF4 의 dequant 와 matmul 이 보통 separate. Unsloth 는 dequant 를 — load 직후 inside-kernel 에서. dequant 결과 텐서를 만들지 않는다.
attention — Flash Attention 그대로
attention 자체는 Tri Dao 의 Flash Attention 호출. Unsloth 가 다시 짜지 않음 — fa 가 충분히 잘 짜여 있다는 인정.
“우리는 PyTorch 의 cross_entropy 가 logit 을 fp32 로 upcast 하는 걸 봤어요. 그런데 그 upcast 가 backward 까지 살아있어야 합니다 — 그게 메모리 spike 입니다. upcast 를 kernel 안에서 하면 그 텐서를 안 만듭니다.”Daniel Han · 32:18
§ 05손으로 유도한 gradient· 매 op 의 backward
“autograd 가 만든 backward 가 — 정확하지만 — 가장 효율적인 코드는 아니다”
Unsloth 의 디자인 결정 중 가장 흥미로운 자리. autograd 가 자동으로 만드는 backward 는 — chain rule 을 그대로 쌓는다. 그게 정확하지만, “이 layer 의 gradient 가 사실 더 짧은 closed-form 으로 표현된다” 면 그게 더 빠르다.
한 사례 — RMSNorm 의 backward. autograd 의 길은 sum, rsqrt, div, mul 의 chain rule. 여러 elementwise 커널 + 임시 텐서.
Daniel 이 손으로 유도한 식 —
y = w · x · rstd (forward)
backward — closed form:
dx = rstd · (dy · w
- x · rstd² · mean(x · dy · w))
dw = sum(dy · x · rstd, axis=batch)
이 식의 핵심 — rstd 가 한 번 계산되면 양쪽 모두 재사용. 한 Triton 커널 안에서 dx, dw 를 동시 산출. 임시 텐서 0.
# Unsloth 의 RMSNorm backward — Triton (간략)
@triton.jit
def rms_norm_backward(X, W, RSTD,
DY, DX, DW_partial,
N: tl.constexpr,
BLOCK: tl.constexpr):
pid = tl.program_id(0)
cols = tl.arange(0, BLOCK)
mask = cols < N
# forward 의 RSTD 재사용
rstd = tl.load(RSTD + pid)
x = tl.load(X + pid*N + cols, mask).to(tl.float32)
w = tl.load(W + cols, mask).to(tl.float32)
dy = tl.load(DY + pid*N + cols, mask).to(tl.float32)
# closed form — 한 줄에서
inner = tl.sum(x * dy * w, axis=0) / N
dx = rstd * (dy * w - x * rstd*rstd * inner)
dw = dy * x * rstd
tl.store(DX + pid*N + cols, dx.to(tl.bfloat16), mask)
tl.store(DW_partial + pid*N + cols, dw, mask)
왜 “손으로 유도” 가 중요한가
autograd 의 backward 는 한 op 단위로 작성된 rule 의 chain. 여러 op 사이의 공유 산출물(예: RSTD, softmax)을 보지 못한다. 손으로 식을 정리하면 — 그 공유를 직접 본다. 그게 메모리 절감 + 속도 향상의 source.
Daniel 이 강의 중 — “Wolfram 으로 매 layer 의 gradient 를 풀고, 손으로 식을 정리하고, 그 다음 Triton 으로 짭니다. 우리 코드의 한 줄마다 이런 작업이 있어요”. loud claim 이지만, GitHub 의 unsloth/kernels 디렉터리를 보면 사실이라는 평가.
§ 06메모리 패턴 최적화· cross-entropy 의 logit cast
“같은 식인데 어디서 cast 하느냐” 가 메모리를 가른다
강의에서 가장 인상적인 디테일 중 하나. PyTorch 의 LM head 는 — logits = hidden @ W_lm.T(BF16) → logits = logits.float()(fp32 upcast) → cross_entropy(logits, labels). 이 두 번째 단계의 fp32 upcast 가 — vocab × seq × batch × 4 byte 의 큰 텐서를 만든다.
FIG · cross_entropy 의 메모리 분포Llama-2 7B · vocab 32K · seq 2048
PyTorch 표준:
hidden BF16: 2048×4096×2 = 16 MB
logits BF16: 2048×32K×2 = 128 MB
logits fp32: 2048×32K×4 = 256 MB ← upcast
total spike: 256 MB · backward 까지 살아있음
Unsloth fused:
hidden BF16: 16 MB
logits BF16: 128 MB
fp32 upcast: inside kernel · register only
total spike: 128 MB (-50%)
vocab 이 큰 모델(Qwen2 152K, Aya 256K)에서는 차이가 더 극적. vocab 큰 모델 fine-tuning 의 enabler.
이 trick 의 일반화 —
- 큰 텐서의 cast 는 가능하면 다음 op 안으로 미룬다. 별도 텐서로 만들지 않는다.
- upcast 가 정확도에 필요한 자리는 명시적으로 분리. softmax 의 exp 안 등.
- partial upcast 의 함정 — Triton 컴파일러가 자동 promote 한다는 점을 알고 있어야. 강의 Q&A 의 핵심 질문.
Daniel 의 Q&A — “두 input 중 하나만 upcast 하면 다른 쪽도 자동 promote 된다는 걸 처음에 모르고 — 한 자리에서 silent fp16 mul 이 일어나서 학습이 한 번 발산했어요. 그 다음부터는 양쪽 모두 명시적으로 cast”.
§ 07정확도 vs 속도· 발산하지 않는 시퀀스
“같은 모델을 더 빠르게 — 결과는 그대로” 가 거짓말이 아닌 이유
QLoRA + Unsloth 의 fine-tuning 결과가 — 같은 hyperparameter 로 — vanilla QLoRA 와 1% 미만 loss diff. 그게 가능한 이유는 손으로 짠 커널의 numerics 가 정확하기 때문.
storage bf16, accum fp32
모든 reduction (RMSNorm 의 sum, softmax 의 sum) 이 fp32 안에서. cast 는 마지막에. Liger 와 같은 패턴.
stochastic rounding (옵션)
master weight cast 시. 작은 grad update 의 누적이 BF16 안에서 정확. 강의 시점에 옵션이지만 default 로 가는 추세.
numerics gate
PR 마다 baseline 과의 output diff 검증. 1e-3 이하. Liger 와 같은 표준.
학습 결과 검증
Alpaca, OpenAssistant 등 표준 fine-tuning 셋의 final loss + downstream eval. vanilla 와 visible 차이 없음.
발산의 자리
attention logit 이 매우 큰 자리(긴 seq 또는 outlier)에서 발산 가능. Unsloth 의 fa 호출 부분이 BF16 fallback 분기 보유.
RoPE 의 long context 함정
long seq 에서 sin/cos 의 정확도가 떨어짐. fp64 로 cos/sin 를 미리 계산 → fp32/bf16 로 cast 의 패턴이 중요.
“우리는 ‘공짜 속도’ 를 약속하지 않습니다. 우리가 약속하는 건 — 같은 학습 결과, 더 적은 시간. 모든 PR 은 발산 검증을 통과해야 머지됩니다.”Daniel Han · 1:01:47
§ 08실측 사례· 속도 · 메모리
“실제 fine-tuning 워크로드에서 어떤 숫자가 나오는가”
FIG · fine-tuning throughput 비교Llama-2 7B QLoRA · seq 2048
vanilla HF + bnb (T4)
1× (ref)
vanilla HF + bnb (4090)
1× (ref)
Colab T4 에서 처음으로 7B QLoRA 가 OOM 없이 도는 자리. 4090 같은 하이엔드 카드에서도 같은 비율의 win.
FIG · 메모리 사용 비교같은 batch · 같은 모델
+ Unsloth long ctx
10 GB (8K seq)
메모리 절감의 source — cross_entropy 의 logit upcast in-kernel. 큰 vocab · 긴 seq 에서 효과 폭발적.
강의에서 강조된 시나리오들 —
- Colab free tier T4 위 7B QLoRA. vanilla 는 OOM, Unsloth 는 동작. “Colab 에서 무료로 LLaMA fine-tune 가능” 의 길.
- 긴 context (8K, 16K, 32K). seq 가 길수록 cross_entropy + RoPE 의 절감 효과가 커짐.
- 큰 vocab (Qwen2, Aya). logit upcast in-kernel 의 dominant 한 자리.
- 다양한 모델 지원. Llama, Mistral, Gemma, Qwen, Phi. 모델별 patch 함수 분리.
Unsloth Pro / Open Source 분리
Unsloth 는 open source 버전(GitHub) + Pro 버전(상업) 으로 나뉨. OSS 가 단일 GPU 시나리오를 cover 하고, Pro 가 multi-GPU + 더 깊은 최적화. 강의의 모든 디자인 결정은 OSS 의 자리. Daniel 의 honest 한 주석 — “우리가 살아남으려면 Pro 도 있어야 합니다, OSS 만으로는 회사가 안 돼요”.
§ 09backend 선택· free vs Pro vs OSS
fine-tuning 의 도구 선택 — Unsloth 가 어디 서 있는가
2024 시점의 fine-tuning 도구 풍경. 각 도구의 자리.
Hugging Face Trainer + bnb
표준 vanilla. 가장 호환성 좋음. perf 는 baseline. 모델 코드 변경 거의 없음.
PEFT (HF)
LoRA / QLoRA 의 abstract layer. Trainer 와 함께. Adapter 관리는 가장 깨끗.
Unsloth
QLoRA 의 가장 빠른 single-GPU 구현. open source 단일 GPU. Triton 커널 + 손 유도 gradient.
Liger Kernel
학습 커널의 fusion. full fine-tune에 더 유효. monkey-patch 한 줄.
Axolotl
YAML config 기반 fine-tuning workflow. Unsloth 와 Liger 둘 다 옵션으로 받음.
torchtune
PyTorch 공식. 간결한 학습 루프. Liger / 일부 최적화 통합.
DeepSpeed ZeRO + LoRA
multi-GPU full fine-tune 또는 LoRA. 큰 모델/큰 batch 의 자리. Unsloth 와 직접 비교 안 됨.
Llama Factory
UI 기반 fine-tuning. Unsloth 를 backend 로 받음. 기능적 wrapping.
언제 Unsloth 를 선택하는가
(1) 단일 GPU + QLoRA. 가장 큰 win 의 자리.
(2) 제한된 메모리 (Colab T4, 8GB 노트북 GPU). vanilla 가 OOM 인 자리.
(3) 긴 context fine-tuning. cross_entropy 절감의 dominant 자리.
Unsloth 가 적합하지 않은 자리 — multi-GPU full fine-tune (DeepSpeed/FSDP 자리), 사전학습(scale 이 다름).
§ 10기억할 메모와 코드· key takeaways
다시 열었을 때 5분 안에 손에 잡혀야 할 것
fine-tuning 사다리
vanilla full → Adafactor → LoRA → QLoRA → Unsloth(QLoRA + 손 짠 커널). 위로 갈수록 메모리/속도 win.
QLoRA 핵심
base weight 4-bit (NF4), adapter BF16. forward 시 4-bit dequant + matmul. backward 는 adapter 만.
Unsloth 의 4 자리
RMSNorm + RoPE + cross_entropy + swiglu 의 Triton 커널. 매 layer 의 backward 를 손으로 유도.
cross_entropy logit cast
PyTorch 의 fp32 upcast 텐서 = vocab×seq×batch×4. Unsloth 는 inside-kernel upcast — 텐서 안 만듦.
손 유도의 의미
autograd 의 chain rule 보다 짧은 closed form 가능. 공유 산출물(RSTD, softmax 결과)을 직접 재사용.
numerics 약속
storage bf16 + accum fp32. PR 마다 output diff 1e-3 검증. 학습 loss 가 baseline 과 visible 일치.
실측 win
Llama 7B QLoRA — 속도 ~2×, 메모리 ~50%. T4 위에서 처음으로 OOM 없이.
위치
단일 GPU + QLoRA + 메모리 제한 환경의 first choice. multi-GPU full FT 는 다른 도구.
손에 새기기 — 실습 시퀀스
- vanilla QLoRA baseline — Hugging Face PEFT + bnb 로 Llama 7B QLoRA fine-tune. T4 또는 4090 위. 시간/메모리 측정.
- Unsloth 적용 —
FastLanguageModel.from_pretrained 한 줄. 같은 hyperparameter 로 같은 작업. 시간/메모리 비교.
- RMSNorm backward 손으로 유도 — 노트에 chain rule 로. closed form 정리. Unsloth 의 Triton 코드와 비교.
- cross_entropy logit upcast 직접 측정 — PyTorch profiler 로 logit 의 fp32 텐서 크기 확인. Unsloth on/off 비교.
- Triton 커널 직접 짜기 — RMSNorm forward 의 한 row 단위 커널. row-per-program, BLOCK = pow2(N). Unsloth 코드를 보지 않고 first principle 로.
- 긴 context 시나리오 — seq 8K, 16K 로 fine-tune. 메모리 곡선의 차이가 어디서 가장 크게 나는지.
- 큰 vocab 모델 — Qwen2 (152K vocab) 로 fine-tune. cross_entropy 절감의 dominant 자리.
- 학습 loss 검증 — vanilla QLoRA vs Unsloth 의 1B token 학습 loss 곡선. visible 차이 없어야 함.
§ 11다른 강의로 이어지는 길· connections
이 강의의 도구가 시리즈 안에 어떻게 다시 등장하는지
§ 12열린 질문· open questions
다음에 다시 들었을 때 직접 검증해야 할 것들
- Unsloth + Liger 의 합집합 — 같은 자리(RMSNorm/RoPE)를 푸는 두 도구가 함께 쓰일 수 있는가? monkey-patch 가 충돌하는가.
- multi-GPU Unsloth (Pro) — open source 의 single-GPU 위 패턴이 multi-GPU 에서도 같은 비율의 win 을 내는가? FSDP/ZeRO 와의 통합.
- FP8 mma 위의 fine-tuning — Hopper 의 FP8 가 Unsloth 의 다음 단계인가? QLoRA 의 NF4 와 mma 의 FP8 의 합집합.
- autograd vs 손 유도의 한계 — 모든 op 의 backward 를 손으로 유도할 수 있는가? attention 의 backward 가 가장 복잡한 자리.
- flash attention v3 통합 — Unsloth 가 fa3 를 받으면 추가 win 이 있는가? 현재는 fa2.
- 긴 context (128K) 지원 — RoPE 의 긴 seq 정확도, cross_entropy 의 메모리 절감의 한계점은 어디인가.
- 다른 모델 아키텍처 — RWKV, Mamba, MoE 같은 비-Transformer 의 fine-tuning 자리.
검증 메모
이 노트의 perf 수치(2× 속도, -50% 메모리 등)는 강의 시점(2024 10월)의 측정. Unsloth 가 빠르게 발전하고 있으니 GitHub 의 README benchmark 를 다시 확인할 것. 모델별 patch 함수의 지원 상태도 변동.