LLM 추론 서버를 여러 대 띄웠을 때, 들어온 요청을 어느 서버로 보낼 것인가.
이 뻔해 보이는 질문이 LLM 에서는 왜 어려운지, 그리고 llm-d 가 이를 머신러닝으로 예측해서 푸는 방식을 확인해봤습니다.


1. 문제: 왜 평범한 로드밸런싱이 안 통하는가

웹 서버라면 로드밸런싱은 단순합니다. 요청 하나하나가 비슷한 비용이니 라운드로빈(순서대로 돌아가며 배분)이나 least-connection(연결이 가장 적은 서버 선택)이면 충분합니다.

LLM 추론은 다릅니다. 요청마다 비용이 수십 배씩 차이 납니다.

요청 A: "안녕"                    입력 2토큰,      출력 10토큰
요청 B: 200페이지 문서 요약해줘     입력 50,000토큰, 출력 2,000토큰

로드밸런서 입장에서 두 요청은 똑같이 "연결 1개"입니다. 그래서 연결 수나 대기열(큐) 길이로 판단하면, 큐가 짧아 보이는 서버가 사실은 거대한 요청 하나를 처리하느라 몇 초씩 막혀 있는 상황이 자주 생깁니다.

LLM 추론의 두 단계

두 용어를 먼저 잡고 가겠습니다. 글 전체에서 계속 나옵니다.

단계 하는 일 지표
Prefill (사전 채움) 입력 프롬프트 전체를 한 번에 읽어 첫 토큰을 만든다 TTFT (Time To First Token) — 요청부터 첫 글자가 나오기까지 걸린 시간
Decode (생성) 그 뒤로 토큰을 하나씩 순차 생성한다 TPOT (Time Per Output Token) — 토큰 하나당 평균 시간

Prefill 은 입력 길이에, Decode 는 출력 길이에 비례합니다. 즉 "이 요청이 서버를 얼마나 오래 붙잡을지"는 입력·출력 길이에 달려 있는데, 큐 길이라는 숫자에는 그 정보가 담겨 있지 않습니다.

캐시라는 두 번째 변수

여기에 KV 캐시가 붙습니다. 모델이 프롬프트를 읽으며 만든 중간 계산 결과를 서버가 메모리에 들고 있는 것입니다. 앞부분이 똑같은 프롬프트가 같은 서버로 다시 오면, 그 부분의 prefill 을 통째로 건너뛸 수 있습니다. 이렇게 "앞부분이 일치하는지"를 보고 재사용하는 것을 프리픽스(prefix) 캐시라고 부릅니다.

그래서 "덜 바쁜 서버"보다 "내 프롬프트의 캐시를 이미 갖고 있는 서버"가 훨씬 빠를 수 있습니다.

정리하면 이렇습니다. 부하 지표(큐 길이)와 캐시 적중은 서로 다른 서버를 가리킬 때가 많습니다. 큐가 짧은 A 와 캐시를 가진 B 중 어디로 보낼지 정하려면 둘에 가중치를 매겨야 하는데, 그 가중치를 사람이 손으로 정하면 트래픽 패턴이 바뀔 때마다 다시 맞춰야 합니다.


2. llm-d 는 무엇인가

llm-d 는 쿠버네티스에서 LLM 추론을 분산 서빙하기 위한 스택입니다. 세 조각으로 이해하면 됩니다.

  Envoy 프록시
     |   "이 요청 어디로 보낼까?" 라고 물어봄 (ext-proc 규약)
     v
  EPP (Endpoint Picker)   <-- llm-d 의 스케줄러
     |   "3번 파드로 보내"
     v
  vLLM 파드들 (실제 모델 서버)
  • vLLM — 실제 모델을 돌리는 추론 엔진
  • Gateway API Inference Extension (GAIE) — 쿠버네티스 표준 규약. 프록시가 라우팅 결정을 스스로 하지 않고 외부 컴포넌트에 물어보는 방식을 정의합니다. 이 "물어보는 통신 규약"을 ext-proc(external processing)이라고 합니다
  • EPP (Endpoint Picker) — 그 질문에 답하는 llm-d 의 스케줄러. 이 글에서 다루는 대상입니다

3. Predicted Latency-Based Scheduling: 개념

기존 스케줄링은 관측된 상태(큐 깊이, KV 캐시 사용률)에 사람이 정한 가중치를 곱해 점수를 냈습니다. Predicted Latency 방식은 질문 자체를 바꿉니다.

기존 : "어느 서버가 덜 바빠 보이나?"
변경 : "이 요청을 각 서버에 보내면 실제로 몇 ms 걸릴까?"

XGBoost(의사결정나무 기반의 머신러닝 알고리즘) 모델로 그 시간을 예측합니다. 특징은 온라인 학습이라는 점입니다. 미리 준비한 학습 데이터셋을 쓰는 것이 아니라, 지금 이 클러스터에서 처리되는 요청의 실제 결과를 계속 학습에 되먹입니다.

구성 요소

예측에 필요한 서버들이 EPP 파드 안에 사이드카로 함께 뜹니다. 사이드카란 같은 파드 안에서 주 컨테이너를 보조하는 별도 컨테이너를 말합니다. 실제로 설치하면 이렇게 보입니다.

$ kubectl get pod <epp-pod> -o jsonpath='{.spec.containers[*].name}'
envoy-proxy  epp  training-server  prediction-server-1
컨테이너 역할
envoy-proxy 요청을 받고, EPP 에 "어디로 보낼까" 질의
epp 스케줄러 본체. 플러그인 파이프라인 실행
training-server 완료된 요청을 모아 XGBoost 모델을 주기적으로 재학습
prediction-server-1 학습된 모델로 TTFT/TPOT 예측. 모든 요청이 지나가는 경로에 있어 여기가 느리면 전체가 느려집니다

요청 한 건의 흐름

1. 요청 도착, Envoy 가 받아서 EPP 에 질의
2. EPP 가 prediction-server 에 물어봄
     "이 요청을 파드 1/2/3에 보내면 각각 TTFT·TPOT 몇 ms?"
3. prediction-server 가 파드별 예측값 반환
4. EPP 의 latency-scorer 가 점수를 매겨 파드 선택
5. 선택된 vLLM 파드가 응답
6. 실제로 걸린 시간을 training-server 로 전송하여 모델 갱신

6번이 이 구조의 핵심입니다. 예측이 틀리면 그 오차가 다음 학습에 반영되므로, 트래픽 패턴이 바뀌어도 모델이 따라갑니다.

실제로 어떤 값을 보고 예측하나

학습 서버 로그에서 모델이 쓰는 입력값(피처) 목록을 그대로 확인할 수 있었습니다. 11개입니다.

is_queued, kv_cache_percentage, input_token_length,
num_request_waiting, num_request_running,
prefill_tokens_in_flight, decode_tokens_in_flight,
prefix_cache_score, effective_input_tokens,
prefill_score_bucket, pod_type_cat

여기서 눈여겨볼 점은 큐 길이(num_request_waiting)가 11개 중 하나에 불과하다는 것입니다. 입력 토큰 수, KV 캐시 사용률, 프리픽스 캐시 점수, 현재 처리 중인 prefill/decode 토큰 수가 함께 들어갑니다. 1절에서 말한 "큐 길이에는 요청 비용 정보가 없다"는 문제를, 나머지 10개 피처로 메우는 구조입니다.


4. 실습 환경

항목
쿠버네티스 v1.36.1
GPU NVIDIA RTX PRO 6000 Blackwell Server Edition, 96GB, 1장
드라이버 / CUDA 580.105.08 / 13.0
GPU 분할 HAMi hami-core 로 5분할 (nvidia.com/gpu: 5)
GPU 노드 16 vCPU / 64GB RAM
모델 Qwen/Qwen3-0.6B (인증 불필요), 복제본 3개
vLLM v0.26.0
llm-d 라우터 차트 oci://ghcr.io/llm-d/charts/llm-d-router-standalone:v0

HAMi 는 물리 GPU 1장을 여러 파드가 나눠 쓰게 해주는 도구입니다. 여기서는 96GB 카드 하나를 5개 슬롯으로 쪼개 썼습니다.

전제 두 가지

1. 모델 서버가 2대 이상이어야 합니다.
스케줄링은 후보 중에서 고르는 일이라, 파드가 1개면 고를 것이 없어 관찰할 수 없습니다. GPU 1장을 HAMi 로 나눠 복제본 3개를 띄웠습니다.

2. 모델 서버들의 구성이 서로 같아야 합니다.
같은 모델·같은 설정의 복제본 묶음을 풀(pool) 이라고 부릅니다. 공식 문서는 GPU 종류나 모델 변형, 서빙 설정이 섞인 풀에서는 예측이 부정확해진다고 명시적으로 경고합니다.


5. 설치

5-1. 리포지토리와 환경변수

git clone https://github.com/llm-d/llm-d.git && cd llm-d
export REPO_ROOT=$(realpath $(git rev-parse --show-toplevel))
source ${REPO_ROOT}/guides/env.sh

export GUIDE_NAME="predicted-latency-routing"
export NAMESPACE=llm-d-predicted-latency
export MODEL_NAME="Qwen/Qwen3-0.6B"

env.sh 가 아래 변수를 채워 줍니다. 직접 정의할 필요 없습니다.

ROUTER_STANDALONE_CHART="oci://ghcr.io/llm-d/charts/llm-d-router-standalone"
ROUTER_CHART_VERSION="v0"
GAIE_URL="releases/latest/download"

5-2. Gateway API Inference Extension CRD 설치

CRD(Custom Resource Definition)는 쿠버네티스에 새로운 리소스 종류를 추가하는 정의입니다. llm-d 의 EPP 는 GAIE 규약 위에서 동작하므로, 그 규약이 정의한 리소스를 클러스터가 먼저 알고 있어야 합니다.

kubectl apply -f https://github.com/kubernetes-sigs/gateway-api-inference-extension/${GAIE_URL}/v1-manifests.yaml

설치되면 다음 리소스가 생깁니다.

$ kubectl get crd | grep -i inference
inferencepools.inference.networking.k8s.io   2026-09-16T11:41:44Z

5-3. 네임스페이스와 토큰

kubectl create namespace ${NAMESPACE}

# 인증이 필요 없는 모델이면 값이 비어 있어도 됩니다.
# 다만 매니페스트가 이 시크릿을 참조하므로 "존재" 자체는 필요합니다.
kubectl create secret generic llm-d-hf-token \
    --from-literal="HF_TOKEN=" \
    --namespace "${NAMESPACE}" \
    --dry-run=client -o yaml | kubectl apply -f -

HuggingFace 의 일부 모델은 사용 동의를 받아야 내려받을 수 있고(이를 gated 모델이라 합니다), 그런 모델은 실제 토큰이 필요합니다. Qwen 계열은 gated 가 아니라 빈 토큰으로도 받아집니다.

5-4. 모델 서버 배포 — 가이드 기본값은 그대로 쓸 수 없습니다

공식 가이드가 제공하는 설정은 이렇습니다.

replicas: 10
args:
  - "Qwen/Qwen3-32B"
  - "--tensor-parallel-size=2"
  - "--max-model-len=131072"
resources:
  limits: { cpu: "16", memory: 128Gi, nvidia.com/gpu: 2 }

복제본 10개가 각각 GPU 2장을 쓰므로 GPU 20장이 필요합니다. 대형 클러스터 기준이라, 실습 환경에서는 설정을 직접 덮어써야 합니다.

쿠버네티스에서 기존 매니페스트의 일부만 바꿔 쓰는 방식을 kustomize 오버레이라고 합니다. 원본은 그대로 두고 바꿀 부분만 따로 적는 것입니다.

overlay/kustomization.yaml — 무엇을 기반으로 무엇을 덮어쓸지 선언합니다.

apiVersion: kustomize.config.k8s.io/v1beta1
kind: Kustomization
resources:
  - <llm-d>/guides/optimized-baseline/modelserver/gpu/vllm/base
patches:
  - path: patch-vllm.yaml

overlay/patch-vllm.yaml — 실제로 바꿀 내용입니다.

apiVersion: apps/v1
kind: Deployment
metadata:
  name: optimized-baseline-nvidia-gpu-vllm-decode
spec:
  replicas: 3
  template:
    spec:
      containers:
        - name: modelserver
          args:
            - "Qwen/Qwen3-0.6B"
            - "--disable-access-log-for-endpoints=/health,/metrics,/v1/models"
            - "--tensor-parallel-size=1"
            - "--max-model-len=8192"
            - "--gpu-memory-utilization=0.85"
          resources:
            limits:
              cpu: "4"
              memory: 12Gi
              nvidia.com/gpu: 1
              nvidia.com/gpumem: 12000
            requests:
              cpu: "2"
              memory: 8Gi
              nvidia.com/gpu: 1
              nvidia.com/gpumem: 12000

주의: --gpu-memory-utilization 을 반드시 낮춰야 합니다.

vLLM 은 기본적으로 이 값이 0.9 인데, 이는 "자기에게 보이는 GPU 메모리의 90%를 미리 잡는다"는 뜻입니다. 슬라이스 제한이 없으면 첫 파드가 96GB 중 86GB 를 선점해 나머지 2개가 뜨지 못합니다.

그래서 두 가지를 같이 걸었습니다. HAMi 의 nvidia.com/gpumem: 12000(단위 MB)으로 파드마다 12GB 만 보이게 자르고, 그 안에서 85%를 쓰게 했습니다. 3개 × 12GB = 36GB 로 96GB 안에 여유 있게 들어갑니다.

kubectl apply -n ${NAMESPACE} -k ./overlay

5-5. 라우터(EPP) 배포

helm install ${GUIDE_NAME} ${ROUTER_STANDALONE_CHART} \
    -f ${REPO_ROOT}/guides/recipes/router/base.values.yaml \
    -f ${REPO_ROOT}/guides/${GUIDE_NAME}/router/predicted-latency.values.yaml \
    -n ${NAMESPACE} --version ${ROUTER_CHART_VERSION}

이 기능을 켜는 설정 자체는 단 두 줄입니다.

router:
  latencyPredictor:
    enabled: true          # 이 한 줄이 training/prediction 사이드카를 띄운다

그리고 EPP 가 요청을 처리할 때 거치는 플러그인 순서가 함께 정의됩니다.

plugins:
- type: approx-prefix-cache-producer
- type: predicted-latency-producer          # 예측 서버 호출
- type: prefix-cache-affinity-filter        # 캐시 보유 파드 우선 (임계값 0.8)
- type: latency-scorer                      # 예측 지연으로 점수화
- type: weighted-random-picker              # 최종 선택

6. 설치하다 막힌 지점 세 가지

한 번에 되지 않았습니다. 같은 데서 막히실 분들을 위해 그대로 적습니다.

막힌 지점 1. Helm 차트 내려받기가 403 거부

Error: INSTALLATION FAILED: failed to perform "FetchReference" on source:
GET "https://ghcr.io/v2/llm-d/charts/llm-d-router-standalone/manifests/v0":
response status code 403: denied: denied

curl 로 직접 시도하면 익명 접근 토큰이 잘 발급되는데 helm 만 실패했습니다. 원인은 ~/.config/helm/registry/config.json 에 남아 있던 만료된 인증 정보였습니다. helm 이 그 자격증명을 우선 사용했고, 만료된 자격증명이라 거절당한 것입니다.

# 빈 registry 설정으로 우회
mkdir -p /tmp/helmreg && echo '{}' > /tmp/helmreg/config.json
export HELM_REGISTRY_CONFIG=/tmp/helmreg/config.json

교훈: 익명 접근이 되는 저장소인데 403 이 난다면, 인증 정보가 없어서가 아니라 잘못된 것이 있어서일 수 있습니다.

막힌 지점 2. EPP 파드가 Pending — 기본 리소스 요청이 18 vCPU

0/2 nodes are available: 1 Insufficient memory, 2 Insufficient cpu

EPP 파드 하나가 요구하는 자원의 합계입니다.

컨테이너 CPU 메모리
envoy-proxy 4 8Gi
epp 8 8Gi
training-server 2 4Gi
prediction-server-1 8 4Gi
합계 18 24Gi

16 vCPU 짜리 노드에는 들어갈 수 없는 크기입니다. 실습용으로 줄였습니다.

막힌 지점 3. prediction-server 가 OOMKilled — 이번엔 너무 줄였다

위를 해결하려고 메모리를 2Gi 로 줄였더니 이번에는 컨테이너가 반복해서 죽었습니다. OOMKilled 는 컨테이너가 할당된 메모리 한도를 넘겨 커널에 강제 종료된 상태를 말합니다(종료 코드 137).

lastState.terminated: reason=OOMKilled exitCode=137

로그:
  "Uvicorn running on http://0.0.0.0:8001"
  "Started parent process [1]"        <- 여기서 멈춤, 워커 로그가 없음

부모 프로세스만 뜨고 워커가 전혀 시작되지 못한 모습입니다. 웹 서버(uvicorn)가 노드의 CPU 개수(16)만큼 워커를 띄우려다 2Gi 한도 안에서 메모리가 부족해진 것으로 보입니다. 정확한 원인은 더 파보지 않았고, 메모리를 올리니 해결됐습니다.

최종 실습용 설정값:

router:
  proxy:
    resources:
      requests: { cpu: "300m", memory: 512Mi }
      limits:   { memory: 2Gi }
  epp:
    resources:
      requests: { cpu: "500m", memory: 1Gi }
      limits:   { memory: 4Gi }
  latencyPredictor:
    trainingServer:
      resources:
        requests: { cpu: "500m", memory: 2Gi }
        limits:   { cpu: "2", memory: 4Gi }
      volumeSize: "5Gi"
    predictionServers:
      resources:
        requests: { cpu: "1", memory: 4Gi }      # 4Gi 미만으로 내리면 OOM 발생
        limits:   { cpu: "4", memory: 8Gi }

기동 완료

$ kubectl -n llm-d-predicted-latency get po
optimized-baseline-nvidia-gpu-vllm-decode-...-8mtld   1/1   Running
optimized-baseline-nvidia-gpu-vllm-decode-...-csbws   1/1   Running
optimized-baseline-nvidia-gpu-vllm-decode-...-ztn5d   1/1   Running
predicted-latency-routing-epp-...-vq6lm               4/4   Running

vLLM 파드가 요청을 받을 수 있는 상태가 되기까지 약 9분 걸렸습니다. 대부분 컨테이너 이미지 내려받기와 CUDA 그래프 준비에 쓰인 시간입니다.


7. 요청 보내기

export IP=$(kubectl get service ${GUIDE_NAME}-epp -n ${NAMESPACE} -o jsonpath='{.spec.clusterIP}')

kubectl -n ${NAMESPACE} run curlbox --image=cfmanteiga/alpine-bash-curl-jq \
  --restart=Never -- sleep 900

EPP 서비스는 포트가 세 개입니다. 추론 요청은 80번으로 보냅니다.

port 9002 -> 9002   ext-proc gRPC (Envoy 가 EPP 에 질의하는 통로)
port 9090 -> 9090   메트릭
port 80   -> 8081   추론 API      <- 요청은 여기로
kubectl -n ${NAMESPACE} exec curlbox -- curl -s -X POST http://${IP}:80/v1/completions \
  -H 'Content-Type: application/json' \
  -d '{
    "model": "Qwen/Qwen3-0.6B",
    "prompt": "Explain the difference between prefill and decode.",
    "max_tokens": 64,
    "temperature": 0,
    "stream": true,
    "stream_options": {"include_usage": true}
  }'

응답:

data: {"id":"cmpl-ee2d838c...","object":"text_completion","model":"Qwen/Qwen3-0.6B",
       "usage":{"prompt_tokens":10,"total_tokens":74,"completion_tokens":64}}
data: [DONE]

"stream": true 는 빼면 안 됩니다. 이 구성은 토큰이 하나씩 도착하는 간격을 재서 TPOT 을 학습하기 때문에, 스트리밍이 아니면 학습에 쓸 샘플이 제대로 쌓이지 않습니다.

부하 생성

예측 모델은 트래픽이 있어야 학습됩니다. 그리고 1절에서 본 것처럼 이 기능은 요청 크기가 제각각일 때 의미가 있으므로, 길이를 세 종류로 섞어 150건을 보냈습니다.

짧은 프롬프트(약 3자)  / max_tokens 32
중간 프롬프트(600자)   / max_tokens 64
긴 프롬프트(4000자)    / max_tokens 128
  세 종류를 무작위로, 동시 5건씩, 총 150건

8. 분석 1: 예측 모델이 정말 학습했는가

학습 서버의 로그가 가장 직접적인 증거입니다.

kubectl -n ${NAMESPACE} logs <epp-pod> -c training-server --tail=200 | grep -i "model trained"

부하를 넣기 전에는 학습할 것이 없다고 나옵니다.

Skipping training: only 0 samples (< 10).

부하를 넣은 뒤 — 10초 간격으로 재학습하는 모습이 그대로 보입니다.

11:57:56  TTFT model trained on  64 samples.  MAE = 18.1990,  RMSE = 28.8657
11:58:06  TTFT model trained on 132 samples.  MAE = 20.3510,  RMSE = 28.3740
11:58:16  TTFT model trained on 132 samples.  MAE = 20.3510,  RMSE = 28.3740
11:58:16  TTFT XGBoost trees saved to /models/ttft_trees.json
11:58:16  TPOT XGBoost trees saved to /models/tpot_trees.json

MAE 와 RMSE 읽는 법. 둘 다 예측이 실제와 얼마나 어긋났는지를 나타내는 오차 지표이고, 작을수록 좋습니다.

  • MAE(평균 절대 오차) — 오차의 크기를 단순 평균한 값. "평균적으로 이만큼 틀렸다"
  • RMSE(제곱근 평균 제곱 오차) — 큰 오차에 더 큰 벌점을 주는 방식. MAE 보다 크게 나오면 가끔 크게 틀리는 경우가 있다는 뜻

여기서 읽어낼 수 있는 것:

  • p90 을 목표로 학습합니다. 로그에 quantile 0.9 로 찍히는데, 평균이 아니라 상위 10% 느린 요청까지 맞추겠다는 뜻입니다. 사용자 체감을 좌우하는 것은 평균이 아니라 느린 쪽이기 때문입니다
  • 재학습 주기는 10초입니다(LATENCY_RETRAINING_INTERVAL_SEC: "10"), 학습을 시작하는 최소 샘플 수는 기본값 100 입니다
  • 64샘플일 때 MAE 18.20, 132샘플일 때 20.35. 샘플이 늘었는데 오차가 커졌습니다. 나빠진 것으로 보이지만, 같은 시점의 RMSE 는 28.87 에서 28.37 로 오히려 줄었습니다. 처음 64건은 비슷한 요청이었다가 이후 긴 프롬프트가 섞이면서 다뤄야 할 범위가 넓어진 것으로 보입니다
  • Insufficient samples for ensemble split, using single model only — 샘플이 부족해 여러 모델을 조합하지 않고 단일 모델만 씁니다. 150건은 이 모델이 제 성능을 내기에는 한참 모자란 양입니다

초반 수치만으로 이 기능의 정확도를 판단해서는 안 됩니다. 온라인 학습이라 트래픽이 쌓여야 정확해지고, 공식 문서는 충분히 수렴했을 때 평균 오차율 약 5% 라고 밝히고 있습니다.


9. 분석 2: 스케줄러가 실제로 서버를 "고르고" 있는가

이 실습에서 가장 흥미로웠던 부분입니다. 150건이 3개 파드에 어떻게 나뉘었는지 vLLM 자체 메트릭으로 확인했습니다.

kubectl -n ${NAMESPACE} exec curlbox -- curl -s http://<pod-ip>:8000/metrics \
  | grep -E 'vllm:request_success_total|vllm:prefix_cache_(hits|queries)_total'
파드 처리한 요청 비율 프리픽스 캐시 적중률
...8mtld 82건 54% 96.7% (122,592 / 126,729)
...ztn5d 36건 24% 94.2% (14,688 / 15,585)
...csbws 33건 22% 0% (0 / 66)
합계 151건

라운드로빈이었다면 50 / 50 / 50 이 나와야 합니다. 실제로는 한 파드에 54%가 몰렸습니다. 스케줄러가 균등 분배가 아니라 선택을 하고 있다는 직접적인 증거입니다.

그리고 가장 많이 받은 파드는 프리픽스 캐시 적중률이 가장 높은 파드였습니다. 캐시가 있으면 prefill 을 건너뛰어 TTFT 가 낮아지고, 예측 모델이 그 상태를 입력으로 받아 낮은 지연을 예측하고, 그래서 더 많은 요청이 그리로 갑니다. 요청이 갈수록 캐시는 더 쌓입니다.

특히 눈에 띄는 것은 csbws 의 캐시 적중률이 0% 라는 점입니다. 캐시를 조회한 횟수 자체가 66건뿐입니다. 스케줄러가 이 파드를 "캐시가 없어 느릴 것"으로 판단해 덜 보냈고, 덜 보냈기 때문에 캐시가 쌓이지 않아 계속 불리해지는 스스로 강화되는 순환이 나타난 것으로 보입니다.

이 동작에는 양면이 있습니다.

  • 얻는 것: 캐시를 가진 서버로 요청을 모아 전체 지연을 낮춥니다
  • 잃는 것: 특정 파드에 부하가 쏠려 자원이 고르게 쓰이지 않습니다. 실제로 22%만 받은 파드가 있습니다

prefix-cache-affinity-filteraffinityThreshold: 0.8 이 이 쏠림의 강도를 정하는 값입니다. 캐시 지역성과 균등 분배 중 어디에 무게를 둘지 여기서 조절하게 됩니다.


10. 분석 3: 메트릭 확인 — 여기서는 실패했습니다

공식 문서가 안내하는 확인 방법은 EPP 가 노출하는 Prometheus 메트릭입니다.

메트릭 의미
llm_d_epp_request_ttft_seconds 실제 TTFT
llm_d_epp_request_predicted_ttft_seconds 예측 TTFT
llm_d_epp_request_ttft_prediction_duration_seconds 예측 계산 자체에 걸린 시간
llm_d_epp_request_streaming_tpot_seconds 실제 TPOT
llm_d_epp_request_predicted_tpot_seconds 예측 TPOT
llm_d_epp_request_slo_violation_total SLO 위반 누적 횟수

예측 정확도는 예측값과 실제값을 같은 구간에서 비교하면 됩니다.

histogram_quantile(0.90, sum by (le) (rate(llm_d_epp_request_predicted_ttft_seconds_bucket[5m])))
histogram_quantile(0.90, sum by (le) (rate(llm_d_epp_request_ttft_seconds_bucket[5m])))

다만 이번 실습에서는 이 엔드포인트에 접근하지 못했습니다. 9090 포트가 인증을 요구합니다.

$ curl http://<epp-pod-ip>:9090/metrics
HTTP 401

# ServiceAccount 토큰과 /metrics 접근 권한을 부여한 뒤에도
HTTP 500  "Authentication failed"

별도의 인증 방식이 필요한 것으로 보이나 원인을 끝까지 규명하지 못했습니다. 그래서 정확도 관련 수치는 8절의 학습 서버 로그(MAE/RMSE)로 대체했습니다. 모니터링 스택(monitoring.values.yaml)을 함께 설치하면 ServiceMonitor 를 통해 Prometheus 가 수집하므로, 그 경로라면 이 문제를 겪지 않을 가능성이 높습니다.

SLO 헤더는 이번 구성에서 동작하지 않습니다

x-llm-d-slo-ttft-msx-llm-d-slo-tpot-ms 헤더로 요청마다 지연 목표를 지정할 수 있고, 어떤 파드도 그 목표를 맞추지 못한다고 예측되면 느린 곳에 억지로 보내는 대신 요청을 거절합니다.

다만 이 동작은 predicted-latency-slo.values.yaml 을 써야 켜집니다. 이번에 사용한 기본 구성(predicted-latency.values.yaml)에서는 헤더를 보내도 무시됩니다. 설정 파일 주석에도 그렇게 적혀 있어, 이번 실습에서는 검증하지 못했습니다.


11. 정리하며 배운 것

1. 스케줄링이 실제로 일어납니다. 54% / 24% / 22% 라는 분배는 라운드로빈에서는 나올 수 없는 숫자입니다. 그리고 그 쏠림이 프리픽스 캐시 적중률 순서와 정확히 일치했습니다.

2. 큐 길이는 11개 입력값 중 하나일 뿐입니다. input_token_length, kv_cache_percentage, prefix_cache_score, prefill_tokens_in_flight 같은 값이 함께 들어갑니다. 요청 비용을 직접 나타내는 값들이 포함된다는 점이 기존 방식과의 실질적인 차이입니다.

3. 학습 목표가 평균이 아니라 p90 입니다. 느린 쪽 요청을 줄이는 데 맞춰진 설계입니다.

4. 작은 실습으로 정확도를 논하기는 어렵습니다. 150건은 로그에 Insufficient samples 가 찍히는 수준입니다. 의미 있는 정확도를 보려면 훨씬 많은 트래픽과 긴 관측 기간이 필요합니다.

5. 기본값은 대형 클러스터 기준입니다. 모델 서버는 GPU 20장을, EPP 는 18 vCPU 를 요구합니다. 실습하려면 자원을 직접 줄여야 하는데, 너무 줄이면 이번처럼 OOM 으로 되돌아옵니다.

언제 도입할 만한가

아래는 워크로드가 좋은지 나쁜지를 가르는 기준이 아니라, 이 스케줄링 기능을 켤지 말지를 정하는 기준입니다.

도입하면 이득이 큰 경우

  • 요청 크기의 편차가 큰 트래픽. 짧은 챗봇 질의와 긴 문서 요약이 한 엔드포인트에 섞여 들어오는 경우가 전형적입니다. 큐 길이가 실제 부하를 대변하지 못하는 바로 그 상황이라, 예측이 값어치를 합니다
  • 대화형 요청과 배치 요청을 같은 GPU 풀에서 처리해야 하고, 요청마다 지연 목표를 다르게 걸고 싶은 경우 (SLO 구성이 별도로 필요합니다)

도입해도 이득이 적은 경우 — 켜지 않아도 됩니다

  • 요청 크기가 고른 트래픽. 이런 워크로드가 나쁘다는 뜻이 전혀 아닙니다. 요청 비용이 비슷하면 큐 길이만으로도 이미 판단이 잘 맞기 때문에, 예측을 얹어도 개선폭이 작다는 의미입니다
  • 비용도 함께 보셔야 합니다. 이 기능을 켜면 EPP 파드에 사이드카 2개(training, prediction)가 추가되고, 차트 기본값 기준으로 그 둘만 CPU 10개와 메모리 8Gi 를 요구합니다. 여기에 요청마다 예측 호출이 한 번씩 더 붙습니다. 그 지연은 llm_d_epp_request_ttft_prediction_duration_seconds 로 잴 수 있지만, 이번 실습에서는 메트릭 접근에 실패해 값을 확인하지 못했습니다

도입하면 안 되는 경우 — 앞의 둘과 성격이 다릅니다

  • GPU 종류나 모델 설정이 섞인 풀. 이것은 "효과가 없다"가 아니라 예측 자체가 틀린다는 문제입니다. 공식 문서가 "will produce inaccurate predictions" 라고 명시적으로 경고합니다. 풀이 이질적이라면 먼저 같은 구성끼리 풀을 나누고 나서 적용해야 합니다

결국 판단은 자기 트래픽의 요청 크기 분포를 먼저 확인하는 것에서 시작하는 것이 맞겠습니다. 편차가 작다면 굳이 얹을 이유가 없고, 크다면 얹을 값어치가 있습니다.


참고 자료

같은 문제를 푸는 네 가지 방식(NVIDIA, AMD, Google TPU, Cerebras)


1. 왜 갈라졌나

이 이야기는 두 가지 법칙이 깨진 데서 시작한다.

무어의 법칙(1965, Gordon Moore)은 칩 하나에 들어가는 트랜지스터 수가 약 2 년마다 두 배가 된다 는 것이다. 물리 법칙이 아니라 경험적·경제적 관찰이고, 중요한 건 이게 개수에 대한 이야기이지 속도에 대한 이야기가 아니라는 점이다. 트랜지스터가 많아진다고 저절로 빨라지지는 않는다.

속도를 담당한 건 다른 법칙이었다. 데나드 스케일링(1974, Robert Dennard)은 트랜지스터를 작게 만들면 단위 면적당 전력 소모가 그대로 유지된다 는 것이다. 치수를 k 배 줄이면 이렇게 된다.

항목 변화 결과
면적 1/k² 트랜지스터가 k² 배 더 들어간다
전압 1/k  
전류 1/k  
정전용량 1/k  
지연시간 1/k 클럭을 k 배 올릴 수 있다
트랜지스터당 전력 (C·V²·f) 1/k²  
면적당 전력 그대로 ★ 더 많이, 더 빠르게, 전력은 그대로

마지막 줄이 핵심이다. 두 법칙이 합쳐지면 이렇게 된다.

  주는 것
무어의 법칙 트랜지스터를 더 많이
데나드 스케일링 그 트랜지스터가 더 빠르고, 전력은 안 늘어남

그래서 1980~2000 년대 초에는 코드를 한 줄도 안 고쳐도 2 년마다 프로그램이 두 배 빨라졌다. 아키텍트가 할 일은 공정이 주는 것을 받아 쓰는 일이었다.

데나드가 먼저 무너졌다 (2005 년 전후)

전압을 계속 낮추려면 문턱 전압도 같이 낮춰야 하는데, 그러면 누설 전류가 폭증한다. 그래서 전압 강하가 멈췄고, 위 공식에서 가 더는 안 줄어드니 면적당 전력이 세대마다 올라가기 시작했다 — 전력 벽이다.

결과가 클럭 정체다. Pentium 4 가 2004 년에 이미 3.8 GHz 였는데 20 년 뒤 주류 CPU 도 5~6 GHz 언저리다. 그 전엔 2 년마다 두 배였던 것이 20 년에 1.5 배가 됐다.

업계의 대응은 멀티코어였다. 코어 하나를 더 빠르게 못 만드니 개수를 늘린 것이다. 하지만 그것도 암달의 법칙에 막힌다. 무어의 법칙은 더 늦게 둔화됐다 — 트랜지스터는 여전히 작아지지만 트랜지스터당 비용이 더는 안 떨어지고 주기가 2 년보다 길어졌다.

단일 스레드 성능 증가율로 보면 가장 선명하다.

  Single-thread performance growth, per year
  (Hennessy & Patterson, approx.)

  1986-2003   52%   ################################################
  2003-2011   23%   #####################
  2011-2015   12%   ###########
  2015-2018  3.5%   ###

데나드 붕괴 이후 멀티코어로 한동안 버텼지만, 그것도 한계에 닿은 모습이 그대로 보인다.

그래서 특화로 간다

공정에서 공짜 속도를 못 얻고 범용 멀티코어도 막히면 남는 레버는 하나뿐이다 — 하나의 워크로드만 아주 잘하는 칩을 설계하는 것(도메인 특화 아키텍처). Hennessy 와 Patterson 이 2018 년에 "아키텍처의 캄브리아 대폭발"을 예고한 게 정확히 이 지점이다.

그리고 AI 는 그 대상으로 이상적이었다. 연산이 행렬곱 하나로 지배되고, 경제적 가치가 막대한 워크로드이기 때문이다.

그런데 그 행렬곱이 두 가지 모양이다

성능을 더 얻는 길이 "트랜지스터를 더 작게"에서 "이 워크로드만을 위한 칩을 설계" 로 옮겨간 뒤, 설계자들이 마주한 문제는 이것이었다 — 같은 행렬곱인데 두 가지 다른 모양으로 나타난다.

  Training / Prefill                    Decode
  ------------------                    ------

    [   A   ] x [   B   ] = [   C   ]     [   A   ] x [v] = [v]
      M x K       K x N       M x N         M x K     K x 1  M x 1

    GEMM  (matrix x matrix)               GEMV  (matrix x vector)

    산술 강도(arithmetic intensity) 높음   산술 강도 낮음
    -> 연산에 묶인다 (compute-bound)       -> 대역폭에 묶인다 (bandwidth-bound)
    -> FLOPs 를 원한다                     -> HBM 대역폭을 원한다

가중치 한 번 읽어서 연산을 많이 하면(GEMM) 연산기가 병목이고, 가중치 한 번 읽어서 벡터 하나만 곱하면(GEMV) 메모리가 병목이다. 같은 모델, 같은 가중치인데 단계에 따라 병목이 반대편에 있다.

여기에 메모리 벽이 겹친다. 연산 성능은 지수적으로 올랐는데 메모리 대역폭은 그만큼 못 따라왔다. 그래서 모든 설계 결정이 결국 하나의 질문으로 수렴한다 — 데이터를 어디에 두고, 어떻게 움직일 것인가.

네 회사가 이 질문에 서로 다르게 답했다.


2. NVIDIA GPU — 프로그래밍 가능성에 건다

철학

수천 개의 스레드를 띄워 놓고, 하나가 메모리를 기다리는 동안 다른 것을 실행해 지연을 숨긴다. 고정 기능 회로 대신 CUDA 로 프로그래밍 가능하게 두고, 그 대가(SIMT 세금)를 지불한다.

구조

+---------------------------------------------------------------+
|  Streaming Multiprocessor        (100+ per package)           |
|                                                               |
|   +----------------+   +----------------+                     |
|   | Warp Scheduler |   | Warp Scheduler |   ... x4            |
|   +-------+--------+   +-------+--------+                     |
|           |                    |                              |
|   +-------v--------+   +-------v--------+                     |
|   |   CUDA Cores   |   |  Tensor Core   |                     |
|   |  (element-wise)|   |    (matmul)    |                     |
|   +----------------+   +----------------+                     |
|                                                               |
|   Registers  ->  L1 / SMEM  ->  TMEM (Blackwell+)             |
+---------------------------------------------------------------+
                              |
                          L2 cache
                              |
                            HBM

SM 하나에 최대 64 개의 워프가 상주한다. 하나가 막히면 다른 워프로 갈아타는 식으로 지연을 덮는다.

세대별 변화

세대 연도 핵심
Volta 2017 최초의 Tensor Core
Hopper 2022 비동기 wgmma, Transformer Engine, thread-block cluster
Blackwell 2024 2-SM MMA (256×256×16), 네이티브 FP4, Tensor Memory(TMEM)
Rubin 2026+ HBM4, prefill 분리, 패키지당 1 TB HBM4e / 600 kW 랙

다섯 가지 베팅

  1. 고정 기능이 아니라 CUDA 로 프로그래밍 가능하게
  2. 대규모 멀티스레딩으로 지연을 숨긴다 (SM 당 64 워프 상주)
  3. 행렬곱을 워프 단위로 감싼다 (mma.syncwgmmatcgen05.mma)
  4. 메모리 계층을 비동기로 (TMA, TMEM, mbarrier)
  5. SIMT 세금을 상각한다 — Tensor Core 를 워프 기계장치가 값어치 있을 만큼 크게 만든다

확장

  • 스케일업: NVLink + NVSwitch. 캐시 일관성이 있고 수동 구리 케이블을 쓴다. NVL72 기준 all-to-all 약 130 TB/s (케이블 5,184 가닥)
  • 스케일아웃: ConnectX NIC + InfiniBand. Blackwell 은 GPU 당 800 Gbps, Rubin 은 1.6 Tbps

소프트웨어

PTX/SASS 인트린식부터 cuBLAS·cuDNN·CUTLASS·TensorRT-LLM, 그 위에 PyTorch·JAX 까지. FlashAttention 은 세대마다 손으로 최적화한다. NVIDIA 는 고객 연구소 안에 엔지니어를 상주시킨다 — 이게 생태계 해자의 실체다.


3. Google TPU — 컴파일러가 전부를 안다

철학

하드웨어 스케줄러가 없다. 캐시가 없다. 스레드도 워프도 없다. XLA 컴파일러가 모든 사이클과 모든 바이트를 미리 계획한다. 학습 워크로드는 정적으로 예측 가능하다는 데 건 것이다.

구조 — 시스톨릭 배열

                  activations 가 흘러 들어온다 ->
                    a3    a2    a1
                     |     |     |
                +----v-----v-----v----+
      w  ->     |  w  |  w  |  w  |   |   가중치는 셀에 머문다
      w  ->     |  w  |  w  |  w  |   |   (weight-stationary)
      w  ->     |  w  |  w  |  w  |   |
                +--+-----+-----+------+   256 x 256 = 65,536 MAC
                   |     |     |          per cycle (Ironwood)
                   v     v     v
                  부분합이 아래로 누적된다

가중치를 배열에 고정해 두고 활성값만 흘려보낸다. 명령어 캐시도, 분기 예측기도, 비순차 실행 엔진도 없다.

주변 장치는 이렇게 붙는다:

유닛 역할
MXU 시스톨릭 배열, 밀집 행렬곱
VPU 2D 벡터 element-wise 연산
Scalar Unit 단일 스레드 스칼라
XLU 레인 교차 리덕션
SparseCore (v4+) 임베딩 워크로드 전용
CAE (v8i) decode 단계 collective 리덕션

메모리는 소프트웨어가 관리하는 스크래치패드뿐이다(VMEM/CMEM/SMEM). 322 비트 VLIW 번들에 사이클당 8 개의 기능 슬롯을 컴파일러가 채운다.

세대별 변화

세대 연도 핵심
v1 2015 INT8 추론 전용
v2 2017 BF16 학습, 듀얼 MXU + HBM
v4 2020 광 회선 교환(Palomar), SparseCore
v5 2023 v4 대비 INT8 3.3 배, v5p 는 8,960 칩 팟
Ironwood(v7) 2024 256×256 MXU, 9,216 칩 슈퍼팟, 42.5 ExaFLOPS FP8
v8 2026 네이티브 FP4, 121 ExaFLOPS FP4 (9,600 칩)

다섯 가지 베팅

  1. 밀집 행렬곱은 시스톨릭 배열로
  2. 캐시 대신 소프트웨어 스크래치패드 — 연산은 싸고 메모리는 비싸다
  3. 컴파일러가 스케줄링 — VLIW, 추측 실행 없음, 동적 디스패치 없음
  4. MAC 만 남기고 다 지운다 — 태그·예측기·리오더 버퍼 전부 삭제
  5. 모양이 안 맞는 워크로드는 전용 엔진으로 (SparseCore, CAE)

확장

  스케일업                          스케일아웃
  --------                          ----------
  3D torus (칩 직결)                Virgo  : TPU 간 east-west
  + Palomar 광 회선 교환             Jupiter: north-south / 스토리지
  144 큐브를 가로질러 재구성          Multislice 가 팟을 넘어 확장
  밀리초급으로 토폴로지 변경          Pathways 가 멀티리전 잡을 조율
  캐시 일관성 없음 (메시지 패싱)      47 Pb/s bisection (134,000+ TPU)

소프트웨어

JAX → StableHLO → HLO → LLO → VLIW. XLA 가 연산 융합, 레이아웃 배정, 버퍼 배정, SPMD 분할을 전부 한다. 커널을 직접 쓰고 싶으면 Pallas(TPU 판 Triton)가 탈출구다. 생태계가 구글 자체 제작으로 중앙집중적이다 — Flax NNX, Optax, MaxText, vLLM, Pathways. PyTorch 경로(torch_xla)는 2등 시민이지만 2025 년부터 JAX→XLA 로 낮추는 경로가 통일되며 나아지고 있다.


4. AMD — 코어는 보수적으로, 포장은 공격적으로

철학

Compute Unit 의 모양을 2012 년 GCN 이후 거의 그대로 유지한다. 대신 3D 적층, CPU·GPU 일관성, 개방 표준에 투자한다. 수직 통합이 아니라 파트너십으로 경쟁한다.

구조

+------------------------------------------+
|  Compute Unit                            |
|                                          |
|   SIMD16   SIMD16   SIMD16   SIMD16      |  <- 벡터 엔진 4 개
|      \        \       /        /         |
|       +--------+-----+--------+          |
|                 |                        |
|          +------v-------+                |
|          | Matrix Core  |  (MFMA)        |  <- wave64 하나가 발행,
|          +--------------+                |     4 SIMD 가 4 사이클 협력
|                                          |
|   Scalar Unit        LDS 64 kB           |
+------------------------------------------+
                   |
          L1 vector cache
                   |
      Infinity Cache (MI300X+ 에서 256 MB)
                   |
                  HBM

트레이드오프가 분명하다. 웨이브 단위 집합 연산인 MFMA 는 NVIDIA 의 비동기 Tensor Core 와 달리 VPU 와 겹쳐 돌릴 수 없다. 그래서 소프트웨어 파이프라인 커널에서 행렬곱과 softmax/마스킹을 깔끔하게 겹치지 못한다. 대신 Infinity Cache 와 HBM 용량으로 벌충한다.

세대별 변화

세대 연도 핵심
CDNA 2020 최초 MFMA 행렬 코어, 그래픽 실리콘 제거, 네이티브 BF16
CDNA 2 2021 듀얼 다이 MCM, 풀레이트 FP64
CDNA 3 2023 TSMC SoIC 최초 3D 적층 칩렛, FP8, Infinity Cache 256 MB, MI300A APU
CDNA 4 2025 네이티브 FP4/FP6 (OCP MX), CU 당 FP64 절반으로 — AI 우선 전환
CDNA Next 2026 HBM4, Helios 랙 (UALink 로 72 GPU 스케일업)

다섯 가지 베팅

  1. Compute Unit 모양은 유지하고, 패키징에 투자한다
  2. HBM 용량 우선 — 2021 년 이후 매 세대 NVIDIA 와 같거나 더 많이
  3. 3D 적층을 일찍 (2023 년 SoIC 하이브리드 본딩)
  4. CPU·GPU 일관 APU (MI300A: Zen 4 CCD + CDNA 3 XCD + HBM3)
  5. 개방 표준 — UALink, UET, OCP MX

확장

  • 스케일업: Infinity Fabric 8-GPU OAM 박스(완전 연결, MI350X 기준 GPU 당 1,075 GB/s) → Helios 랙(72 GPU, 출시 시 UALoE 로 260 TB/s, 2027 년부터 네이티브 UALink)
  • 스케일아웃: Ultra Ethernet(UET RDMA) + Pensando Vulcano 800 NIC + Broadcom Tomahawk 6

소프트웨어 — 격차를 솔직히

ROCm 스택(HIP, rocBLAS, MIOpen, RCCL, Composable Kernel). XLA 에 해당하는 것이 없다. PyTorch 는 Triton ROCm 백엔드 + torch.compile 로 1급 지원. FlashAttention 은 Composable Kernel 로 FA2 가 프로덕션, FA3/FA4 는 NVIDIA 를 몇 달 뒤따른다.

동일 워크로드에서 CUDA 대비 10–25% 느리고, FlashAttention-4 꼬리는 여전히 NVIDIA 의 해자 영역이다.


5. Cerebras WSE — 웨이퍼를 자르지 않는다

철학

메모리 벽은 물리 법칙이 아니라 설계 선택이다. HBM 을 아예 버리고 SRAM 만 쓴다. 웨이퍼를 칩으로 자르지 않고 46,225 mm² 한 덩어리로 쓴다.

구조

+=====================================================+
|  One wafer : 46,225 mm2 , 84 reticles stitched      |
|                                                     |
|   [c][c][c][c][c][c][c][c][c][c][c][c][c][c][c]     |
|   [c][c][c][c][c][c][c][c][c][c][c][c][c][c][c]     |
|   [c][c][c][c][c][c][c][c][c][c][c][c][c][c][c]     |
|   [c][c][c][c][c][c][c][c][c][c][c][c][c][c][c]     |
|            ... 900,000 cores in 2D mesh ...         |
|                                                     |
|   per core: ~38k um2 , SRAM 48 kB , FP16 SIMD       |
|   no shared cache / no global mem / no matrix unit  |
+=====================================================+
     총 SRAM 44 GB  ·  온칩 집계 21 PB/s
                        |
              12 x 100 GbE  =  1.2 Tb/s   <- 절벽

코어는 데이터가 오기 전까지 놀고 있다. wavelet 이 도착하면 깨어나고, 8 개의 하드웨어 마이크로스레드가 사이클 단위로 전환한다. 행렬 유닛이 없어서 행렬곱을 브로드캐스트 + FMAC + 메시 리덕션으로 조립한다.

이 구조의 부수 효과가 재밌다 — 행렬 유닛이 없으니 송신 측에서 0 을 건너뛰는 것이 공짜다. 비정형 희소성으로 FP16 밀집 15.8 PFLOPS 가 희소 125 PFLOPS 까지 간다.

숫자 하나로 보는 차이

  바이트당 FLOP 비율
WSE-3 (온웨이퍼 SRAM) 약 1.3
B200 (HBM 에서) 약 0.002

세 자릿수 차이다. 이게 "메모리 벽을 설계로 없앴다"는 주장의 근거다.

Weight Streaming — 학습을 어떻게 하나

44 GB 안에 모델이 안 들어가는데 어떻게 학습하나? 활성값을 SRAM 에 고정하고 가중치를 흘려보낸다.

  MemoryX (DRAM/flash 어플라이언스)
        |  weights stream in
        v
  +-----------------+
  |   WSE wafer     |   활성값은 여기 상주
  |  (activations)  |
  +-----------------+
        |  gradients stream out
        v
  옵티마이저는 CPU 에서

  => 모델 크기는 44 GB 가 아니라 MemoryX 가 결정한다
  => 샤딩 없이 단일 디바이스처럼 프로그래밍
  => 스케일아웃은 SwarmX 로 순수 데이터 병렬

다섯 가지 베팅

  1. 웨이퍼를 자르지 않는다 (다이 이음매당 2,880 GB/s 로 금속 스티칭)
  2. SRAM 만 (44 GB, 업계 최고 대역폭/용량 비)
  3. 데이터플로 코어, 행렬 유닛 없음 — 0 건너뛰기가 공짜로 따라온다
  4. 가중치가 움직이고 활성값은 머문다 — 모델 크기를 칩 메모리에서 분리
  5. 처리량이 아니라 지연을 판다 — 토큰당 가장 빠른 추론, 프리미엄 가격

현실 — 원문이 짚는 한계

원문은 여기서 솔직하다.

항목 상태
공개된 최대 클러스터 64 시스템 (Condor Galaxy 3)
공개된 최대 모델 70B (Jais 2, G42, 2.6T 토큰)
70B 초과 고객 사례 없음
MFU(실효 활용률) 미공개
시스템당 처리량 미공개

독립 측정(Artificial Analysis, 2024-08)에서 Llama 70B decode 가 446–969 tokens/sec 로 B200 의 약 2.4 배지만, 처리량이 아니라 지연 우위다. 프런티어 모델을 웨이퍼당 44 GB 로 서빙하려면 토큰당 비용이 높을 것으로 추정된다.


6. 한눈에 비교

NVIDIA Google TPU AMD Cerebras
철학 대규모 병렬·프로그래밍 가능 행렬 기계·컴파일러 스케줄 보수적 CU·공격적 패키징 단일 웨이퍼·SRAM 전용
연산 추상 워프/스레드 + Tensor Core VLIW 스칼라 + VPU + MXU 웨이브프런트 + MFMA 데이터플로 코어 + wavelet
행렬곱 비동기 (Hopper+) 아니오 (고정 스케줄) 아니오 (웨이브 집합) 아니오 (브로드캐스트+리덕션)
메모리 모델 HW 캐시 + SMEM + TMEM SW 스크래치패드 Infinity Cache + LDS 온웨이퍼 SRAM 만
스케일업 NVLink+NVSwitch (일관성) ICI torus + 광 스위치 (메시지 패싱) Infinity Fabric (IOD 통해 일관성) 웨이퍼 메시 (고정)
수치 형식 FP32→FP16→FP8→FP4 BF16→FP8→FP4 FP16→FP8→FP4/FP6 FP16/BF16 (FP8 아직)
소프트웨어 커널 주도 (CUDA) 컴파일러 주도 (JAX/XLA) 하이브리드 (PyTorch+ROCm) 단일 디바이스 / 레이어 병렬
최대 배포 576 GPU (NVL576) 9,216 TPU (Ironwood) 72 GPU (Helios) 64 시스템

 

스케일업 패브릭을 그림으로

네 회사가 "칩 여러 개를 하나처럼 쓰는" 방법이 근본적으로 다르다.

 

 

NVIDIA NVL72 — 스위치를 통과한다

     G    G    G    G    G    G
      \    \   |    |   /    /
       +------------------+
       |     NVSwitch     |
       +------------------+
      /    /   |    |   \    \
     G    G    G    G    G    G

GPU 끼리 직접 잇지 않는다. 전부 스위치로 들어가고, 스위치가 all-to-all 을 만든다. 그래서 캐시 일관성이 성립하고 — 다른 GPU 의 메모리를 그냥 주소로 읽을 수 있다 — 랙 안 배선이 수동 구리 5,184 가닥으로 끝난다.

 

 

Google Ironwood — 스위치 없이 칩끼리 직결

     +--- c --- c --- c ---+
     |    |     |     |    |
     +--- c --- c --- c ---+
     |    |     |     |    |
     +--- c --- c --- c ---+

이웃한 칩끼리만 연결되고, 양 끝이 감겨 torus 를 이룬다(그림은 3D 중 한 면). 스위치가 없으니 캐시 일관성도 없다 — 대신 메시지 패싱으로 collective 를 돈다. 팟을 넘어갈 때만 광 회선 교환(Palomar)이 끼어들고, 그 스위치는 밀리초급으로 토폴로지 자체를 바꾼다.

 

 

AMD Helios — 작게 완전 연결하고, 그걸 쌓는다

     +---------------------------+
     |  G  G  G  G  G  G  G  G   |
     |  \__\__\__|__/__/__/__/   |
     |        all-to-all         |     OAM box : 8 GPU 완전 연결
     +-------------+-------------+
                   |
           [ UALoE / UALink ]
                   |
              x 9  =  72 GPU

8-GPU 박스 안에서는 서로 완전 연결이고, 박스 바깥은 패브릭으로 잇는다. 일관성은 IOD(I/O 다이)를 통해 유지된다.

 

 

Cerebras WSE-3 — 패브릭이 곧 칩이다

     +=================================+
     |                                 |
     |   900,000 cores  ·  2D mesh     |
     |   one wafer  ·  46,225 mm2      |
     |   no switch  ·  no off-chip hop |
     |                                 |
     +=================================+

연결할 칩이 없다. 웨이퍼 하나가 곧 패브릭이고, 크기가 고정이다. 늘리려면 웨이퍼를 더 사서 이더넷으로 묶는 수밖에 없다.

 

 

  NVIDIA NVL72 Google Ironwood AMD Helios Cerebras WSE-3
묶는 단위 72 GPU 9,216 칩 72 GPU 웨이퍼 1 장
대역폭 130 TB/s all-to-all 68 PB/s aggregate 260 TB/s 21 PB/s on-chip
연결 방식 스위치 경유 칩 직결 (torus) 박스 내 완전 연결 온웨이퍼 메시
일관성 캐시 일관성 없음 (메시지 패싱) IOD 통해 일관성 해당 없음
물리 매체 수동 구리 5,184 가닥 3D torus + 광 회선 교환 UALoE → UALink 금속 스티칭
확장성 랙 단위로 증설 팟·슬라이스로 증설 랙 단위로 증설 고정 (웨이퍼 크기)

7. 그래서 무엇을 고르나

원문의 결론은 "최적 플랫폼은 하나로 수렴하지 않는다" 이다.

  • 칩 하나의 FP8 성능은 약 2 배 범위 안으로 모였다. 이 층위에서는 이미 비슷해졌다.
  • 그런데 시스템 규모는 자릿수가 다르다. 9,216 칩 슈퍼팟과 72 GPU 랙과 64 시스템은 같은 축에 놓을 수 없다.
  • 따라서 최적 플랫폼은 하나로 수렴하지 않고, 무엇을 우선하느냐가 답을 정한다.

네 회사의 베팅이 모두 프런티어 워크로드에서 실제 배포에 성공했다는 점이 이 주장을 뒷받침한다 — Gemini, Llama, xAI, Jais 가 각각 다른 실리콘 위에서 돈다.

소프트웨어 쪽은 좁혀지는 중이다. Triton 과 torch.compile 이 CUDA 와 컴파일러 주도 진영 사이의 간격을 줄이고 있다. 다만 철학의 극은 남는다 — 원문의 표현을 빌리면:

TPU 에서는 컴파일러가 유일하게 의미 있는 인터페이스이고,
GPU 에서는 컴파일러가 여러 인터페이스 중 하나다.


읽고 남는 것

개인적으로 인상적이었던 건 TPU 가 무엇을 "삭제"했는가 였다. 명령어 캐시, 분기 예측기, 리오더 버퍼, 하드웨어 스케줄러 — CPU 가 수십 년에 걸쳐 쌓아온 것을 전부 지우고 MAC 만 남겼다. 그게 가능한 이유가 "학습 워크로드는 정적으로 예측 가능하다"는 단 하나의 가정이고, 그 가정에 회사 전체를 걸었다는 점이 흥미롭다.

반대로 Cerebras 는 모두가 물리적 제약이라고 부르던 것을 설계 선택이라고 부른다. 메모리 벽은 HBM 을 쓰기로 했기 때문에 존재하는 것이고, SRAM 만 쓰면 바이트당 FLOP 비율이 세 자릿수 달라진다 — 대신 44 GB 라는 벽이 새로 생긴다. 벽을 없앤 게 아니라 다른 벽과 바꾼 것에 가깝다.

결국 네 회사 모두 같은 것을 하고 있다. 어떤 제약을 받아들이고 어떤 제약과 싸울지를 고르는 일.


출처

LLM을 한 대만 굴릴 때는 없던 문제들이, 두 대가 되는 순간 한꺼번에 생깁니다.

어느 쪽으로 보낼지, 한 대가 죽으면 어떻게 할지, 프롬프트에 섞여 들어온 개인정보를 어디서 막을지 여러 문제가 발생하죠.

이걸 애플리케이션마다 구현하는 대신 게이트웨이 한 겹으로 모으는 게 LiteLLM Proxy가 하는 일입니다.

이 글은 그 게이트웨이를 Kubernetes 위에 직접 세우고, 마스킹 · 라우팅 · Failover · Auto Routing + 모니터링이 정말 그렇게 동작하는지 하나씩 확인한 기록입니다.

환경 Kubernetes · NVIDIA L40S 1장(HAMi 분할)
구성 LiteLLM Proxy + vLLM ×3 + Presidio ×2
모델 Qwen2.5-0.5B / Qwen2.5-1.5B

최종 구성

                    ┌──────────────────────────────┐
   Client ────────▶ │      LiteLLM Proxy :4000     │
                    │                              │
                    │  마스킹 · 라우팅 · Failover     │
                    │        · Auto Routing        │
                    └───┬──────────────────────┬───┘
                        │                      │
       ┌────────────────┴──────┐               │
       ▼                       ▼               ▼
  vllm-a :8000          vllm-b :8000     vllm-c :8000
  Qwen2.5-0.5B          Qwen2.5-0.5B     Qwen2.5-1.5B
       └──── 그룹: qwen-fast ─────┘     그룹: qwen-quality

                        │
                   Presidio ×2  (분석 + 치환, GPU 불필요)

용어부터

처음 보는 이름이 여럿이라 먼저 짧게 정리합니다.

이름 무엇인가
vLLM LLM 추론 서버. OpenAI 호환 API(/v1/chat/completions)를 제공합니다. 실제로 GPU를 쓰는 쪽
LiteLLM Proxy 여러 LLM 앞에 서는 게이트웨이. 역시 OpenAI 호환이라 클라이언트는 바뀔 게 없습니다
HAMi GPU 한 장을 여러 파드가 나눠 쓰게 해주는 스케줄러. 물리 1장을 N개 슬라이스로 광고합니다
Presidio Microsoft의 PII 탐지·치환 엔진. LiteLLM이 마스킹을 여기에 위임합니다
모델 그룹 LiteLLM에서 같은 이름으로 묶인 백엔드 묶음. 라우팅·Failover의 단위

왜 백엔드가 2 + 1인가

같은 모델을 서빙하는 두 대(vllm-a, vllm-b)는 로드밸런싱과 Failover를 관찰할 대상입니다. 한 대만 있으면 "분산됐다"를 볼 수 없고, 죽였을 때 갈 곳도 없습니다.

더 큰 모델 한 대(vllm-c)는 Auto Routing이 고를 상위 티어입니다. 티어가 하나뿐이면 자동 라우팅은 고를 게 없어 볼 것도 없습니다.


0. 사전 준비

시작 전에 네 가지를 확인합니다. 하나라도 안 되면 뒤에서 막힙니다.

GPU와 분할 방식

kubectl get node -o json | python3 -c "
import json,sys
for n in json.load(sys.stdin)['items']:
    a = n['status']['allocatable']
    g = {k:v for k,v in a.items() if 'nvidia' in k}
    if g: print(n['metadata']['name'], g)
"
gpu2-gpu-pool-mk7jk-hgbdd {'nvidia.com/gpu': '10'}

물리 GPU는 L40S 한 장(45 GiB) 인데 nvidia.com/gpu10으로 나옵니다. HAMi가 10분할해서 광고하고 있다는 뜻입니다.

HAMi가 없다면
순정 NVIDIA device plugin은 GPU를 통째로만 줍니다(nvidia.com/gpu: 1 = 물리 1장). 이 경우 GPU가 3장 있어야 이 실습을 그대로 할 수 있고, 1장뿐이라면 vLLM을 1대만 띄우고 라우팅·Failover 부분은 건너뛰어야 합니다. 뒤에 나오는 nvidia.com/gpumem 줄도 HAMi 전용이라 지워야 합니다.

모니터링 스택

kubectl get prometheus -A
kubectl -n monitoring get svc | grep prometheus

kube-prometheus-stack이 이미 있다고 가정합니다. ServiceMonitor CRD가 없으면 05장이 동작하지 않습니다.

Prometheus가 어떤 ServiceMonitor를 주워가는지도 미리 봅니다. 이게 비어 있으면 전 네임스페이스를 다 수집한다는 뜻이라 편합니다.

kubectl -n monitoring get prometheus \
  -o jsonpath='{.items[0].spec.serviceMonitorSelector}'
{}      ← 비어 있음 = 라벨 조건 없이 전부 수집

비어 있지 않다면, 뒤에서 만들 ServiceMonitor거기 맞는 라벨을 달아야 합니다. 흔한 예로 release: kube-prometheus-stack.

인터넷 접근

모델을 Hugging Face에서 받고, 컨테이너 이미지도 밖에서 당겨옵니다.

kubectl create ns llm-lab
kubectl -n llm-lab run nettest --image=busybox:1.36 --restart=Never -- \
  sh -c 'wget -q -T 8 --spider https://huggingface.co && echo HF-OK'
sleep 15 && kubectl -n llm-lab logs nettest
kubectl -n llm-lab delete pod nettest

폐쇄망이라면 모델을 미리 받아 PVC에 넣고 --model에 로컬 경로를 주는 방식으로 바꿔야 합니다.

디스크

vLLM 이미지가 10 GB 남짓입니다. 모델까지 하면 노드에 여유가 20 GB는 있어야 합니다.

kubectl get node <GPU노드> -o jsonpath='{.status.allocatable.ephemeral-storage}'

1. vLLM 백엔드 세 대 띄우기

가장 오래 걸리는 단계라 먼저 시작해두고 나머지를 준비하는 게 좋습니다.

매니페스트

01-vllm.yaml로 저장합니다.

apiVersion: apps/v1
kind: Deployment
metadata:
  name: vllm-a
  namespace: llm-lab
  labels: {app: vllm, backend: a}
spec:
  replicas: 1
  selector: {matchLabels: {app: vllm, backend: a}}
  template:
    metadata:
      labels: {app: vllm, backend: a}
    spec:
      containers:
      - name: vllm
        image: vllm/vllm-openai:v0.10.1.1
        args:
          - --model=Qwen/Qwen2.5-0.5B-Instruct
          - --served-model-name=qwen-small
          - --max-model-len=4096
          - --gpu-memory-utilization=0.70
          - --disable-log-requests
        env:
        - {name: HF_HOME, value: /models}
        ports:
        - {name: http, containerPort: 8000}
        resources:
          limits:
            nvidia.com/gpu: "1"        # 슬라이스 1개
            nvidia.com/gpumem: "8000"  # 그 슬라이스에 8000 MiB (HAMi 전용)
        volumeMounts:
        - {name: models, mountPath: /models}
        - {name: shm, mountPath: /dev/shm}
        readinessProbe:
          httpGet: {path: /health, port: 8000}
          initialDelaySeconds: 60
          periodSeconds: 10
          failureThreshold: 60      # 모델 다운로드가 길어 넉넉히
      volumes:
      - name: models
        emptyDir: {sizeLimit: 20Gi}
      - name: shm
        emptyDir: {medium: Memory, sizeLimit: 2Gi}   # vLLM에 필요
---
apiVersion: v1
kind: Service
metadata:
  name: vllm-a
  namespace: llm-lab
  labels: {app: vllm, backend: a}
spec:
  selector: {app: vllm, backend: a}
  ports: [{name: http, port: 8000, targetPort: 8000}]

각 줄이 하는 일

  • --served-model-name — 클라이언트에게 보여줄 모델 이름. 실제 HF 경로(Qwen/Qwen2.5-...)와 분리해두면 나중에 모델을 갈아끼워도 호출부가 안 바뀝니다.
  • --max-model-len — 컨텍스트 길이. 길수록 KV 캐시를 많이 먹으므로 슬라이스가 작으면 줄여야 합니다.
  • --gpu-memory-utilization=0.70여기가 헷갈리는 지점. 이건 컨테이너에 보이는 메모리 기준입니다. 8 GiB 슬라이스에 0.70이면 약 5.6 GiB를 씁니다. 물리 45 GiB의 70%가 아닙니다.
  • emptyDir on /dev/shm — vLLM이 공유 메모리를 씁니다. 기본 64 MB로는 부족해 터집니다.
  • failureThreshold: 60 — 모델을 처음 받으면 수 분이 걸립니다. 기본값이면 readiness가 먼저 포기합니다.

b와 c 만들기

ba와 완전히 같고 이름·라벨만 다릅니다. c는 더 큰 모델이라 메모리도 더 줍니다.

# b: a를 그대로 복제
sed 's/vllm-a/vllm-b/g; s/backend: a/backend: b/g' 01-vllm.yaml > 01-vllm-b.yaml

# c: 1.5B 모델 + 12 GiB
sed 's/vllm-a/vllm-c/g; s/backend: a/backend: c/g;
     s|Qwen/Qwen2.5-0.5B-Instruct|Qwen/Qwen2.5-1.5B-Instruct|;
     s/served-model-name=qwen-small/served-model-name=qwen-large/;
     s/gpumem: "8000"/gpumem: "12000"/' 01-vllm.yaml > 01-vllm-c.yaml

kubectl apply -f 01-vllm.yaml -f 01-vllm-b.yaml -f 01-vllm-c.yaml

진행 확인

kubectl -n llm-lab get pod -w

ContainerCreating이 한동안 이어집니다. 이미지 10 GB를 받는 중입니다. HAMi가 스케줄에 관여하는 것도 이벤트로 볼 수 있습니다.

kubectl -n llm-lab get events --sort-by=.lastTimestamp | tail -5
Normal  FilteringSucceed   pod/vllm-b-…  find fit node(gpu2-gpu-pool-…), 1 nodes fit
Normal  BindingSucceed     pod/vllm-b-…  Successfully binding node […]
Normal  Pulling            pod/vllm-b-…  Pulling image "vllm/vllm-openai:v0.10.1.1"

FilteringSucceed / BindingSucceed가 HAMi 스케줄러가 남기는 로그입니다. 이게 안 보이면 HAMi가 아니라 기본 스케줄러가 붙었다는 뜻입니다.

슬라이스가 제대로 배정됐나

kubectl -n llm-lab get pod -o json | python3 -c "
import json,sys
for p in json.load(sys.stdin)['items']:
    for c in p['spec']['containers']:
        g = {k:v for k,v in c.get('resources',{}).get('limits',{}).items() if 'nvidia' in k}
        if g: print(p['metadata']['name'][:30], g)
"
vllm-a-8b74859f7-z4wb7   {'nvidia.com/gpu': '1', 'nvidia.com/gpumem': '8k'}
vllm-b-99bc4dfc6-9mfng   {'nvidia.com/gpu': '1', 'nvidia.com/gpumem': '8k'}
vllm-c-696b7f95dd-npk2f  {'nvidia.com/gpu': '1', 'nvidia.com/gpumem': '12k'}

물리 GPU 1장에 세 개의 추론 서버가 올라갔습니다. 8 + 8 + 12 = 28 GiB로 45 GiB 안에 들어갑니다.


2. 마스킹용 Presidio 올리기

LiteLLM 자체에는 정규식 마스킹이 없습니다. Microsoft Presidio에 위임하는 구조입니다.

  • analyzer — "어디에 무엇이 있는지" 찾습니다 (이름? 카드번호? 이메일?)
  • anonymizer — 찾은 자리를 치환합니다

둘 다 CPU만 쓰므로 GPU 자원이 필요 없습니다.

02-presidio.yaml:

apiVersion: apps/v1
kind: Deployment
metadata: {name: presidio-analyzer, namespace: llm-lab}
spec:
  replicas: 1
  selector: {matchLabels: {app: presidio-analyzer}}
  template:
    metadata: {labels: {app: presidio-analyzer}}
    spec:
      containers:
      - name: analyzer
        image: mcr.microsoft.com/presidio-analyzer:latest
        ports: [{containerPort: 3000}]
        resources:
          requests: {cpu: 200m, memory: 1Gi}
          limits:   {cpu: "2",  memory: 3Gi}   # NLP 모델을 메모리에 올린다
        readinessProbe:
          httpGet: {path: /health, port: 3000}
          initialDelaySeconds: 30
          failureThreshold: 30
---
apiVersion: v1
kind: Service
metadata: {name: presidio-analyzer, namespace: llm-lab}
spec:
  selector: {app: presidio-analyzer}
  ports: [{port: 3000, targetPort: 3000}]
---
apiVersion: apps/v1
kind: Deployment
metadata: {name: presidio-anonymizer, namespace: llm-lab}
spec:
  replicas: 1
  selector: {matchLabels: {app: presidio-anonymizer}}
  template:
    metadata: {labels: {app: presidio-anonymizer}}
    spec:
      containers:
      - name: anonymizer
        image: mcr.microsoft.com/presidio-anonymizer:latest
        ports: [{containerPort: 3000}]
        resources:
          requests: {cpu: 100m, memory: 256Mi}
          limits:   {cpu: "1",  memory: 1Gi}
        readinessProbe:
          httpGet: {path: /health, port: 3000}
          initialDelaySeconds: 15
          failureThreshold: 30
---
apiVersion: v1
kind: Service
metadata: {name: presidio-anonymizer, namespace: llm-lab}
spec:
  selector: {app: presidio-anonymizer}
  ports: [{port: 3000, targetPort: 3000}]
kubectl apply -f 02-presidio.yaml

analyzer는 NLP 모델을 로드하느라 1분쯤 걸립니다.


3. LiteLLM 게이트웨이 — 네 기능이 전부 여기 들어간다

이 글의 핵심 설정 파일입니다. 길지만 블록별로 나눠 설명하겠습니다.

03-litellm.yaml:

apiVersion: v1
kind: Secret
metadata: {name: litellm-secret, namespace: llm-lab}
stringData:
  master-key: "sk-lab-master-key"     # 실습용. 운영에서는 반드시 교체
---
apiVersion: v1
kind: ConfigMap
metadata: {name: litellm-config, namespace: llm-lab}
data:
  config.yaml: |
    model_list:
      # ── 그룹 1: qwen-fast (백엔드 2대) ──
      - model_name: qwen-fast
        litellm_params:
          model: hosted_vllm/qwen-small
          api_base: http://vllm-a.llm-lab.svc.cluster.local:8000/v1
          api_key: "dummy"
        model_info: {id: vllm-a}
      - model_name: qwen-fast
        litellm_params:
          model: hosted_vllm/qwen-small
          api_base: http://vllm-b.llm-lab.svc.cluster.local:8000/v1
          api_key: "dummy"
        model_info: {id: vllm-b}

      # ── 그룹 2: qwen-quality (백엔드 1대) ──
      - model_name: qwen-quality
        litellm_params:
          model: hosted_vllm/qwen-large
          api_base: http://vllm-c.llm-lab.svc.cluster.local:8000/v1
          api_key: "dummy"
        model_info: {id: vllm-c}

      # ── Auto Routing 진입점 ──
      - model_name: smart
        litellm_params:
          model: auto_router/complexity_router
          drop_params: true
          complexity_router_config:
            tiers:
              SIMPLE:    qwen-fast
              MEDIUM:    qwen-fast
              COMPLEX:   qwen-quality
              REASONING: qwen-quality
            complexity_router_default_model: qwen-fast

    router_settings:
      routing_strategy: latency-based-routing
      num_retries: 2
      allowed_fails: 2
      cooldown_time: 20
      fallbacks:
        - qwen-fast: ["qwen-quality"]

    guardrails:
      - guardrail_name: "pii-mask"
        litellm_params:
          guardrail: presidio
          mode: "pre_call"
          default_on: true
          presidio_language: "en"
          pii_entities_config:
            EMAIL_ADDRESS: "MASK"
            PHONE_NUMBER:  "MASK"
            CREDIT_CARD:   "MASK"
            PERSON:        "MASK"
            IP_ADDRESS:    "MASK"

    litellm_settings:
      callbacks: ["prometheus"]
      drop_params: true

    general_settings:
      master_key: os.environ/LITELLM_MASTER_KEY

블록별 해설

model_list — 모델 그룹이 만들어지는 곳

model_name같으면 같은 그룹입니다. 위에서 qwen-fast가 두 번 나오는데, 이게 오타가 아니라 핵심입니다. 클라이언트는 qwen-fast만 부르고, 두 백엔드 중 어디로 갈지는 라우터가 정합니다.

  • model: hosted_vllm/... — LiteLLM에게 "이건 vLLM 서버다"라고 알려주는 접두사
  • api_key: "dummy" — vLLM은 키를 검사하지 않지만 OpenAI SDK가 값을 요구합니다
  • model_info.id — 로그·지표에서 백엔드를 구분할 이름

router_settings — 라우팅과 Failover

설정 하는 일
routing_strategy 그룹 안에서 어느 백엔드를 고를지
num_retries 실패 시 같은 그룹의 다른 백엔드로 재시도
allowed_fails / cooldown_time 연속 실패한 백엔드를 잠시 그룹에서 뺌 (회로 차단)
fallbacks 그룹 전체가 죽었을 때 다른 그룹으로

재시도와 폴백은 층이 다릅니다. 앞은 그룹 , 뒤는 그룹 입니다.

guardrails — 마스킹

mode: pre_call이 핵심입니다. 요청이 vLLM으로 나가기 전에 치환하므로 원문 PII가 백엔드에 도달하지 않습니다. post_call이면 이미 모델이 본 뒤라 늦습니다.

litellm_settings.callbacks: ["prometheus"]

이 한 줄이 /metrics 엔드포인트를 켭니다. 05장에서 씁니다.

프록시 배포

04-litellm-deploy.yaml:

apiVersion: apps/v1
kind: Deployment
metadata:
  name: litellm
  namespace: llm-lab
  labels: {app: litellm}
spec:
  replicas: 1
  selector: {matchLabels: {app: litellm}}
  template:
    metadata: {labels: {app: litellm}}
    spec:
      containers:
      - name: litellm
        image: ghcr.io/berriai/litellm:main-stable
        args: ["--config", "/config/config.yaml", "--port", "4000"]
        env:
        - name: LITELLM_MASTER_KEY
          valueFrom: {secretKeyRef: {name: litellm-secret, key: master-key}}
        # 마스킹 위임 대상 — 이 두 환경변수가 없으면 guardrail이 동작하지 않는다
        - {name: PRESIDIO_ANALYZER_API_BASE,   value: "http://presidio-analyzer.llm-lab.svc.cluster.local:3000"}
        - {name: PRESIDIO_ANONYMIZER_API_BASE, value: "http://presidio-anonymizer.llm-lab.svc.cluster.local:3000"}
        ports: [{name: http, containerPort: 4000}]
        volumeMounts: [{name: config, mountPath: /config}]
        resources:
          requests: {cpu: 200m, memory: 512Mi}
          limits:   {cpu: "2",  memory: 2Gi}
        readinessProbe:
          httpGet: {path: /health/liveliness, port: 4000}
          initialDelaySeconds: 20
          failureThreshold: 30
      volumes:
      - name: config
        configMap: {name: litellm-config}
---
apiVersion: v1
kind: Service
metadata:
  name: litellm
  namespace: llm-lab
  labels: {app: litellm}
spec:
  selector: {app: litellm}
  ports: [{name: http, port: 4000, targetPort: 4000}]
kubectl apply -f 03-litellm.yaml -f 04-litellm-deploy.yaml

전부 뜰 때까지

kubectl -n llm-lab get pod -w
# 6개가 모두 1/1 Running이 될 때까지 (5~10분)

테스트용 클라이언트 파드

앞으로의 모든 테스트를 클러스터 안에서 실행합니다.

kubectl -n llm-lab run cli --image=curlimages/curl:8.10.1 --restart=Never -- sleep 7200

첫 확인 — 모델 그룹이 보이는지.

kubectl -n llm-lab exec cli -- \
  curl -s http://litellm:4000/v1/models \
  -H "Authorization: Bearer sk-lab-master-key"
{"data":[{"id":"qwen-fast",...},{"id":"qwen-quality",...},{"id":"smart",...}]}

백엔드는 3대인데 모델은 3개로 보입니다. qwen-fast가 두 백엔드를 감춘 하나의 그룹이기 때문입니다.


4. 마스킹 검증 — PII가 정말 안 나가나

"마스킹이 켜졌다"와 "원문이 모델에 안 갔다"는 다른 말입니다. 후자를 증명하려면 모델에게 받은 걸 그대로 되풀이하게 시키면 됩니다.

kubectl -n llm-lab exec cli -- sh -c 'cat > /tmp/req.json <<JSON
{"model":"qwen-fast","max_tokens":80,"temperature":0,
 "messages":[{"role":"user","content":"Echo the following back verbatim: My name is John Smith, email john.smith@acme.com, phone 555-123-4567, card 4111-1111-1111-1111, ip 192.168.1.42."}]}
JSON
curl -s http://litellm:4000/v1/chat/completions \
  -H "Authorization: Bearer sk-lab-master-key" \
  -H "Content-Type: application/json" --data @/tmp/req.json'

모델이 되돌려준 문장

My name is [PERSON], email [EMAIL_ADDRESS], phone [PHONE_NUMBER],
card [CREDIT_CARD], and IP address [IP_ADDRESS].

"그대로 되풀이하라"고 했는데 플레이스홀더가 돌아왔습니다. 모델이 본 것 자체가 이미 치환된 문장이었다는 뜻입니다.

응답만 검사하는 방식이었다면 이 증명이 안 됩니다. 응답에서 지우는 건 이미 모델이 본 뒤니까요 — pre_callpost_call의 차이가 여기서 갈립니다.

잘 안 될 때
마스킹이 안 먹으면 LiteLLM 로그에 Presidio 연결 오류가 찍힙니다.
kubectl -n llm-lab logs deploy/litellm | grep -i presidio
대개 PRESIDIO_*_API_BASE 환경변수 오타이거나 analyzer가 아직 안 떴을 때입니다.


5. 라우팅 — 전략마다 분배가 다르다

어디로 갔는지 아는 법

먼저 함정 하나. 응답으로는 어느 백엔드가 처리했는지 알 수 없습니다. LiteLLM은 클라이언트가 요청한 모델명(qwen-fast)을 그대로 돌려줍니다.

그래서 vLLM 각 파드의 카운터를 직접 세는 방식으로 측정합니다.

kubectl -n llm-lab exec cli -- sh -c '
count() {
  curl -s "http://vllm-$1:8000/metrics" \
    | awk "/^vllm:request_success_total/ {s+=\$2} END {printf \"%d\", s+0}"
}

A0=$(count a); B0=$(count b)
echo "시작:  vllm-a=$A0  vllm-b=$B0"

i=1
while [ $i -le 12 ]; do
  curl -s -o /dev/null http://litellm:4000/v1/chat/completions \
    -H "Authorization: Bearer sk-lab-master-key" \
    -H "Content-Type: application/json" \
    -d "{\"model\":\"qwen-fast\",\"max_tokens\":8,\"temperature\":0,\"messages\":[{\"role\":\"user\",\"content\":\"say hi\"}]}"
  i=$((i+1))
done

A1=$(count a); B1=$(count b)
echo "증분:  vllm-a=$((A1-A0))  vllm-b=$((B1-B0))"
'

latency-based-routing · 12건

증분:  vllm-a = 0    vllm-b = 12

오해하기 쉬운 지점
"라우팅이 안 되는 것"처럼 보이지만 정상입니다. latency-based-routing은 최근 응답이 빠른 쪽으로 몰아주는 전략입니다. 순차 요청에서 한쪽이 계속 빠르면 계속 그쪽으로 갑니다. 균등 분배를 기대하고 이 전략을 쓰면 오해합니다.

전략을 바꿔 다시 재보기

# ConfigMap을 꺼내 전략만 바꾸고 다시 넣는다
kubectl -n llm-lab get cm litellm-config -o jsonpath='{.data.config\.yaml}' > /tmp/cfg.yaml
sed -i 's/routing_strategy: latency-based-routing/routing_strategy: simple-shuffle/' /tmp/cfg.yaml
kubectl -n llm-lab create cm litellm-config --from-file=config.yaml=/tmp/cfg.yaml \
  --dry-run=client -o yaml | kubectl apply -f -

# ConfigMap을 바꿔도 프로세스는 다시 읽지 않는다 — 재시작 필요
kubectl -n llm-lab rollout restart deploy/litellm
kubectl -n llm-lab rollout status deploy/litellm

같은 측정을 다시 돌리면 결과가 뚜렷이 갈립니다.

simple-shuffle · 12건

증분:  vllm-a = 9    vllm-b = 3

완벽한 반반은 아닙니다 — 가중 랜덤이라 12건 표본에서 9:3은 충분히 나올 수 있습니다. 중요한 건 양쪽에 퍼졌다는 것입니다.

전략 기준 쓸 자리
simple-shuffle 가중 랜덤 백엔드 성능이 균질할 때
latency-based-routing 최근 응답 지연 성능이 들쭉날쭉할 때
least-busy 진행 중 요청 수 요청별 처리시간 편차가 클 때
usage-based-routing TPM/RPM 여유 쿼터가 걸린 상용 API와 섞어 쓸 때

6. Failover — 한 대, 그리고 그룹 전체

두 층으로 방어합니다. 앞서 본 num_retries(그룹 안)와 fallbacks(그룹 간)입니다. 하나씩 확인합니다.

관측용 스크립트

이번엔 응답의 model 필드를 봅니다. 폴백이 일어나면 요청한 것과 다른 이름이 돌아옵니다.

kubectl -n llm-lab exec cli -- sh -c '
i=1
while [ $i -le 6 ]; do
  m=$(curl -s http://litellm:4000/v1/chat/completions \
    -H "Authorization: Bearer sk-lab-master-key" \
    -H "Content-Type: application/json" \
    -d "{\"model\":\"qwen-fast\",\"max_tokens\":8,\"temperature\":0,\"messages\":[{\"role\":\"user\",\"content\":\"say hi\"}]}" \
    | sed -n "s/.*\"model\":\"\([^\"]*\)\".*/\1/p" | head -1)
  echo "요청 $i → ${m:-실패}"
  i=$((i+1))
done'

1단계 — 그룹 안에서 한 대 제거

kubectl -n llm-lab scale deploy/vllm-a --replicas=0
sleep 15
# 위 스크립트 다시 실행
요청 1 → qwen-fast    요청 4 → qwen-fast
요청 2 → qwen-fast    요청 5 → qwen-fast
요청 3 → qwen-fast    요청 6 → qwen-fast

남은 vllm-b가 전부 받았습니다. 클라이언트 입장에서는 아무 일도 없었습니다.

2단계 — 그룹 전체 제거

kubectl -n llm-lab scale deploy/vllm-b --replicas=0
sleep 15
# 이제 qwen-fast 그룹에 살아있는 백엔드가 없다. 그래도 qwen-fast를 부른다
요청 1 → qwen-large   요청 4 → qwen-large
요청 2 → qwen-large   요청 5 → qwen-large
요청 3 → qwen-large   요청 6 → qwen-large

클라이언트는 qwen-fast를 불렀는데 qwen-large가 응답했습니다. 에러 없이 그룹을 넘어간 것입니다. 이게 fallbacks가 하는 일입니다.

# 복구
kubectl -n llm-lab scale deploy/vllm-a deploy/vllm-b --replicas=1

설계상 유의
폴백은 조용히 다른 모델로 바꿔칩니다. 응답 품질·비용·토큰 한도가 달라지는데 클라이언트는 모릅니다. 상용에서는 폴백 대상을 아무거나 두지 말고, 뒤에 나오는 litellm_deployment_state 지표로 폴백이 발동한 사실 자체를 관측할 수 있게 해두는 편이 낫습니다.


7. Auto Routing — 프롬프트를 보고 티어를 고른다

같은 smart라는 이름으로 요청하되, 프롬프트의 복잡도를 채점해 티어를 고릅니다. 기본 채점기는 휴리스틱이라 외부 임베딩 서비스가 필요 없습니다 — 토큰 수, 코드 포함 여부, 추론 표지 같은 신호를 봅니다.

앞의 설정에서 이 부분입니다.

tiers:
  SIMPLE:    qwen-fast
  MEDIUM:    qwen-fast      # ← 이 매핑이 정책이다
  COMPLEX:   qwen-quality
  REASONING: qwen-quality

티어를 어떻게 확인하나

여기도 함정이 있습니다. 응답의 model 필드는 요청한 이름(smart)을 그대로 돌려줍니다. 어느 티어로 갔는지 응답만으로는 알 수 없습니다.

라우팅 때처럼 백엔드 카운터를 세면 됩니다.

kubectl -n llm-lab exec cli -- sh -c '
mkdir -p /tmp/auto
cat > /tmp/auto/simple.json <<JSON
{"model":"smart","max_tokens":16,"temperature":0,"messages":[{"role":"user","content":"What is 2+2?"}]}
JSON
cat > /tmp/auto/complex.json <<JSON
{"model":"smart","max_tokens":16,"temperature":0,"messages":[{"role":"user","content":"Refactor this Python function to be thread-safe and explain your reasoning step by step: def add(self, k, v): self.cache[k] = v; if len(self.cache) > self.cap: self.cache.pop(next(iter(self.cache)))"}]}
JSON

cnt() { curl -s "http://vllm-$1:8000/metrics" \
  | awk "/^vllm:request_success_total/ {s+=\$2} END {printf \"%d\", s+0}"; }

for f in simple complex; do
  A0=$(cnt a); B0=$(cnt b); C0=$(cnt c)
  curl -s -o /dev/null http://litellm:4000/v1/chat/completions \
    -H "Authorization: Bearer sk-lab-master-key" \
    -H "Content-Type: application/json" --data @/tmp/auto/$f.json
  A1=$(cnt a); B1=$(cnt b); C1=$(cnt c)
  fast=$(( (A1-A0)+(B1-B0) )); qual=$(( C1-C0 ))
  [ $qual -gt 0 ] && t="qwen-quality (1.5B)" || t="qwen-fast (0.5B)"
  echo "$f → $t   [fast+$fast quality+$qual]"
done'

네 가지 프롬프트로 돌려본 결과입니다.

프롬프트 tier score 도착지
"What is 2+2?" SIMPLE −0.150 qwen-fast
"hi" SIMPLE −0.150 qwen-fast
분산 rate limiter 설계 + 실패 모드 추론 MEDIUM 0.325 qwen-fast
코드 스레드 안전 리팩터링 + 단계별 설명 REASONING 0.565 qwen-quality

예상과 달랐던 것

세 번째가 이상합니다. 꽤 복잡한 설계 질문인데 상위 티어로 가지 않았습니다.

DEBUG 로그를 켜면 왜인지 그대로 나옵니다.

kubectl -n llm-lab set env deploy/litellm LITELLM_LOG=DEBUG
kubectl -n llm-lab rollout status deploy/litellm
# 다시 요청한 뒤
kubectl -n llm-lab logs deploy/litellm | grep -i "routing decision"
ComplexityRouter: routing decision cause=heuristic_scorer,
  tier=MEDIUM, score=0.325, signals=('code (api)', 'reasoning (step by ste…')

ComplexityRouter: routing decision cause=reasoning_override,
  tier=REASONING, score=0.565,
  signals=('code (function, def, refactor)', 'reasoning (step by step,
           explain your reasoning)', 'multi-step')

thresholds:  simple_medium=0.15   medium_complex=0.35   …

0.325medium_complex 임계값 0.350.025 모자랐습니다. 그래서 MEDIUM으로 분류됐고, 제 설정에서 MEDIUM은 qwen-fast를 가리킵니다.

라우터가 틀린 게 아니라 제가 쓴 정책이 그렇게 동작한 것입니다. 바꾸고 싶다면 손댈 곳은 둘 중 하나입니다.

  • 매핑 바꾸기: MEDIUM: qwen-quality — 애매하면 좋은 모델로
  • 임계값 낮추기: medium_complex를 낮춰 COMPLEX 판정을 넓힌다

확인이 끝나면 DEBUG는 끕니다. 로그가 매우 많아집니다.

kubectl -n llm-lab set env deploy/litellm LITELLM_LOG-

8. 모니터링 — 두 층을 다 본다

두 층의 지표는 겹치지 않습니다.

  • vLLM — KV 캐시 사용률, 대기 큐 같은 엔진 상태
  • LiteLLM — 어느 백엔드로 몇 건 갔는지, 가드레일이 몇 번 돌았는지

둘 다 /metrics를 내보내므로 ServiceMonitor 두 개면 됩니다.

05-monitoring.yaml:

apiVersion: monitoring.coreos.com/v1
kind: ServiceMonitor
metadata:
  name: litellm
  namespace: llm-lab
  # Prometheus의 serviceMonitorSelector가 비어있지 않다면
  # 여기에 그 라벨을 달아야 한다 (예: release: kube-prometheus-stack)
spec:
  selector: {matchLabels: {app: litellm}}
  endpoints:
    - port: http
      path: /metrics/          # ← 끝 슬래시 필수
      interval: 15s
      authorization:           # ← LiteLLM은 인증 필요
        type: Bearer
        credentials: {name: litellm-secret, key: master-key}
---
apiVersion: monitoring.coreos.com/v1
kind: ServiceMonitor
metadata: {name: vllm, namespace: llm-lab}
spec:
  selector: {matchLabels: {app: vllm}}   # a·b·c 공통 라벨로 3대를 한 번에
  endpoints:
    - port: http
      path: /metrics
      interval: 15s
kubectl apply -f 05-monitoring.yaml

여기서 두 번 막혔습니다

① 인증 — LiteLLM의 /metricsmaster_key로 보호됩니다. 인증 없이 긁으면 401입니다. vLLM 쪽은 열려 있어서 이 블록이 필요 없습니다.

② 끝 슬래시 — 인증을 붙여도 307 → /metrics/ 리다이렉트가 나옵니다. curl-L로 따라가지만 Prometheus 스크레이퍼는 리다이렉트를 따라가지 않습니다. 슬래시를 빼면 타깃은 up인데 시계열만 비는, 진단하기 나쁜 상태가 됩니다.

직접 확인해보면 이렇습니다.

# 인증 없이 → 401
kubectl -n llm-lab exec cli -- curl -s -o /dev/null -w "%{http_code}\n" \
  http://litellm:4000/metrics

# 인증만 → 307
kubectl -n llm-lab exec cli -- curl -s -o /dev/null -w "%{http_code}\n" \
  http://litellm:4000/metrics -H "Authorization: Bearer sk-lab-master-key"

# 인증 + 슬래시 → 200
kubectl -n llm-lab exec cli -- curl -s -o /dev/null -w "%{http_code}\n" \
  http://litellm:4000/metrics/ -H "Authorization: Bearer sk-lab-master-key"

타깃이 붙었는지 확인

kubectl -n monitoring exec prometheus-kube-prometheus-stack-prometheus-0 -c prometheus -- \
  sh -c 'wget -qO- "http://localhost:9090/api/v1/targets?state=active"' \
  | python3 -c "
import json,sys
for t in json.load(sys.stdin)['data']['activeTargets']:
    j = t['labels'].get('job','')
    if 'litellm' in j or 'vllm' in j:
        print(f\"{j:10s} {t['health']:8s} {t.get('lastError','')[:50]}\")
"
litellm    up
vllm-a     up
vllm-b     up
vllm-c     up

실제로 뽑아본 것

Prometheus UI(/graph)나 아래 방식으로 쿼리합니다.

kubectl -n monitoring exec prometheus-kube-prometheus-stack-prometheus-0 -c prometheus -- \
  sh -c "wget -qO- 'http://localhost:9090/api/v1/query?query=<쿼리>'"

라우팅이 분산됐나sum by(api_base)(litellm_deployment_success_responses_total)

http://vllm-a…:8000/v1/chat/completions    9
http://vllm-b…:8000/v1/chat/completions    6
http://vllm-c…:8000/v1/chat/completions    1

마스킹이 돌았나sum(litellm_guardrail_requests_total)

32

백엔드가 살아있나litellm_deployment_state (0 = 정상, Failover 관측용)

qwen-small  vllm-a   0
qwen-small  vllm-b   0
qwen-large  vllm-c   0

엔진 상태vllm:gpu_cache_usage_perc · vllm:num_requests_waiting
vLLM 쪽에서만 보이는 지표입니다.

qwen-large   0.0000892…
qwen-small   0.0000456…
qwen-small   0.0000456…

이 실습은 부하가 거의 없어 KV 캐시가 사실상 0입니다. 실제 운영에서 gpu_cache_usage_perc가 1에 가까워지고 num_requests_waiting이 쌓이기 시작하면 그 백엔드는 이미 포화 상태입니다 — 라우팅 전략을 least-busy로 바꾸거나 슬라이스를 키울 시점입니다.


9. 정리(cleanup)

네임스페이스 하나만 지우면 전부 사라집니다. ServiceMonitor도 함께 지워지므로 Prometheus 쪽에 따로 손댈 게 없습니다.

kubectl delete ns llm-lab

GPU 슬라이스가 회수됐는지 확인합니다.

kubectl get pod -A -o json | python3 -c "
import json,sys
u=0
for p in json.load(sys.stdin)['items']:
    for c in p['spec']['containers']:
        v=c.get('resources',{}).get('limits',{}).get('nvidia.com/gpu')
        if v: u+=int(v)
print('사용 중 슬라이스:', u)
"

막혔을 때 보는 곳

증상 확인할 것
vLLM 파드가 Pending kubectl describe pod → GPU 슬라이스 부족 or HAMi 스케줄러 미동작
vLLM이 OOM으로 죽음 --gpu-memory-utilization 낮추기, --max-model-len 줄이기
vLLM이 /dev/shm 오류 shm emptyDir를 빠뜨렸는지
마스킹이 안 됨 logs deploy/litellm | grep -i presidio, 환경변수 2개 확인
라우팅이 한쪽으로만 정상일 수 있음 — latency-based인지 확인
폴백이 안 됨 fallbacks 대상 그룹이 살아있는지
Prometheus 타깃은 up인데 데이터 없음 /metrics/ 끝 슬래시
타깃이 아예 안 뜸 serviceMonitorSelector 라벨 조건
설정을 바꿨는데 반영 안 됨 ConfigMap 수정 후 rollout restart 했는지

정리

기능 확인된 것 주의할 것
마스킹 원문 PII가 백엔드에 도달하지 않음 Presidio 2개 파드가 별도로 필요
라우팅 전략에 따라 12:0 ↔ 9:3으로 갈림 latency-based는 균등 분배가 아님
Failover 그룹 내 재시도 + 그룹 간 폴백 모두 동작 폴백은 조용히 다른 모델로 바꿔침
Auto Routing 휴리스틱만으로 티어 분류, 외부 의존 없음 임계값 부근에서 갈리며 응답으로는 못 봄
모니터링 두 층 지표를 기존 Prometheus에 그대로 편입 /metrics/ 슬래시와 인증 두 군데

관통하는 교훈이 하나 있습니다. 네 기능 전부 "동작한다"와 "의도대로 동작한다" 사이에 관측이 필요했습니다. 마스킹은 되풀이 프롬프트로, 라우팅과 Auto Routing은 백엔드 카운터로, Failover는 응답 모델명으로 확인했습니다. 응답 본문만 봤다면 라우팅도 Auto Routing도 전부 똑같아 보였을 것입니다.

게이트웨이를 세우는 일의 절반은 설정이고, 나머지 절반은 그 설정이 실제로 무엇을 하는지 볼 수 있게 만드는 일이었습니다.


실습 환경: Kubernetes 테넌트 클러스터 · NVIDIA L40S 1장(HAMi 분할) · vLLM v0.10.1.1 · LiteLLM main-stable · Presidio latest · kube-prometheus-stack 67.3.1

NVIDIA L40S 1장(쿠버네티스 테넌트 클러스터, KubeVirt 워커 VM + HAMi) 위에서 TensorRT-LLM 1.2.1을 처음부터 굴려본 기록입니다.
1부 — NGC 컨테이너 설치 → LLM API quickstart → trtllm-serve OpenAI 호환 서빙
2부 — trtllm-bench 네이티브 벤치 + 엔진 중립 하네스로 성능 실측
3부ModelOpt FP8 PTQ (L40S = Ada = FP8 네이티브)
보너스 — 동일 하드웨어·동일 모델·동일 하네스로 vLLM 0.28.0과 2×2 대조

질문답

TRT-LLM이 vLLM보다 빠른가? 워크로드에 따라 다르다. 짧은 입력(1K)에선 처리량 **+1.8%**로 사실상 동률, 긴 입력(4K)에선 +10%. FP8을 켜면 격차가 **+15% / +30%**로 벌어진다.
가장 큰 레버는? 엔진 선택이 아니라 FP8 양자화. 같은 엔진에서 FP8이 처리량을 +49~51% 올린다. 엔진 교체 효과(+2~10%)보다 한 자릿수 배 크다.
첫 토큰 지연(TTFT)은? 정반대로 갈린다. 짧은 입력은 vLLM이 중앙값 1.8배 빠르고, 긴 입력은 TRT-LLM이 2.9배 빠르다. 기본값 chunked prefill이 서로 반대(vLLM ON / TRT-LLM OFF)인 것이 핵심.
꼬리 지연은? TRT-LLM이 전 구간 우세 (P99 TTFT 13~26% 낮음).
FP8 양자화 비용은? 66.7초. 캘리브레이션 256샘플 포함. 가중치 15GB→8.2GB(−45%), KV 캐시 +24%.
Int8과 비교하면? GPTQ-Int8은 FP16보다 느렸다(커널 미성숙). 반면 FP8은 Ada 텐서코어 네이티브라 전 지표에서 이득. "비트를 줄이면 빨라진다"가 아니라 **"하드웨어가 그 포맷을 네이티브로 아느냐"**가 갈랐다.
 

한 줄 요약: L40S에서 7B를 굴린다면 엔진 고민보다 FP8을 먼저 켜라. 그다음 트래픽이 프리필 무겁고(긴 입력) 꼬리 지연이 중요하면 TRT-LLM, 짧은 입력에 중앙값 TTFT가 중요하면 vLLM.

0. 개념 — TensorRT-LLM은 무엇이 다른가

0.1 두 개의 백엔드

TensorRT-LLM은 이름 때문에 오해받는데, 1.x부터는 TensorRT 엔진을 굽지 않는 경로가 기본이다.

백엔드흐름상태
PyTorch 백엔드 (--backend pytorch) HF 체크포인트 → 곧바로 실행. 커스텀 CUDA 커널 + CUDA Graph + in-flight batching 1.x 기본값, 본 실습에서 사용
TensorRT 엔진 백엔드 trtllm-build로 AOT 컴파일 → 직렬화된 .engine → 로드 레거시. 형상(batch/seq len)을 미리 고정해야 함
 

즉 "TensorRT-LLM = 엔진 빌드가 오래 걸린다"는 인상은 구버전 이야기다. 실습에서 trtllm-serve가 모델 로드부터 서빙 준비까지 약 100초에 끝난다.

0.2 핵심 메커니즘

  • In-flight batching (= continuous batching): 요청이 끝나는 대로 배치 슬롯을 즉시 회수. vLLM의 continuous batching과 같은 개념.
  • Paged KV cache: 블록 단위 KV 관리(기본 tokens_per_block=32). vLLM의 PagedAttention과 대응.
  • Block reuse (enable_block_reuse=True): 프리픽스 캐싱. vLLM의 enable_prefix_caching과 대응.
  • attn_backend='TRTLLM': 자체 fused MHA 커널.
  • CUDA Graph: 배치 크기별로 그래프를 미리 캡처(기본 1~32, 64, 128)해 디코드 스텝의 런치 오버헤드 제거.

0.3 L40S에서의 양자화 — 왜 FP8인가

L40S는 Ada Lovelace, compute capability 8.9다.

포맷L40S 지원비고
BF16/FP16 ✅ 네이티브 기준선
FP8 (E4M3) 텐서코어 네이티브 Ada부터 지원. 본 실습의 주인공
NVFP4 Blackwell(sm100+) 세대 필요
INT8/INT4 ✅ (커널 성숙도 편차) Int8이 FP16보다 느렸던 이유
 

"비트 반 줄이면 반 빨라진다가 아니다"Int8이 느렸던 건 비트 폭 때문이 아니라 커널이 덜 여물어서였고, FP8은 하드웨어가 네이티브로 아는 포맷이라 전 지표에서 이득이 난다.

# 내 GPU가 FP8을 네이티브로 아는지 확인 (8.9 이상이면 FP8 OK)
nvidia-smi --query-gpu=name,compute_cap --format=csv,noheader
# NVIDIA L40S, 8.9

1. 환경 & 설치

1.1 환경

  • 쿠버네티스 테넌트 클러스터(KubeVirt 워커 VM), NVIDIA L40S 1장HAMi로 할당 ([nvidia.com/gpu](https://nvidia.com/gpu): 1, [nvidia.com/gpumem](https://nvidia.com/gpumem): 44000, [nvidia.com/gpucores](https://nvidia.com/gpucores): 100)
  • 드라이버 580.105.08 / CUDA 13.0
  • TensorRT-LLM 1.2.1 (nvcr.io/nvidia/tensorrt-llm/release:1.2.1), 대조군 vLLM 0.28.0 (vllm/vllm-openai:latest)
  • 모델 Qwen/Qwen2.5-7B-Instruct (BF16, TP=1)
  • 노드: 16 vCPU / 64GB RAM / 디스크 150GB

⚠️ 트러블슈팅 ① — 디스크 100GB로는 시작조차 못 한다

첫 시도에서 워커 노드(디스크 100GB)가 DiskPressure로 파드를 전부 Evict했다.

항목크기
기존 OS/k8s 18GB
TRT-LLM 이미지 (압축 19.5GB → 언팩) ~50GB
Qwen2.5-7B BF16 15GB
kubelet eviction 임계 ~15GB
합계 ~98GB / 100GB
 

Evicted: The node was low on resource: ephemeral-storage → 노드 디스크를 150GB로 리사이즈해서 해결. FP8 체크포인트(8.2GB)까지 얹어도 여유 30GB 남는다. TRT-LLM 컨테이너는 vLLM(~16GB)의 3배라는 점을 용량 산정에 반드시 반영해야 한다.

1.2 이미지 확인 — 문서의 태그를 그대로 믿지 말 것

공식 설치 가이드는 nvcr.io/nvidia/tensorrt-llm/release:1.3.0rc25를 안내하지만, 실제 레지스트리엔 그 태그가 없었다. 익명으로 태그를 직접 조회하는 편이 확실하다:

R=nvidia/tensorrt-llm/release
T=$(curl -s "https://nvcr.io/proxy_auth?scope=repository:$R:pull" \
    | python3 -c 'import sys,json; print(json.load(sys.stdin)["token"])')

# 버전 태그만 추출
curl -s -H "Authorization: Bearer $T" "https://nvcr.io/v2/$R/tags/list" \
 | python3 -c 'import sys,json,re; t=json.load(sys.stdin)["tags"]; print("\n".join([x for x in t if re.match(r"^\d+\.\d+\.\d+$",x)]))'
# 0.20.0 / 0.21.0 / 1.0.0 / 1.1.0 / 1.2.0 / 1.2.1   ← 최신 stable = 1.2.1

# 이미지 크기(압축 기준) 미리 확인
curl -s -H "Authorization: Bearer $T" \
  -H 'Accept: application/vnd.docker.distribution.manifest.v2+json' \
  "https://nvcr.io/v2/$R/manifests/1.2.1" \
 | python3 -c 'import sys,json; d=json.load(sys.stdin); print(f"{sum(l[\"size\"] for l in d[\"layers\"])/2**30:.1f} GB")'
# 19.5 GB

NGC의 TensorRT-LLM release 저장소는 NGC API 키 없이 익명 pull이 된다(위 proxy_auth가 그 증거).

1.3 워크벤치 파드

# trtllm-shell.yaml — 실습용 셸 파드
apiVersion: v1
kind: Pod
metadata: { name: trtllm-shell, namespace: trtllm }
spec:
  restartPolicy: Never
  nodeSelector: { nvidia.com/gpu.present: "true" }   # ← nodeName 쓰지 말 것 (아래 트러블슈팅 ②)
  containers:
  - name: trtllm
    image: nvcr.io/nvidia/tensorrt-llm/release:1.2.1
    command: ["bash","-lc","sleep infinity"]          # ← -l 필수 (아래 트러블슈팅 ③)
    env: [ { name: HF_HOME, value: /hf } ]
    resources:
      limits: { nvidia.com/gpu: "1", nvidia.com/gpumem: "44000", nvidia.com/gpucores: "100", cpu: "12", memory: 48Gi }
      requests: { cpu: "4", memory: 16Gi }
    volumeMounts:
      - { name: hf,  mountPath: /hf }
      - { name: shm, mountPath: /dev/shm }            # ← MPI/NCCL용
  volumes:
    - { name: hf,  hostPath: { path: /var/lib/hf-cache, type: DirectoryOrCreate } }
    - { name: shm, emptyDir: { medium: Memory, sizeLimit: 16Gi } }
kubectl create ns trtllm
kubectl apply -f trtllm-shell.yaml
kubectl -n trtllm exec trtllm-shell -- bash -lc 'python3 -c "import tensorrt_llm; print(tensorrt_llm.__version__)"'
# TensorRT-LLM 1.2.1

⚠️ 트러블슈팅 ② — HAMi는 nodeName 직접 지정을 거부한다

Error from server (Forbidden): admission webhook "vgpu.hami.io" denied the request: pod has node assigned

spec.nodeName으로 노드를 못 박으면 기본 스케줄러를 우회하게 되고, HAMi는 자기 스케줄러(hami-scheduler)를 거치지 않은 GPU 파드를 막는다. GPU를 쓰는 파드는 반드시 nodeSelector/affinity로 유도하고 배치는 스케줄러에 맡겨야 한다. (GPU를 안 쓰는 보조 파드는 nodeName을 써도 무방하다.)

⚠️ 트러블슈팅 ③ — command:로 바이너리를 직접 부르면 라이브러리를 못 찾는다

Deployment에 command: ["trtllm-serve"]로 적었더니 기동 즉시 죽었다: ImportError: libnvonnxparser.so.10: cannot open shared object file: No such file or directory

이 라이브러리는 /usr/local/tensorrt/lib에 분명히 존재한다. 문제는 그 경로를 LD_LIBRARY_PATH에 넣어주는 주체가 이미지의 ENTRYPOINT/프로파일이라는 것. 쿠버네티스에서 command:를 주면 ENTRYPOINT가 교체되면서 그 환경 설정이 통째로 사라진다.

# ❌ 라이브러리 경로 소실
command: ["trtllm-serve"]
args: ["Qwen/Qwen2.5-7B-Instruct", "--port=8000"]

# ✅ 로그인 셸을 거쳐 프로파일을 로드
command: ["bash","-lc"]
args: ["exec trtllm-serve Qwen/Qwen2.5-7B-Instruct --host 0.0.0.0 --port 8000 --backend pytorch"]

exec를 붙여야 셸이 PID 1을 넘겨줘서 SIGTERM이 서버까지 전달된다(graceful shutdown).

2. Quickstart — LLM API

# quickstart.py
from tensorrt_llm import LLM, SamplingParams

def main():                                  # ← 이 감싸기가 필수 (아래 트러블슈팅 ④)
    prompts = [
        "Hello, my name is",
        "The capital of France is",
        "The future of AI is",
    ]
    sampling_params = SamplingParams(temperature=0.8, top_p=0.95)

    llm = LLM(model="TinyLlama/TinyLlama-1.1B-Chat-v1.0")
    outputs = llm.generate(prompts, sampling_params)

    for output in outputs:
        print(f"Prompt: {output.prompt!r}, Generated text: {output.outputs[0].text!r}")

if __name__ == "__main__":
    main()
kubectl -n trtllm exec trtllm-shell -- bash -lc 'python3 /tmp/quickstart.py'
[TensorRT-LLM][INFO] Number of tokens per block: 32.
[TensorRT-LLM][INFO] [MemUsageChange] Allocated 35.84 GiB for max tokens in paged KV cache (1708256).
Prompt: 'Hello, my name is', Generated text: 'John Smith. I am a male and I love to play video games...'
Prompt: 'The capital of France is', Generated text: 'Paris.'
Prompt: 'The future of AI is', Generated text: "something that everyone will see in their own lives..."

1.1B 모델인데 KV 캐시로 35.84 GiB(170만 토큰)를 잡는 데 주목. 기본 free_gpu_memory_fraction=0.9가 남는 메모리를 거의 다 KV로 가져간다. 한 GPU에 여러 모델을 얹을 생각이라면 이 값을 반드시 낮춰야 한다.

⚠️ 트러블슈팅 ④ — 노트북에선 되는데 스크립트로 옮기면 터진다

The main script or module attempted to spawn new MPI worker processes...

LLM 생성자가 MPI 워커 프로세스를 spawn하는데, 가드가 없으면 자식이 모듈을 다시 import하며 무한 재귀 spawn이 된다. 주피터 셀에선 __name__이 "__main__"이라 안 걸리고, .py로 옮기는 순간 터진다. 노트북 예제를 프로덕션 스크립트로 옮길 때 가장 먼저 밟는 지뢰.

3. trtllm-serve — OpenAI 호환 서빙

# engine.yaml — 서버 + Service (실험마다 args만 교체)
apiVersion: v1
kind: Service
metadata: { name: engine, namespace: trtllm }
spec:
  selector: { app: engine }
  ports: [ { port: 8000, targetPort: 8000 } ]
---
apiVersion: apps/v1
kind: Deployment
metadata: { name: engine, namespace: trtllm }
spec:
  replicas: 1
  strategy: { type: Recreate }          # GPU 1장 → 롤링 불가
  selector: { matchLabels: { app: engine } }
  template:
    metadata: { labels: { app: engine } }
    spec:
      nodeSelector: { nvidia.com/gpu.present: "true" }
      containers:
      - name: engine
        image: nvcr.io/nvidia/tensorrt-llm/release:1.2.1
        command: ["bash","-lc"]
        args: ["exec trtllm-serve Qwen/Qwen2.5-7B-Instruct --host 0.0.0.0 --port 8000 --backend pytorch"]
        env: [ { name: HF_HOME, value: /hf } ]
        ports: [ { containerPort: 8000 } ]
        resources:
          limits: { nvidia.com/gpu: "1", nvidia.com/gpumem: "44000", nvidia.com/gpucores: "100", cpu: "12", memory: 48Gi }
          requests: { cpu: "4", memory: 16Gi }
        readinessProbe: { httpGet: { path: /health, port: 8000 }, initialDelaySeconds: 30, periodSeconds: 10, failureThreshold: 180 }
        volumeMounts:
          - { name: hf,  mountPath: /hf }
          - { name: shm, mountPath: /dev/shm }
      volumes:
        - { name: hf,  hostPath: { path: /var/lib/hf-cache, type: DirectoryOrCreate } }
        - { name: shm, emptyDir: { medium: Memory, sizeLimit: 16Gi } }
kubectl apply -f engine.yaml
kubectl -n trtllm rollout status deploy/engine --timeout=1200s   # 최초 기동 ~100초 (모델 캐시 있을 때)

kubectl -n trtllm exec bench -- curl -s http://engine:8000/v1/models
# {"object":"list","data":[{"id":"Qwen/Qwen2.5-7B-Instruct","object":"model","owned_by":"tensorrt_llm"}]}

kubectl -n trtllm exec bench -- curl -s http://engine:8000/v1/chat/completions \
  -H 'Content-Type: application/json' -d '{
    "model":"Qwen/Qwen2.5-7B-Instruct",
    "messages":[{"role":"user","content":"What is model quantization? Answer in one sentence."}],
    "max_tokens":60,"temperature":0}'

엔드포인트는 /v1/chat/completions, /v1/completions, /health, /metrics, /version.

3.1 실제 적용된 기본값 — 로그가 정답이다

trtllm-serve는 기동 시 최종 확정된 설정 객체를 통째로 로그에 찍는다. 문서보다 이쪽이 정확하다.

kubectl -n trtllm logs -l app=engine --tail=1200 | grep -oE \
  "max_batch_size=[0-9]+|max_num_tokens=[0-9]+|free_gpu_memory_fraction=[0-9.]+|max_tokens=[0-9]+|enable_chunked_prefill=[A-Za-z]+|enable_block_reuse=[A-Za-z]+|tokens_per_block=[0-9]+|attn_backend='[A-Z]+'"

 

파라미터 TRT-LLM 1.2.1 기본값 vLLM 0.28.0 기본값
KV 메모리 노브 free_gpu_memory_fraction=0.9

(가중치 적재 후 남은 메모리의 90%)
gpu_memory_utilization=0.92

(전체 GPU 메모리의 92%, 가중치 포함)
KV 캐시 크기 (실측) 446,681 tok (25.6GB) 431,024 tok
tokens_per_block / block_size 32 auto (기본 None)
enable_chunked_prefill False ⚠️ True ⚠️
프리픽스 캐싱 enable_block_reuse=True enable_prefix_caching=True
max_batch_size 2048 auto
max_num_tokens 8192 auto
어텐션 백엔드 TRTLLM FlashAttention/FlashInfer
스케줄러 정책 GUARANTEED_NO_EVICT
 

메모리 노브는 이름도 의미도 다르다. TRT-LLM의 0.9는 가중치를 올리고 남은 메모리 기준이고, vLLM의 0.92는 전체 메모리 기준(가중치 포함)이다.

결과적으로 잡힌 KV 캐시가 446,681 vs 431,024 토큰으로 3.6% 차이에 그쳤다. 운이 좋았던 것이며, 덕분에 별도 보정 없이 "기본값 대 기본값"으로 붙여도 메모리 예산이 대등했다. 반면 chunked prefill 기본값이 정반대다. 이 한 줄이 뒤에서 TTFT 결과를 완전히 갈라놓는다.

4. 벤치마크 방법론 — 엔진이 달라도 공정하게

부하 생성기를 vllm bench serve 하나로 통일했다. TRT-LLM도 OpenAI 호환 서버를 제공하므로 같은 클라이언트로 테스트를 수행할 수 있다.

┌──────────────┐    OpenAI /v1/completions    ┌──────────────────────────┐
│  bench 파드   │ ───────────────────────────► │  engine Service :8000     │
│ vllm bench    │                              │  ├ trtllm-serve  (TRT-LLM)│
│   serve       │ ◄─────────────────────────── │  └ vllm serve    (vLLM)   │
│ (GPU 미점유)  │      동일 지표 정의            └──────────────────────────┘
└──────────────┘

핵심은 클라이언트를 서버 파드 바깥에 두는 것. 양쪽 모두 별도 클라이언트 파드로 조건을 맞췄다.

# 공통 벤치 커맨드 — 엔진과 무관하게 동일
vllm bench serve --backend openai --endpoint /v1/completions \
  --base-url http://engine:8000 --model Qwen/Qwen2.5-7B-Instruct \
  --dataset-name random --random-input-len 1024 --random-output-len 256 \
  --num-prompts 300 --max-concurrency 128 \
  --ignore-eos --temperature 0 --seed 12345 \
  --percentile-metrics ttft,tpot,itl,e2el --metric-percentiles 90,99

공정 비교 장치:

  • --ignore-eos → 출력이 항상 정확히 256토큰. (20요청 스모크 테스트에서 Total generated tokens: 5120 = 20 × 256 확인)
  • --temperature 0 + --seed 12345 → 결정적·동일 프롬프트
  • --max-concurrency N → closed-loop(항상 N개 in-flight)
  • 두 워크로드로 프리필/디코드 지배 상황을 분리:
워크로드 입력 출력 요청 동시성 성격
W1 1024 256 300 128 디코드 지배
W2 4096 256 200 64 프리필 무거움
 

5. TRT-LLM vs vLLM — BF16

W1 (in 1024 / out 256, 300요청, 동시성 128)

엔진 req/s out tok/s 총 tok/s TTFT(중앙) TTFT(p99) TPOT(중앙) ITL(중앙)
TRT-LLM 1.2.1 6.15 1574 7872 1683ms 8696ms 70.1ms 38.7ms
vLLM 0.28.0 6.04 1547 7733 925ms 10496ms 70.9ms 37.4ms
차이 +1.8% +1.8% +1.8% 1.82배 느림 −17% −1% +4%
 

W2 (in 4096 / out 256, 200요청, 동시성 64)

엔진 req/s out tok/s 총 tok/s TTFT(중앙) TTFT(p99) TPOT(중앙) ITL(중앙)
TRT-LLM 1.2.1 1.96 501 8521 1041ms 19233ms 119.6ms 49.0ms
vLLM 0.28.0 1.78 456 7751 3007ms 22130ms 126.7ms 46.6ms
차이 +10.1% +9.9% +9.9% 2.89배 빠름 −13% −5.6% +5%
 

해석

  1. 처리량은 프리필이 무거울수록 TRT-LLM이 유리하다. W1에서 +1.8%(측정 노이즈 수준)였던 격차가 W2에서 +10%로 벌어진다. 긴 컨텍스트 프리필에서 TRT-LLM의 fused MHA 커널이 이득을 낸다.
  2. TTFT 중앙값이 워크로드에 따라 완전히 뒤집힌다.
    • W1: vLLM이 1.82배 빠름 (925 vs 1683ms)
    • W2: TRT-LLM이 2.89배 빠름 (1041 vs 3007ms) 원인은 chunked prefill 기본값이 반대라는 것이다. vLLM은 기본 ON이라 프리필을 잘게 쪼개 새 요청의 첫 토큰을 빨리 뽑는다(입력이 짧은 1K일 때 유리). 하지만 입력이 4K로 커지면 청크가 여러 스텝에 걸쳐 디코드와 계속 섞이며 오히려 첫 토큰이 늦어진다. TRT-LLM은 기본 OFF라 4K 프리필을 한 번에 처리해 중앙값이 훨씬 좋다.
  3. 꼬리(P99 TTFT)는 TRT-LLM이 전 구간 우세 (−13~17%). GUARANTEED_NO_EVICT 스케줄러 정책이 한번 받은 요청을 선점당하지 않게 보장하는 것과 맞닿아 있다. SLO를 P99로 잡는 서비스라면 중앙값만 보고 vLLM을 고르면 안 된다.
  4. TPOT/ITL은 BF16에서 사실상 동률(±5%).

6. FP8 양자화 — ModelOpt PTQ

L40S가 FP8을 네이티브로 지원하므로, TensorRT Model Optimizer(ModelOpt)로 직접 양자화한다. 컨테이너에 이미 nvidia-modelopt 0.37.0이 들어있다.

# ptq_fp8.py — Qwen2.5-7B-Instruct → FP8 (weight+activation per-tensor)
import torch, time
from transformers import AutoModelForCausalLM, AutoTokenizer
import modelopt.torch.quantization as mtq
from modelopt.torch.export import export_hf_checkpoint

MODEL = "Qwen/Qwen2.5-7B-Instruct"
OUT   = "/hf/qwen2.5-7b-fp8"
NCALIB, SEQLEN = 256, 512

def main():
    tok   = AutoTokenizer.from_pretrained(MODEL)
    model = AutoModelForCausalLM.from_pretrained(MODEL, torch_dtype=torch.bfloat16, device_map="cuda")

    # 캘리브레이션: 활성값 스케일(amax)을 관측하기 위한 실제 데이터가 필요
    from datasets import load_dataset
    ds = load_dataset("cnn_dailymail", "3.0.0", split=f"train[:{NCALIB}]")
    texts = [x["article"] for x in ds]

    def forward_loop(m):
        for t in texts:
            ids = tok(t, return_tensors="pt", truncation=True, max_length=SEQLEN).input_ids.cuda()
            m(ids)

    model = mtq.quantize(model, mtq.FP8_DEFAULT_CFG, forward_loop)

    with torch.inference_mode():
        export_hf_checkpoint(model, export_dir=OUT)   # ModelOpt 통합 HF 체크포인트
    tok.save_pretrained(OUT)

if __name__ == "__main__":
    main()
kubectl -n trtllm exec quant -- bash -lc 'python3 /tmp/ptq_fp8.py'
Registered Qwen2Attention to _QuantAttention for KV Cache quantization
Inserted 675 quantizers
  calib 256/256
[quantize] 31.0s
[export]    8.5s → /hf/qwen2.5-7b-fp8
[total]    66.7s

총 66.7초. 7B 모델 FP8 PTQ가 1분대에 끝난다. FP8은 per-tensor 스케일만 구하면 되고 가중치 재구성 최적화가 따로 필요 없다.

# 산출물 확인
du -sh /hf/qwen2.5-7b-fp8          # 8.2G  (BF16 15G 대비 −45%)
cat /hf/qwen2.5-7b-fp8/hf_quant_config.json
{
    "producer": { "name": "modelopt", "version": "0.37.0" },
    "quantization": {
        "quant_algo": "FP8",
        "kv_cache_quant_algo": null,
        "exclude_modules": ["lm_head"]
    }
}

kv_cache_quant_algo: null — KV 캐시는 양자화하지 않았다(가중치·활성값만 FP8). lm_head는 출력 분포 왜곡을 피하려 제외한다.

서빙

# TRT-LLM
trtllm-serve /hf/qwen2.5-7b-fp8 --host 0.0.0.0 --port 8000 --backend pytorch

# vLLM — hf_quant_config.json 을 읽어 quantization=modelopt 로 자동 감지
vllm serve /hf/qwen2.5-7b-fp8 --served-model-name qwen2.5-7b-fp8 --host 0.0.0.0 --port 8000

같은 체크포인트를 두 엔진이 모두 그대로 사용할 수 있다.

메모리 효과

가중치KV 캐시(TRT-LLM)KV 캐시(vLLM)
BF16 15GB 446,681 tok (25.6GB) 431,024 tok
FP8 8.2GB (−45%) 553,718 tok (+24%) 563,312 tok (+31%)
 

고정된 free_gpu_memory_fraction=0.9 아래에서 가중치가 줄어든 만큼 KV로 전환되어 동시 수용력이 늘어난다.

7. 2×2 종합 — 엔진 × 정밀도

W1 (in 1024 / out 256, 300요청, 동시성 128)

엔진 정밀도 req/s out tok/s TTFT(중앙) TTFT(p99) TPOT(중앙) ITL(중앙) KV(tok)
TRT-LLM BF16 6.15 1574 1683ms 8696ms 70.1ms 38.7ms 446,681
vLLM BF16 6.04 1547 925ms 10496ms 70.9ms 37.4ms 431,024
TRT-LLM FP8 9.27 2372 1339ms 5985ms 45.2ms 28.2ms 553,718
vLLM FP8 8.03 2055 735ms 8119ms 54.5ms 27.7ms 563,312

W2 (in 4096 / out 256, 200요청, 동시성 64)

엔진 정밀도 req/s 총 tok/s TTFT(중앙) TTFT(p99) TPOT(중앙) ITL(중앙)
TRT-LLM BF16 1.96 8521 1041ms 19233ms 119.6ms 49.0ms
vLLM BF16 1.78 7751 3007ms 22130ms 126.7ms 46.6ms
TRT-LLM FP8 2.92 12724 693ms 10784ms 79.5ms 37.2ms
vLLM FP8 2.25 9783 2326ms 14401ms 92.4ms 36.1ms
 

핵심 분석

① 정밀도 > 엔진. 처리량 개선 폭을 보면 명확하다.


W1 W2
FP8 전환 (TRT-LLM 내) +50.7% +49.0%
FP8 전환 (vLLM 내) +32.9% +26.4%
엔진 교체 (BF16, vLLM→TRT) +1.8% +10.1%
엔진 교체 (FP8, vLLM→TRT) +15.4% +29.8%
 

엔진을 바꿀지 고민하기 전에 FP8부터 켜는 것이 유리하다. 단일 변경으로 얻는 이득이 압도적으로 크고, 체크포인트 하나로 두 엔진 모두에서 바로 활용할 수 있다.

② FP8 이득은 TRT-LLM에서 더 크다 (+50.7% vs +32.9%). FP8 커널 최적화 성숙도 차이로 보이며, 그 결과 FP8을 켜면 엔진 간 격차가 오히려 벌어진다(+1.8%→+15.4%, +10.1%→+29.8%).

③ Int8 vs FP8 대조. GPTQ-Int8은 FP16보다 11% 느렸다(auto_gptq Int8 커널의 dequant 오버헤드). 반면 FP8은 전 지표에서 이득이다.


  처리량 디코드(TPOT) 메모리
GPTQ-Int8 −11% −18% ❌ +35% ✅
FP8 (본 실측) +51% +36% +24%

양자화 이득은 비트 폭이 아니라 "그 포맷을 하드웨어 텐서코어가 네이티브로 아느냐"가 결정한다. Ada는 FP8을 알고, Int8 GPTQ 커널은 그 경로를 못 탄다.

④ TTFT 중앙값은 여전히 vLLM(W1) / TRT-LLM(W2)로 갈리지만, P99는 FP8에서 TRT-LLM이 더 벌린다 (W1 −26%, W2 −25%).

8. 네이티브 trtllm-bench — 교차 검증

서버를 경유한 숫자가 HTTP/토크나이즈 오버헤드로 왜곡되지 않았는지, TRT-LLM 자체 오프라인 벤치로 대조한다.

# 1) 합성 데이터셋 생성 (W1 형상)
trtllm-bench --model Qwen/Qwen2.5-7B-Instruct prepare-dataset \
  --output /tmp/ds_w1.jsonl \
  token-norm-dist --num-requests 300 \
  --input-mean 1024 --input-stdev 0 --output-mean 256 --output-stdev 0

# 2) BF16
trtllm-bench --model Qwen/Qwen2.5-7B-Instruct \
  throughput --dataset /tmp/ds_w1.jsonl --backend pytorch \
  --concurrency 128 --num_requests 300

# 3) FP8 — 로컬 체크포인트는 --model_path 로 지정 (아래 트러블슈팅 ⑤)
trtllm-bench --model Qwen/Qwen2.5-7B-Instruct --model_path /hf/qwen2.5-7b-fp8 \
  throughput --dataset /tmp/ds_w1.jsonl --backend pytorch \
  --concurrency 128 --num_requests 300

 

측정 방식 BF16 req/s  BF16 out tok/s FP8 req/s FP8 out tok/s
서버 경유 (vllm bench serve) 6.15 1574 9.27 2372
네이티브 (trtllm-bench) 6.33 1622 9.60 2457
차이 +2.9% +3.0% +3.5% +3.6%

네이티브가 일관되게 3% 정도 높다. HTTP 왕복 + 서버 측 토크나이즈/디토크나이즈 비용에 해당하며, 두 정밀도에서 같은 폭으로 나타나므로 서버 경유 측정이 편향 없이 유효함을 뒷받침한다.

⚠️ 트러블슈팅 ⑤ — trtllm-bench의 문서/CLI 드리프트
문서에 나오는 prepare-dataset --stdout은 1.2.1에 없다(Error: No such option: --stdout). → --output <파일> 사용.

--tokenizer도 prepare-dataset 레벨엔 없다. 상위 --model이 그 역할을 한다.

--model에 로컬 경로를 주면 HF repo id로 파싱해서 실패한다. HFValidationError: Repo id must be in the form 'repo_name' or 'namespace/repo_name': '/hf/qwen2.5-7b-fp8' → --model <HF이름>(토크나이저·설정용)과 --model_path <로컬경로>(가중치용)를 함께 지정한다.

9. 튜닝 치트시트 (L40S / 7B 기준)

  • 무조건 먼저: FP8 양자화. ModelOpt PTQ 1분, 처리량 +26~51%, 메모리 −45%, 두 엔진 공용. Ada 이상 하드웨어에서 필수 적용 권장.
  • 긴 입력(RAG·문서요약, 프리필 무거움): TRT-LLM. 처리량 +10~30%, TTFT 중앙값 2.9~3.4배 우위.
  • 짧은 입력 대화형(중앙값 TTFT 중요): vLLM. 기본 chunked prefill ON이 유리.
  • P99 SLO가 계약 조건: TRT-LLM. 전 조합에서 꼬리 TTFT 13~26% 낮음.
  • 한 GPU에 여러 모델: --kv_cache_free_gpu_memory_fraction을 반드시 낮출 것. 기본 0.9는 남는 메모리를 거의 다 가져간다(1.1B 모델이 KV로 35.8GB를 잡음).
  • TRT-LLM에서 TTFT를 줄이고 싶다면: enable_chunked_prefill이 기본 OFF임을 기억하고 명시적으로 켜서 재측정해볼 수 있다. 단, 긴 입력에선 켜는 게 손해일 수 있음.
  • 용량 산정: TRT-LLM 컨테이너는 압축 19.5GB / 언팩 ~50GB. 노드 디스크 150GB 이상 권장.

10. 트러블슈팅 총정리

# 증상 원인 해결
파드 전부 Evicted, DiskPressure TRT-LLM 이미지 ~50GB + 모델 15GB > 100GB 노드 노드 디스크 150GB로 리사이즈
admission webhook "vgpu.hami.io" denied: pod has node assigned nodeName이 HAMi 스케줄러를 우회 GPU 파드는 nodeSelector로 유도
ImportError: libnvonnxparser.so.10 command:가 ENTRYPOINT를 교체해 LD_LIBRARY_PATH 소실 command: ["bash","-lc"] + exec ...
MPI ... forgotten to use the proper idiom LLM()이 MPI 워커 spawn, 가드 없으면 재귀 if __name__ == "__main__":
No such option: --stdout / HFValidationError 문서·CLI 버전 간 차이 --output, --model_path
문서의 1.3.0rc25 태그 pull 실패 레지스트리에 존재하지 않음 태그 목록 직접 조회 → 1.2.1 사용
CephFS PVC Pending rados: ret=-22 (테넌트 Ceph 연결) 벤치엔 영속성 불필요 → hostPath 우회
Kyverno PolicyViolation 경고 hostPath/non-root 위반 audit 모드라 통과. 프로덕션이면 정책 예외 필요
 

11. 한계 / 주의사항

  • 각 조합 1회 측정(반복 없음) → 분산·신뢰구간 미반영. W1 BF16의 +1.8% 차이는 노이즈 범위 내일 수 있다.
  • random 데이터셋 사용 → 실트래픽과 프리픽스 캐싱 효과가 다르다(양쪽 다 프리픽스 캐싱 기본 ON이라 랜덤 토큰에선 거의 히트하지 않음).
  • 동시성 1점만 측정 → latency–throughput 곡선 전체를 커버하지는 못함.
  • FP8 품질을 정량 검증하지 않았다. 서빙 sanity(동일 프롬프트 정상 응답)만 확인했고 perplexity/lm-eval 미실시. 채택 전 정확도 하락 측정 필요.
  • 기본값 대 기본값 비교다. 양쪽 다 튜닝하면 순위가 바뀔 수 있다.
  • HAMi 소프트 파티션 위 측정이나, 두 엔진 모두 44GB를 온전히 보므로 통짜 GPU와 동등하게 동작함.

부록 — 재현 스크립트

# 전체 벤치 스윕
./run_trtllm_bench.sh <label> <in> <out> <num_prompts> <concurrency> <outfile>

# 엔진 교체는 args 만 patch (Recreate 전략)
kubectl -n trtllm patch deploy engine --type=json -p '[
 {"op":"replace","path":"/spec/template/spec/containers/0/image","value":"vllm/vllm-openai:latest"},
 {"op":"replace","path":"/spec/template/spec/containers/0/args",
  "value":["exec vllm serve Qwen/Qwen2.5-7B-Instruct --host 0.0.0.0 --port 8000"]}]'
kubectl -n trtllm rollout status deploy/engine --timeout=1200s

# 정리
kubectl delete ns trtllm            # GPU 반납
# 모델 캐시는 노드 hostPath /var/lib/hf-cache 에 잔존

Pre-quantized 모델 4종(Qwen2.5-7B-Instruct FP16, GPTQ-Int8, GPTQ-Int4, AWQ-Int4)을 vLLM V1 엔진에 배포하여 성능을 실측 비교했습니다. 별도의 로컬 양자화 과정 없이 Hugging Face의 pre-quantized 모델 ID를 지정하면 vLLM이 config.json의 quantization_config를 참조해 연산 커널을 자동으로 설정합니다.

변형 사용 모델 ID vLLM 감지 방식
FP16 Qwen/Qwen2.5-7B-Instruct None

GPTQ-Int8 Qwen/Qwen2.5-7B-Instruct-GPTQ-Int8 auto_gptq

GPTQ-Int4 Qwen/Qwen2.5-7B-Instruct-GPTQ-Int4 auto_gptq

AWQ-Int4 Qwen/Qwen2.5-7B-Instruct-AWQ auto_awq

실측 결과 (워크로드: input 1024 / output 256, 300 요청, 동시성 128)

변형 KV 캐시 (tok) 최대 동시성 req/s out tok/s Med TTFT Med TPOT Med ITL
FP16 374,704 11.44x 6.04 1,546 933ms 70.7ms 37.3ms
GPTQ-Int8 505,440 (+35%) 15.42x 5.38 (-11%) 1,377 1,292ms 83.8ms 30.4ms
GPTQ-Int4 560,816 (+50%) 17.11x 7.29 (+21%) 1,867 941ms 62.0ms 24.7ms
AWQ-Int4 561,040 (+50%) 17.12x 7.12 (+18%) 1,822 978ms 63.5ms 24.9ms

동일 프롬프트 샘플링 테스트(Greedy, temp=0) 결과 4가지 모델 모두 문맥 붕괴 없이 일관되고 정확한 정의 문장을 생성했습니다.

심층 분석

  • 메모리 확보 및 동시 수용력 증대: 모델 가중치가 작아진 만큼 동일 GPU 메모리 비율(gpu-memory-utilization=0.85) 환경에서 KV 캐시 영역이 확장됩니다. Int4 적용 시 KV 캐시 용량이 374k에서 561k tokens로 50% 증가하여 최대 동시성이 11.4x에서 17.1x로 대폭 상승합니다.
  • Int4의 처리량 역전 상승: LLM 디코드 연산은 메모리 대역폭 병목(Memory Bandwidth Bound) 환경입니다. Int4 도입으로 가중치 전송량이 줄어들면서 토큰당 처리 속도가 빨라졌고, 전체 처리량이 21% 상향되었습니다.
  • Int8 성능 저하 원인: GPTQ-Int8은 FP16보다 정밀도가 높음에도 불구하고 처리량이 11% 감소했습니다. 이는 vLLM 내 auto_gptq Int8 커널의 역양자화(Dequantization) 오버헤드가 메모리 대역폭 이득을 상쇄했기 때문입니다. L40S 환경에서는 Int4 Marlin 커널이 훨씬 최적화되어 있습니다.

튜닝 치트시트

  • 실시간 대화형 서비스 (TTFT 최우선)
    • Chunked Prefill ON + --max-num-batched-tokens=2048 설정.
    • --max-num-seqs는 너무 작게 설정하지 않아 큐 대기 시간 증가 방지.
  • 오프라인 대량 배치 처리 (Throughput 최우선)
    • --max-num-batched-tokens=16384 이상 설정 또는 Chunked Prefill OFF (--no-enable-chunked-prefill).
  • 단기 컨텍스트 전용 서빙
    • --max-model-len을 실제 최대 응답 길이로 축소하여 KV 캐시 수용력(Concurrency) 확보.
  • 7B 모델 기준 최적 조합 (NVIDIA L40S)
    • GPTQ-Int4 / AWQ-Int4 양자화 적용을 기본 옵션으로 채택 (메모리 9GB 절감, 처리량 20% 이상 상승).

## 스크립트 및 결과

run_vllm_quant.sh
0.00MB
vllm-quant-results.txt
0.02MB

NVIDIA L40S 1장(HAMi vGPU) 환경에서 vLLM V1 엔진의 주요 서빙 파라미터를 변경하며 처리량(Throughput)과 지연시간(TTFT·TPOT)의 트레이드오프를 실측 분석한 기록입니다. 모든 재현 명령어를 순서대로 포함했습니다.

핵심 요약

파라미터 크게(↑) 작게(↓)
--max-num-seqs (배치 상한) 처리량↑ (auto 이상은 KV캐시 한계로 무의미) 디코드 지연(TPOT)↓, 부하 시 TTFT 큐잉 폭증
--max-model-len 긴 컨텍스트 수용 (단기 성능 영향 없음) 동시 수용력(Concurrency) 대폭 증가, 긴 요청 거부
Chunked Prefill & --max-num-batched-tokens 크게/OFF → 총 처리량↑ · TPOT↓, TTFT↑ 작게/ON → 첫 토큰 지연(TTFT)↓, 처리량↓ · TPOT↑
 

0. 테스트 환경 및 GPU 상태 확인

KubeVirt VM 기반 Kubernetes 노드에서 HAMi를 통해 L40S 1장을 할당 받았습니다.

# 1. GPU 노드 및 HAMi 할당 슬롯 확인
kubectl get nodes -o custom-columns='NODE:.metadata.name,GPU:.status.capacity.nvidia\.com/gpu'

# 2. 물리 GPU 세부 사양 및 HAMi 분할 상태 확인 (44GB 메타데이터 확인)
kubectl get node <gpu-node-name> -o jsonpath='{.metadata.annotations.hami\.io/node-nvidia-register}'

1. vLLM 서버 배포

Recreate 전략의 Deployment로 배포하고, 노드 hostPath를 사용하여 모델 캐시를 재사용합니다.

# vllm.yaml
apiVersion: apps/v1
kind: Deployment
metadata:
  name: vllm
  namespace: vllm-bench
  labels:
    app: vllm
spec:
  replicas: 1
  strategy:
    type: Recreate
  selector:
    matchLabels:
      app: vllm
  template:
    metadata:
      labels:
        app: vllm
    spec:
      containers:
      - name: vllm
        image: vllm/vllm-openai:latest
        args:
          - "--model=Qwen/Qwen2.5-7B-Instruct"
          - "--port=8000"
          - "--gpu-memory-utilization=0.85"
        ports:
        - containerPort: 8000
        env:
        - name: HF_HOME
          value: /root/.cache/huggingface
        resources:
          limits:
            nvidia.com/gpu: "1"
            nvidia.com/gpumem: "44000"
            nvidia.com/gpucores: "100"
            cpu: "10"
            memory: 40Gi
          requests:
            cpu: "4"
            memory: 16Gi
        readinessProbe:
          httpGet:
            path: /health
            port: 8000
          initialDelaySeconds: 30
          periodSeconds: 5
          failureThreshold: 240
        volumeMounts:
        - name: hf-cache
          mountPath: /root/.cache/huggingface
        - name: shm
          mountPath: /dev/shm
      volumes:
      - name: hf-cache
        hostPath:
          path: /var/lib/vllm-hf-cache
          type: DirectoryOrCreate
      - name: shm
        emptyDir:
          medium: Memory
          sizeLimit: 8Gi
# 네임스페이스 생성 및 vLLM 배포
kubectl create ns vllm-bench
kubectl apply -f vllm.yaml

# 롤아웃 대기 (최초 배포 시 모델 다운로드로 수 분 소요)
kubectl -n vllm-bench rollout status deploy/vllm --timeout=1800s

# 서빙 상태 확인
kubectl -n vllm-bench exec deploy/vllm -- curl -s localhost:8000/v1/models

트러블슈팅 Note: vLLM 0.27.1 기준 --disable-log-requests 옵션은 삭제되었으므로 제외해야 컨테이너 크래시(unrecognized arguments)를 방지할 수 있습니다.

2. 엔진 기본값 확인 명령어

vLLM CLI 도움말 및 실행 로그를 조회하여 자동 산출된 기본값을 확인합니다.

POD=$(kubectl -n vllm-bench get pod -l app=vllm -o name | head -1)

# CLI 기본 도움말 조회
kubectl -n vllm-bench exec $POD -- sh -c 'vllm serve --help=max-num-seqs'
kubectl -n vllm-bench exec $POD -- sh -c 'vllm serve --help=max-num-batched-tokens'
kubectl -n vllm-bench exec $POD -- sh -c 'vllm serve --help=max-model-len'
kubectl -n vllm-bench exec $POD -- sh -c 'vllm serve --help=enable-chunked-prefill'

# 기동 로그에서 실제 산출된 값 확인
kubectl -n vllm-bench logs $POD | grep -iE "max_seq_len|enable_chunked_prefill|enable_prefix_caching|GPU KV cache size|Maximum concurrency"

 

확인된 기본 산출 값:

  • max_model_len: 32,768 (모델 Config 기반 자동)
  • max_num_seqs: auto
  • max_num_batched_tokens: auto (~8,192 수준)
  • enable_chunked_prefill: True (V1 Engine 기본 ON)
  • GPU KV Cache Size: 374,704 tokens (동시성 11.44x)

3. 벤치마크 실행 명령어

컨테이너 내부에서 내장 vllm bench serve를 실행하여 런 간 공정한 실측 데이터를 수집합니다.

POD=$(kubectl -n vllm-bench get pod -l app=vllm -o name | head -1)

kubectl -n vllm-bench exec $POD -- sh -c 'vllm bench serve \
  --model Qwen/Qwen2.5-7B-Instruct --host localhost --port 8000 \
  --dataset-name random --random-input-len 1024 --random-output-len 256 \
  --num-prompts 300 --max-concurrency 128 \
  --ignore-eos --temperature 0 --seed 12345 \
  --percentile-metrics ttft,tpot,itl,e2el --metric-percentiles 90,99'

4. 과제 1: max_num_seqs 및 max_model_len 실측

Deployment의 args를 패치하여 재배포 후 실측합니다. (input 1024 / output 256 / concurrency 128)

# 예시: max-num-seqs=32 설정 패치 명령
kubectl -n vllm-bench patch deploy vllm --type=json -p \
 '[{"op":"replace","path":"/spec/template/spec/containers/0/args",
    "value":["--model=Qwen/Qwen2.5-7B-Instruct","--port=8000","--gpu-memory-utilization=0.85","--max-num-seqs=32"]}]'

kubectl -n vllm-bench rollout status deploy/vllm --timeout=420s

실측 결과

Config req/s out tok/s TTFT (중앙) TTFT (p99) TPOT (중앙) 최대 동시성
baseline (auto) 5.96 1,525 951ms 10.8s 71.9ms 11.44x
--max-num-seqs=32 3.41 (-43%) 873 27,470ms 29.3s 34.0ms 12.09x
--max-num-seqs=512 5.96 (=) 1,526 943ms 10.8s 71.9ms 11.34x
--max-model-len=8192 5.96 (=) 1,525 946ms 10.8s 71.9ms 45.74x
 

분석:

  • max-num-seqs=32: 동시 요청 128개 중 96개가 큐 대기 상태로 밀려나며 TTFT가 27초로 폭증하고 전체 처리량이 43% 급감합니다.
  • max-model-len=8192: 긴 입력이 없는 워크로드에서는 성능 변화가 없으나, 엔진 내부 KV 캐시를 적게 할당하여 최대 동시 수용력이 11.44x에서 45.74x로 4배 급증합니다.

5. 과제 2: Chunked Prefill 및 max_num_batched_tokens 실측

Prefill 부하를 높여 청크 효과를 명확히 측정합니다. (input 4096 / output 256 / concurrency 64)

# 1. Chunked Prefill ON & batched tokens=2048 패치
kubectl -n vllm-bench patch deploy vllm --type=json -p \
 '[{"op":"replace","path":"/spec/template/spec/containers/0/args",
    "value":["--model=Qwen/Qwen2.5-7B-Instruct","--port=8000","--gpu-memory-utilization=0.85","--max-num-batched-tokens=2048"]}]'

# 2. Chunked Prefill OFF 패치
kubectl -n vllm-bench patch deploy vllm --type=json -p \
 '[{"op":"replace","path":"/spec/template/spec/containers/0/args",
    "value":["--model=Qwen/Qwen2.5-7B-Instruct","--port=8000","--gpu-memory-utilization=0.85","--no-enable-chunked-prefill"]}]'

실측 결과

Config req/s 총 tok/s TTFT (중앙) TTFT (p99) TPOT (중앙)
baseline (chunked ON, auto) 1.73 7,534 2,310ms 23.5s 132ms
--max-num-batched-tokens=2048 1.70 7,410 1,404ms 24.3s 138ms
--max-num-batched-tokens=16384 1.80 7,855 4,333ms 21.8s 116ms
--no-enable-chunked-prefill 1.83 7,971 7,297ms 21.3s 104ms
 

분석:

  • TTFT 개선: 청크 크기를 2048로 줄이면 Prefill을 잘게 쪼개어 스케줄링하므로 첫 토큰 지연이 1404ms로 가장 우수해집니다.
  • 처리량 최적화: Chunked Prefill을 끄거나(--no-enable-chunked-prefill) 청크를 대형화(16384)하면 GPU Batching 효율이 높아져 총 처리량이 최고(1.83 req/s)가 되며 TPOT(104ms)이 가장 빠릅니다.

6. 전체 테스트 자동화 스크립트 및 cleanup

반복 테스트를 진행한 전체 스크립트 흐름 및 cleanup 명령어입니다.

# 튜닝 자동화 스크립트 예시
apply_cfg() {
  ARG_STR=$1
  kubectl -n vllm-bench patch deploy vllm --type=json -p \
   "[{\"op\":\"replace\",\"path\":\"/spec/template/spec/containers/0/args\",
      \"value\":[\"--model=Qwen/Qwen2.5-7B-Instruct\",\"--port=8000\",\"--gpu-memory-utilization=0.85\" ${ARG_STR:+, $ARG_STR}]}]"
  kubectl -n vllm-bench rollout status deploy/vllm --timeout=420s
}

bench() {
  LABEL=$1; IN_LEN=$2; OUT_LEN=$3; NUM_P=$4; CONC=$5
  POD=$(kubectl -n vllm-bench get pod -l app=vllm -o name | head -1)
  echo "=== Running $LABEL ==="
  kubectl -n vllm-bench exec $POD -- sh -c "vllm bench serve \
    --model Qwen/Qwen2.5-7B-Instruct --host localhost --port 8000 \
    --dataset-name random --random-input-len $IN_LEN --random-output-len $OUT_LEN \
    --num-prompts $NUM_P --max-concurrency $CONC \
    --ignore-eos --temperature 0 --seed 12345 \
    --percentile-metrics ttft,tpot,itl,e2el --metric-percentiles 90,99"
}

# 과제 1 실행
apply_cfg ""
bench C1_baseline 1024 256 300 128

apply_cfg '"--max-num-seqs=32"'
bench C1_seqs32 1024 256 300 128

apply_cfg '"--max-num-seqs=512"'
bench C1_seqs512 1024 256 300 128

apply_cfg '"--max-model-len=8192"'
bench C1_modellen8192 1024 256 300 128

# 과제 2 실행
apply_cfg ""
bench C2_baseline 4096 256 200 64

apply_cfg '"--max-num-batched-tokens=2048"'
bench C2_mnbt2048 4096 256 200 64

apply_cfg '"--max-num-batched-tokens=16384"'
bench C2_mnbt16384 4096 256 200 64

apply_cfg '"--no-enable-chunked-prefill"'
bench C2_nochunk 4096 256 200 64

# 자원 자원 해제 (네임스페이스 삭제)
kubectl delete ns vllm-bench

 

### bench 스크립트 및 성능 체크 결과

run_vllm_bench.sh
0.00MB
vllm-bench-results.txt
0.05MB

 

+ Recent posts