Beginners Guide to Metal

강의의 출발점은 — Apple Silicon 의 GPU 가 사실 매우 강력한데(M2 Max 의 GPU 는 이론적 13.6 TFLOPS FP32) 대부분의 ML tutorial 이 이 자리를 가르치지 않는다. Nikita 의 입장은 명시적으로 “agnostic 하게 GPU 를 쓰는 길의 일부”.

강의가 답하려는 두 줄 —

1. PyTorch 안에서 Apple GPU 를 어떻게 쓰는가 — device="mps" 가 무엇이고 어떤 길로 도는가.
2. 그 위에서 직접 커널을 짜야 할 때 어떻게 하는가 — Metal shader 를 PyTorch extension 으로 박는 길, 디버깅 가능한 자리.

강의의 frame 은 정직하다 — “나도 advanced engineer 가 아니다, 그저 코드를 쓰는 사람이다”(Nikita 의 자기소개). 그래서 강의의 톤이 입문자 시점의 정직한 재구성. CUDA 와 mapping 되는 자리는 명시적으로 짚되, “여기서는 cudaProfiler 같은 게 없습니다” 같이 갭도 가린 적 없다.

M1 부터 M3 까지의 GPU 는 NVIDIA 와는 다른 단위 분해를 쓴다. 강의에서 Nikita 가 깐 매핑 — 이름은 다르지만 같은 자리에서 같은 일을 한다.

Metal 은 graphics + compute 를 한 API 로 묶은 Apple 의 native GPU 추상. 강의에서 Nikita 가 한 화면씩 보여준 객체들 — 각 자리에서 무엇이 일어나는지.

// Metal compute — vector add 의 Objective-C++ (간략)
id<MTLDevice> device = MTLCreateSystemDefaultDevice();
id<MTLCommandQueue> queue = [device newCommandQueue];

// 1. shader 컴파일 — 문자열에서 jit
NSString* src = @"kernel void add(device float* a, ...) { ... }";
id<MTLLibrary> lib = [device newLibraryWithSource:src options:nil error:&err];
id<MTLFunction> fn  = [lib newFunctionWithName:@"add"];
id<MTLComputePipelineState> pipe = [device newComputePipelineStateWithFunction:fn error:&err];

// 2. buffer — unified memory 면 CPU 에서 직접 fill 가능
id<MTLBuffer> bufA = [device newBufferWithLength:N*sizeof(float) options:MTLResourceStorageModeShared];

// 3. encode + dispatch
id<MTLCommandBuffer> cb = [queue commandBuffer];
id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder];
[enc setComputePipelineState:pipe];
[enc setBuffer:bufA offset:0 atIndex:0];
[enc dispatchThreadgroups:MTLSizeMake(N/256,1,1)
    threadsPerThreadgroup:MTLSizeMake(256,1,1)];
[enc endEncoding];
[cb commit];
[cb waitUntilCompleted];          // ≡ cudaDeviceSynchronize

PyTorch 는 1.12 부터 MPS backend 를 가진다. 이름의 미묘한 점 — MPS 는 사실 Metal Performance Shaders(Apple 의 vendor 라이브러리)이지만, PyTorch 의 device="mps" 안에는 Metal 직접 호출과 MPS 호출이 모두 섞여 있다.

# PyTorch MPS — 사용자 시점
import torch

assert torch.backends.mps.is_available()
device = "mps"

x = torch.randn(8192, 8192, device=device)
y = torch.randn(8192, 8192, device=device)

# MPS backend dispatch — MPSMatrixMultiplication
z = x @ y

torch.mps.synchronize()

# 직접 shader 끼워넣기 — torch.utils.cpp_extension 의 MPS 버전
from torch.utils.cpp_extension import load
ext = load(name="my_mps",
            sources=["my_kernel.mm"],
            extra_cflags=["-fobjc-arc"])

Apple Silicon 의 가장 큰 차이. NVIDIA GPU 는 별도 VRAM(HBM)이 있고 PCIe 를 통해 host RAM 과 통신. M-series 는 — CPU 와 GPU 가 같은 LPDDR5 메모리를 공유. 이게 디자인의 거의 모든 결정에 영향.

강의에서 Nikita 가 가장 정직하게 깐 자리. CUDA 의 ncu, nsys, compute-sanitizer 의 자리에 — Metal 위에는 무엇이 있고 무엇이 없는가.

강의의 큰 메시지 — “이 표만 옆에 두고 보면, 일주일 안에 CUDA 사용자가 Metal 커널을 읽을 수 있다”. 단어가 다를 뿐 — 모델은 거의 그대로.

강의 후반부의 본론. Apple Silicon 위 LLM inference 가 — 2024 시점에 — 어떻게 구체적으로 도는지. 두 큰 stack 이 있다 — llama.cpp(C++) 와 MLX(Apple 의 ML 프레임워크).

강의의 정직한 끝맺음. Apple GPU 가 빛나는 자리와 그렇지 않은 자리.

손에 새기기 — 실습 시퀀스

1. PyTorch MPS 첫 코드 — Mac 에서 device="mps" 로 LLaMA-2 7B inference. 토큰/초 측정.
2. storage mode 의 영향 측정 — 같은 buffer 를 Shared vs Private 로 — 작은 sync 워크로드에서 cache coherence 비용 직접 관찰.
3. vector add Metal shader 직접 짜기 — gemm_perf_studies.mm 같은 형식. kernel void add(...), jit 컴파일, dispatch.
4. tiled GEMM 짜기 — 강의 repo 의 gemm_perf_studies.mm 을 따라서. naive → threadgroup tiled → SIMD-group matrix(M3+).
5. llama.cpp Metal 코드 읽기 — ggml-metal.m 한 파일을 정독. Q4_K 의 dequant + GEMM 이 한 shader 안에 어떻게 짜여 있는지.
6. MLX 로 같은 모델 — 같은 LLaMA-2 7B 를 MLX 로. PyTorch MPS 와 throughput 비교.
7. Frame Capture 로 timeline — Xcode 의 “Capture GPU Frame”. PyTorch profiler 와 합쳐서 하나의 step 의 분포 확인.
8. os_signpost instrumentation — 자기 shader 안 핵심 지점에 marker. Instruments 의 timeline 에서 보기.

• M3 의 SIMD-group matrix 의 perf — 8×8 fp16 mma 가 NVIDIA mma 와 같은 자리의 throughput 비교. 실제로 GEMM perf 가 어디까지 따라오는가.
• PyTorch MPS 의 op coverage — 강의 시점 이후 어디까지 늘었는지. 특히 학습용 op (autograd, FSDP)의 지원 상태.
• MLX 와 PyTorch MPS 의 perf 갭 — Apple-native 디자인이 정확히 어디서 win 인가. lazy graph compile 의 효과.
• multi-Mac 학습 가능성 — Thunderbolt 5 + 향후 unified compute 발전이 multi-node 학습을 가능하게 할지. 현재는 사실상 단일 노드.
• Vulkan/MoltenVK 의 자리 — Apple 위에 Vulkan layer 를 깔면 portable 코드가 Metal 로 변환됨. 손실은 어디서?
• Apple Neural Engine (ANE) 의 자리 — GPU 와 별개 가속기. CoreML 만이 직접 노출. PyTorch 가 ANE 를 backend 로 쓰는 길이 있는가.
• Apple 의 ML 투자 방향 — Apple Intelligence 가 발표된 이후 Metal 의 ML 워크로드 도구가 개선되는 속도.