더 빠른 컴파일러 생성 GEMM 및 FlashAttention 커널을 사용하여 LLM 추론을 최적화한 작업 로그와 벤치마크 결과를 제시합니다. 주의할 점은 토큰 생성 속도가 빨라지는 것이 아니라, 단지 지연 시간(Latency)이 낮아진다는 것입니다. TPOT (생성)는 메모리 대역폭 제한(Memory-bound)을 받기 때문에, 더 빠른 커널이 도움이 되지 않습니다. 또한, 이 추론 프레임워크는 아직 실험적인 단계입니다. 통합은 vLLM 플러그인을 통해 이루어지는데, 불행히도 이는 추가적인 오버헤드를 발생시켜 순정 vLLM에 비해 TPOT가 약간 느려지는 결과를 초래합니다. 그럼에도 불구하고, 이 결과가 LLM 서빙 최적화에 관심 있는 개인들에게 도움이 될 수 있을 것입니다. 또한, 통합 과정의 특이 사항들이 해결되면 RTX GPU에서 전반적인 성능 향상이 있을 것으로 기대합니다. 5090에서 테스트해보려면 다음 명령어를 실행하세요:

git clone https://github.com/cloudrift-ai/emmy
cd emmy && make setup
./venv/bin/emmy deploy local --recipe recipes/gemma-4-12B-it

미리 빌드된 Docker 이미지도 있습니다:

docker run --gpus all --ipc=host -p 8000:8000 cloudriftai/vllm-emmy-gemma-4-12b-it:latest

End-to-End 벤치마크

출력 토큰 처리량 (tok/s):

¹ llama.cpp의 64-슬롯 구성은 OOM(Out of Memory)이 발생합니다: 이는 64×640-토큰 슬롯을 위해 약 15GB의 KV 캐시 전체를 사전에 할당하기 때문입니다. vLLM의 페이지 할당기(Paged allocator)와 달리, 여기서는 Gemma의 슬라이딩 윈도우(Sliding-window) 레이어를 활용하지 못하며, 이는 23GB의 FP16 가중치와 병행하여 사용하기에 적합하지 않습니다. 긴 시퀀스(16K, 24K, 32K, 64K)에 대해서도 테스트를 진행했으나, 간결함을 위해 생략했습니다. 결과는 4K/4K 벤치마크 지점의 경향을 따릅니다.

중앙값 지연 시간(Median latencies), TTFT / TPOT (ms)는 동일한 실행에서 측정되었습니다:

² c=64의 TTFT는 단일 64개 요청 웨이브( --num-prompts 64 )로 측정되었습니다: 각 요청의 프리필(prefill)은 하나의 웨이브에서 허용되므로, 중앙값은 포화 상태에서의 프리필을 측정합니다. ( --num-prompts 256의 경우 중앙값은 두 번째 및 세 번째 웨이브 요청에 도달하며, 이는 프리필이 아닌 대기열 소진(queue drain) — TPOT×256에 진입 허용량 포함 — 을 측정합니다. 엔진당 평균/중앙값 역전 값들이 이를 보여줍니다.) 행의 TPOT 절반은 --num-prompts 256의 정상 상태 실행에서 가져온 것입니다.

Speculative Decoding (MTP) Smoke-test
이 섹션에서는 Emmy가 MTP를 사용하여 Gemma 4를 실행할 때 예상치 못한 성능 또는 품질 저하가 발생하지 않는지 확인합니다. 아래 표들은 무작위 토큰 입력에 대해 MTP를 사용한 Gemma 4의 처리량(tok/s)을 보여주며, 각 추측 깊이별로 하나의 표가 제공되고 각각은 기본 vLLM 기준과 비교됩니다.

emmy bench experiments/gemma-4-12B/serving_mtp_rtx5090 --local 2 speculated tokens (tok/s) : Tokens in Tokens out Concurrency vLLM (stock) stock + MTP Emmy + MTP 256 256 64 1434.7 1423.5 882.8 4096 4096 1 57.3 109.3 107.2 4096 4096 4 216.6 355.8 352.1 4096 4096 8 384.1 609.9 446.2 8192 256 4 112.9 114.8 102.5 3 speculated tokens : Tokens in Tokens out Concurrency vLLM (stock) stock + MTP Emmy + MTP 256 256 64 1434.7 — — 4096 4096 1 57.3 145.9 135.8 4096 4096 4 216.6 443.7 436.3 4096 4096 8 384.1 716.5 496.0 8192 256 4 112.9 124.1 103.3 5 speculated tokens : Tokens in Tokens out Concurrency vLLM (stock) stock + MTP Emmy + MTP 256 256 64 1434.7 — — 4096 4096 1 57.3 197.2 187.4 4096 4096 4 216.6 586.5 404.2 4096 4096 8 384.1 728.2 632.7 8192 256 4 112.9 129.3 113.9 커널 최적화는 주로 Emmy 컴파일러가 RTX 5090 및 RTX 4090에서 일반적인 GEMM 및 FA 형태에 대해 cuBLAS를 능가할 수 있도록 한 두 가지 최적화를 활용한 것입니다: TMA 전송(Blackwell 전용)과 FP32 그림 누적 레지스터를 사용한 전체 FP16 텐서 코어 활용입니다. Matmul 및 Flash 커널을 위한 TMA 전송은 소비자 Blackwell 다이에서도 여전히 동일한 cp.async 전송을 사용합니다 (실제로 소비자 Blackwell의 대부분 cuBLAS GEMM 커널, FP16 tensorop 경로를 포함하여,는 순방향 포팅된 Ampere 시대의 cutlass_80_* 커널입니다). cp.async를 TMA로 교체하면 필요한 명령어 수를 줄일 수 있고, TMA는 공유 메모리 뱅크 충돌을 무료로 스위즐(swizzle) 방식으로 처리합니다. 커널 형태 M×N×K cuBLAS (cp.async) Emmy (TMA) 속도 향상 q_proj 512×4096×3840 97.6 84.2 1.16× kv_proj 512×2048×3840 48.8 45.3 1.08× o_proj 512×3840×4096 103.3 82.0 1.26× gate_up (융합 게이트+업) 512×30720×3840 573.3 561.7 1.02× down_proj 512×3840×15360 286.3 286.2 1.00× 하이브리드 FP16/FP32 누적 프로덕션 스택의 기본 matmul 경로는 FP32 누적을 사용하는 FP16 텐서 코어를 사용합니다.

하지만 소비자용 다이(consumer dies)에서는 FP16 입력/FP32 누적(FP16-input/FP32-accumulate) HMMA가 FP16 입력/FP16 누적(FP16-input/FP16-accumulate) 속도의 정확히 절반으로 실행됩니다. 이를 해결하기 위해, 저는 빠른 atom인 mma_m16n8k16_f16_f16을 사용하되, FP16 부분합(partials)을 FP32 레지스터로 승격(promoting)시켜 전역 누적(global accumulation)을 FP32에서 정확하게 수행함으로써 정확도를 유지합니다. 이를 통해 모든 개별 mma마다 FP32 누적 비용(tax)을 지불하는 대신, FP32 누적을 사용하면서도 FP16 텐서 코어(tensor-core)의 속도를 얻을 수 있습니다. 이와 유사한 트릭이 Hopper 아키텍처에서 FP8 텐서 코어의 제한된 정밀도 누적 문제를 해결하기 위해 DeepSeek의 DeepGEMM에서 사용된 바 있습니다. 저는 이를 FP16에 맞게 조정하였으며, 컴파일러는 Gemma 모델의 추론에 사용되는 수많은 GEMM 커널 변형(variants)을 생성했습니다. Gemma 프로젝션(projections) 전반에 걸쳐 측정된 이득(RTX 5090, seq_len 512, µs):

| Kernel Shape | M×N×K | cuBLAS HGEMM | Emmy FP32 | Emmy Hybrid | Hybrid vs cuBLAS |
| :--- | :--- | :--- | :--- | :--- |
| q_proj | 512×4096×3840 | 97.6 | 84.2 | 61.5 | 1.59× |
| kv_proj | 512×2048×3840 | 48.8 | 45.3 | 36.4 | 1.34× |
| o_proj | 512×3840×4096 | 103.3 | 82.0 | 64.1 | 1.61× |
| gate_up (fused gate+up) | 512×30720×3840 | 573.3 | 561.7 | 362.4 | 1.58× |
| down_proj | 512×3840×15360 | 286.3 | 286.2 | 214.2 | 1.34× |

정확도 검증 (Accuracy Verification)
C = A@B의 오차는 N(0,1)에서 추출된 FP16 입력을 사용하여 측정되었으며, 동일한 FP16 반올림 피연산자(FP16-rounded operands)에 대해 각 누적 전략을 FP64 참조값과 비교했습니다 (이를 통해 입력 반올림이 아닌 누적 오차만을 분리하여 측정합니다).

K (누적 깊이) FP32-accum FP16-accum Emmy Hybrid 256 9.2e-8 6.0e-4 3.3e-4 3840 ( qkv / gate K) 2.8e-7 2.3e-3 3.3e-4 4096 ( o_proj K) 3.0e-7 2.4e-3 3.3e-4 15360 ( down_proj K) 5.6e-7 4.5e-3 3.3e-4 32768 8.1e-7 6.6e-3 3.3e-4 정확성 검사: 200 GSM8K 문제, few-shot 프롬프트, 동일 시드 (lm-eval 0.4.12, 엄격한 정확 일치): Lane GSM8K exact-match vLLM (stock) 0.685 ± 0.033 vLLM + Emmy 0.670 ± 0.033 vLLM + Emmy FAST_MATH 0.695 ± 0.033 llama.cpp 0.665 ± 0.033 네 가지 구성 모두 하나의 표준 오차 범위 내에 있습니다: FAST_MATH의 하이브리드 누적은 작업 품질을 저하시키지 않으며, 커널 수준 오류 분석과 일치합니다 (그 상대 오차 ~3.3×10⁻⁴는 FP16 자체 표현 노이즈 안에 위치합니다). Flash Attention Numbers 동일한 최적화가 FA에 활용됩니다. 전체 FA 최적화 이야기는 이전 게시물에서 다루었으므로, 점수판만 올립니다 (causal, seq 512, 16 heads, head_dim 256): Card torch SDPA (FA-2) Emmy FP32 Emmy Hybrid best vs SDPA RTX 5090 30.7 31.7 29.7 1.03× RTX 4090 41.0 37.1 33.8 1.21× 전체 기사: https://medium.com/ai-advances/beating-vllm-and-llama-cpp-on-gemma4-12b-68a8412f57ea /u/NoVibeCoding이 r/LocalLLaMA에 제출함 [링크] [댓글]