FreeToken은 개인 PC 한 대에서 프런티어급 오픈웨이트 MoE 모델을 서빙하는 추론 엔진입니다. 게이밍 데스크탑, 노트북, 워크스테이션처럼 GPU가 한 장뿐이고 그 VRAM이 모델 전체를 담지 못하는 환경이 대상입니다. 버클리·MIT 계열 연구진이 만들었고 논문(arXiv:2608.16157)과 코드가 함께 공개돼 있습니다.
왜 이게 가능한지는 MoE라는 모델 구조에서 나옵니다. 논문이 드는 예를 그대로 옮기면, DeepSeek-V4-Flash는 43개 레이어 각각에서 256개 전문가 중 6개만 골라 씁니다. 파라미터는 284B인데 토큰 하나당 실제로 계산에 참여하는 건 13B뿐입니다. 13B는 RTX 5090의 32GB에 들어갑니다. 문제는 어느 6개가 뽑힐지 미리 모른다는 것, 그래서 256개 전부를 어딘가에 갖고 있어야 한다는 것입니다.
작업대(GPU) 옆 공구함(VRAM)은 작습니다. 전체 공구(전문가 256개)는 멀리 있는 창고(호스트 RAM)에 있고, 왕복 통로(PCIe)는 하나뿐이라 여러 개를 한꺼번에 나르면 줄이 섭니다.
기존 엔진들은 "없으면 무조건 가져온다"였습니다. 그래서 통로 폭이 곧 속도 상한이 됐고, 그동안 옆에 놀고 있는 CPU는 아무 일도 안 했습니다. FreeToken은 여기서 갈래를 하나 더 냅니다 — 없는 공구 중 일부는 가져오지 않고 창고 옆 작업대(CPU)에서 그 자리에서 깎아 씁니다. 나르는 시간과 깎는 시간이 같이 끝나도록 개수를 나누는 게 이 프로젝트의 정체입니다.
284B-A13B 같은 표기가 이 둘을 뜻합니다.공개 첫 커밋이 2026년 8월 11일, 그로부터 2주도 안 돼 별 3.2k. 속도의 이유는 세 가지가 겹친 것입니다 — 실측 수치가 크고, 저자진이 화려하고(Ion Stoica · Matei Zaharia · Song Han · Kurt Keutzer), 대상이 개인 PC라 확인해 볼 사람이 많습니다.
① 프리필이 희소성을 파괴한다. 프롬프트가 길어지면 토큰마다 다른 전문가가 뽑히고, 그 합집합이 결국 전문가 전체를 덮습니다. "토큰당 6개"라는 희소성이 프롬프트 단위에서는 사라집니다.
② 디코드 미스가 전부 PCIe로 몰린다. 캐시에 없는 전문가는 매번 PCIe로 실어 나르므로, 속도가 링크 대역폭에 그대로 묶입니다. 그동안 호스트 CPU와 RAM 대역폭은 놀고 있습니다.
③ 엣지 GPU는 전용이 아니다. 게임이 켜지고 브라우저가 VRAM을 먹으면 가용량이 세션 중에도 변합니다. 기동 시점에 메모리 배분을 고정하는 설계는 여기서 깨집니다.
① q* 정책 — 디코드 스텝마다 캐시 미스를 "PCIe로 가져올 것"과 "CPU에서 계산할 것"으로 쪼갭니다. 공식은 q* = m × B_P / B_H (m=미스 수, B_P=PCIe 대역폭, B_H=호스트 메모리 대역폭). 두 경로가 동시에 끝나도록 나눈 값입니다.
② 의미 기반 캐싱 — 에이전트가 컨텍스트를 고쳐 쓰는 지점(툴콜·thinking 블록)에 KV 스냅샷을 앵커로 박아 둡니다. 대화가 편집돼도 그 지점까지는 재사용하고 새로 붙은 부분만 다시 프리필합니다.
③ 탄력적 메모리 — 전문가 캐시와 KV 캐시의 VRAM 배분을 엔진 재시작 없이 스케줄러의 안전 지점에서 재조정합니다.
흔한 오해 하나를 먼저 정리해야 합니다. FreeToken이 겨루는 상대는 vLLM이나 SGLang이 아닙니다. 그쪽은 여러 장의 데이터센터 GPU에 모델이 다 들어가는 상황을 전제합니다. 실제로 README의 감사 문단은 vLLM · SGLang · FlashInfer · llama.cpp를 코드와 설계를 빌려온 상류(上流)로 명시합니다. 논문이 실험에서 비교한 baseline은 llama.cpp · Ollama · KTransformers · MoE-Infinity — 전부 "엣지에서 오프로딩으로 버티는" 진영입니다.
| 엔진 | 전문가를 어떻게 다루나 | 한계 |
|---|---|---|
| llama.cpp / Ollama | 로드 시점에 레이어를 GPU/CPU로 정적 분배 | 라우팅을 안 봄 → 미스율 높음 |
| KTransformers | AMX 최적화 CPU 실행 + 프리필 시점 배치 결정 | 컨텍스트가 바뀌면 배치가 낡음 |
| MoE-Infinity | 요청별 활성화 패턴 추적 → 프리페치 | 예측이 빗나가면 그대로 미스 |
| FreeToken | 미스는 인정하고, 미스를 어떻게 서빙할지를 바꿈 | x86 + NVIDIA Ampere 이상 전용 |
논문의 표현이 이 차이를 정확히 짚습니다 — "라우팅 계산은 정확하게, 모델은 손대지 않은 채 둔다. 바꾸는 것은 예측을 얼마나 잘하느냐가 아니라 남은 미스를 어떻게 서빙하느냐다."
실험은 모델 3종(Qwen3.6-35B-A3B · DeepSeek-V4-Flash · GLM-5.2) × 머신 6종(8GB RTX 4060 노트북부터 RTX PRO 6000 워크스테이션까지) × 실제 에이전트 워크로드 4종으로 돌렸습니다. 가중치 포맷은 모든 엔진이 비트 단위로 동일한 체크포인트를 쓰도록 맞췄습니다.
| 측정 항목 | 결과 |
|---|---|
| RTX 5090 / 35B | 77~83 tok/s — 최고 baseline 대비 1.8~2.3배 |
| RTX 5090 / 284B | 22~25 tok/s — 1.5~1.9배 |
| 8GB RTX 4060 노트북 | 39.3 tok/s (NVFP4) — RTX 4090의 92% |
| RTX PRO 6000 / 753B | 14.9 tok/s vs llama.cpp 7.3 (2.0배) |
| 에이전트 대화 안정성 | 1턴 대비 12% 이내 유지 (KTransformers는 2번째 워크로드에서 이미 31% 손실) |
| 최악 TTFT | 44초 이하 (모든 baseline은 어딘가에서 150초 초과) |
| 서버→소비자 메인보드 | FreeToken은 4% 손실, llama.cpp는 20% 손실 (듀얼채널 DDR5에서 CPU 전문가가 굶음) |
기준선으로 삼을 숫자가 하나 있습니다. 논문은 프로덕션 코딩 에이전트 트레이스의 중앙 디코드 속도 33 tok/s를 "대화 가능"의 임계선으로 잡습니다. 8GB 노트북의 39.3 tok/s가 이 선을 넘긴다는 게 이 프로젝트의 헤드라인입니다.
수치는 저자들이 직접 측정한 값이고 아직 제3자 재현 결과가 널리 나오지 않았습니다. 또 README는 "290B+ 프런티어 MoE"라고 쓰지만 논문 본문의 실제 모델은 284B와 753B입니다. 벤치마크를 인용할 때는 같은 워크로드·같은 양자화 포맷인지 확인하는 게 안전합니다.
언어 비율만 보면 Python 88.4% / CUDA 5.5% / C++ 2.8% / C 2.5%로 "파이썬 프로젝트"처럼 보입니다. 하지만 실제 성능은 그 안의 Triton 커널 44개와 C++/CUDA 12,419줄에서 나옵니다. Python은 조립과 스케줄링을 맡습니다.
| 층 | 구성 |
|---|---|
| API 서버 | FastAPI + uvicorn. OpenAI /v1/chat/completions · Anthropic /v1/messages · Responses /v1/responses 3종 호환 |
| 프로세스 간 통신 | pyzmq + msgpack (프런트엔드 ↔ 스케줄러) |
| 추론 런타임 | PyTorch 2.11 (cu130 빌드, >=2.11,<2.12로 좁게 고정) |
| 커널 | Triton 3.6.0 (정확 핀) 자체 커널 44개 + CUDA/C++ 확장 2개 + flashlib==0.3.0의 슬롯 캐시 커널 |
| 선택 가속 | [fi] flashinfer · [sgl] sglang-kernel 0.4.5 → 권장 설치는 [accel] (둘 다) |
| 커널 JIT | apache-tvm-ffi==0.1.13.post3 + 별도 freetoken-kernel-cache 휠(미리 빌드된 .so) |
| 모델 로딩 | transformers 5.x · huggingface_hub · modelscope · safetensors · gguf |
| 빌드 | setuptools + BuildExtension(use_ninja=True). Dockerfile은 없음 |
| 배포 CI | GitHub Actions 2개 — nightly 휠(self-hosted EC2) + 태그 릴리스(manylinux, cp310~cp313). 테스트 CI는 없음 |
triton==3.6.0, apache-tvm-ffi==0.1.13.post3, flashlib==0.3.0처럼 정확 핀이 여럿이고 numpy<2.5 같은 상한도 있습니다. pyproject.toml에 사유가 주석으로 달려 있습니다 — numpy 상한은 flashlib이 끌고 오는 numba 때문, torch 하한은 sglang-kernel 0.4.5가 2.11에만 있는 심볼을 링크하기 때문.
가장 성가신 건 Marlin NVFP4 경로입니다. vLLM의 미리 빌드된 커널을 빌려 쓰는데 vLLM이 transformers<5를 핀해서 본체의 transformers>=5.5와 충돌합니다. 그래서 이 경로만 별도 가상환경에 수동 설치해야 합니다. sm_80~99(RTX 30/40 세대)에서 NVFP4를 쓰려면 이 함정을 먼저 만납니다.
어텐션 구현이 레지스트리로 관리되고, 모델과 GPU에 따라 자동으로 선택됩니다. --attention-backend에 prefill,decode 쌍을 주면 두 단계에 서로 다른 백엔드를 쓰는 하이브리드로 합성됩니다.
| 이름 | 필요 조건 | 비고 |
|---|---|---|
trtllm | flashinfer + sm100 (Blackwell) | 페이지 크기 16/32/64 |
fi | flashinfer | 범용 |
fa | sgl_kernel | FlashAttention 계열 |
triton | 없음 | 순수 Triton 폴백 — 어디서든 돎 |
dsa / dsv4_sparse / m3_sparse | 없음 | DeepSeek·MiniMax 전용 희소 어텐션 |
@triton.jit을 붙이면 GPU 코드로 컴파일됩니다. FreeToken의 캐시 관리 로직 같은 "복잡하지만 데이터 크기는 작은" 커널이 전부 Triton으로 쓰여 있습니다.구조도에서 눈여겨볼 곳은 맨 아래 세 갈래입니다. 전문가 가중치의 원본은 항상 호스트 RAM에 있고, GPU에 있는 건 캐시일 뿐입니다. 그래서 캐시가 작든 크든 결과는 같고 속도만 달라집니다. 이 성질이 나중에 "실행 중 메모리 재분배"를 안전하게 만들어 줍니다.
프롬프트를 읽는 단계(프리필)에서는 어차피 전문가 대부분이 쓰입니다. 그래서 캐싱을 포기하고 레이어 순서대로 전문가 풀 전체를 스트리밍합니다. 대신 버퍼를 두 개 두고 layer_id % 2로 번갈아 씁니다 — 지금 레이어를 계산하는 동안 다음 레이어를 미리 실어 옵니다.
영리한 부분은 이 두 버퍼가 별도 할당이 아니라는 점입니다. 슬롯 캐시의 앞쪽 2 × 전문가수 칸을 빌려 쓰는 뷰입니다. 그래서 프리필이 끝나면 거기 남은 전문가가 그대로 디코드 캐시의 씨앗이 됩니다. 논문 표현으로는 "별도 프리필 캐시도, 단계 간 인수인계도 없다". 다만 대가가 있어서, --moe-cache-size가 2 × 전문가수보다 작으면 오버랩이 자동으로 꺼집니다.
여기가 이 프로젝트의 심장입니다. 토큰을 하나씩 뱉는 디코드 단계에서, 캐시에 없는 전문가 m개를 만나면 세 갈래로 갈립니다.
이 계산을 파이썬이 아니라 GPU 커널 안에서 한다는 게 핵심 제약입니다. 미스 개수 m은 그 스텝의 실제 라우팅에 따라 달라지는데, 디코드 경로 전체가 CUDA Graph로 캡처돼 있어서 그 값을 호스트로 가져올 수 없습니다. 그래서 분할 결정 자체를 디바이스로 내렸습니다.
moe/offload_kernels.py — Triton 커널 안에서 Q16 고정소수점으로 q*를 푸는 부분입니다.
# 페치 개수에 상한을 건다. 넘치는 미스는 CPU가 계산한다(GPU에 올리지 않음).
if fetch_frac_q16 > 0:
# 대역폭 정합 분할 (fetch_frac = pcie_bw / cpu_bw):
# 페치 시간은 F * (1 - frac)에, CPU 시간은 (M - F) * frac에 비례.
# F = frac * M 에서 균형. 정수가 아니므로 양옆 정수 중
# "느린 쪽(max)"이 더 작아지는 값을 고른다.
lo = (num_missing * fetch_frac_q16) >> 16
cost_lo = tl.maximum(lo * ((1 << 16) - fetch_frac_q16),
(num_missing - lo) * fetch_frac_q16)
cost_hi = tl.maximum((lo + 1) * ((1 << 16) - fetch_frac_q16),
(num_missing - lo - 1) * fetch_frac_q16)
max_fetch = tl.where(cost_lo <= cost_hi, lo, lo + 1)
num_fetch = tl.minimum(num_missing, max_fetch)
fetch_frac_q16은 B_P / B_H를 65536을 곱한 정수로 표현한 값입니다(GPU 커널에서 부동소수점 나눗셈을 피하려는 것). 이 값의 출처는 ft bench bw가 남긴 프로파일인데, 여기서도 세심한 선택이 하나 있습니다 — 단독 측정값이 아니라 두 경로를 동시에 돌렸을 때의 측정값을 우선합니다. 실제 상황에서는 PCIe 전송과 CPU 연산이 메모리 컨트롤러를 놓고 경쟁하니까요.
같은 로직의 순수 파이썬 레퍼런스 구현이 같은 파일에 나란히 있습니다. 테스트가 GPU 커널과 이 레퍼런스의 결과를 비트 단위로 대조합니다.
if frac_q16 > 0:
m, q = len(missing), 1 << 16
lo = (m * frac_q16) >> 16
cost = lambda f: max(f * (q - frac_q16), (m - f) * frac_q16)
max_fetch = lo if cost(lo) <= cost(lo + 1) else lo + 1
GPU와 CPU가 동시에 돌면 "누가 무엇을 맡았는지"를 표시해야 합니다. 그런데 CUDA Graph 안에서는 배열 크기가 바뀌면 안 되므로 리스트를 잘라내는 방식을 쓸 수 없습니다. 해법은 -1이라는 sentinel 값 하나와 마스킹입니다.
raw = topk_ids.clone() # 원본 전문가 id 보관
cache.ensure_experts_hybrid(layer_id, topk_ids) # → 슬롯 번호, 아니면 -1
on_gpu = topk_ids >= 0
# CPU 쪽: GPU가 맡은 자리는 id를 -1로 (C++ 커널이 id<0을 건너뜀)
cpu_ids = torch.where(on_gpu, raw.new_full((), -1), raw)
pending = executor.decode_submit(layer_id, hidden_states, topk_weights, cpu_ids)
↑ CPU 작업을 먼저 던져 놓는다
cache.copy_missing() # ↓ 그 사이 PCIe 페치 진행
# GPU 쪽: CPU가 맡은 자리는 가중치를 0으로 (계산해도 결과에 안 더해짐)
gpu_slots = topk_ids.clamp_min(0) # -1 → 슬롯 0
gpu_w = torch.where(on_gpu, topk_weights, topk_weights.new_zeros(()))
gpu_routed = self._expert_gemm(cache, hidden_states, gpu_w, gpu_slots, ...)
cpu_routed = executor.decode_sync(pending) # CPU 결과 회수
return gpu_routed + cpu_routed # 부분합 병합
분기(if) 없이 elementwise 연산만으로 두 경로를 갈랐습니다. clamp_min(0)이 -1을 슬롯 0으로 밀어 넣지만 그 자리의 가중치가 0이라 결과에 영향이 없습니다. 배열 모양이 고정되니 CUDA Graph 캡처가 그대로 통과합니다.
GPU가 CPU 결과를 기다려야 하는데, 그래프 안에서는 일반적인 동기화를 쓸 수 없습니다. FreeToken은 두 가지를 구현해 두고 골라 씁니다.
| 방식 | 동작 | 비용 |
|---|---|---|
cudaLaunchHostFunc | 호스트 함수를 그래프의 노드로 삽입 | 호출당 30~50μs. MoE 레이어당 2회 × 75레이어 = 스텝당 약 6ms |
| 스트림 메모리 연산 기본값 | cuStreamWriteValue64 / WaitValue64로 플래그 핸드셰이크 | 거의 0. 단 Windows WDDM · vGPU · 구형 드라이버에서 불가 → 기동 시 프로브 후 폴백 |
여기 딸린 각주가 하나 있는데, 이게 엣지 환경 최적화의 성격을 잘 보여줍니다. "GPU에서 spin-wait 커널을 돌려 기다리면 안 되냐"에 대한 코드 주석의 답은 — 그렇게 하면 GPU 사용률이 99%로 보고되고, 노트북의 전력 스케줄러가 그걸 보고 CPU 클럭을 낮춰서 결과적으로 더 느려진다는 것입니다. 데이터센터에서는 나오지 않을 종류의 제약입니다.
여러 전문가를 한 번에 복사하려고 cudaMemcpyBatchAsync를 쓰는데, 배치 안에 256KB 미만 항목과 큰 항목이 섞이면 이 API가 말없이 동기 복사로 강등됩니다(H100 + CUDA 13.0에서 이분 탐색으로 확인). 실측 피해는 gpt-oss 2048토큰에서 전체 -22%. 그래서 코드에 _SMALL_BANK_FEAT_BYTES = 256 * 1024라는 상수와 함께 긴 주석이 남아 있습니다.
전문가 캐시는 LRU(가장 오래 안 쓴 것부터 버리기)로 관리되는데, 실은 LRU가 두 개입니다.
usage라는 int64 타임스탬프 배열과 전역 step 카운터. 자리가 없으면 argmin(usage)로 희생자를 고릅니다. 전문가 id는 layer_id × 전문가수 + expert라는 평평한 하나의 번호 공간으로 관리해서, 버릴 때 (레이어, 전문가)로 되돌릴 필요가 없습니다.미스율 실측이 이 설계의 효과를 보여줍니다. 같은 라우팅 트레이스를 재생했을 때, RTX 5090이 담을 수 있는 용량(Qwen3.6 전문가 풀의 37%)에서:
| 배치 전략 | Qwen3.6 미스율 | DSV4-Flash 미스율 |
|---|---|---|
| FreeToken 전역 LRU | 16% | 39% |
| KTransformers (프리필 시점 배치) | 41% | 59% |
| llama.cpp (정적 분배) | 62% | 89% |
ft daemon은 별도 프로세스이고, torch · CUDA · flashinfer · sgl_kernel을 절대 임포트하지 않습니다. 이유가 README에 명시돼 있습니다 — "CUDA 오류나 네이티브 확장의 세그폴트는 그것을 로드한 프로세스를 죽인다. 그리고 데몬은 그중 아무것도 로드하지 않는다."
더 인상적인 건 이 원칙을 테스트로 강제한다는 점입니다. tests/daemon/test_daemon_import_safety.py가 import sentinel을 심어 두고, 데몬 모듈을 임포트했을 때 금지 모듈이 딸려 들어오면 실패합니다. 규칙을 문서에만 적어 두면 6개월 뒤 누군가 무심코 어깁니다.
실행 중 캐시 재분배(ft ctl cache --moe 4096 --kv 32k)는 위험한 연산입니다. 기존 캐시를 해제한 뒤 새 크기로 잡는데, 중간에 실패하면 서빙 불가 상태로 영구히 떨어집니다. 그래서 순서를 뒤집었습니다 — 해제 전에 기하 구조와 메모리 예산 적합성을 전부 검증하고, 안 맞으면 CacheRebuildRejected로 거절합니다. 거절은 복구 가능하고, 반쯤 해제된 상태는 복구 불가능하니까요.
이 트리에서 읽어 낼 것이 몇 가지 있습니다.
kernel/이 가장 크다(16k줄). 파이썬 프로젝트로 보이지만 무게중심은 커널입니다. 그중 csrc/gguf/ 5,733줄은 llama.cpp의 ggml 커널을 그대로 이식한 것 — 남의 검증된 커널을 다시 안 짜고 가져왔습니다.function_call_parser.py 하나가 3,735줄. 모델 패밀리마다 툴콜을 표현하는 형식이 제각각이라, 그걸 전부 파싱하는 코드입니다. "에이전트를 제대로 지원한다"는 말의 실제 비용이 여기 있습니다.tests/README.md: GPU 박스에서 2~4분).deepseek_v4/는 10개 파일(args · attention · compress · config · layers · model · moe · ops · weight)로 쪼개져 있습니다. 공통 추상화로 억지로 묶지 않고 모델별로 펼친 구조입니다.코드를 따라 읽을 거라면 이 순서가 가장 빠릅니다: moe/offload_kernels.py(q* 커널) → layers/moe.py(하이브리드 디코드) → moe/offload_cache.py(더블 버퍼 · 뱅크 스키마) → moe/cpu_executor.py(CPU 쪽) → engine/cache_budget.py(메모리 배분 정책). 이 5개 파일이면 논문의 주장 대부분이 코드로 확인됩니다.
추론 엔진을 만들 일이 없어도 배울 게 남는 저장소입니다. 여기 있는 기법 대부분은 "제약이 심한 환경에서 어떻게 설계하나"에 대한 답이라 범용성이 있습니다.
"배열 크기가 고정이어야 하고 중간 결과를 CPU로 못 가져온다"는 제약 아래서 LRU 캐시를 굴리는 방법 전체가 여기 있습니다. 히트 판정 · 희생자 선택 · 슬롯 재매핑 · 인덱스 덮어쓰기를 단일 블록 Triton 커널 하나에서 3단계로 처리합니다. sentinel(-1) + clamp_min + 0 가중치로 분기를 없애는 패턴은 다른 곳에도 그대로 쓸 수 있습니다.
볼 파일: moe/offload_kernels.py
"두 개의 서로 다른 처리기에 일을 나눠 둘 다 동시에 끝내기"는 고전 스케줄링 문제입니다. FreeToken은 그 답 q* = m·B_P/B_H를 유도하고 → Q16 고정소수점으로 정수화하고 → 정수가 아닌 최적점 주변의 두 정수 중 느린 쪽이 더 작아지는 것을 고르는 데까지 갑니다. 이론식이 실제 GPU 커널 열 줄로 착지하는 과정을 통째로 볼 수 있는 흔치 않은 사례입니다.
볼 파일: moe/offload_kernels.py + moe/bench_profile.py
전용 copy 스트림 + release_event / ready_event 한 쌍으로 더블 버퍼를 안전하게 돌리는 정석. "다음 사용자가 버퍼를 덮어쓰기 전에 이전 사용자가 다 읽었는지"를 이벤트로만 보장합니다. 스레드·락 없이 동시성을 다루는 방식이라 GPU 밖에서도 응용됩니다.
볼 파일: moe/offload_cache.py · layers/moe.py
모든 텐서를 4096바이트에 정렬하고, 샤드 경계도 4096에 맞춰서, 텐서가 샤드를 걸쳐도 각 조각이 정렬을 유지하게 만든 포맷입니다. 여기에 preadv의 short read 루프 처리가 붙는데 — 이걸 빼먹으면 가중치 중간에 에러 없이 쓰레기 바이트가 들어갑니다. 파일 I/O에서 가장 무서운 종류의 버그입니다.
볼 파일: checkpoint/ftw.py
큰 호스트 버퍼를 준비할 때 ① lazy 익명 mmap 할당 → ② 실제 데이터로 채움 → ③ 그 다음에 cudaHostRegister로 pin. 순서를 ①③②로 바꾸면 pin이 전 페이지에 fault를 내고 0으로 채우는데, 그 0이 곧바로 데이터로 덮어써져 통째로 낭비됩니다. DSV4의 137GiB에서 약 47초가 여기서 갈립니다.
볼 파일: moe/host_banks.py
KV 페이지와 GDN recurrent state라는 단위가 다른 두 자원을 같은 접두사 트리에 매달아 두고, 각자의 LRU로 축출하되 full_ref ≥ mamba_ref라는 불변식을 유지합니다. 내부 노드는 tombstone으로 처리. 트랜스포머와 상태공간 모델이 섞인 하이브리드 아키텍처를 지원하려면 반드시 마주치는 문제입니다.
볼 파일: kvcache/hybrid_radix_cache.py · swa_radix_cache.py
코딩 에이전트는 대화 중간을 계속 고쳐 씁니다(툴 결과를 끼워 넣고, thinking 블록을 지우고). 일반적인 접두사 캐시는 편집 지점 이후를 통째로 버리는데, FreeToken은 툴콜 오프너 토큰 위치에 상태 스냅샷을 ping-pong 슬롯으로 얼려 둡니다. 이 도메인 지식이 앞의 "에이전트 워크로드에서 12% 이내 유지"라는 수치를 만듭니다.
볼 파일: scheduler/cache.py
8종 양자화 포맷(bf16 · fp8_block · q4_0 · nvfp4 · mxfp4 · …)을 뱅크 이름 튜플로만 선언해 두면, 캐시 이동 기계는 뱅크가 몇 개인지 레이아웃이 어떤지 전혀 몰라도 동작합니다. 포맷을 아는 건 맨 마지막 GEMM 디스패치뿐입니다. 새 포맷 추가가 표에 한 줄 넣는 일이 됩니다.
볼 파일: moe/offload_cache.py
축출·페치 결정을 순수 파이썬으로 한 번 더 구현해 두고 비트 단위로 대조합니다. GPU 커널은 디버깅이 어려우니 "정답"을 옆에 두는 방식입니다. tests/README.md의 한 줄이 이 저장소에서 가장 인용할 만합니다 — "기댓값을 무엇에 대조해 검사하는지 말할 수 없다면, 그 테스트는 자리값을 못 한다."
볼 파일: moe/offload_kernels.py · tests/README.md
"제어 프로세스는 torch를 임포트하지 않는다"는 규칙을 import sentinel 테스트로 강제합니다. 아키텍처 문서에 적어 두는 것과 테스트가 막아 주는 것은 6개월 뒤에 전혀 다른 결과를 냅니다.
볼 파일: daemon/README.md · tests/daemon/test_daemon_import_safety.py
| 항목 | 요구사항 |
|---|---|
| CPU 아키텍처 | x86_64 전용 (aarch64 · DGX Spark은 로드맵) |
| GPU | NVIDIA Ampere 이상 — RTX 30 / 40 / 50 시리즈. 1080 · 2080 등은 미지원 |
| 드라이버 | r580 이상 (CUDA 13) |
| OS | pip 설치는 Linux. 데스크탑 앱은 Windows + Linux. macOS 미지원 |
| Python | 3.10 이상 (휠은 cp310~cp313) |
| CUDA 툴킷 | 커널 JIT 컴파일용 nvcc(CUDA 13)가 PATH에. kernel-cache 휠이 있으면 대부분 생략 가능 |
| 호스트 RAM | 전문가 가중치 크기만큼 여유. 예: Qwen3.6-35B-A3B(bf16)에 약 70GB. NVFP4 체크포인트는 훨씬 적음 |
| VRAM | 공식 하한 명시 없음. 실측으로 8GB에서 35B(NVFP4) 39.3 tok/s 확인 |
많은 사람이 "8GB GPU로 35B이 돌아간다"만 보고 시도했다가 여기서 막힙니다. 전문가 전체가 호스트 RAM에 상주해야 하므로, bf16 35B이면 RAM 70GB가 필요합니다. 게이밍 PC의 32GB로는 부족합니다. 대안은 NVFP4 같은 저비트 체크포인트를 쓰는 것 — 그래서 8GB 노트북 실험도 NVFP4 빌드였습니다.
CPU 쪽도 조건이 붙습니다. CPU MoE 실행기는 실행 시점에 명령어 집합을 골라 씁니다 — AVX-512-BF16 → AVX-512F → AVX2+FMA → 스칼라 순. 스칼라로 떨어지면 하이브리드 모드의 이점이 사실상 사라집니다. 스레드는 물리 코어당 하나만 고정하는데, 이유가 명확합니다 — MoE 디코드는 메모리 대역폭에 묶인 작업이라 하이퍼스레딩 형제 스레드는 같은 코어의 로드 포트만 두고 다툴 뿐 대역폭을 늘려 주지 않습니다.
| 제약 | 상태 |
|---|---|
| 멀티 GPU / 텐서 병렬 | 미지원 (TP=1 고정). 이슈 #29 · #15 · #43 · #54 · #95 |
| AMD ROCm · Apple Silicon | 미지원. 커뮤니티 RDNA4 포팅 시도가 있음(비공식) |
| 멀티모달 입력 | 체크포인트는 로드되지만 텍스트 전용으로 서빙 |
| speculative decoding | 미구현 (로드맵) |
| GGUF | Gemma-4에만 지원 |
공개 2주차 프로젝트답게 열린 이슈가 96개이고 그중 치명적인 것들이 있습니다 — #87 48k 토큰 초과 프리필이 서버 전체를 크래시(fla/l2norm 커널에서 device not ready), #110 프리필 워크스페이스의 CUDA OOM이 처리되지 않아 엔진이 죽는데 서버는 계속 요청을 받음, #111 KV 풀보다 긴 요청이 에러 없이 영원히 큐에 남음. 프로덕션 투입 전에 이슈 트래커를 한 번 훑는 편이 좋습니다.
커널 코드를 읽지 말고 공식만 보고 직접 계산해 보는 겁니다. 미스 m=20, PCIe 대역폭 32GB/s, 호스트 메모리 대역폭 80GB/s일 때 몇 개를 페치해야 할까요? 그리고 페치 개수를 0부터 20까지 바꿔 가며 max(페치시간, CPU시간)을 그래프로 그려 보세요. V자 골짜기의 바닥이 공식값과 일치하는지 확인합니다.
확인 포인트: 최적점이 정수가 아닐 때 양옆 중 어느 쪽이 나은지 — 그게 코드의 cost_lo vs cost_hi 비교입니다.
ft bench bw가 하는 일을 직접 해 봅니다. ⓐ numpy로 큰 배열 복사를 반복해 호스트 메모리 대역폭을 재고, ⓑ GPU가 있다면 torch의 pinned 텐서 → GPU 복사로 PCIe 대역폭을 잽니다. 두 값의 비가 곧 fetch_frac입니다.
확인 포인트: 둘을 따로 잴 때와 동시에 돌리며 잴 때의 값이 다를 겁니다. FreeToken이 동시 측정값을 우선하는 이유를 여기서 직접 보게 됩니다.
가짜 라우팅 트레이스를 만듭니다 — 레이어 32개 × 전문가 64개, 토큰마다 각 레이어에서 6개를 (균등하지 않게, 지수 분포로) 뽑습니다. 그 트레이스에 대해 세 가지 배치 전략의 미스율을 재 보세요: ⓐ 전역 LRU ⓑ 처음 N개를 고정 배치 ⓒ 레이어 앞쪽 절반만 GPU. 캐시 용량을 전체의 20% / 40% / 60%로 바꿔 가며 표를 만듭니다.
확인 포인트: 논문의 16% vs 41% vs 62%와 비슷한 서열이 나오는지. 절댓값은 트레이스에 따라 다르지만 순서는 재현될 겁니다. 안 나온다면 라우팅 분포의 편향(전문가마다 인기 차이)을 더 세게 줘 보세요.
FreeToken은 측정 전용 환경변수를 프로덕션 코드에 남겨 뒀습니다. FREETOKEN_HYBRID_OVERLAP=0을 주면 CPU 풀을 PCIe 페치 전에 동기화해서 두 경로를 직렬화합니다. 같은 프롬프트로 켠 상태 / 끈 상태의 tok/s를 재면 겹치기가 만든 이득만 분리됩니다.
같은 방식으로 FREETOKEN_CPU_MOE_FLAG_SYNC=0(스트림 메모리 연산 대신 host-func)도 재 보세요. 레이어 수가 많은 모델일수록 차이가 커야 합니다.
확인 포인트: 자기 하드웨어에서의 이득이 논문 수치와 다를 수 있습니다. 그게 정상입니다 — q*는 머신마다 다른 값이니까요. 오히려 왜 다른지를 대역폭 비로 설명할 수 있으면 성공입니다.
attention/__init__.py의 BackendInfo 구조를 읽고, 순수 Triton 백엔드(triton)를 베껴 일부러 느린 백엔드를 하나 추가해 보세요. 목표는 성능이 아니라 확장 지점의 계약을 이해하는 것입니다 — 어떤 능력(FULL / SWA / MLA)을 선언해야 하는지, 페이지 크기 제약이 어디서 검증되는지, --attention-backend prefill,decode로 합성될 때 무엇이 요구되는지.
다음 단계: 여기까지 왔으면 moe/__init__.py의 MoE 백엔드 레지스트리도 같은 구조입니다. --moe-backend의 auto 선택 규칙(engine/engine.py)을 읽어 보면, 왜 auto가 절대 fused를 고르지 않는지 알게 됩니다 — 잘못 고르면 느려지는 게 아니라 가중치 로딩 OOM이라서입니다.
| 주차 | 주제 | 구체적으로 할 것 |
|---|---|---|
| 1주차 | MoE 구조 이해 | Switch Transformer · Mixtral 논문을 읽고, 라우터 + top-k 게이팅을 PyTorch로 직접 구현. 총 파라미터와 활성 파라미터가 왜 다른지 손으로 계산해 볼 것. 여기가 안 되면 나머지가 다 안 됩니다. |
| 2주차 | 추론 엔진의 기본기 | KV 캐시 · 프리필 vs 디코드 · 연속 배칭(continuous batching) · PagedAttention. vLLM 논문 한 편이면 충분합니다. FreeToken의 scheduler/가 여기서부터 읽힙니다. |
| 3주차 | 메모리 계층과 대역폭 | pinned memory · cudaMemcpyAsync · CUDA 스트림과 이벤트. 직접 벤치를 짜서 pageable vs pinned, 동기 vs 비동기 복사의 차이를 재 볼 것. 과제 2가 여기 해당합니다. |
| 4주차 | Triton 커널 작성 | 공식 튜토리얼(vector add → fused softmax → matmul)을 끝까지. 그다음 moe/offload_kernels.py를 열면 읽을 수 있는 상태가 됩니다. 단일 블록 grid=(1,) 커널이 왜 이상해 보이지 않게 되는 지점. |
| 5주차 | CUDA Graph | 캡처와 재생, 그리고 무엇이 캡처를 깨뜨리는지. 작은 예제로 일부러 깨뜨려 보는 게 가장 빠릅니다(동적 shape · .item() 호출 · 호스트 동기화). 이걸 겪고 나면 FreeToken의 sentinel 트릭이 "당연한 해법"으로 보입니다. |
| 6주차 | 양자화 포맷 | FP8 · MXFP4 · NVFP4 · GGUF Q4_0의 비트 레이아웃과 스케일 처리. moe/offload_cache.py의 뱅크 스키마 8종이 각각 어떤 포맷에 대응하는지 매핑해 보면 정리됩니다. |
6주가 부담스럽다면 1주차 + 4주차 + 5주차만 해도 이 저장소의 핵심(q* 커널과 하이브리드 디코드)은 읽힙니다. 반대로 시간이 남으면 llama.cpp의 MoE 오프로딩 코드를 같이 읽어 보세요 — 같은 문제를 정적 분배로 푼 쪽을 보면 FreeToken의 선택이 왜 다른지가 선명해집니다.
| 용어 | 뜻 |
|---|---|
| MoE | Mixture of Experts. 레이어를 여러 전문가로 쪼개고 토큰마다 일부만 통과시키는 구조 |
| 활성 파라미터 | 토큰 하나를 처리할 때 실제로 계산에 참여하는 파라미터. 284B-A13B의 뒤쪽 숫자 |
| 라우터 (게이트) | 토큰마다 어느 전문가로 보낼지 고르는 작은 신경망. 이 선택이 미리 예측 불가한 게 문제의 출발점 |
| 프리필 (prefill) | 프롬프트 전체를 한 번에 읽어 KV 캐시를 채우는 단계. 계산량이 많고 병렬성이 큼 |
| 디코드 (decode) | 토큰을 하나씩 생성하는 단계. 계산량은 적고 메모리 대역폭에 묶임 |
| TTFT | Time To First Token. 첫 글자가 나오기까지의 지연. 프리필 성능의 지표 |
| KV 캐시 | 이미 처리한 토큰들의 어텐션 중간값 저장소. 안 두면 매 토큰마다 전체를 다시 계산해야 함 |
| 오프로딩 | GPU에 안 들어가는 가중치를 호스트 RAM·디스크에 두고 필요할 때 옮기는 기법 |
| pinned memory | OS가 스왑하지 못하게 고정한 호스트 메모리. GPU가 DMA로 직접 읽을 수 있어 전송이 빠름 |
| PCIe | GPU와 메인보드를 잇는 통로. 5.0 x16의 실측 상한이 약 52.7GB/s — 이 숫자가 오프로딩 성능의 천장 |
| q* (q-star) | 이 프로젝트의 핵심 개념. 미스 중 PCIe로 가져올 개수. q* = m × B_P / B_H |
| makespan | 병렬로 돌린 작업들 중 가장 늦게 끝나는 것의 완료 시각. q*가 최소화하는 대상 |
| CUDA Graph | GPU 명령 묶음을 녹화해 재생하는 기능. 오버헤드를 줄이지만 동적 shape 금지 |
| Triton | GPU 커널을 파이썬 문법으로 작성하는 언어이자 컴파일러 |
| LRU | Least Recently Used. 가장 오래 안 쓴 항목부터 버리는 캐시 축출 정책 |
| radix prefix 캐시 | 대화 접두사를 트리로 공유해 재계산을 피하는 KV 캐시 구조 |
| SWA | Sliding Window Attention. 최근 N개 토큰만 보는 어텐션. KV 캐시가 무한히 안 커짐 |
| GDN | Gated DeltaNet. 선형 어텐션 계열의 recurrent 상태. KV 페이지와 단위가 달라 별도 관리 필요 |
| NVFP4 / MXFP4 | 4비트 부동소수점 양자화 포맷. 가중치 크기를 크게 줄여 작은 RAM에서 큰 모델을 가능하게 함 |
| FTW | FreeToken Weights. 4096바이트 정렬로 O_DIRECT 읽기에 최적화한 자체 체크포인트 포맷 |
| AVX-512-BF16 | CPU의 bfloat16 곱셈누산 명령어. 이게 있어야 CPU MoE 실행이 쓸 만한 속도가 나옴 |
| sentinel 값 | "해당 없음"을 뜻하는 특수 값(여기서는 -1). 배열 크기를 고정한 채 분기를 표현하는 수단 |