Metal, MPS, Tensor API: GPU의 공개된 문들

"애플 GPU로 들어가는 길은 정확히 하나고, 나머지 전부는 그 앞에 지은 문이야."

들어가는 단 하나의 길

애플 GPU에서 도는 모든 프로그램은 Metal을 거쳐. 애플의 그래픽·연산 API고, 하드웨어가 노출하는 유일한 API야. Metal 컴퓨트 커널은 Metal Shading Language로 쓴 작은 함수고, 애플 컴파일러가 컴파일하고, 커맨드 큐가 스레드 격자 위에 디스패치해. 그게 문이야. CUDA도 없고, 애플이 아직 유지하는 OpenCL도 없고, 애플이 주는 Vulkan도 없어. 프레임워크가 '애플 실리콘 지원'이라고 하면, 필요한 연산에 대해 Metal 커널을 썼거나 빌렸다는 뜻이야. 코드 블록은 MLX의 fast.metal_kernel로 그 커널 하나를 써. Python에서 Metal Shading Language 소스를 GPU에 넘기고 배열을 돌려받게 해주는 거야.

그 앞에 지은 문들

문	뭔지	누가 지나가나	증거
Metal Performance Shaders (MPS)	애플이 미리 써둔 Metal 커널 라이브러리. 행렬 곱셈, 컨볼루션, 이미지 연산	밑에서 Core ML, 그리고 여기서 이름을 딴 PyTorch 백엔드	벤더 (애플 문서)
MLX	애플의 배열 프레임워크. 자체 Metal 커널(과 CPU 백엔드), 지연 평가, 통합 메모리 네이티브	mlx-lm, Ollama의 MLX 엔진, 이 퀘스트의 실험, 이 집 로컬 추론 대부분	벤더 (MLX 문서)
PyTorch MPS 백엔드	`'mps'`에 놓인 PyTorch 텐서를 Metal로 디스패치 (처음엔 MPS 경유, 지금은 대부분 커스텀 커널)	Hugging Face 파이프라인, diffusers, CUDA 먼저 쓴 모든 것. 이 집의 이미지 생성 엔진	벤더 (PyTorch 문서)
llama.cpp의 Metal 백엔드	양자화 추론용 손으로 쓴 Metal 커널	llama.cpp, LM Studio, Ollama의 원래 엔진	벤더 (llama.cpp README: "애플 실리콘은 일급 시민 — ARM NEON, Accelerate, Metal 프레임워크로 최적화")
Core ML	레이어마다 CPU, GPU, 뉴럴 엔진을 고르는 컴파일된 모델 런타임	앱에 내장된 모델들. 이 집의 온디바이스 음성 인식기	벤더 (애플 문서)
Metal 4 Tensor API	Metal 4(macOS 26+)의 `MTLTensor`와 텐서 연산. M5 GPU의 Neural Accelerator를 직접 프로그래밍하는 길	MLX 0.30부터 ("M5의 Neural Accelerator 지원 (macOS >= 26.2)")	벤더 (애플, MLX 릴리스 노트)

이 표엔 눈여겨볼 모양이 있어. 로컬 언어 모델 추론을 지배하는 문들, MLX랑 llama.cpp는 양자화된 행렬-벡터 곱을 위해 손으로 쓴 커널을 가진 쪽이야. 디코드는 대역폭에 묶이고, 가중치 하나하나를 딱 한 번, 전폭으로 읽는 커널이 게임의 전부니까. 나머지 전부를 지배하는 문들, PyTorch랑 Core ML은 다른 데서 프로그래밍 모델을 물려받아 그 번역 비용을 치르는 쪽이야. Ollama의 2026년 이동이 바람이 어디로 부는지 제일 분명하게 보여줘. "Ollama는 이제 애플 실리콘에서 MLX로 구동됩니다, 프리뷰". llama.cpp는 옆에 남겨둔 채로.

M5의 문, 주장으로 적어두기

M5 세대에서 애플은 모든 GPU 코어에 Neural Accelerator를 넣고 Metal 4로 열었어. "개발자는 Metal 4의 Tensor API로 Neural Accelerator를 직접 프로그래밍해 앱을 위한 솔루션을 만들 수도 있다." MLX 자체 글은 "Metal 4와 함께 도입된 Tensor Operations(TensorOps)와 Metal Performance Primitives 프레임워크를 활용한다"고 하고, 그러려면 macOS 26.2 이상이 필요해. 효과에 대한 애플 숫자는 1번 트랙의 프리필 숫자야. "M4 기준 대비 첫 토큰까지 시간 최대 4배 향상". 그리고 디코드는 "더 큰 메모리 대역폭 덕분에 … 19-27% 성능 향상". 두 번째 절을 잘 읽어. 애플 스스로 디코드 향상을 가속기가 아니라 대역폭 덕으로 돌려. 이 레슨의 M5 문장은 전부 벤더 주장이야. 이 퀘스트는 M3 너머의 GPU를 재지 않았고(집의 M5 Max 노트북은 플랜의 M5 미측정 규칙 아래 있어), 가속기는 이 퀘스트가 지나가 보지 않은 문이야.

Code

metal_door.py — MLX를 통해 Python에서 디스패치하는 Metal Shading Language 커널·python

#!/usr/bin/env python3
"""The one real door: a Metal compute kernel. MLX compiles the MSL body below
into a kernel, dispatches it over a grid, and hands back an MLX array in
unified memory. Requires mlx on Apple silicon."""
import time
import mlx.core as mx

SOURCE = """
    uint i = thread_position_in_grid.x;      // one GPU thread per element
    out[i] = 2.0f * inp[i] + 1.0f;           // the whole kernel: an affine map
"""
affine = mx.fast.metal_kernel(name="affine", input_names=["inp"], output_names=["out"], source=SOURCE)


def run(x: mx.array) -> mx.array:
    return affine(inputs=[x], grid=(x.size, 1, 1), threadgroup=(256, 1, 1),
                  output_shapes=[x.shape], output_dtypes=[mx.float32])[0]


x = mx.arange(8, dtype=mx.float32)
print(run(x))                                   # array([1, 3, 5, ..., 11, 13, 15])

n = 1 << 28                                      # 1 GiB in, 1 GiB out
big = mx.ones((n,), dtype=mx.float32); mx.eval(big)
mx.eval(run(big))                                # compile + warm
best = 1e9
for _ in range(3):
    t = time.perf_counter(); mx.eval(run(big)); best = min(best, time.perf_counter() - t)
print(f"custom Metal kernel, 1 GiB in + 1 GiB out: {best*1e3:.1f} ms -> {2*n*4/best/1e9:.0f} GB/s")
# office, M3 Ultra, mlx 0.32.2, 2026-09-15: 10.0 ms -> 215 GB/s (a naive one-element-per-thread kernel;
# MLX's own add reached 635 GB/s in stream.py — writing a fast kernel is its own craft)

Exercise

metal_door.py를 돌린 다음 커널 본문을 입력 둘의 원소별 곱을 계산하게 바꿔봐(두 번째 입력 이름을 추가해). 작은 배열에서 결과를 확인해. 그다음 네 맥에 설치된, GPU에 닿는 프레임워크를 전부 나열하고 각각 표의 어느 문을 쓰는지 이름 붙여. 그중 네가 방금 쓴 커널을 쓰게 해주는 건 뭐야?

Hint

Python이나 Swift에서 자기 커널을 쓰게 해주는 건 MLX(fast.metal_kernel)랑 날것의 Metal뿐이야. PyTorch MPS와 Core ML은 커널을 대신 골라주고, llama.cpp는 Metal 소스를 편집하게는 해주지만 런타임에 주입은 안 돼. 네 코드로 지나갈 수 있는 문이 알아둘 가치가 있는 문이야.