동시 접속자 100명 몰리자마자 H100 8대가 얼어붙었다: vLLM 청크 프리필(Chunked Prefill)과 KV 캐시 분할 내부 동작 원리

NVIDIA H100 80GB SXM5 GPU의 공식 데이터시트를 책상 위에 올려두고 공학용 계산기를 두드리며 한참을 멍하니 앉아 있었습니다. 장당 5,000만 원이 넘는 이 하드웨어는 FP16 기준 초당 약 2,000조 번($1,979 \text{ TFLOPS}$)의 행렬 곱셈을 수행할 수 있는 연산 괴물입니다. 그런데 사내 클러스터에 Meta-Llama-3-70B-Instruct 모델을 4-way Tensor Parallel로 띄우고 DCGM(Data Center GPU Manager) 텔레메트리를 열어보았을 때, 저희 팀이 마주한 실측 숫자는 기괴했습니다.

[H100 SXM5 Hardware Telemetry Snapshot]
• HBM3 메모리 대역폭 점유율: 98.4% (3.35 TB/s 메모리 버스 포화 🔥)
• 텐서 코어 SM 연산 활성도:   14.2% (유휴 텐서 코어 85.8% 💤)
• 실시간 소비 전력:          312 W / 700 W (TDP 대비 절반 이하)
• 병목 원인:                 디코드(GEMV) 단계의 산술 강도가 1~2 FLOP/Byte에 불과함

HBM3 메모리 버스는 98.4%까지 차올라 비명을 지르고 있는데, 정작 값비싼 텐서 코어(Tensor Core) 연산 유닛의 85.8%는 아무런 일도 하지 못한 채 클럭만 축내고 있었습니다.

더 심각한 문제는 프로덕션 프론트엔드에서 날아온 사용자 체감 장애 리포트였습니다. 초당 30글자씩 매끄럽게 타이핑되던 실시간 챗봇 스트리밍 화면이, 아무런 예고도 없이 불규칙하게 2초 넘게 꽁꽁 얼어붙었다가(Freezing) 밀려 있던 글자들이 한꺼번에 쏟아져 나오는 현상이 반복되고 있었습니다.

[실시간 챗봇 SSE 스트리밍 구간의 실제 패킷 타임스탬프 로그]
00.00s: "고객님, 문의하신 2026년 해외 출장 여비 정산 규정에 대..." (정상 스트리밍)
00.45s: "...해 안내해 드리겠습니다. 우선 제출하셔야 할 증빙 서류는..." (정상 스트리밍)
00.52s: 🚨 [여기서 정확히 2.14초간 SSE 청크 유입 완전 정지 / 화면 프리징]
02.66s: "항공권 e-티켓 영수증과 법인카드 매출전표이며..." (밀린 토큰이 순식간에 쏟아짐)

처음에는 백엔드 FastAPI 게이트웨이의 비동기 이벤트 루프 블로킹이나 Envoy 리버스 프록시의 버퍼링 설정을 의심했습니다. 하지만 네트워크와 웹 서버 계층을 모두 우회하여 vLLM 엔진 프로세스에 직접 프로파일러 심침을 꽂았을 때, 범인은 소프트웨어가 아니라 H100 실리콘 내부의 루프라인 물리학(Roofline Model)과 초기 연속 배치(Continuous Batching) 스케줄러의 프리필 독점 구조에 있다는 사실이 드러났습니다.

H100의 손익분기점(590.7 FLOPS/Byte)과 GEMV 디코딩의 물리적 모순

LLM 추론 워크로드가 왜 GPU 내부에서 극단적인 불균형을 일으키는지 이해하려면, 하드웨어의 연산 능력과 메모리 대역폭 사이의 비율인 산술 강도(Arithmetic Intensity)를 계산해 보아야 합니다. NVIDIA H100 80GB SXM5의 FP16 텐서 코어 피크 연산력은 $1,979 \text{ TFLOPS}$이고, HBM3 고대역폭 메모리의 최대 대역폭은 $3.35 \text{ TB/s}$입니다. 텐서 코어와 메모리 버스가 어느 한쪽도 놀지 않고 100% 맞물려 돌아가기 위해 필요한 최소 연산 밀도, 즉 루프라인 모델의 손익분기점(Ridge Point)은 다음과 같이 도출됩니다.

$$\text{Ridge Point} = \frac{\text{Peak Computation (FLOPS)}}{\text{Memory Bandwidth (Bytes/s)}} = \frac{1,979 \times 10^{12}}{3.35 \times 10^{12}} \approx 590.7 \text{ FLOPS/Byte}$$

HBM3 메모리에서 가중치나 KV 캐시 데이터를 1바이트 퍼 올릴 때마다 최소 591번 이상의 부동소수점 곱셈·누산(FMA) 연산을 수행해야만 텐서 코어가 메모리를 기다리지 않고 풀가동된다는 뜻입니다.

문제는 트랜스포머 추론이 물리적 성격이 완전히 정반대인 두 단계로 나뉜다는 점입니다. 사용자가 입력한 프롬프트(예: 4,000토큰)를 처음 받아 초기 KV 캐시를 생성하는 프리필(Prefill) 단계는 4,000개의 토큰 행렬 전체를 가중치 행렬과 곱하는 대규모 행렬-행렬 곱셈(GEMM, General Matrix-Matrix Multiplication)입니다. 이 구간에서는 한 번 SRAM으로 불러온 모델 가중치를 4,000개의 토큰이 공유하여 연산하므로 산술 강도가 800~1,500 FLOPS/Byte까지 치솟습니다. 철저한 연산 병목(Compute-bound) 구간이며 텐서 코어가 100% 가동됩니다.

반면 프리필이 끝나고 사용자 화면에 글자를 한 자씩 찍어내는 디코드(Decode) 단계는 직전까지 생성된 문맥을 바탕으로 오직 단 1개의 다음 토큰만 생성하는 행렬-벡터 곱셈(GEMV, General Matrix-Vector Multiplication)입니다. Llama-3-70B 모델의 FP16 가중치는 약 140GB이며, 이를 4장의 H100으로 텐서 병렬화(TP=4)하면 GPU 한 장당 35GB의 가중치를 분담합니다. 동시 접속자 32명의 다음 글자 1개를 뽑아내기 위해, GPU는 매 스텝마다 HBM3에서 35GB의 모델 가중치 전체와 32명분의 KV 캐시를 온칩 SRAM으로 전부 읽어들여야 합니다. 그렇게 수십 기가바이트를 퍼 올리고도 정작 수행하는 연산은 벡터 내적 몇 번이 전부라, 배치가 32일 때의 산술 강도는 고작 16~32 FLOPS/Byte 수준에 그칩니다. 손익분기점인 590.7에 턱없이 모자라기 때문에, 디코드 구간에서는 HBM3 메모리 버스만 98%로 포화되고 텐서 코어의 85%는 데이터가 도착하기만을 기다리며 공회전(Stall)하는 것입니다.

Nsight Systems(nsys) 타임라인에 찍힌 1,882ms의 디코드 기아(Decode Starvation)

그렇다면 평소 30ms 간격으로 글자를 잘 찍어내던 챗봇이 왜 불규칙하게 2초씩 멈춰 섰을까요? 저희는 컨테이너 내부에 NVIDIA Nsight Systems(nsys)를 붙여 CUDA 스트림 위에서 실행되는 커널 타임라인을 마이크로초 단위로 캡처했습니다.

nsys profile --trace=cuda,nvtx,osrt --output=vllm_hw_trace \
  python3 -m vllm.entrypoints.openai.api_server \
    --model meta-llama/Meta-Llama-3-70B-Instruct --tensor-parallel-size 4

추출된 트레이스 파일(.nsys-rep)을 열어 스트리밍이 멈췄던 바로 그 시각의 SM(Streaming Multiprocessor) 점유 상태를 확대한 결과, 결정적인 장면이 포착되었습니다.

[CUDA Stream Kernel Execution Timeline]
00.00s ~ 00.35s: ── GEMV 커널 (기존 32명 디코드 토큰 1개씩 생성: 이터레이션당 31.8ms) ──
00.35s ~ 00.38s: ── GEMV 커널 (기존 32명 디코드 토큰 1개씩 생성: 이터레이션당 32.1ms) ──
[🚨 사내 위키 검색 봇에서 16,384 토큰 크기의 장문 RAG 프롬프트 1건 유입]
00.38s ~ 02.26s: [=== flash_fwd_kernel + cutlass_gemm (16K 프리필 실행 / 1,882.4ms 독점) ===]
                 ➔ 이 1.88초 동안 기존 32명 세션의 디코드 커널 스케줄링 전면 중단!
02.26s ~ 02.29s: ── GEMV 커널 (대기열에 묶여 있던 32명 디코드 연산 재개) ──────────────

범인은 동시 접속자 수의 폭증이 아니라 단 한 건의 거대한 RAG 프롬프트 난입이었습니다. 사내 개발자 중 한 명이 50페이지 분량의 아키텍처 설계서와 사규 문서를 통째로 긁어모은 16,384토큰짜리 RAG 질의를 던진 순간, vLLM의 기본 스케줄러(Prefill-prioritizing Continuous Batching)는 새로 유입된 요청의 KV 캐시를 먼저 생성하기 위해 프리필 우선 스케줄링을 발동했습니다.

16,384토큰의 대규모 GEMM 연산과 FlashAttention 순방향 커널이 GPU의 132개 SM 전체를 1,882.4밀리초(약 1.88초) 동안 독점해 버린 것입니다. CUDA 스트림은 이 거대한 프리필 커널이 끝날 때까지 기존에 초당 30자씩 토큰을 받아보고 있던 32명의 디코드(GEMV) 스텝을 단 한 번도 스케줄링하지 못했습니다. 왕복 1차선 도로에서 시속 100km로 달리고 있던 승용차 32대 앞으로 25톤 대형 덤프트럭 한 대가 끼어들어 1.88초 동안 차선을 완전히 막아버린 셈입니다. 이를 추론 시스템 엔지니어링에서는 디코딩 기아(Decode Starvation) 혹은 Head-of-Line Blocking 현상이라고 부릅니다.

유휴 텐서 코어에 프리필 조각을 얹는 청크 피기배킹(Chunked Piggybacking)

덤프트럭이 기존 차량들의 흐름을 가로막지 않게 만들 해법은, 앞서 루프라인 모델에서 확인했던 디코드 단계의 하드웨어 유휴 자원에 숨어 있었습니다. 디코딩이 돌아갈 때 HBM3 메모리 대역폭은 포화 상태이지만 텐서 코어 연산기는 85% 이상 놀고 있다는 사실에 착안한 것이 바로 청크 프리필(Chunked Prefill)과 피기배킹(Piggybacking) 아키텍처입니다.

16,384토큰짜리 거대 프롬프트를 한 번의 이터레이션에 통째로 밀어 넣는 대신, 예를 들어 512토큰 단위의 작은 청크(Chunk) 32개로 쪼갭니다. 그리고 기존 32명의 사용자가 다음 디코드 토큰 1개를 생성하기 위해 가중치 행렬(35GB)을 HBM3에서 SRAM으로 퍼 올리는 바로 그 이터레이션에, 512토큰짜리 프리필 청크 하나를 같은 배치 텐서에 슬쩍 합쳐서(Piggyback) 단일 패스로 함께 연산해 버리는 것입니다.

flowchart LR
    subgraph ASIS ["AS-IS: 기존 연속 배치 (1.88초 화면 동결)"]
        direction TB
        D1["32명 디코드 (32ms)"]:::decNode --> Prefill["🚨 16K RAG 프리필 독점<br/>(1,882ms 동안 디코드 완전 정지)"]:::freezeNode --> D2["디코드 재개"]:::decNode
    end

    subgraph TOBE ["TO-BE: 청크 프리필 + 피기배킹 (41ms 균등 실행)"]
        direction TB
        C1["32명 디코드 + 프리필 청크<br/>(41ms 혼합 배치)"]:::chunkNode --> C2["32명 디코드 + 다음 청크<br/>(41ms 혼합 배치)"]:::chunkNode --> C3["연속 균등 스트리밍 유지<br/>(프리징 0초)"]:::chunkNode
    end

    classDef freezeNode fill:#3f0f1c,stroke:#f43f5e,stroke-width:2px,color:#ffe4e6;
    classDef decNode fill:#0c2340,stroke:#38bdf8,stroke-width:1.5px,color:#e0f2fe;
    classDef chunkNode fill:#073b61,stroke:#0ea5e9,stroke-width:1.5px,color:#e0f2fe;
    style ASIS fill:#0f172a,stroke:#ef4444,stroke-width:2px,color:#fee2e2;
    style TOBE fill:#031a2e,stroke:#38bdf8,stroke-width:2px,color:#e0f2fe;

실리콘 내부에서 일어나는 연산을 수식으로 뜯어보면 왜 이 결합이 기적 같은 효율을 내는지 단번에 이해할 수 있습니다. 32명의 디코드 요청이 35GB의 가중치 행렬 $W$를 HBM3에서 온칩 SRAM으로 이미 한 번 로드해 놓았기 때문에, 512토큰의 프리필 청크는 단 1바이트의 추가 가중치 메모리 읽기(HBM Read)도 발생시키지 않고 이미 SRAM에 올라와 있는 $W$를 그대로 재사용하여 텐서 코어에서 행렬 곱셈을 수행합니다.

기존에는 32토큰(디코드 32명)만 곱하고 버려지던 가중치 로드 작업이 $32 + 480 = 512$토큰 규모의 결합 행렬 곱셈으로 변환되면서, 배치가 갖는 유효 산술 강도가 32 FLOPS/Byte에서 약 450~510 FLOPS/Byte로 훌쩍 뛰어오릅니다. H100의 Ridge Point(590.7) 부근까지 연산 밀도가 차오르면서, 메모리 대역폭 병목에 가려져 놀고 있던 85%의 유휴 텐서 코어가 사실상 공짜에 가까운 시간 오버헤드(+9ms)만으로 프리필 청크를 소화해 내는 연산 은폐(Latency Hiding)가 성립하는 것입니다.

Hopper TMA와 워프 특화(Warp Specialization), 그리고 NVLink-4의 정속 주행

특히 H100(Hopper) 아키텍처에서는 이 혼합 배치(Fused Decode + Chunked Prefill)가 전 세대 A100(Ampere) 대비 훨씬 매끄럽게 동작하는데, 그 물리적 배경에는 TMA(Tensor Memory Accelerator)와 워프 특화(Warp Specialization) 하드웨어 파이프라인이 자리 잡고 있습니다.

기존 A100에서는 HBM 글로벌 메모리에서 온칩 공유 메모리(Shared Memory, SRAM)로 KV 캐시 페이지와 가중치 타일을 복사할 때, SM 내부의 연산 스레드들이 직접 주소 오프셋을 계산하고 ld.global 인스트럭션을 발행해야 했습니다. 이 과정에서 레지스터 파일이 주소 계산에 소모되고 연산 파이프라인이 메모리 지연을 직접 뒤집어썼습니다. 반면 H100은 SM마다 독립된 하드웨어 비동기 DMA 엔진인 TMA를 탑재하여, 스레드 워프들의 역할을 물리적으로 양분합니다.

프로듀서 워프(Producer Warps)는 오직 TMA 유닛에 텐서 디스크립터 지시만 내려 다음 타일의 가중치와 비연속적인 PagedAttention KV 블록들을 백그라운드에서 SRAM의 핑퐁 버퍼(Ping-Pong Buffer)로 선반입(Prefetch)하는 작업만 수행합니다. 동시에 컨슈머 워프(Consumer Warps)는 HBM 지연시간에 전혀 간섭받지 않고 이미 SRAM에 도착해 있는 직전 타일에 대해 wgmma.mma_async(Warp-Group Matrix Multiply-Accumulate) 명령어를 쉼 없이 텐서 코어에 밀어 넣습니다.

여기에 더해 4장의 H100을 묶는 NVLink-4 인터커넥트(양방향 900 GB/s) 통신 패턴도 극적으로 안정화됩니다. 트랜스포머의 각 Attention 및 FFN 레이어가 끝날 때마다 4장의 GPU는 부분합 텐서를 동기화하는 All-Reduce 집합 통신을 수행해야 합니다. 청크 프리필이 없던 시절에는 16K 프리필이 터질 때 수백 메가바이트의 거대 텐서가 NVLink로 쏟아져 나와 마이크로 버스트(Micro-burst) 정체를 일으키다가, 디코드 구간에서는 수백 킬로바이트짜리 소형 패킷만 오가며 커널 런치 및 배리어(Barrier) 오버헤드만 낭비했습니다. 반면 한 스텝의 최대 배치 토큰 예산을 512로 고정하자, 매 이터레이션마다 NVLink 버스를 통과하는 All-Reduce 페이로드 크기가 약 $512 \times 8192 \times 2\text{ Bytes} \approx 8.4\text{ MB}$로 일정하게 정규화되어 GPU 간 배리어 대기 시간(Barrier Stall)이 84% 감소했습니다.

프로덕션 vLLM 쿠버네티스 배포 명세와 TTFT 대 ITL의 엔지니어링 등가교환

사내 H100 클러스터에 실제 적용하여 2초 프리징 현상을 박멸한 쿠버네티스 Deployment 매니페스트와 vLLM 엔진 튜닝 파라미터입니다.

# vllm-llama3-70b-chunked.yaml
apiVersion: apps/v1
kind: Deployment
metadata:
  name: vllm-llama3-70b-serving
  namespace: llm-serving
spec:
  replicas: 2
  template:
    spec:
      containers:
      - name: vllm-engine
        image: vllm/vllm-openai:v0.6.3
        args:
        - "--model=meta-llama/Meta-Llama-3-70B-Instruct"
        - "--tensor-parallel-size=4"
        - "--gpu-memory-utilization=0.92"
        - "--max-model-len=32768"
        - "--enable-chunked-prefill=true"
        - "--max-num-batched-tokens=512"
        - "--block-size=32"
        - "--enable-prefix-caching"
        resources:
          limits:
            nvidia.com/gpu: 4

여기서 핵심 조절 레버는 --max-num-batched-tokens=512와 --block-size=32의 조합입니다. --max-num-batched-tokens는 한 번의 스케줄러 스텝(Iteration)에서 디코드 토큰과 프리필 청크 토큰을 합쳐 최대 몇 개까지 묶을 것인지를 결정하는 상한선입니다. 만약 현재 32명의 사용자가 실시간 스트리밍 디코딩 중이라면 해당 스텝에서 32개의 토큰 슬롯이 먼저 디코드에 할당되고, 남은 $512 - 32 = 480$개의 토큰 예산만큼만 대기 중인 16K RAG 문서의 프리필 청크를 잘라내어 한 배치로 묶습니다. 또한 --block-size=32로 PagedAttention의 물리 KV 블록 단위를 기본값(16)의 두 배로 넓혀 FlashAttention-2 커널의 내부 타일 크기와 정렬을 맞춤으로써, 청크 분할 시 수반되는 KV 캐시 페이지 테이블 조회 오버헤드를 상쇄했습니다.

동일한 4-way H100 SXM5 노드에서 32명의 동시 스트리밍 세션이 돌아가는 와중에 15초 주기로 16,384토큰짜리 장문 RAG 프롬프트를 난입시키며 측정한 프로덕션 A/B 벤치마크 결과입니다.

벤치마크 핵심 지표	기존 연속 배치 (AS-IS)	청크 프리필 512 (TO-BE)	엔지니어링 임팩트
토큰 간 지연시간 (ITL P99)	1,882.4 ms	41.2 ms	-97.8% 단축 (스트리밍 멈춤 완전 소멸 ⚡)
토큰 간 지연시간 (ITL P50)	32.1 ms	38.4 ms	+6.3 ms (인간이 인지 불가능한 미세 증가)
클러스터 전체 처리량 (Throughput)	842 tok/s	1,086 tok/s	+28.9% 향상 (유휴 텐서 코어 85% 회수 🚀)
16K RAG 문서 첫 응답 (TTFT P99)	1,882 ms	2,190 ms	+16.3% 증가 (실시간 UX 보호를 위한 등가교환)
텐서 코어 SM 활성도 (DCGM)	14.2%	78.6%	+64.4%p 상승 (실리콘 연산 효율 극대화)

표를 자세히 들여다보면 시스템 엔지니어링의 냉정한 등가교환 법칙이 드러납니다. 16,384토큰짜리 문서를 던진 요청 당사자의 첫 토큰 응답 시간(TTFT)은 기존 1,882ms에서 2,190ms로 약 308ms(16.3%) 늘어났습니다. 한 번의 거대 커널로 몰아쳐서 끝내던 프리필 연산을 30여 번의 이터레이션으로 잘게 쪼개어 기존 디코드 배치에 얹어가며 처리했으니, 개별 프리필 요청의 완료 시점은 소폭 뒤로 밀릴 수밖에 없습니다.

그러나 서비스 전체의 사용자 경험 관점에서 이 트레이드오프는 비교조차 할 수 없는 압승이었습니다. 50페이지짜리 방대한 사규 문서를 첨부해 질문한 사용자는 첫 답변이 시작되기까지 1.8초가 걸리든 2.1초가 걸리든 문서 분량 때문에 자연스럽게 대기 시간으로 수용합니다. 반면 이미 답변을 실시간으로 읽어나가고 있던 32명의 다른 사용자들에게 화면이 2초 동안 뚝 멈추는 현상은 서비스 장애로 직결됩니다. 300ms의 TTFT를 지불하고 ITL P99를 1,882.4ms에서 41.2ms로 97.8% 끌어내린 것, 그리고 놀고 있던 텐서 코어를 깨워 전체 클러스터의 초당 토큰 처리량을 28.9% 끌어올린 것이야말로 nvidia-smi의 99% 착시 너머에서 하드웨어 물리학을 이해한 엔지니어만이 거둘 수 있는 전리품입니다.