5분 AI.
TRENDSHIFT #2 딥다이브 · 2026-08-24 분석

FreeToken 딥다이브
— 게이밍 PC 한 대로 284B 모델을 돌리는 엔진

VRAM에 안 들어가는 거대 MoE 모델을, 집에 있는 그래픽카드 한 장으로 대화 가능한 속도로 서빙합니다. 핵심 아이디어는 단순합니다 — 부족한 전문가를 PCIe로 "가져올 것"과 CPU에서 "그냥 계산할 것"으로 쪼개서 동시에 굴린다. (저장소: FlashML-org/FreeToken · ⭐3.2k · Python · Apache-2.0 · v0.1.2 · 2026-08-24 기준)
목차
  1. 프로젝트 한줄 요약
  2. 왜 주목받는가
  3. 기술 스택 전체 지도
  4. 아키텍처 심화 분석
  5. 디렉토리 구조 해부
  6. 학습 포인트 (기술별)
  7. 하드웨어 / 시스템 요구사항
  8. 직접 해볼 수 있는 실습 과제
  9. 관련 기술 심화 학습 로드맵
  10. 핵심 키워드 사전
  11. 참고 링크

1프로젝트 한줄 요약

FreeToken이 정확히 무엇을 하는 물건인가

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)에서 그 자리에서 깎아 씁니다. 나르는 시간과 깎는 시간이 같이 끝나도록 개수를 나누는 게 이 프로젝트의 정체입니다.

용어
MoE (Mixture of Experts, 전문가 혼합)
신경망 레이어 하나를 여러 개의 작은 전문가로 쪼개 두고, 토큰마다 라우터가 그중 몇 개만 골라 통과시키는 구조. 파라미터는 크게 늘리면서 토큰당 계산량은 작게 유지하는 게 목적입니다. 전체 파라미터를 "총 파라미터", 실제로 쓰이는 부분을 "활성 파라미터"라 부르고, 284B-A13B 같은 표기가 이 둘을 뜻합니다.
용어
오프로딩 (offloading)
GPU 메모리에 다 안 들어가는 가중치를 호스트 RAM이나 디스크에 두고, 필요할 때만 GPU로 옮겨 쓰는 기법. FreeToken에서 전문가 가중치의 원본은 항상 호스트 RAM의 고정 메모리(pinned memory)이고, GPU에 있는 건 캐시일 뿐입니다. 그래서 GPU 메모리가 얼마든 정답은 안 바뀌고 속도만 바뀝니다.

2왜 주목받는가

트렌딩 이유와 경쟁 제품 대비 강점

공개 첫 커밋이 2026년 8월 11일, 그로부터 2주도 안 돼 별 3.2k. 속도의 이유는 세 가지가 겹친 것입니다 — 실측 수치가 크고, 저자진이 화려하고(Ion Stoica · Matei Zaharia · Song Han · Kurt Keutzer), 대상이 개인 PC라 확인해 볼 사람이 많습니다.

기존 방식의 문제
엣지 MoE 서빙이 막히는 지점 3가지

① 프리필이 희소성을 파괴한다. 프롬프트가 길어지면 토큰마다 다른 전문가가 뽑히고, 그 합집합이 결국 전문가 전체를 덮습니다. "토큰당 6개"라는 희소성이 프롬프트 단위에서는 사라집니다.

② 디코드 미스가 전부 PCIe로 몰린다. 캐시에 없는 전문가는 매번 PCIe로 실어 나르므로, 속도가 링크 대역폭에 그대로 묶입니다. 그동안 호스트 CPU와 RAM 대역폭은 놀고 있습니다.

③ 엣지 GPU는 전용이 아니다. 게임이 켜지고 브라우저가 VRAM을 먹으면 가용량이 세션 중에도 변합니다. 기동 시점에 메모리 배분을 고정하는 설계는 여기서 깨집니다.

FreeToken의 해법
q* 분할 · 의미 기반 캐싱 · 탄력적 메모리

① q* 정책 — 디코드 스텝마다 캐시 미스를 "PCIe로 가져올 것"과 "CPU에서 계산할 것"으로 쪼갭니다. 공식은 q* = m × B_P / B_H (m=미스 수, B_P=PCIe 대역폭, B_H=호스트 메모리 대역폭). 두 경로가 동시에 끝나도록 나눈 값입니다.

② 의미 기반 캐싱 — 에이전트가 컨텍스트를 고쳐 쓰는 지점(툴콜·thinking 블록)에 KV 스냅샷을 앵커로 박아 둡니다. 대화가 편집돼도 그 지점까지는 재사용하고 새로 붙은 부분만 다시 프리필합니다.

③ 탄력적 메모리 — 전문가 캐시와 KV 캐시의 VRAM 배분을 엔진 재시작 없이 스케줄러의 안전 지점에서 재조정합니다.

비교 대상은 vLLM이 아니다

흔한 오해 하나를 먼저 정리해야 합니다. FreeToken이 겨루는 상대는 vLLM이나 SGLang이 아닙니다. 그쪽은 여러 장의 데이터센터 GPU에 모델이 다 들어가는 상황을 전제합니다. 실제로 README의 감사 문단은 vLLM · SGLang · FlashInfer · llama.cpp를 코드와 설계를 빌려온 상류(上流)로 명시합니다. 논문이 실험에서 비교한 baseline은 llama.cpp · Ollama · KTransformers · MoE-Infinity — 전부 "엣지에서 오프로딩으로 버티는" 진영입니다.

엔진전문가를 어떻게 다루나한계
llama.cpp / Ollama로드 시점에 레이어를 GPU/CPU로 정적 분배라우팅을 안 봄 → 미스율 높음
KTransformersAMX 최적화 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 / 35B77~83 tok/s — 최고 baseline 대비 1.8~2.3배
RTX 5090 / 284B22~25 tok/s — 1.5~1.9배
8GB RTX 4060 노트북39.3 tok/s (NVFP4) — RTX 4090의 92%
RTX PRO 6000 / 753B14.9 tok/s vs llama.cpp 7.3 (2.0배)
에이전트 대화 안정성1턴 대비 12% 이내 유지 (KTransformers는 2번째 워크로드에서 이미 31% 손실)
최악 TTFT44초 이하 (모든 baseline은 어딘가에서 150초 초과)
서버→소비자 메인보드FreeToken은 4% 손실, llama.cpp는 20% 손실 (듀얼채널 DDR5에서 CPU 전문가가 굶음)

기준선으로 삼을 숫자가 하나 있습니다. 논문은 프로덕션 코딩 에이전트 트레이스의 중앙 디코드 속도 33 tok/s를 "대화 가능"의 임계선으로 잡습니다. 8GB 노트북의 39.3 tok/s가 이 선을 넘긴다는 게 이 프로젝트의 헤드라인입니다.

주의해서 읽을 것

수치는 저자들이 직접 측정한 값이고 아직 제3자 재현 결과가 널리 나오지 않았습니다. 또 README는 "290B+ 프런티어 MoE"라고 쓰지만 논문 본문의 실제 모델은 284B와 753B입니다. 벤치마크를 인용할 때는 같은 워크로드·같은 양자화 포맷인지 확인하는 게 안전합니다.

3기술 스택 전체 지도

무엇으로 만들어졌나 — 파이썬 껍데기 안의 커널 뭉치

언어 비율만 보면 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] (둘 다)
커널 JITapache-tvm-ffi==0.1.13.post3 + 별도 freetoken-kernel-cache 휠(미리 빌드된 .so)
모델 로딩transformers 5.x · huggingface_hub · modelscope · safetensors · gguf
빌드setuptools + BuildExtension(use_ninja=True). Dockerfile은 없음
배포 CIGitHub 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-backendprefill,decode 쌍을 주면 두 단계에 서로 다른 백엔드를 쓰는 하이브리드로 합성됩니다.

이름필요 조건비고
trtllmflashinfer + sm100 (Blackwell)페이지 크기 16/32/64
fiflashinfer범용
fasgl_kernelFlashAttention 계열
triton없음순수 Triton 폴백 — 어디서든 돎
dsa / dsv4_sparse / m3_sparse없음DeepSeek·MiniMax 전용 희소 어텐션
용어
Triton
GPU 커널을 파이썬 문법으로 쓰는 언어이자 컴파일러(OpenAI 개발). CUDA C++로 짜야 했던 커널을 파이썬 함수처럼 작성하고 @triton.jit을 붙이면 GPU 코드로 컴파일됩니다. FreeToken의 캐시 관리 로직 같은 "복잡하지만 데이터 크기는 작은" 커널이 전부 Triton으로 쓰여 있습니다.

4아키텍처 심화 분석

진입점부터 전문가 하나가 계산되기까지

전체 구조도

ft (콘솔 스크립트) ─ python/freetoken/cli.py:main() │ ※ 모든 서브커맨드를 함수 안에서 지연 임포트 → ft --version 이 torch 없이 뜬다 ├─ serve 엔진 + API 서버 기동 ├─ shell 터미널 대화 UI ├─ ctl 실행 중 엔진 제어 (통계 · 캐시 리사이즈) ├─ daemon ★ torch를 임포트하지 않는 감시 프로세스 (포트 1900) ├─ launch 코딩 에이전트(claude/codex/opencode…) 설정 + 기동 ├─ checkpoint HuggingFace → FTW 포맷 변환 └─ bench bw 이 머신의 CPU vs PCIe 대역폭 측정 (q* 의 입력값) ┌──────────────────────────────────────────────────────────────┐ │ FastAPI server/api_server.py 포트 1919 │ │ OpenAI /v1/chat · Anthropic /v1/messages · /v1/responses │ └───────────────────────────┬──────────────────────────────────┘ │ pyzmq + msgpack ┌───────────────────────────▼──────────────────────────────────┐ │ Scheduler scheduler/scheduler.py │ │ ├ TableManager 페이지 테이블 슬롯 │ │ ├ CacheManager prefix 캐시 + KV/SWA/GDN 3통화 할당 │ │ ├ PrefillManager 청크 단위 프리필 │ │ └ DecodeManager 디코드 배치 │ └───────────────────────────┬──────────────────────────────────┘ ┌───────────────────────────▼──────────────────────────────────┐ │ Engine engine/engine.py │ │ ├ Context (전역): page_table · attn_backend · moe_backend │ │ ├ OffloadMoeCache moe/offload_cache.py ★ 핵심 │ │ ├ CpuMoeExecutor moe/cpu_executor.py ★ 핵심 │ │ └ GraphRunner CUDA Graph 캡처 │ └───────────────────────────┬──────────────────────────────────┘ │ ┌───────────────────────┼───────────────────────┐ ▼ ▼ ▼ GPU VRAM PCIe 링크 호스트 RAM ├ KV 캐시 (52.7 GB/s @5.0 x16) └ 전문가 원본 └ 전문가 캐시 슬롯 ←── 미스 페치 ────── (pinned, 진실의 원천) │ CPU MoE 실행기 (AVX-512 / AVX2)

구조도에서 눈여겨볼 곳은 맨 아래 세 갈래입니다. 전문가 가중치의 원본은 항상 호스트 RAM에 있고, GPU에 있는 건 캐시일 뿐입니다. 그래서 캐시가 작든 크든 결과는 같고 속도만 달라집니다. 이 성질이 나중에 "실행 중 메모리 재분배"를 안전하게 만들어 줍니다.

설계 패턴 ① — 프리필: 전 레이어 더블 버퍼 스트리밍

프롬프트를 읽는 단계(프리필)에서는 어차피 전문가 대부분이 쓰입니다. 그래서 캐싱을 포기하고 레이어 순서대로 전문가 풀 전체를 스트리밍합니다. 대신 버퍼를 두 개 두고 layer_id % 2로 번갈아 씁니다 — 지금 레이어를 계산하는 동안 다음 레이어를 미리 실어 옵니다.

버퍼 A: [레이어0 전송]──[레이어0 계산] [레이어2 전송]──[레이어2 계산] 버퍼 B: [레이어1 전송]──[레이어1 계산] [레이어3 전송]──… └── 전송이 계산 뒤에 완전히 숨는다 ──┘ 실측 (RTX 5090, Qwen3.6 BF16): 8,192토큰 청크 하나 = 1.19~1.22초 = 64.4GB 전문가 풀을 52.7GB/s로 한 번 흘리는 시간 = PCIe 5.0 x16의 실질 상한 → 전문가 계산이 전송 뒤에 완전히 가려짐 → 16k 토큰에서 6.7k tok/s 버퍼를 하나로 줄이면: 4k에서 -19% 8k에서 -25% 16k에서 -26%

영리한 부분은 이 두 버퍼가 별도 할당이 아니라는 점입니다. 슬롯 캐시의 앞쪽 2 × 전문가수 칸을 빌려 쓰는 뷰입니다. 그래서 프리필이 끝나면 거기 남은 전문가가 그대로 디코드 캐시의 씨앗이 됩니다. 논문 표현으로는 "별도 프리필 캐시도, 단계 간 인수인계도 없다". 다만 대가가 있어서, --moe-cache-size2 × 전문가수보다 작으면 오버랩이 자동으로 꺼집니다.

설계 패턴 ② — 디코드: q* 대역폭 정합 분할

여기가 이 프로젝트의 심장입니다. 토큰을 하나씩 뱉는 디코드 단계에서, 캐시에 없는 전문가 m개를 만나면 세 갈래로 갈립니다.

이번 스텝의 라우팅 결과: 전문가 12개 필요 │ ├── 캐시에 있음 (7개) ────────────────→ GPU가 바로 계산 │ └── 캐시에 없음 = 미스 (5개) │ ├── q* 개 (3개) ── PCIe로 페치 ──→ GPU가 계산 │ ↑ ↑ │ └─ 동시 진행 ─┘ └── 나머지 (2개) ─── CPU가 호스트 RAM에서 그 자리에서 계산 ↓ 부분합을 더해서 최종 출력 q* = m × (PCIe 대역폭 / 호스트 메모리 대역폭) → 페치에 걸리는 시간과 CPU 계산에 걸리는 시간이 같아지는 지점 → 둘 중 늦게 끝나는 쪽(makespan)이 최소가 된다

이 계산을 파이썬이 아니라 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_q16B_P / B_H65536을 곱한 정수로 표현한 값입니다(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 캡처가 그대로 통과합니다.

용어
CUDA Graph
GPU에 보낼 명령들을 미리 한 덩어리로 녹화해 두고 재생하는 기능. 디코드는 같은 모양의 연산을 수천 번 반복하므로, 매번 커널을 하나씩 실행 요청하면 그 오버헤드만으로 상당한 시간이 나갑니다. 대신 제약이 큽니다 — 배열 크기가 매번 같아야 하고, 중간에 결과를 CPU로 가져와 판단할 수 없습니다. FreeToken의 설계 난이도 대부분이 여기서 나옵니다.

설계 패턴 ④ — CUDA Graph 안에서 CPU를 기다리는 두 가지 방법

GPU가 CPU 결과를 기다려야 하는데, 그래프 안에서는 일반적인 동기화를 쓸 수 없습니다. FreeToken은 두 가지를 구현해 두고 골라 씁니다.

방식동작비용
cudaLaunchHostFunc호스트 함수를 그래프의 노드로 삽입호출당 30~50μs. MoE 레이어당 2회 × 75레이어 = 스텝당 약 6ms
스트림 메모리 연산 기본값cuStreamWriteValue64 / WaitValue64로 플래그 핸드셰이크거의 0. 단 Windows WDDM · vGPU · 구형 드라이버에서 불가 → 기동 시 프로브 후 폴백

여기 딸린 각주가 하나 있는데, 이게 엣지 환경 최적화의 성격을 잘 보여줍니다. "GPU에서 spin-wait 커널을 돌려 기다리면 안 되냐"에 대한 코드 주석의 답은 — 그렇게 하면 GPU 사용률이 99%로 보고되고, 노트북의 전력 스케줄러가 그걸 보고 CPU 클럭을 낮춰서 결과적으로 더 느려진다는 것입니다. 데이터센터에서는 나오지 않을 종류의 제약입니다.

코드에 박제된 함정
cudaMemcpyBatchAsync의 조용한 성능 붕괴

여러 전문가를 한 번에 복사하려고 cudaMemcpyBatchAsync를 쓰는데, 배치 안에 256KB 미만 항목과 큰 항목이 섞이면 이 API가 말없이 동기 복사로 강등됩니다(H100 + CUDA 13.0에서 이분 탐색으로 확인). 실측 피해는 gpt-oss 2048토큰에서 전체 -22%. 그래서 코드에 _SMALL_BANK_FEAT_BYTES = 256 * 1024라는 상수와 함께 긴 주석이 남아 있습니다.

설계 패턴 ⑤ — 두 겹의 LRU

전문가 캐시는 LRU(가장 오래 안 쓴 것부터 버리기)로 관리되는데, 실은 LRU가 두 개입니다.

미스율 실측이 이 설계의 효과를 보여줍니다. 같은 라우팅 트레이스를 재생했을 때, RTX 5090이 담을 수 있는 용량(Qwen3.6 전문가 풀의 37%)에서:

배치 전략Qwen3.6 미스율DSV4-Flash 미스율
FreeToken 전역 LRU16%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로 거절합니다. 거절은 복구 가능하고, 반쯤 해제된 상태는 복구 불가능하니까요.

5디렉토리 구조 해부

어디에 무엇이 있고, 어디가 무거운가
FreeToken/ ├── .github/workflows/ 휠 빌드 CI 2개 (테스트 CI 없음) ├── benchmarks/ 벤치 스크립트 3개 ├── docs/ install · quickstart · models · cli ├── freetoken-kernel-cache/ 별도 휠 프로젝트 (미리 빌드된 .so) ├── scripts/ci/ manylinux 빌드 · 리태그 ├── python/freetoken/ ★ 본체 │ ├── kernel/ 16,045줄 / 67파일 ← 최대. Triton + CUDA/C++ + FLA │ │ ├── csrc/cpu_moe/ 2,150줄 CPU MoE 실행기 (C++) │ │ ├── csrc/gguf/ 5,733줄 llama.cpp ggml 커널 이식 │ │ ├── fla/ 13파일 flash-linear-attention 포팅 │ │ └── triton/ (+dsv4/) 44파일 자체 Triton 커널 │ ├── models/ 14,828줄 /110파일 아키텍처 13종 · 등록 문자열 22개 │ ├── server/ 10,876줄 / 20파일 ← function_call_parser.py 혼자 3,735줄 │ ├── moe/ 6,290줄 / 18파일 ★ 오프로드 코어 │ ├── kvcache/ 3,869줄 / 14파일 풀 + radix 캐시 3변종 │ ├── daemon/ 3,007줄 / 16파일 torch-free 슈퍼바이저 + systemd │ ├── attention/ 2,850줄 / 14파일 백엔드 7종 │ ├── scheduler/ 2,300줄 / 10파일 scheduler(905) · cache(667) │ ├── engine/ 1,933줄 / 6파일 engine.py(1,394) · cache_budget · graph │ ├── layers/ 1,623줄 / 9파일 MoELayer / OffloadMoELayer │ ├── shell/ 1,515줄 / 5파일 터미널 TUI │ ├── checkpoint/ 953줄 / 4파일 FTW 포맷 (ftw.py 577) │ └── tokenizer · utils · distributed · llm · message · benchmark └── tests/ 24,708줄 / 79파일 / 테스트 함수 854개 └── daemon · dsv4 · e2e · engine · kernels · models · moe · scheduler · server · tokenizer · kvcache/radix 전체: Python 434파일 97,345줄 + CUDA/C++ 23파일 12,419줄

이 트리에서 읽어 낼 것이 몇 가지 있습니다.

읽는 순서 추천

코드를 따라 읽을 거라면 이 순서가 가장 빠릅니다: moe/offload_kernels.py(q* 커널) → layers/moe.py(하이브리드 디코드) → moe/offload_cache.py(더블 버퍼 · 뱅크 스키마) → moe/cpu_executor.py(CPU 쪽) → engine/cache_budget.py(메모리 배분 정책). 이 5개 파일이면 논문의 주장 대부분이 코드로 확인됩니다.

6학습 포인트 (기술별)

이 저장소에서 실제로 가져갈 수 있는 것

추론 엔진을 만들 일이 없어도 배울 게 남는 저장소입니다. 여기 있는 기법 대부분은 "제약이 심한 환경에서 어떻게 설계하나"에 대한 답이라 범용성이 있습니다.

01 · GPU 프로그래밍

CUDA Graph와 호환되는 동적 자료구조 만들기

"배열 크기가 고정이어야 하고 중간 결과를 CPU로 못 가져온다"는 제약 아래서 LRU 캐시를 굴리는 방법 전체가 여기 있습니다. 히트 판정 · 희생자 선택 · 슬롯 재매핑 · 인덱스 덮어쓰기를 단일 블록 Triton 커널 하나에서 3단계로 처리합니다. sentinel(-1) + clamp_min + 0 가중치로 분기를 없애는 패턴은 다른 곳에도 그대로 쓸 수 있습니다.

볼 파일: moe/offload_kernels.py

02 · 스케줄링 이론 → 코드

이종 자원의 makespan 최소화를 커널 안에서 푸는 법

"두 개의 서로 다른 처리기에 일을 나눠 둘 다 동시에 끝내기"는 고전 스케줄링 문제입니다. FreeToken은 그 답 q* = m·B_P/B_H를 유도하고 → Q16 고정소수점으로 정수화하고 → 정수가 아닌 최적점 주변의 두 정수 중 느린 쪽이 더 작아지는 것을 고르는 데까지 갑니다. 이론식이 실제 GPU 커널 열 줄로 착지하는 과정을 통째로 볼 수 있는 흔치 않은 사례입니다.

볼 파일: moe/offload_kernels.py + moe/bench_profile.py

03 · 시스템 프로그래밍

비동기 복사의 이벤트 규율

전용 copy 스트림 + release_event / ready_event 한 쌍으로 더블 버퍼를 안전하게 돌리는 정석. "다음 사용자가 버퍼를 덮어쓰기 전에 이전 사용자가 다 읽었는지"를 이벤트로만 보장합니다. 스레드·락 없이 동시성을 다루는 방식이라 GPU 밖에서도 응용됩니다.

볼 파일: moe/offload_cache.py · layers/moe.py

04 · 파일 포맷 설계

O_DIRECT 친화 가중치 포맷 (FTW)

모든 텐서를 4096바이트에 정렬하고, 샤드 경계도 4096에 맞춰서, 텐서가 샤드를 걸쳐도 각 조각이 정렬을 유지하게 만든 포맷입니다. 여기에 preadv의 short read 루프 처리가 붙는데 — 이걸 빼먹으면 가중치 중간에 에러 없이 쓰레기 바이트가 들어갑니다. 파일 I/O에서 가장 무서운 종류의 버그입니다.

볼 파일: checkpoint/ftw.py

05 · 메모리 관리

pin-after-fill — 순서 하나로 47초를 버는 법

큰 호스트 버퍼를 준비할 때 ① lazy 익명 mmap 할당 → ② 실제 데이터로 채움 → ③ 그 다음에 cudaHostRegister로 pin. 순서를 ①③②로 바꾸면 pin이 전 페이지에 fault를 내고 0으로 채우는데, 그 0이 곧바로 데이터로 덮어써져 통째로 낭비됩니다. DSV4의 137GiB에서 약 47초가 여기서 갈립니다.

볼 파일: moe/host_banks.py

06 · 자료구조

다중 통화(multi-currency) radix prefix 캐시

KV 페이지와 GDN recurrent state라는 단위가 다른 두 자원을 같은 접두사 트리에 매달아 두고, 각자의 LRU로 축출하되 full_ref ≥ mamba_ref라는 불변식을 유지합니다. 내부 노드는 tombstone으로 처리. 트랜스포머와 상태공간 모델이 섞인 하이브리드 아키텍처를 지원하려면 반드시 마주치는 문제입니다.

볼 파일: kvcache/hybrid_radix_cache.py · swa_radix_cache.py

07 · 도메인 특화 최적화

"의미 경계 앵커" — 에이전트를 아는 캐싱

코딩 에이전트는 대화 중간을 계속 고쳐 씁니다(툴 결과를 끼워 넣고, thinking 블록을 지우고). 일반적인 접두사 캐시는 편집 지점 이후를 통째로 버리는데, FreeToken은 툴콜 오프너 토큰 위치에 상태 스냅샷을 ping-pong 슬롯으로 얼려 둡니다. 이 도메인 지식이 앞의 "에이전트 워크로드에서 12% 이내 유지"라는 수치를 만듭니다.

볼 파일: scheduler/cache.py

08 · 추상화 설계

양자화 포맷을 "뱅크 스키마"로 밀어내기

8종 양자화 포맷(bf16 · fp8_block · q4_0 · nvfp4 · mxfp4 · …)을 뱅크 이름 튜플로만 선언해 두면, 캐시 이동 기계는 뱅크가 몇 개인지 레이아웃이 어떤지 전혀 몰라도 동작합니다. 포맷을 아는 건 맨 마지막 GEMM 디스패치뿐입니다. 새 포맷 추가가 표에 한 줄 넣는 일이 됩니다.

볼 파일: moe/offload_cache.py

09 · 테스트 전략

GPU 커널의 CPU 레퍼런스 미러

축출·페치 결정을 순수 파이썬으로 한 번 더 구현해 두고 비트 단위로 대조합니다. GPU 커널은 디버깅이 어려우니 "정답"을 옆에 두는 방식입니다. tests/README.md의 한 줄이 이 저장소에서 가장 인용할 만합니다 — "기댓값을 무엇에 대조해 검사하는지 말할 수 없다면, 그 테스트는 자리값을 못 한다."

볼 파일: moe/offload_kernels.py · tests/README.md

10 · 운영 설계

크래시 격리를 문서가 아니라 테스트로 지키기

"제어 프로세스는 torch를 임포트하지 않는다"는 규칙을 import sentinel 테스트로 강제합니다. 아키텍처 문서에 적어 두는 것과 테스트가 막아 주는 것은 6개월 뒤에 전혀 다른 결과를 냅니다.

볼 파일: daemon/README.md · tests/daemon/test_daemon_import_safety.py

7하드웨어 / 시스템 요구사항

내 컴퓨터에서 돌아가는지 먼저 확인
항목요구사항
CPU 아키텍처x86_64 전용 (aarch64 · DGX Spark은 로드맵)
GPUNVIDIA Ampere 이상 — RTX 30 / 40 / 50 시리즈. 1080 · 2080 등은 미지원
드라이버r580 이상 (CUDA 13)
OSpip 설치는 Linux. 데스크탑 앱은 Windows + Linux. macOS 미지원
Python3.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 확인
가장 먼저 걸리는 벽
VRAM이 아니라 호스트 RAM

많은 사람이 "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미구현 (로드맵)
GGUFGemma-4에만 지원
v0.1.2의 현실
아직 크래시하는 열린 이슈들

공개 2주차 프로젝트답게 열린 이슈가 96개이고 그중 치명적인 것들이 있습니다 — #87 48k 토큰 초과 프리필이 서버 전체를 크래시(fla/l2norm 커널에서 device not ready), #110 프리필 워크스페이스의 CUDA OOM이 처리되지 않아 엔진이 죽는데 서버는 계속 요청을 받음, #111 KV 풀보다 긴 요청이 에러 없이 영원히 큐에 남음. 프로덕션 투입 전에 이슈 트래커를 한 번 훑는 편이 좋습니다.

8직접 해볼 수 있는 실습 과제

난이도별 5개 — GPU 없이 할 수 있는 것부터
과제 1난이도 ★☆☆☆☆ · GPU 불필요

q* 공식을 스프레드시트로 재현하기

커널 코드를 읽지 말고 공식만 보고 직접 계산해 보는 겁니다. 미스 m=20, PCIe 대역폭 32GB/s, 호스트 메모리 대역폭 80GB/s일 때 몇 개를 페치해야 할까요? 그리고 페치 개수를 0부터 20까지 바꿔 가며 max(페치시간, CPU시간)을 그래프로 그려 보세요. V자 골짜기의 바닥이 공식값과 일치하는지 확인합니다.

확인 포인트: 최적점이 정수가 아닐 때 양옆 중 어느 쪽이 나은지 — 그게 코드의 cost_lo vs cost_hi 비교입니다.

과제 2난이도 ★★☆☆☆

내 컴퓨터의 대역폭 비율 측정하기

ft bench bw가 하는 일을 직접 해 봅니다. ⓐ numpy로 큰 배열 복사를 반복해 호스트 메모리 대역폭을 재고, ⓑ GPU가 있다면 torch의 pinned 텐서 → GPU 복사로 PCIe 대역폭을 잽니다. 두 값의 비가 곧 fetch_frac입니다.

확인 포인트: 둘을 따로 잴 때동시에 돌리며 잴 때의 값이 다를 겁니다. FreeToken이 동시 측정값을 우선하는 이유를 여기서 직접 보게 됩니다.

과제 3난이도 ★★★☆☆ · GPU 불필요

LRU 캐시 시뮬레이터로 미스율 재현하기

가짜 라우팅 트레이스를 만듭니다 — 레이어 32개 × 전문가 64개, 토큰마다 각 레이어에서 6개를 (균등하지 않게, 지수 분포로) 뽑습니다. 그 트레이스에 대해 세 가지 배치 전략의 미스율을 재 보세요: ⓐ 전역 LRU ⓑ 처음 N개를 고정 배치 ⓒ 레이어 앞쪽 절반만 GPU. 캐시 용량을 전체의 20% / 40% / 60%로 바꿔 가며 표를 만듭니다.

확인 포인트: 논문의 16% vs 41% vs 62%와 비슷한 서열이 나오는지. 절댓값은 트레이스에 따라 다르지만 순서는 재현될 겁니다. 안 나온다면 라우팅 분포의 편향(전문가마다 인기 차이)을 더 세게 줘 보세요.

과제 4난이도 ★★★★☆ · GPU 필요

오버랩 노브를 껐다 켜며 A/B 측정하기

FreeToken은 측정 전용 환경변수를 프로덕션 코드에 남겨 뒀습니다. FREETOKEN_HYBRID_OVERLAP=0을 주면 CPU 풀을 PCIe 페치 전에 동기화해서 두 경로를 직렬화합니다. 같은 프롬프트로 켠 상태 / 끈 상태의 tok/s를 재면 겹치기가 만든 이득만 분리됩니다.

같은 방식으로 FREETOKEN_CPU_MOE_FLAG_SYNC=0(스트림 메모리 연산 대신 host-func)도 재 보세요. 레이어 수가 많은 모델일수록 차이가 커야 합니다.

확인 포인트: 자기 하드웨어에서의 이득이 논문 수치와 다를 수 있습니다. 그게 정상입니다 — q*는 머신마다 다른 값이니까요. 오히려 왜 다른지를 대역폭 비로 설명할 수 있으면 성공입니다.

과제 5난이도 ★★★★★ · GPU 필요

새 어텐션 백엔드를 레지스트리에 등록하기

attention/__init__.pyBackendInfo 구조를 읽고, 순수 Triton 백엔드(triton)를 베껴 일부러 느린 백엔드를 하나 추가해 보세요. 목표는 성능이 아니라 확장 지점의 계약을 이해하는 것입니다 — 어떤 능력(FULL / SWA / MLA)을 선언해야 하는지, 페이지 크기 제약이 어디서 검증되는지, --attention-backend prefill,decode로 합성될 때 무엇이 요구되는지.

다음 단계: 여기까지 왔으면 moe/__init__.py의 MoE 백엔드 레지스트리도 같은 구조입니다. --moe-backend의 auto 선택 규칙(engine/engine.py)을 읽어 보면, 왜 auto가 절대 fused를 고르지 않는지 알게 됩니다 — 잘못 고르면 느려지는 게 아니라 가중치 로딩 OOM이라서입니다.

9관련 기술 심화 학습 로드맵

6주 코스 — 이 저장소를 읽을 수 있는 상태까지
주차주제구체적으로 할 것
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의 선택이 왜 다른지가 선명해집니다.

10핵심 키워드 사전

이 문서와 저장소에 반복해서 나오는 말들
용어
MoEMixture of Experts. 레이어를 여러 전문가로 쪼개고 토큰마다 일부만 통과시키는 구조
활성 파라미터토큰 하나를 처리할 때 실제로 계산에 참여하는 파라미터. 284B-A13B의 뒤쪽 숫자
라우터 (게이트)토큰마다 어느 전문가로 보낼지 고르는 작은 신경망. 이 선택이 미리 예측 불가한 게 문제의 출발점
프리필 (prefill)프롬프트 전체를 한 번에 읽어 KV 캐시를 채우는 단계. 계산량이 많고 병렬성이 큼
디코드 (decode)토큰을 하나씩 생성하는 단계. 계산량은 적고 메모리 대역폭에 묶임
TTFTTime To First Token. 첫 글자가 나오기까지의 지연. 프리필 성능의 지표
KV 캐시이미 처리한 토큰들의 어텐션 중간값 저장소. 안 두면 매 토큰마다 전체를 다시 계산해야 함
오프로딩GPU에 안 들어가는 가중치를 호스트 RAM·디스크에 두고 필요할 때 옮기는 기법
pinned memoryOS가 스왑하지 못하게 고정한 호스트 메모리. GPU가 DMA로 직접 읽을 수 있어 전송이 빠름
PCIeGPU와 메인보드를 잇는 통로. 5.0 x16의 실측 상한이 약 52.7GB/s — 이 숫자가 오프로딩 성능의 천장
q* (q-star)이 프로젝트의 핵심 개념. 미스 중 PCIe로 가져올 개수. q* = m × B_P / B_H
makespan병렬로 돌린 작업들 중 가장 늦게 끝나는 것의 완료 시각. q*가 최소화하는 대상
CUDA GraphGPU 명령 묶음을 녹화해 재생하는 기능. 오버헤드를 줄이지만 동적 shape 금지
TritonGPU 커널을 파이썬 문법으로 작성하는 언어이자 컴파일러
LRULeast Recently Used. 가장 오래 안 쓴 항목부터 버리는 캐시 축출 정책
radix prefix 캐시대화 접두사를 트리로 공유해 재계산을 피하는 KV 캐시 구조
SWASliding Window Attention. 최근 N개 토큰만 보는 어텐션. KV 캐시가 무한히 안 커짐
GDNGated DeltaNet. 선형 어텐션 계열의 recurrent 상태. KV 페이지와 단위가 달라 별도 관리 필요
NVFP4 / MXFP44비트 부동소수점 양자화 포맷. 가중치 크기를 크게 줄여 작은 RAM에서 큰 모델을 가능하게 함
FTWFreeToken Weights. 4096바이트 정렬로 O_DIRECT 읽기에 최적화한 자체 체크포인트 포맷
AVX-512-BF16CPU의 bfloat16 곱셈누산 명령어. 이게 있어야 CPU MoE 실행이 쓸 만한 속도가 나옴
sentinel 값"해당 없음"을 뜻하는 특수 값(여기서는 -1). 배열 크기를 고정한 채 분기를 표현하는 수단

11참고 링크

원본과 이어서 볼 것들

FreeToken 자체

비교해 볼 만한 저장소

배경 지식