Insights
AI가 자동으로 큐레이션·번역·정리하는 기술 동향 피드입니다.
본 글은 LLM 추론 과정에서 발생하는 메모리 비효율성을 다루며, 레거시 시스템이 GPU HBM의 상당 부분을 낭비하는 문제를 지적합니다. vLLM이 도입한 PagedAttention은 운영체제의 가상 메모리 페이징 개념을 차용하여, 연속 할당의 한계를 극복하고 메모리 단편화 및 예약 낭비를 제거함으로써 추론 효율성을 혁신적으로 개선했습니다.
AI 자동 생성 콘텐츠
본 콘텐츠는 Dev.to AI tag의 원문을 AI가 자동으로 요약·번역·분석한 것입니다. 원 저작권은 원저작자에게 있으며, 정확한 내용은 반드시 원문을 확인해 주세요.
원문 바로가기이전 심층 분석(From model.generate() to the Roofline Model: The Physics of Prefill and Decode Divergence)에서 우리는 자가회귀(autoregressive) 디코드 단계가 근본적으로 **메모리 대역폭에 의해 제한됨(memory-bandwidth bound)**을 1원칙부터 도출했습니다. 효과적인 GPU 컴퓨팅 활용률을 극대화하고 메모리 트래픽을 상각하기 위한 유일한 방법은 가능한 한 높은 동시 배치 크기(concurrent batch sizes)로 확장하는 것입니다.
하지만 초기 서빙 스택(예: 초기 FasterTransformer 또는 순수 PyTorch 서빙 스크립트)이 동시성을 확장하려고 할 때, 시스템 엔지니어들은 극복할 수 없는 벽에 부딪혔습니다:
80GB A100 또는 H100 가속기에서, 단 12개의 동시 요청만으로도 치명적인 CUDA Out of Memory (OOM) 오류로 인해 시스템이 자주 충돌했습니다. 이상하게도, 실제로 생성되는 토큰의 KV 캐시 물리적 크기를 계산하면 실제 메모리 사용량은 20GB 미만이었습니다.
나머지 60GB에 달하는 프리미엄 고대역폭 메모리(HBM)는 어디로 사라진 것일까요?
UC Berkeley 연구원들이 SOSP 2023에서 발표한 기념비적인 논문, _"Efficient Memory Management for Large Language Model Serving with PagedAttention,"_에 따르면, 레거시 추론 런타임은 총 GPU 메모리의 60%에서 80%를 낭비했습니다.
이 물리적 장벽을 극복하기 위해, vLLM은 PagedAttention을 도입했습니다. 이 글에서는 운영 체제 설계의 핵심 개념인 **가상 메모리 페이징(virtual memory paging)**이 어떻게 GPU 가속기에 적용되어 현대 LLM 추론 아키텍처를 재편했는지 탐구합니다.
PagedAttention 이전에는 딥러닝 프레임워크(PyTorch 및 TensorFlow 포함)가 근본적인 가정에 의존했습니다: 텐서는 반드시 연속적인 물리 메모리를 차지해야 한다.
동적이고 토큰별 자가회귀 생성 과정에서, 이 연속 메모리 모델은 세 가지 주요 메모리 낭비 원인을 만듭니다:
graph TD
subgraph Waste["연속 메모리 할당의 세 가지 결함 (60%-80% 낭비)"]
W1["<b>내부 단편화(Internal Fragmentation)</b><br/>max_model_len에 대한 정적 사전 할당 (예: 4096).<br/>사용자가 50 토큰에서 완료되면, 나머지는 잠금 상태로 유지됨.]
...
대규모 언어 모델(LLM)은 <EOS> 토큰을 방출할 때 비결정적으로 종료되므로, 서비스 엔진은 주어진 요청이 몇 개의 토큰을 생성할지 예측할 수 없습니다.
중간 생성 중 충돌을 방지하기 위해 기존 시스템들은 모델의 최대 컨텍스트 창(max_model_len, 예: 4096 또는 8192 토큰) 크기로 연속적인 메모리 청크를 사전 할당했습니다.
엔진이 점진적 할당(progressive allocation)을 시도하더라도(예: 128 토큰 단위로 메모리를 예약), 연속적인 텐서 확장은 다른 곳에 더 큰 버퍼를 할당하고 기존 캐시를 복사해야 합니다.
잦은 메모리 복사로 인한 처리량 페널티를 피하기 위해, 메모리 할당자들은 대부분 유휴 상태로 남아 있는 큰 안전 버퍼를 유지했습니다.
동시 서비스 환경에서 요청들은 비동기적으로 도착하고 완료됩니다. 요청 A가 256 토큰을 소비하고, 요청 B가 2048 토큰을, 그리고 요청 C가 512 토큰을 소비할 수 있습니다. 크기가 다른 연속적인 버퍼들이 할당되고 해제되면서, GPU 물리 메모리는 고립된 청크들로 단편화됩니다.
GPU 전체에 걸쳐 총 15GB의 집계 메모리가 사용 가능하더라도, 2GB 연속적 버퍼를 요구하는 들어오는 요청은 해당 크기의 연속적인 블록이 존재하지 않으면 OOM(Out of Memory) 오류를 유발할 수 있습니다.
⚠️ [TIP]
인터랙티브 계산기: 프로덕션 KV 캐시 할당량 및 최대 동시성 제한을 산정해야 한다면, 내장된 LLM VRAM 계산기를 사용하여 다양한 컨텍스트 길이 하에서 동적 메모리 풀 동작을 시뮬레이션하십시오.
메모리 단편화를 제거하기 위해 컴퓨터 과학은 이미 운영 체제 설계에서 해결책을 개발했습니다: 가상 메모리 페이징(virtual memory paging).
현대 운영 체제에서 프로세스는 연속적인 가상 주소 공간 내에서 작동하며, 물리적 메모리는 이산적인 페이지 프레임(page frames) (일반적으로 4KB)으로 나뉩니다. 하드웨어 메모리 관리 장치(MMU)는 **페이지 테이블(page tables)**을 사용하여 연속적인 가상 주소를 비연속적인 물리 페이지에 투명하게 매핑합니다.
vLLM은 이 아키텍처를 GPU KV 캐시에 적용하여 PagedAttention을 만들었습니다:
flowchart LR
subgraph LogicalSpace["논리적 시퀀스 뷰 (연속적)"]
L0["논리적 블록 0<br/>토큰 0 ~ 15"]
...
block_number를 가집니다.filled_count)을 추적합니다.PagedAttention 설계 하에서:
– 외부 단편화(External fragmentation) 제거: 모든 물리적 블록이 동일한 치수(예: 정확히 16개의 토큰 슬롯)를 공유하기 때문에, 메모리는 불일치한 간극을 만들지 않고 균일한 단위로 할당 및 해제됩니다.
– 내부 단편화(Internal fragmentation) 최소화: 메모리 낭비는 시퀀스의 마지막 논리적 블록에서만 발생합니다. 블록 크기가 $B = 16$인 경우, 요청당 최대 낭비는 15개의 토큰 슬롯입니다: $$\text{평균 낭비} = \frac{B}{2} \times \text{1 토큰 KV 캐시 크기}$$ 요청당 메모리 낭비가 60%–80%에서 4% 미만으로 떨어져, 엔진이 동일한 GPU 메모리 풀에 2배에서 4배 더 많은 동시 시퀀스를 수용할 수 있게 합니다.
vLLM과 같은 프로덕션 엔진에서는 호스트 측의 **블록 관리자(Block Manager)**가 메모리 관리를 조정합니다:
sequenceDiagram
autonumber
actor Client as Client Request
...
엔진은 수천 개의 개별 GPU 텐서 할당을 관리하는 대신, 시작 시점에 두 개의 크고 연속적인 물리적 텐서인 key_cache와 value_cache를 할당합니다.
반정밀도(BF16/FP16)에서, 물리적 키 캐시 텐서는 **응집된 메모리 접근(coalesced memory access)**을 가능하도록 배치됩니다:
# vLLM의 물리적 키 캐시 텐서 레이아웃
# 차원은 128바이트 워프 메모리 정렬을 보장하기 위해 순열됩니다
key_cache = torch.empty(
...
블록을 할당하는 것은 CPU 메모리에서 정수 인덱스를 업데이트하는 것만 필요합니다. 생성 과정 중에는 런타임 cudaMalloc 호출이 발생하지 않아, 서브 마이크로초(sub-microsecond) 단위의 블록 관리가 가능해집니다.
PagedAttention은 또한 시퀀스 간에 **Copy-on-Write (CoW)**를 통한 제로 카피 메모리 공유를 가능하게 합니다.
이는 여러 출력이 공통된 프롬프트 접두사(prompt prefixes)를 공유하는 워크로드에서 특히 유용합니다:
연속 할당 모델(contiguous allocation models) 하에서는 프롬프트의 KV 캐시가 각 분기(branch)마다 중복되어야 했습니다. PagedAttention은 **참조 카운팅 (reference counting)**을 통해 이를 관리합니다:
flowchart TD
Prompt["입력 프롬프트 (2개의 물리 블록 소비)"] --> PB0["물리 블록 #10<br/>(참조 횟수 = 2)"]
Prompt --> PB1["물리 블록 #11<br/>(참조 횟수 = 2)"]
...
#10과 #11을 참조합니다. 이 블록들의 참조 카운터는 2로 증가합니다.#10과 #11에서 두 분기 모두에 걸쳐 동시에 읽기를 수행합니다.ref_count > 1인 부분적으로 채워진 블록에 새로운 토큰을 쓸 때, 엔진은 새로운 물리 블록(#45)을 할당하고, 기존 내용을 복사하며, Branch A의 매핑 포인터를 재지정하고, 원래 블록의 참조 카운터를 감소시킵니다.이 메커니즘은 SGLang의 RadixAttention tree cache와 같은 고급 접두사 캐싱 시스템을 위한 아키텍처적 기반 역할을 합니다.
표준 어텐션 커널(예: cuBLAS GEMM 또는 표준 FlashAttention)은 키와 값(keys and values)이 연속적인 물리 메모리 버퍼에 존재해야 합니다.
데이터를 복사하지 않고 비연속적인 물리 블록에 걸쳐 어텐션을 계산하기 위해, PagedAttention은 주소를 동적으로 해결하는 사용자 정의 CUDA 커널을 사용합니다:
디스패치(dispatch) 과정에서 스케줄러는 block_tables를 2D 정수 텐서로 GPU 커널에 전달합니다:
// PagedAttention CUDA Kernel 실행 개념
template<typename scalar_t, int BLOCK_SIZE>
__global__ void paged_attention_kernel(...)
엔지니어들은 종종 블록 조회 시 발생하는 포인터 간접 참조(pointer indirection)가 캐시 미스(cache misses)를 통해 GPU 메모리 효율성을 떨어뜨리는지 질문합니다:
GPU 소프트웨어에서 가상 메모리 개념을 구현함으로써 PagedAttention은 초기 서빙 엔진의 주요 메모리 비효율성을 해결했습니다:
하지만 메모리 효율성은 퍼즐의 첫 번째 조각일 뿐입니다. 메모리 할당이 최적화되면, 스케줄러는 파이프라인 정지(pipeline stalls)를 유발하지 않으면서 다양한 길이의 요청을 관리해야 합니다.
다음 장에서는 실행 스케줄링에 대해 살펴볼 것입니다: Continuous Batching and Chunked Prefill: Eliminating Head-of-Line Blocking.
블록 크기(Block size)는 **내부 단편화(internal fragmentation)**와 GPU 메모리 버스 정렬(GPU memory bus alignment) 사이의 엔지니어링 트레이드오프를 나타냅니다: