gpumode · 강의 아카이브
《GPU Mode》 L038 2024 · NOV High priority transcript · available

Low Bit ARM kernels

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 와 모바일은 왜 모양이 다른가” 의 답.

ARM NEON SVE low-bit ExecuTorch torchao tile unpacking
S
Speaker
Scott Roy
Meta · github.com/metascroy · ExecuTorch / torchao
강의 번호
L038
스피커
Scott Roy
학습 우선순위
High · 정독
다시 볼 때
macOS · iPhone 측정
§ 01강의가 풀려는 문제· on-device LLM

“데이터센터 GPU 가 아닌 device 에서 LLM 을 돌릴 수 있는가” 의 답을 ARM 위에서

데스크탑 / 폰 / 임베디드 — 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 으로 동작.

강의의 인지적 frame

“Triton/CUDA 의 mental model 이 그대로 옮겨오지 않는다” — 이게 강의의 출발. ARM 의 vector ISA(NEON/SVE)는 GPU SIMT 와 다르다. tile 의 단위, register 의 양, cache 의 layer, 병렬화의 방향 이 통째로 다르다. 그래서 같은 “4-bit weight-only GEMM” 라도 NVIDIA 위 코드와 ARM 위 코드는 닮은 점이 거의 없다.

“같은 알고리즘이지만 hardware 가 너무 달라서 — kernel 코드는 거의 다른 우주에서 살게 된다.”학습 노트

강의 끝에 손에 잡혀야 할 자산 — (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 와 어떻게 다른지.

§ 02모바일 / edge 의 compute 모델· CPU 위주, GPU 부수

왜 모바일에서 GPU 가 “부수적” 인가

데스크탑/모바일 환경에서 LLM 의 자연스러운 compute target 이 CPU 인 이유는 단순하다 — 가장 통합된, 가장 큰, 가장 호환되는 unit 이기 때문.

CPU (ARM)
표준 통합 unit
모든 device 에 있고 메모리 일관 (unified memory). NEON/SVE 가 vector compute. LLM weight 를 그대로 적재하기 가장 쉽다.
GPU (Adreno · Mali · Apple)
크지만 dispatch 비용 큼
제조사마다 다른 ISA. compute API 도 Metal/Vulkan/OpenCL 분열. GPU/CPU 메모리 공유 양상도 device 마다 다름.
NPU / DSP
전용 가속기
전력 효율 best 이지만 SDK 별 호환성 모자람. 모델 형태가 NPU 가 받아들일 수 있어야 하고, 일반 LLM 은 아직 잘 안 맞는 경우 많음.

강의에서 Scott 이 명시한 한 줄 — “CPU 가 가장 portable 하고 가장 통합된 target. 모바일 GPU/NPU 는 좋지만 fragmentation 이 심해서 한 stack 으로 모든 device 를 cover 하기 어렵다.” 그래서 ExecuTorch 는 CPU backend 를 main 으로 두고, GPU/NPU 는 가능한 device 에서 추가로 넘긴다.

unified memory 의 의미

Apple Silicon, 최근 ARM SoC 들은 CPU 와 GPU 가 같은 메모리를 본다. NVIDIA 처럼 PCIe 위 분리된 HBM 이 아니다. 그래서 “memcpy H→D” 같은 비용이 없거나 낮다. 그 대신 memory bandwidth 가 절대량으로 작다 — M1 Pro 가 ~200GB/s, A100 이 ~2TB/s 로 10배 차이. 이 차이가 모든 후속 결정을 끌고 간다.

§ 03ARM NEON / SVE· 128-bit · 가변 vector

SIMT 가 아니라 SIMD — 한 instruction 이 여러 lane 을 동시에 친다

NVIDIA GPU 는 SIMT (한 thread 가 한 element, 32 thread = warp). ARM 은 SIMD (한 thread 가 여러 element 동시). 이 차이가 mental model 의 첫 번째 갈림길.

smmla 가 왜 결정적인가

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 로 던진다.

“NVIDIA 는 thread 가 element 한 개. ARM 은 thread 가 vector 한 개. 그 차이가 코드 구조 전체를 다르게 만든다.”학습 노트
§ 04INT8 GEMM / GEMV 패턴· smmla / sdot

decode 의 batch=1 GEMV 가 사실상 모바일 LLM 의 hot path

decode 는 batch=1 → activation 이 vector 한 줄, weight 가 matrix. GEMV (matrix-vector). prefill 만 GEMM. 모바일 LLM 의 latency 는 거의 GEMV 가 결정.

FIG · GEMV tile 한 개 · M=1 (decode)NEON · BLOCK_N=8, BLOCK_K=16
a0
a1
a2
a3
a4
a5
a6
a7
w00
w01
w02
w03
w04
w05
w06
w07
w10
w11
w12
w13
w14
w15
w16
w17
acc0
acc1
acc2
acc3
M=1 의 decode GEMV. activation row 한 개 + weight tile 8×16 → output 4 개. SMMLA 한 두 발행으로 inner step. tile 크기 = 8 이 NEON 128-bit / int8 16-lane 의 자연스러운 단위.

강의에서 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 의 다른 tile

prefill 은 M 이 큰 GEMM. tile 이 M=8 / N=8 / K=16 같은 직사각형 형태. 코드는 비슷한데 outer loop 가 추가. ARM low-bit kernel 보통 GEMV / GEMM 두 본을 따로 둔다 — Scott 이 강의에서 “24개 GEMV kernel + 별도 GEMM kernel” 이라고 표현한 자리.

§ 05cache-friendly tiling· L1 · register

register tile / L1 tile / threading tile — 세 단계 분할

CPU kernel 의 tiling 은 GPU 와 다르게 명확한 hierarchy 가 있다. register tile (가장 안쪽), L1 cache tile (중간), thread tile (가장 바깥). Scott 의 ARM kernel 도 같은 구조.

L0 · 표준 layer linear (M × K weight, K × 1 activation)모델 안 한 layer 의 일 PyTorch op
L1 · thread tile N 차원 분할 → 각 thread 가 자기 row blockpthreadpool 로 dispatch. parallel 1D N/cores 행
L2 · L1 cache tile N=수십, K=수백 단위 sub-blockweight 와 activation 의 working set 이 L1 (~64KB) 안에 메모리 밴드 절약
L3 · register tile M=1 / N=8 / K=16 inner kernelSMMLA 와 dequant 가 register 위에서 vectorized inner

이 hierarchy 가 GPU 의 “grid → block → warp → thread” 와 닮아 보이지만 의미가 다르다 — CPU 에서는 cache layer 의 명시성 이 강하다. L1 / L2 / L3 의 크기가 분명하고, working set 이 cache 안에 들어가는지 안 들어가는지가 latency 를 직접 결정.

tile 크기 결정의 휴리스틱

강의에서 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 으로 적어둔다.

§ 06정확도 처리· group · zero point

weight 1~8 bit 모두 같은 패턴 — group + scale + (옵션) zero point

L034 의 GPU 와 마찬가지로, ARM low-bit 도 핵심은 group-wise scale. 다만 group_size 의 단위가 더 작다 (32~128). 그리고 ARM 쪽은 zero point (asymmetric quant) 도 더 자주 쓴다.

unpacking routine 의 범위

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 는 비슷하거나 약간 더 느림.

§ 07ExecuTorch 통합· custom op · pthreadpool

kernel 한 본을 PyTorch 모델 안의 한 op 로 만드는 길

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.

  • kernel — register tile 의 SMMLA loop. asm intrinsic 직접.
  • unpack — packed weight 를 int8 register 로 펴는 함수. bitwidth 별.
  • parallel_1d — pthreadpool 위에서 N 차원 thread 분할.

모델 layer (Linear) 한 개를 quantize → packing → custom op 호출 의 시퀀스로 바꾼다. PyTorch 의 graph mode (torch.export) 가 이 op 을 그대로 받아들여 ExecuTorch binary 로 export.

heterogeneous core 처리는 단순

강의 Q&A 에서 — “모바일 SoC 의 P-core / E-core 비대칭성을 parallel_1d 가 handle 하는가?” 답은 “아니, parallel_1d 는 그냥 N 등분.” 더 정교한 work-stealing 은 향후 작업. 현재는 균등 분할이 충분히 좋게 나오는 일이 많음 — weight tile 자체가 동일하므로.

§ 08LLM 모바일 사례· M1 Pro · iPhone

같은 stack 으로 맥북 · iPhone · Linux 위 LLM

강의에서 보여준 측정 — 같은 모델 (Llama 3.1 8B), 같은 quant (4-bit weight + dynamic INT8 act, group=32), 같은 stack (torchao + ExecuTorch + ARM kernels) 의 device 별 latency.

device · contexttok/s decodememory · model참고
M1 Pro · macOS~20 tok/s~5 GBCPU only
M2 Ultra · macOS~30 tok/s~5 GBCPU only
iPhone 15 Pro~6 tok/s~5 GB8GB RAM 한계 근처
Raspberry Pi 5~1 tok/s~5 GB참고용
A100 (참고)~120 tok/s~5 GBbatch=1

두 가지 사실. 첫째 — 맥북에서도 사용 가능한 속도가 나온다. 20 tok/s 면 사람이 읽기 따라잡는 속도. 둘째 — iPhone 은 메모리가 hard limit. 8GB RAM 의 디바이스에서 5GB 모델은 OS 와의 경쟁에서 swapping 위험.

“같은 stack 한 본으로 맥북 · iPhone · Linux 위 LLM 추론. 그게 이 프로젝트의 한 줄 결과.”Scott Roy · 강의 후반

3B 클래스 모델 (Llama 3.2 3B) 으로 가면 iPhone 에서도 20+ tok/s 가 나오는 영역. 모바일 LLM 의 sweet spot 이 1B~4B 인 이유 — 메모리 limit + latency 의 두 제약이 합쳐서.

§ 09NVIDIA 와 trade-off· cloud vs edge

같은 알고리즘이 두 hardware 에서 “같은 일을 하지 않는다”

NVIDIA H100ARM (M1/iPhone)의미
memory bw~3 TB/s~200 GB/s15× 차
tile core 단위WGMMA (warp group)SMMLA (vector)단위 100× 차
parallel 단위SM × SIMTcore × SIMDcores 8~16개
memory model분리 HBMunifiedmemcpy free
전력~700W~10Wedge 전용
target latencythroughputsingle-user latencybatch=1 만

표가 보여주는 것 — 두 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 가 이 일을 부분적으로 함.

§ 10기억할 메모와 코드· torchao · llama.cpp

다시 열었을 때 5분 안에 손에 잡혀야 할 것

CPU 의 자리
모바일/edge 의 “표준 통합 unit”. unified memory + 모든 device 에 존재.
NEON / SVE
128-bit SIMD / 가변 폭. SMMLA 가 int8 matrix mac instruction.
register tile
M=1 N=8 K=16 — GEMV 의 표준. NEON 128-bit + int8 16-lane 의 자연스러운 단위.
cache-friendly tiling
register / L1 / thread 세 단계. weight reuse 가 cache 친화성의 중심.
1~8 bit 모두
scale + (옵션) zero point + group_size = 32 표준. 8개 unpacking routine.
ExecuTorch 통합
C++ kernel 이 그대로 PyTorch custom op. pthreadpool 로 병렬.
모바일 sweet spot
1B~4B 모델. 8B 도 맥북은 가능. iPhone 은 메모리 hard limit.
GPU 와 trade-off
batch=1 latency 게임. throughput 안 따짐. unified memory · 적은 bandwidth.
Code · ExecuTorchpytorch/executorch
Referencellama.cpp — 모바일 LLM 추론의 사실상 표준 비교

손에 새기기 — 실습 시퀀스

  1. NEON intrinsic 한 줄vfmaq_f32, vmmlaq_s32. C++ 한 함수로 vectorized fp32 dot product. 결과를 일반 코드와 비교.
  2. SMMLA loop — int8 weight + int8 act 의 8×16 tile 한 발행. 결과를 reference 와 비교. 이게 ARM low-bit 의 “tensor core hello world”.
  3. register tile 만 단독 — pthreadpool 없이 single thread 에서 GEMV inner kernel. M1 Pro 에서 ~10 tok/s 는 single thread 에서 이미 가능.
  4. parallel_1d 추가 — pthreadpool 띄우고 N 차원 분할. core 수만큼 throughput 곱 — 8 core 면 ~80 tok/s 까지 toy GEMV 에서.
  5. torchao 의 quantize → packing — Llama 3 한 layer 를 quantize 해서 ARM kernel 에 넘긴다. round-trip 으로 정확도 검증.
  6. ExecuTorch exporttorch.export + ExecuTorch binary. 그 binary 를 iPhone 또는 raspberry pi 로 옮겨 동작 확인.
  7. llama.cpp 와 같은 모델 비교 — 같은 weight, 같은 quant 에 대해 두 stack 의 latency. ARM kernel 이 어디서 더 빠른지/느린지.
§ 11다른 강의로 이어지는 길· connections

같은 도구가 시리즈 안에서 어디에 다시 등장하는가

§ 12열린 질문· open questions

다음에 다시 들었을 때 직접 검증해야 할 것들

검증 메모

이 노트의 latency 수치는 강의 시점의 measurement 를 재구성한 예시. ARM hardware 가 빠르게 갱신되고 있고 (M3, M4, A18 …) torchao 의 kernel 도 진행형. 자기 device 에서 직접 측정해야 baseline.

← Lecture 037 SASS & uarch Lecture 039 → Torchtitan — Mark · Tianyu 가 깐 PyTorch 분산 학습 스택