Llama 3 8B 가 맥북에서 토큰 당 50ms 로 도는 시점이 왔다. NVIDIA GPU 가 아니라 ARM CPU 위에서. Scott Roy 가 ExecuTorch / torchao 의 일부로 짜고 있는 1~8 bit ARM kernels 의 내부 — NEON · SVE 의 활용, K=8 / 16 / 32 group quant 의 packing 이 register 와 cache 위에서 어떻게 펴지는가, 그리고 “같은 weight-only quant 인데 NVIDIA 와 모바일은 왜 모양이 다른가” 의 답.
데스크탑 / 폰 / 임베디드 — NVIDIA GPU 가 없는 환경. 이 자리에서 LLM 추론을 돌리려면 거의 모든 결정이 다르게 내려진다. CPU 위주 compute, 적은 메모리, 강력한 cache hierarchy, 저전력 제약. Scott 의 일은 그 자리에 맞는 low-bit kernel set 을 짜는 것.
강의의 첫 측정 — Llama 3.1 8B 가 M1 Pro 맥북에서 토큰 당 약 50ms. 4-bit weight + INT8 activation. 이게 llama.cpp 의 길을 따라가지 않고 PyTorch 생태계 안에서 (torchao + ExecuTorch) 직접 짠 결과. macOS / Linux / iOS 의 같은 stack 으로 동작.
“Triton/CUDA 의 mental model 이 그대로 옮겨오지 않는다” — 이게 강의의 출발. ARM 의 vector ISA(NEON/SVE)는 GPU SIMT 와 다르다. tile 의 단위, register 의 양, cache 의 layer, 병렬화의 방향 이 통째로 다르다. 그래서 같은 “4-bit weight-only GEMM” 라도 NVIDIA 위 코드와 ARM 위 코드는 닮은 점이 거의 없다.
강의 끝에 손에 잡혀야 할 자산 — (1) ARM CPU 의 compute 단위(NEON / SVE) 와 GPU SM 의 직관적 매핑, (2) low-bit kernel 의 “register-level tiling + cache-level tiling” 이라는 두 단계 구조, (3) ExecuTorch 의 custom op 가 어떻게 모델 안에 들어가는지, (4) 같은 1~8 bit 범위에서 정확도와 latency 의 trade-off 가 GPU 와 어떻게 다른지.
데스크탑/모바일 환경에서 LLM 의 자연스러운 compute target 이 CPU 인 이유는 단순하다 — 가장 통합된, 가장 큰, 가장 호환되는 unit 이기 때문.
강의에서 Scott 이 명시한 한 줄 — “CPU 가 가장 portable 하고 가장 통합된 target. 모바일 GPU/NPU 는 좋지만 fragmentation 이 심해서 한 stack 으로 모든 device 를 cover 하기 어렵다.” 그래서 ExecuTorch 는 CPU backend 를 main 으로 두고, GPU/NPU 는 가능한 device 에서 추가로 넘긴다.
Apple Silicon, 최근 ARM SoC 들은 CPU 와 GPU 가 같은 메모리를 본다. NVIDIA 처럼 PCIe 위 분리된 HBM 이 아니다. 그래서 “memcpy H→D” 같은 비용이 없거나 낮다. 그 대신 memory bandwidth 가 절대량으로 작다 — M1 Pro 가 ~200GB/s, A100 이 ~2TB/s 로 10배 차이. 이 차이가 모든 후속 결정을 끌고 간다.
NVIDIA GPU 는 SIMT (한 thread 가 한 element, 32 thread = warp). ARM 은 SIMD (한 thread 가 여러 element 동시). 이 차이가 mental model 의 첫 번째 갈림길.
SMMLA instruction 은 두 8×16 int8 matrix 를 곱해 4×4 int32 acc 에 누적. 한 줄로 16개 mac 연산. NVIDIA 의 mma.sync 와 비슷한 자리지만 단위가 더 작고 thread 단위에서 매번 발행. 이 instruction 이 ARM low-bit GEMM 의 “tensor core”.
강의에서 Scott 이 짚은 사실 — FMLA (vector FMA) 는 fp16/fp32 lane 별 곱셈+누적이고, SMMLA 는 int8 matrix 곱. low-bit kernel 의 inner loop 는 거의 SMMLA. dequant 결과를 int8 형태로 쥐고 있다가 SMMLA 로 던진다.
decode 는 batch=1 → activation 이 vector 한 줄, weight 가 matrix. GEMV (matrix-vector). prefill 만 GEMM. 모바일 LLM 의 latency 는 거의 GEMV 가 결정.
강의에서 Scott 이 명시한 — “mcal=1, ncal=8, kcal=16 가 GEMV 의 표준 tile.” M=1 인 이유는 decode batch=1, N=8 은 NEON int32 acc 4 개 × 2 SMMLA, K=16 은 SMMLA 한 발행이 받는 단위.
prefill 은 M 이 큰 GEMM. tile 이 M=8 / N=8 / K=16 같은 직사각형 형태. 코드는 비슷한데 outer loop 가 추가. ARM low-bit kernel 보통 GEMV / GEMM 두 본을 따로 둔다 — Scott 이 강의에서 “24개 GEMV kernel + 별도 GEMM kernel” 이라고 표현한 자리.
CPU kernel 의 tiling 은 GPU 와 다르게 명확한 hierarchy 가 있다. register tile (가장 안쪽), L1 cache tile (중간), thread tile (가장 바깥). Scott 의 ARM kernel 도 같은 구조.
이 hierarchy 가 GPU 의 “grid → block → warp → thread” 와 닮아 보이지만 의미가 다르다 — CPU 에서는 cache layer 의 명시성 이 강하다. L1 / L2 / L3 의 크기가 분명하고, working set 이 cache 안에 들어가는지 안 들어가는지가 latency 를 직접 결정.
강의에서 Scott 이 명시한 한 줄 — “우리는 tile size 를 cache size 에 맞춰 정밀하게 tuning 하지 않았다. 가장 큰 효과는 register tile 의 SIMD 활용에서 왔다.” L1 의 fragmentation 보다 register 의 SMMLA 활용이 우선. 그게 의외로 모바일 LLM 의 sweet spot.
weight 는 보통 한 번 load 되면 그 thread 동안 여러 번 reuse. activation 은 한 번 load + 한 번 use. 그래서 weight 의 cache 친화성 이 더 결정적. weight 를 미리 packing 단계에서 cache-friendly layout 으로 적어둔다.
L034 의 GPU 와 마찬가지로, ARM low-bit 도 핵심은 group-wise scale. 다만 group_size 의 단위가 더 작다 (32~128). 그리고 ARM 쪽은 zero point (asymmetric quant) 도 더 자주 쓴다.
Scott 의 라이브러리 — 1, 2, 3, 4, 5, 6, 7, 8 bit 모두 대응. 각 bit-width 마다 별도 unpacking routine. 3-bit / 5-bit / 7-bit 처럼 byte alignment 가 안 맞는 bit-width 도 다룰 수 있게 byte-stream packing. 8 개 unpacking routine × N tile size = 24개의 inner kernel.
강의에서 흥미로운 사실 — 3-bit 가 4-bit 와 거의 같은 정확도를 주는 케이스가 의외로 많다. weight 의 effective bits 를 1 bit 줄이는 게 모바일 메모리 절약에 큰 의미. 다만 unpacking 코드가 4-bit 보다 복잡해서 latency 는 비슷하거나 약간 더 느림.
kernel 이 빠른 것만으로는 의미가 없다 — 실제 모델 안에 들어가야. 강의 후반부는 ExecuTorch 의 custom op pipeline 위에 어떻게 이 kernel 들이 박히는지를 깐다.
핵심은 두 가지 — (a) kernel 이 일반 C++ 함수, 즉 추가 컴파일러나 DSL 안 필요. torch.ops.{namespace} 로 등록되며 PyTorch / ExecuTorch 모두에서 호출 가능. (b) 병렬화는 pthreadpool (XNNPACK 의 thread pool wrapper) 를 그대로 빌려 씀. ExecuTorch / PyTorch / XNNPACK 가 같은 thread pool 을 공유.
// torchao 의 ARM kernel — C++ 그대로
struct UKernelConfig {
int mcal, ncal, kcal; // register tile
int n_tiles_per_thread;
KernelFn kernel; // 함수 포인터
UnpackFn unpack; // 1~8 bit 별
};
// 24 개 GEMV kernel + 8 개 unpacking routine
// 의 lookup table — bitwidth × tile
extern UKernelConfig kernels[8][3];
이 layout 의 의미 — kernel 의 모양이 (bitwidth, tile size, asymmetric 여부) 의 cartesian 곱. lookup table 한 개로 모든 dispatch.
모델 layer (Linear) 한 개를 quantize → packing → custom op 호출 의 시퀀스로 바꾼다. PyTorch 의 graph mode (torch.export) 가 이 op 을 그대로 받아들여 ExecuTorch binary 로 export.
강의 Q&A 에서 — “모바일 SoC 의 P-core / E-core 비대칭성을 parallel_1d 가 handle 하는가?” 답은 “아니, parallel_1d 는 그냥 N 등분.” 더 정교한 work-stealing 은 향후 작업. 현재는 균등 분할이 충분히 좋게 나오는 일이 많음 — weight tile 자체가 동일하므로.
강의에서 보여준 측정 — 같은 모델 (Llama 3.1 8B), 같은 quant (4-bit weight + dynamic INT8 act, group=32), 같은 stack (torchao + ExecuTorch + ARM kernels) 의 device 별 latency.
두 가지 사실. 첫째 — 맥북에서도 사용 가능한 속도가 나온다. 20 tok/s 면 사람이 읽기 따라잡는 속도. 둘째 — iPhone 은 메모리가 hard limit. 8GB RAM 의 디바이스에서 5GB 모델은 OS 와의 경쟁에서 swapping 위험.
3B 클래스 모델 (Llama 3.2 3B) 으로 가면 iPhone 에서도 20+ tok/s 가 나오는 영역. 모바일 LLM 의 sweet spot 이 1B~4B 인 이유 — 메모리 limit + latency 의 두 제약이 합쳐서.
표가 보여주는 것 — 두 hardware 는 완전히 다른 trade-off space 위에 있다. cloud GPU 는 throughput 게임(batching, prefix sharing, 대량 동시 요청), edge 는 latency / privacy / offline 게임(batch=1, 모델 작아야 함). 같은 quantization 이라도 hardware 의 다른 제약이 코드를 다르게 펴게 만든다.
Apple Silicon 의 unified memory + 강력한 NPU + 점점 빨라지는 GPU. 모바일 SoC 가 작은 LLM 을 GPU 에서 크게, 큰 LLM 을 CPU 에서 일부 — 같은 device 안에서 layer 별 다른 unit 으로 dispatch 하는 방향이 자연스러움. 이미 ExecuTorch 가 이 일을 부분적으로 함.
vfmaq_f32, vmmlaq_s32. C++ 한 함수로 vectorized fp32 dot product. 결과를 일반 코드와 비교.torch.export + ExecuTorch binary. 그 binary 를 iPhone 또는 raspberry pi 로 옮겨 동작 확인.이 노트의 latency 수치는 강의 시점의 measurement 를 재구성한 예시. ARM hardware 가 빠르게 갱신되고 있고 (M3, M4, A18 …) torchao 의 kernel 도 진행형. 자기 device 에서 직접 측정해야 baseline.