Low Bit Metal Kernels

2024–2025 의 큰 trend 중 하나 — on-device LLM. 데이터 보안, latency, network 의존성 제거. 그러나 7–8B 모델이 iPhone 16 Pro (8 GB unified memory) 안에 들어가려면 — quantization 이 필수. 그리고 Apple Silicon 의 GPU 는 NVIDIA 와 다른 architecture — Metal 위에 직접 짜야 한다.

이 강의의 자리 — NVIDIA 시리즈의 코어 기법(low bit, dequant fusion, tile 최적화) 가 Apple Silicon 위에서 어떻게 다시 깨지는가. 같은 idea, 다른 hardware, 다른 도구.

강의 transcript 가 누락되었지만 ExecuTorch (Meta 의 on-device 추론 framework) + Metal Performance Shaders + torchao 가 모두 open source. 본 노트는 이 자료들의 공개 정보 + Apple Silicon GPU architecture 도메인 지식 위에서 정리한다.

Apple Silicon GPU (M-series) 의 architecture 는 NVIDIA 와 다르다. 알아야 할 자리들.

Apple 의 GPU 프로그래밍은 Metal. compute shader 를 Metal Shading Language (MSL) — C++ 기반 — 로 작성. 추상은 CUDA 와 거의 일대일 대응.

// Metal MSL — INT4 dequant + matmul (개념)
#include <metal_stdlib>
using namespace metal;

kernel void int4_matmul(
    device const uint8_t* W_packed [[buffer(0)]],   // INT4 packed
    device const half*    scales   [[buffer(1)]],   // per-group scale
    device const half*    X         [[buffer(2)]],   // activation FP16
    device       half*    Y         [[buffer(3)]],
    constant uint& M [[buffer(4)]],
    constant uint& N [[buffer(5)]],
    constant uint& K [[buffer(6)]],
    uint2 gid [[threadgroup_position_in_grid]],
    uint2 tid [[thread_position_in_threadgroup]],
    uint  sid [[simdgroup_index_in_threadgroup]])
{
    threadgroup half X_tile[BLOCK_M][BLOCK_K];
    threadgroup half W_tile[BLOCK_K][BLOCK_N];

    half acc = 0;
    for (uint k = 0; k < K; k += BLOCK_K) {
        // 1) X 를 threadgroup memory 로
        X_tile[tid.y][tid.x] = X[...];

        // 2) INT4 dequant: 1 byte = 2 nibble
        uint8_t packed = W_packed[...];
        half low  = ((int)(packed & 0x0F) - 8) * scales[group];
        half high = ((int)(packed >> 4)   - 8) * scales[group];
        W_tile[tid.y * 2    ][tid.x] = low;
        W_tile[tid.y * 2 + 1][tid.x] = high;

        threadgroup_barrier(mem_flags::mem_threadgroup);

        // 3) matmul tile
        for (uint kk = 0; kk < BLOCK_K; ++kk) {
            acc += X_tile[tid.y][kk] * W_tile[kk][tid.x];
        }
        threadgroup_barrier(mem_flags::mem_threadgroup);
    }

    Y[...] = acc;
}

low-bit quantization 의 핵심 — weight 가 INT4 로 packed 되어 있다. 한 byte 안에 두 weight. 사용 직전에 dequant.

Apple Silicon 의 가장 큰 architecture 차이 — CPU 와 GPU 가 같은 DRAM. cudaMemcpy 같은 H↔D copy 가 없다. 같은 buffer 의 같은 메모리를 두 processor 가 본다.

이 차이가 만드는 디자인 결정 — weight loading 이 거의 무료. file → mmap → MTLBuffer 한 줄로 GPU 가 직접 read 가능. 7B INT4 = 3.5 GB 의 weight 를 disk → memory → GPU 의 단계 없이 한 번에. iPhone 의 cold start latency 의 핵심.

ExecuTorch 는 Meta 가 만든 PyTorch 의 on-device 추론 framework. PyTorch 모델을 .pte (PyTorch Edge) 파일로 export 후 mobile / embedded 에서 실행.

이 단계 중 backend partition 가 본 강의의 자리. weight matmul 같은 op 가 Metal backend 로 delegated. 그 Metal backend 가 INT4 dequant + matmul 의 custom Metal 커널 호출. partition 결정이 latency 를 결정 — 일부 op 만 Metal 가면 CPU↔GPU dispatch overhead 증가.

Meta + Apple 의 협업으로 demonstrate 된 자리 — Llama 3 8B 가 iPhone 16 Pro 에서 실시간 token 생성. 10–15 tokens/sec 수준. 같은 hardware, 다른 implementation 으로는 그 절반이 일반적.

NVIDIA GPU 위 INT4 LLM kernel 의 표준 구조 (예: GPTQ, AWQ + tinygemm) 와 Apple Silicon 위 같은 일의 구현이 어디서 갈리는지.

강의가 솔직하게 다뤘을 자리. on-device LLM 이 모든 use case 에 적합하지 않다.

손에 새기기 — 실습 시퀀스

1. Metal Hello kernel — Xcode 또는 metalcpp 로 가장 단순한 vector_add 커널. CUDA 와의 추상 1:1 대응 손에 잡기.
2. FP16 matmul — Metal MSL 로 simple matmul. threadgroup memory tiling. CUDA 의 tiled matmul 직역.
3. INT4 dequant 단독 커널 — packed byte → 2 FP16 의 dequant 만. CPU 와 정확성 비교.
4. INT4 dequant + matmul fusion — 위 두 단계 합치기. inner loop 안 dequant.
5. ExecuTorch 의 INT4 backend 코드 읽기 — executorch/backends/apple/metal 의 실제 구현. 본 노트의 추상이 어떻게 코드화 되어 있는지.
6. iPhone 위 Llama 빌드 — ExecuTorch 의 example 따라가기. .pte 만들고 iOS app 에서 호출.
7. throughput 측정 — token/sec, memory bandwidth 사용률, thermal throttle 시작 시점.

transcript 가 없어 강의의 정확한 강조점은 추정. 영상 직접 확인이 필요한 사항들.

• 구체적 throughput 수치 — 본 노트의 token/sec 추정은 도메인 지식 기반. 강의가 보여준 정확한 측정값은 영상 직접 확인 필요.
• 구체적 Metal 코드 — 본 노트의 INT4 커널은 ExecuTorch 공개 코드 + 일반 패턴 조합. 강의의 정확한 코드는 다를 수 있음.
• Apple Neural Engine (ANE) 의 자리 — Apple 의 specialized matmul accelerator. 본 강의가 GPU (Metal) 만 다뤘는지, ANE 와의 비교를 했는지 확인 필요.
• M3 Dynamic Caching — M3 부터 새로 들어간 GPU feature. 본 강의가 활용했는지 확인 필요.
• KV cache quantization — long context 에서 KV cache 도 4-bit 양자화. 본 강의가 다뤘는지 미확인.
• prefill 최적화의 자리 — decode 에 집중했는지, prefill 의 large batch matmul 도 다뤘는지.
• 다른 quantization recipe — GPTQ vs AWQ vs SmoothQuant 의 비교. 본 강의가 어느 알고리즘을 사용했는지 확인.
• ExecuTorch 의 backend partition 디테일 — 어느 op 가 Metal, 어느 op 가 CPU 인지의 결정. 본 강의가 이 자리를 다뤘는지.