LLM inference 의 latency 곡선에서 보이지 않는 두 도둑 — kernel launch overhead 와 HBM 왕복. CMU 의 Mirage Persistent Kernel 은 LLM forward 전체를 한 개의 대형 fused 커널로 컴파일해서 두 비용을 동시에 절감한다. 1.2-6.7× latency 개선이라는 숫자가 의미하는 것 — search space, super-optimizer, mega-kernel 의 인지적 모델까지 학습 노트.
"하나의 inference token = 수백 개 kernel launch" 의 자리를 다시 본다
LLM 의 inference 를 표준 PyTorch / vLLM 으로 돌리면 한 token 의 forward 가 보통 200-500 개의 GPU kernel 로 쪼개진다. attention, GEMM, layer norm, residual, RMSNorm, rotary embedding... 각자가 별도 launch. small batch (특히 batch=1) 에서는 kernel launch overhead 가 wallclock 의 30-50% 가 되는 자리.
강의의 frame.
각 kernel 사이의 HBM 왕복 — 한 layer 의 output 이 HBM 에 써지고 다음 kernel 이 다시 읽는다. 같은 데이터를 두 번 옮긴다. memory bandwidth 의 낭비.
launch overhead — 각 launch 가 ~5μs 의 host-device 통신. 200 launch 면 1ms. small-batch decode latency 에 직접 박힌다.
이걸 한 커널로 묶을 수 있다 — Mirage Persistent Kernel (MPK). "한 번 launch 된 mega-kernel 이 layer 들을 inside 에서 sequencing".
강의의 인지적 frame
vLLM / SGLang 같은 inference engine 의 다음 자리는 커널 갯수 자체를 줄이는 것. fusion 한두 개가 아니라 전체 forward 를 단일 커널로 컴파일. 그것이 Mirage 의 출발점.
"우리가 본 inference 의 진짜 도둑은 compute 가 아니라 launch 와 HBM 왕복이다 — 둘 다 한 커널로 묶어서 죽인다."Mirage 저자 · 확인 필요
§ 02mega-kernel 의 동기· launch overhead 절감
many-kernel vs mega-kernel — timeline 의 차이
왜 mega-kernel 이 의미가 있는지를 timeline 한 장으로 본다.
FIG · LLM 한 token forward 의 timeline 비교vLLM many-kernel vs Mirage mega-kernel
vLLM
L
k1
L
k2
L
k3
L
k4
L
k5
L
attn
L
k6
L
k7
L
k8
Mirage MPK
L
single mega-kernel · all layers in-kernel scheduled
절감
latency saving (launch + HBM)
L = launch overhead bar, k = kernel. vLLM 의 갭마다 launch 가 박혀 있고, kernel 사이에 HBM round-trip 의 implicit cost. Mirage 는 한 launch + on-chip scheduling.
mega-kernel 이 왜 단순히 fusion 의 한 사례 가 아닌가. 일반 fusion (예: torch.compile) 은 인접한 op 만 묶는다 (matmul + bias + relu). Mirage 는 layer 의 경계를 넘어서 attention 과 다음 layer 의 GEMM 까지 한 커널 안에 들어간다. 이게 가능한 이유 — § 03 의 super-optimizer 가 dependency 를 분석해 inside-kernel scheduling 으로 표현하기 때문.
"persistent" 의 의미
"Persistent Kernel" 은 한 번 launch 된 후 모든 SM 위에 머무르며 작업을 처리하는 패턴. work-queue 에서 task 를 fetch 해서 실행. forward 의 모든 layer 가 task 로 분해되어 한 persistent kernel 안에서 도는 것이 MPK.
§ 03super-optimizer· why "super"
"여러 단계의 IR 을 동시에 search" 가 핵심 차별점
Mirage 가 부른 표어 — multi-level superoptimizer for tensor programs. 일반 컴파일러가 한 IR level 에서만 최적화한다면, Mirage 는 graph-level / tile-level / SM-level 을 동시에 search.
L0graph IRmodel 의 computational graph — matmul, attention, norm 같은 op 들의 DAG
L1block / tile IRgraph 의 op 를 tile 단위로 분해. 어떤 op 끼리 fuse 할지, tile 크기는 얼마인지
L2thread / warp IRtile 안의 thread 분배. shared memory layout. WGMMA 의 매핑
L3PTX → SASS최종 코드 생성. nvcc 가 담당
"super" 의 진짜 의미 — L0/L1/L2 의 결정을 동시에 search. graph 단계에서 fusion 후보를 먼저 정하고 → tile 단계에서 최적 size 를 정하고 → thread 단계에서 layout 을 정하는 sequential pipeline 이 아니다. 세 단계의 결정이 맞물려 있어 분리하면 local optimum 에 갇힌다.
예시 — attention 의 KV cache load 와 GEMM 의 weight load 를 동시에 보면.
graph 단계 — KV load 와 weight load 가 같은 tile group 에 들어가면 fusion 가능.
tile 단계 — 두 load 가 같은 SM 위에서 돌면 SMEM 공유로 bandwidth 절감.
thread 단계 — thread 분배가 두 load 의 latency 를 가리도록 hand-tuned.
이 셋이 분리되면 — graph 단계에서 fusion 결정 시 SMEM 비용을 모르고, tile 단계에서 thread 를 모르고. super-optimizer 는 세 결정을 joint search.
왜 "super-optimization" 이라 부르는가
Massalin 의 1987 superoptimizer 의 정신 — 모든 가능한 program 을 enumerate 하고 검증한다. Mirage 는 LLM-scale 에서 같은 정신을 살리되 search space 를 학습한 heuristic 으로 줄이고, 검증을 자동화. 결과 = "사람이 짤 수 없는 형태의 mega-kernel".
"우리는 graph + tile + thread 를 함께 search 한다 — 분리하면 좋은 mega-kernel 이 나오지 않는다."Mirage 논문 · 확인 필요
§ 04search space· what is searched
scheduler 가 결정해야 하는 4 차원
Mirage 의 search space 를 명시적으로 풀어본다. 4 개 축이 있다.
TAB · search space 의 4 차원모든 결정의 joint optimization
차원
결정
예시
typical 변형
graph fusion
어떤 op 끼리 묶는가
matmul + add + GeLU vs 별도
O(2^N) — pruning 필요
tile shape
각 op 의 BLOCK_M/N/K
(128, 256, 32) vs (64, 128, 64)
~10-20 candidates
SM mapping
persistent task 를 어떤 SM 에
round-robin vs work-stealing
3-4 strategy
memory placement
어떤 tensor 를 SMEM/register/HBM
K cache 를 SMEM 에 keep
capacity 제약 안
총 search space 는 천만 단위 후보. cost model 로 빠르게 prune, 살아남은 후보만 실측. 한 model 당 컴파일 시간은 분-시간 단위.
FIG · search space sketchgraph × tile × thread
vLLM 의 default 는 좌하단 (작은 tile, 얕은 fusion). Mirage 의 best 는 우상단 (큰 tile, 깊은 fusion). 그 영역에 mega-kernel 이 산다.
§ 05검증· verification
"super-optimizer 가 짠 코드가 진짜 맞는가" 의 답
super-optimizer 의 위험 — 컴파일러가 짜낸 코드가 사람이 검토하기 어려운 형태. silent wrong result 가 production 에 풀려 나갈 risk. Mirage 의 답.
numerical 검증 — 같은 input 에 대해 reference (PyTorch eager) 와 mega-kernel 의 output 을 atol/rtol 비교. § L076 의 BackendBench 와 같은 정신.
differential testing — random input 으로 두 backend 비교. NaN/Inf mask 까지 확인.
timing profile — 각 mega-kernel 후보의 실제 wallclock 측정. cost model 의 prediction 과 비교 — error 가 큰 후보는 버린다.
compiler-level invariant — mega-kernel 합성 시 dependency graph 의 분명한 위반 (예: write-before-read) 은 컴파일러 단계에서 reject.
이 검증 layer 가 산업적 채택 가능성의 핵심. 빠른 mega-kernel 이 silent wrong 이면 production 에 풀어둘 수 없다.
test-time vs build-time
Mirage 의 "컴파일" 은 작은 sample input 으로 build-time 검증을 강제. 그러나 production 의 모든 input shape 을 cover 하지 못한다 — 그래서 inference-time 에 shape 가 새로 등장하면 fallback 또는 재컴파일. dynamic shape 의 대응이 § 08 의 한계.
§ 06LLM 추론 결과· latency numbers
1.2× ~ 6.7× — 어떤 자리에서 큰 속도가 나는가
Mirage 가 보고하는 latency 개선 범위는 1.2× 부터 6.7× 까지 넓다. 이 변동의 원인을 보면 어떤 자리에 mega-kernel 이 효과적인지가 보인다.
TAB · model × scenario 별 speedup (개념적)baseline = vLLM many-kernel
scenario
설명
speedup
이유
batch=1, decode
single token 의 generation
~6.7×
launch overhead 가 dominant — mega-kernel 의 sweet spot
batch=16, decode
moderate batch decode
~3-4×
여전히 launch overhead, HBM bandwidth 도
batch=64, decode
large batch decode
~2×
compute 가 dominant 시작 — 효과 줄어듦
prefill (long sequence)
large compute kernel
~1.2-1.5×
compute bound — fusion 효과 작음
training step
backward 까지 포함
미상
강의 시점에 training 은 not yet (확인 필요)
큰 speedup 은 batch=1 decode 에 집중 — 즉 latency-critical inference, 즉 chat 응답 / coding assistance 같은 자리. throughput-critical 자리 (큰 batch prefill) 에서는 효과 약함.
모델별로도 변동이 있다. 강의에서 보고한 패턴 (확인 필요).
Llama 3 8B / 70B — 잘 알려진 모델, search space 가 좁음. speedup ~3-4×.
Mixtral / DeepSeek MoE — expert routing 의 sparsity 가 mega-kernel 에 잘 맞음. ~5×.
Qwen 2.5 32B — non-standard layer norm 등이 fusion 의 새 자리를 열어줌. ~4×.
"Mirage 의 speedup 곡선은 batch size 와 prefill/decode 의 함수 — chat application 의 latency-critical 자리에서 가장 빛난다."학습 노트
§ 07hardware 활용· SM occupation
persistent kernel 이 SM 위에 오래 머무르는 의미
mega-kernel 의 hardware 측면 효과. 단순히 launch 절감 외에 SM 활용도 자체가 변한다.
persistent kernel 의 work-queue — 한 번 launch 된 커널이 SM 위에 머무르며 task 를 fetch. SM 이 idle 한 시간이 줄어든다.
L2 cache locality — 같은 mega-kernel 안의 layer 들이 SM 의 L2 안에서 데이터를 공유. layer 간 HBM 왕복이 사라진다.
register reuse across layers — 한 layer 의 결과가 다음 layer 의 input 으로 register 안에서 직접 흐른다. SMEM 도 안 거침.
task scheduling 의 사용자 통제 — Mirage 가 task graph 를 SM 위에 분배. 일부 SM 은 attention, 다른 SM 은 GEMM 에 specialized 가능 (확인 필요).
이 모델의 trade-off — register / SMEM pressure 가 매우 커진다. mega-kernel 안에 모든 데이터가 들어와야 하므로 occupancy 감소 risk. Mirage 의 search 가 이를 cost 로 반영해 균형점을 찾는다.
CUDA Graphs 와의 관계
CUDA Graphs (cudaGraph) 도 launch overhead 를 절감한다. 차이는 — Graphs 는 여러 kernel 을 미리 묶어 한 번에 launch, kernel 자체는 분리. Mirage 는 코드 자체를 한 커널로 합성. fusion 효과가 추가됨.
§ 08한계· caveats
모든 자리에서 mega-kernel 이 답은 아니다
compile time — search 와 검증을 거치므로 한 model 당 분 ~ 시간. 빠른 prototyping 에는 비싸다. cache 가 핵심.
dynamic shape 약점 — sequence length, batch size 의 변동이 매번 재컴파일을 유발할 수 있다. real-world serving 에서 shape sweep 사전 컴파일 필요.
register / SMEM saturation — 큰 mega-kernel 은 occupancy 감소. 작은 batch 는 OK 지만 큰 batch 에서는 일반 kernel 이 더 빠를 수 있다.
large compute 자리에서의 효과 작음 — prefill 같은 compute-bound 자리에서는 cuBLAS / FA3 의 hand-tuned 코드가 더 빠르다.
training 미지원 (강의 시점) — backward pass 의 mega-kernel 합성은 forward 보다 훨씬 어렵다. autograd 의 graph 가 동적이므로.
multi-GPU 의 위치 — TP / EP / PP 환경에서 mega-kernel 이 어떻게 들어가는지가 확인 필요. Iris (L078) 와 결합 가능성.
새 hardware 흡수 비용 — Hopper / Blackwell 의 새 명령을 super-optimizer 가 흡수하려면 cost model 업데이트 + 새 search 변수.
§ 09다음 단계· future work
강의에서 명시적으로 던진 다음 자리
training 으로의 확장 — backward pass 의 mega-kernel 합성. autograd 의 dynamic graph 와 충돌 어떻게 풀 것인가.
multi-GPU mega-kernel — Iris (L078) 의 distributed Triton 추상과 결합. TP forward 전체를 한 mega-kernel 로.
online recompilation — runtime 의 shape 변동에 대한 fast-path. shape bucketing + lazy compile.
Blackwell 적응 — tcgen05.mma, FP4 의 super-optimizer 흡수.
open serving stack 통합 — vLLM / SGLang 의 backend 로 Mirage 가 들어가는 길.
cost model 의 학습 — 더 정교한 cost prediction. ML-driven cost model.
§ 10기억할 메모· key takeaways
다시 열었을 때 손에 잡혀야 할 것
launch overhead 의 진짜 비용
batch=1 decode 에서 wallclock 의 30-50%. 200-500 launch 가 한 token 마다.
mega-kernel
forward 전체를 한 커널로 합성. persistent + on-chip scheduling. fusion 의 한 단계 위.
super-optimizer
graph + tile + thread 를 joint search. multi-level. cost model + 실측 검증.
search space 4 차원
graph fusion / tile shape / SM mapping / memory placement. 천만 단위 candidate, pruning 후 실측.