Tiny-vLLM: C++와 CUDA로 만드는 고성능 LLM 추론 엔진

TL;DR Highlight

vLLM의 핵심 기능을 C++와 CUDA로 직접 구현하며 배울 수 있는 교육용 LLM 추론 엔진 프로젝트로, 소스코드와 단계별 강의가 함께 제공된다.

Who Should Read

LLM 추론 엔진 내부 동작 원리를 깊이 이해하고 싶은 시스템 개발자나 ML 엔지니어. 특히 CUDA 커널 작성이나 고성능 추론 서버 구현에 관심 있는 개발자에게 적합하다.

Core Mechanics

이 프로젝트는 단순한 코드 저장소가 아니라, LLM 추론 엔진을 처음부터 직접 구현하는 과정을 안내하는 강의 커리큘럼과 소스코드를 동시에 제공한다.
실제 모델인 Llama 3.2 1B Instruct를 Safetensors 형식에서 로드해 전체 forward pass(프리필 + 디코드)를 구현한다. 장난감 예제가 아닌 실제 모델로 동작하는 추론 엔진이다.
모든 연산은 CUDA 커널로 직접 구현된다. GPU와 CPU 메모리 관리, bfloat16 부동소수점 처리까지 다루기 때문에 하드웨어 레벨의 이해를 쌓을 수 있다.
KV cache(키-밸류 캐시, 이전에 계산한 어텐션 값을 저장해 재계산을 피하는 기법)를 구현하며, static batching과 continuous batching(요청이 들어오는 대로 배치를 동적으로 구성하는 방식)을 모두 다룬다.
FlashAttention과 유사한 Online Softmax, PagedAttention(메모리를 페이지 단위로 쪼개 KV cache를 관리하는 vLLM의 핵심 기법)을 직접 구현하는 내용이 포함된다.
RMSNorm, RoPE(위치 인코딩 방식), GQA(Grouped Query Attention, 멀티헤드 어텐션의 효율화 버전), SiLU 활성화 함수, Causal Mask 등 현대 LLM의 핵심 구성요소를 하나씩 CUDA로 구현한다.
cublasGemmEx(NVIDIA의 행렬 연산 라이브러리)를 사용하는 방법과 column-major에서 row-major로 행렬을 전치하는 트릭 같은 실전 CUDA 최적화 기법도 포함된다.
README 자체가 LLM, vLLM, 모델, 추론 서버의 개념을 처음부터 설명하는 강의 형태로 작성됐으며, 저자는 코드를 안 읽어도 스스로 재구현할 수 있는 멘탈 모델을 갖추는 것이 목표라고 밝혔다.

Evidence

저자 본인이 댓글에서 'README가 가장 중요한 부분'이라고 강조했다. 코드 자체보다 개념을 이해하고 스스로 재구현할 수 있도록 설계된 학습 자료라는 점을 명확히 했다.
커뮤니티에서 초창기 llama.cpp와 비슷한 느낌이지만 문서화가 훨씬 잘 되어 있다는 평가가 나왔다. 강의 형식으로 포맷된 문서가 특히 호평을 받았다.
CUDA API 반환값 체크 코드가 없다는 지적이 있었다. CUDA API는 오류를 반환값으로 알려주는데 이를 무시하면 디버깅이 어렵기 때문에, '학습용이라도 에러 핸들링은 넣어야 하는 것 아니냐'는 비판적 의견이 있었다.
저자의 블로그(jedrzej.maczan.pl)에 AI 논문 관련 글이 많다는 정보가 댓글에서 공유됐다. 이 프로젝트 외에도 관련 배경 지식을 얻을 수 있는 자료가 있다는 점이 추가로 언급됐다.

How to Apply

vLLM이나 TensorRT-LLM 같은 추론 엔진을 사용하면서 내부 동작이 궁금했던 경우, 이 프로젝트의 README 강의를 순서대로 따라가면 KV cache, PagedAttention, Continuous Batching이 왜 필요한지 바닥부터 이해할 수 있다.
대학교 ML 시스템 강의나 스터디 그룹에서 LLM 추론 최적화를 가르쳐야 한다면, 이 프로젝트를 교재로 활용할 수 있다. 저자가 명시적으로 대학 강의 자료로 사용해도 된다고 허용하고 있다.
CUDA 커널을 처음 작성해보고 싶은 개발자라면, 임베딩 레이어부터 RMSNorm, 어텐션까지 단계적으로 구성된 커리큘럼을 따라가면서 실제 동작하는 LLM 위에서 CUDA 프로그래밍을 익힐 수 있다.
운영 중인 LLM 서비스에서 메모리 병목이나 배치 처리 비효율 문제를 디버깅할 때, 이 프로젝트에서 Static Batching과 Continuous Batching의 차이를 직접 구현해보면 어떤 상황에서 어떤 전략이 유리한지 실감할 수 있다.

Code Example

snippet

# 빌드 및 실행 방법 (build.sh, run.sh 참조)
# CMake 기반 빌드
bash build.sh

# 추론 실행
bash run.sh

# NVIDIA Nsight Systems로 프로파일링
bash nsys.sh

# NVIDIA NCU(Nsight Compute)로 커널 분석
bash ncu.sh

# 테스트 실행
bash test.sh
bash full_test.sh

Terminology

KV cache어텐션 연산에서 이전 토큰들의 Key/Value 행렬을 메모리에 저장해두는 기법. 토큰을 생성할 때마다 전체를 재계산하지 않아 속도가 크게 빨라진다.

PagedAttentionKV cache 메모리를 OS의 가상 메모리처럼 고정 크기 페이지 단위로 나눠 관리하는 기법. vLLM이 개발해 메모리 낭비를 크게 줄였다.

Continuous BatchingLLM 추론 서버에서 요청이 완료되는 즉시 새 요청을 같은 배치에 끼워 넣는 방식. 기존 Static Batching은 모든 요청이 끝날 때까지 기다려야 해서 GPU가 놀았다.

RoPERotary Position Embedding의 약자. 트랜스포머 모델에서 토큰의 위치 정보를 회전 행렬로 인코딩하는 방식으로, Llama 계열 모델에서 사용된다.

GQAGrouped Query Attention의 약자. 멀티헤드 어텐션에서 Query 헤드 여러 개가 하나의 Key/Value 헤드를 공유하도록 해 메모리와 연산량을 줄이는 기법이다.

SafetensorsHugging Face가 만든 모델 가중치 저장 포맷. 기존 PyTorch pickle 형식보다 빠르고 안전하게 모델을 로드할 수 있다.