NPU·GPU에서의 HBM 활용 — Weight·Activation·KV Cache 배치 분석

#한 줄 요약

“LLM inference의 HBM 점유는 weight·activation·KV cache 세 부분으로 갈립니다.” — 70B 모델 weight 140 GB, batch 128 activation 40 GB, KV cache는 sequence 길이에 비례해 50~200 GB까지 늘어납니다. KV cache가 진짜 메모리 폭증의 원인입니다. 시리즈를 마무리하며 카드급 사례와 CXL·UALink와의 결합까지 봅니다.

Ch 7에서 컨트롤러의 내부를 봤습니다. 이번 마지막 장은 AI workload가 HBM을 어떻게 채우는지입니다. 시리즈를 마무리하면서 현세대 AI 가속기 사례와 HBM 너머의 메모리 tiering까지 정리합니다.

#LLM inference의 메모리 분해

LLaMA 70B 모델을 batch 128, sequence 2048로 서빙한다고 합시다.

1
LLaMA 70B inference (FP16, batch 128, seq 2048)
2

3
Weight (정적, 모델 자체):
4
  70 B parameter × 2 byte = 140 GB
5

6
Activation (per-token, batch 분담):
7
  hidden_size 8192 × layer 80 × batch 128 × seq 1
8
  = 8192 × 80 × 128 × 2 / 1024³
9
  = 167 MB per token (decode)
10

11
KV cache (sequence에 비례):
12
  per token: hidden × 2 (K, V) × layer × dtype
13
           = 8192 × 2 × 80 × 2 byte = 2.5 MB per token
14

15
  batch 128 × seq 2048 × 2.5 MB
16
  = 128 × 2048 × 2.5 MB / 1024
17
  = 640 GB
18

19
총 메모리: 140 + 0.17 + 640 = 약 780 GB

KV cache가 weight보다 4배 큽니다. 이게 LLM serving의 메모리 폭증입니다.

H100 80 GB로 LLaMA 70B를 serving하려면 모델 4분할 + KV cache 별도 호스트가 필요합니다. H200·B200으로 가야 한 장에 weight가 들어갑니다.

#Weight 저장

weight는 학습 후 정적이고 모든 추론에서 그대로 읽힙니다.

Weight layout in HBM — 전형적 매핑은 한 layer를 여러 stack에 분산. Layer 0의 Q/K/V·FFN을 stack 0..7로 쪼개고, Layer 1도 같은 방식.

Layout 규칙

NVIDIA TensorRT-LLM과 vLLM은 weight를 자동으로 channel-aware하게 layout합니다. 수동 tuning은 거의 사라졌습니다.

#Activation — batch와 sequence

activation은 forward pass 중 layer 입출력입니다. batch와 sequence가 곱입니다.

Activation memory

prefill은 throughput bound, decode는 latency bound입니다. activation memory 자체는 KV cache에 비하면 작습니다.

#KV cache — 메모리 폭증의 원인

attention 연산은 과거 모든 토큰의 K, V를 현재 query가 참조합니다. 이것을 재계산 안 하려고 KV cache에 저장합니다.

KV cache size

GQA가 KV cache를 8배 줄였습니다. LLaMA 2 70B부터 GQA가 표준입니다.

KV cache access pattern — decode 단계의 매 토큰

1
for layer in 80:
2
    Q = compute(activation)
3
    K, V = compute(activation)               # new token push
4

5
    # read past K, V for attention
6
    for past_token in past_tokens:
7
        attention += Q @ K[past_token].T
8

9
    # weighted sum with V
10
    output = sum(attention * V[past_token])

Memory traffic per layer per token

Sequence가 길어질수록 KV traffic이 증가 → long context inference는 KV bound.

8K context: KV traffic이 weight traffic의 2배. 128K context: 32배. long context가 memory bound의 극단입니다.

#Memory layout — tiling 전략

GPU/NPU kernel은 HBM access pattern을 명시적으로 tiling합니다.

1
// CUDA matmul tiling 의사 코드
2
__global__ void matmul_tiled(
3
    half *A,    // M × K, in HBM
4
    half *B,    // K × N, in HBM
5
    half *C,    // M × N, in HBM
6
    int M, int N, int K
7
) {
8
    // 1. block-level tile
9
    __shared__ half tileA[128][32];
10
    __shared__ half tileB[32][128];
11

12
    int row = blockIdx.y * 128 + threadIdx.y;
13
    int col = blockIdx.x * 128 + threadIdx.x;
14

15
    float sum = 0;
16

17
    // 2. tile loop
18
    for (int t = 0; t < K; t += 32) {
19
        // load tile A from HBM → shared
20
        tileA[threadIdx.y][threadIdx.x] = A[row * K + t + threadIdx.x];
21
        // load tile B from HBM → shared
22
        tileB[threadIdx.y][threadIdx.x] = B[(t + threadIdx.y) * N + col];
23
        __syncthreads();
24

25
        // compute on shared (fast)
26
        for (int k = 0; k < 32; k++) {
27
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
28
        }
29
        __syncthreads();
30
    }
31

32
    C[row * N + col] = sum;
33
}

핵심은 shared memory(on-chip)에 tile 단위로 load하고 재사용해 HBM access를 최소화하는 것입니다.

1
tile size 선택
2

3
tile 1024 byte (32×32 half):
4
  - shared memory 사용: 작음
5
  - HBM access 빈번
6
  - bandwidth bound
7

8
tile 16 KB (128×64 half):
9
  - shared memory 한계 근접
10
  - HBM access 적음
11
  - compute bound로 이동 가능
12

13
H100 shared memory: 228 KB per SM
14
  → tile 64 KB까지 안전

Flash Attention은 attention 자체를 tile해서 KV cache를 fully on-chip에 올립니다. long context의 game changer입니다.

#MFU·MBU의 현실

대표적인 inference workload의 측정치입니다.

LLaMA 70B inference on H100(FP16)의 측정치는 다음과 같습니다.

batch가 클수록 MFU 상승합니다. weight read가 batch 전체에 amortize되기 때문입니다. batch 1에서는 weight 140 GB를 읽기만 하면 42 ms인데 batch 64에서는 그대로 42 ms입니다 (capacity 한계까지).

1
batch scaling
2

3
batch 1:    weight read 1회 × 42 ms = 42 ms / token
4
batch 64:   weight read 1회 × 42 ms = 0.65 ms / token (이론)
5
            실제 50 ms × token 도달 (limited by FLOPS)
6

7
batch sweet spot:
8
  H100 80 GB: batch 16~32
9
  H200 141 GB: batch 32~64
10
  B200 192 GB: batch 64~128

#카드급 사례

현세대 AI 가속기 카드의 HBM 구성입니다.

한국 NPU의 경우입니다.

#CXL과의 결합 — Memory Tiering

HBM 192 GB로도 대형 LLM weight + 전체 KV cache를 담지 못합니다. CXL이 낮은 tier의 메모리를 제공합니다.

LLM serving 매핑은 다음과 같습니다.

• hot KV cache → HBM
• warm KV cache → CXL
• cold KV cache → SSD
• weight (정적) → HBM (전부 캐시)

vLLM·SGLang 같은 현세대 serving framework는 KV cache를 tier 1~3에 자동 분산합니다. paged attention이라고 부르는 기법이 대표적입니다.

#UALink — GPU 간 메모리 공유

여러 GPU가 서로의 HBM을 직접 access하는 UALink는 2024년 컨소시엄 결성, 2025년 4월 200G 1.0 spec 비준으로 이어졌습니다.

UALink vs NVLink

UALink 컨소시엄: AMD, Broadcom, Cisco, Google, HPE, Intel, Meta, Microsoft.

GPU 간 HBM 공유 효과 — weight 한 GPU에 두고 나머지 GPU가 직접 access. weight replication을 줄여 더 큰 모델 서빙 가능.

UALink 시리즈에서 자세히 다룹니다.

#시리즈 마무리

HBM·GDDR 8개 장을 지났습니다. 시리즈를 한 줄씩 정리합니다.

#추천 후속 시리즈

이 시리즈와 직접 연결되는 세 시리즈를 추천합니다.

#자주 하는 실수

#”HBM capacity로 모든 LLM이 수용된다”

70B 모델 weight 140 GB는 H200(141 GB)에 겨우 들어갑니다. KV cache 80 GB까지 합치면 어느 H200도 못 담습니다. 최소 2장 분할이 필수입니다. 카드 capacity만 보고 수용 가능을 판단하면 오답입니다.

#KV cache를 작다고 가정

GQA로 줄여도 *64 GB+*가 나오기 일쑤입니다. long context는 더 늘어납니다. paged attention 같은 동적 관리가 필수입니다.

#tile size를 수동 tuning하려는 시도

NVIDIA TensorRT, vLLM, AMD ROCm 모두 autotuner가 들어 있습니다. 수동 tile은 autotuner를 이기기 어렵습니다. 새 모델·새 GPU에서 autotuner를 신뢰하는 게 안전합니다.

#”batch 1 inference에 H100을 사주면 빠르다”

batch 1은 *MFU 5%*입니다. H100의 1000 TFLOPS 중 50 TFLOPS만 사용입니다. latency는 H200과 비슷하기도 합니다. batch가 작은 사용처에서는 L40·L4 같은 GDDR 카드가 비용 효율적입니다.

#HBM capacity를 늘리면 무조건 throughput이 늘어난다는 가정

capacity가 늘면 batch를 키울 수 있고 throughput이 늘긴 합니다. 하지만 batch가 어느 점을 넘으면 latency가 SLA를 깨고, MBU 상한에 닿아 gain이 멈춥니다. capacity와 bandwidth의 균형이 핵심입니다.

#정리

• LLM inference 메모리는 weight·activation·KV cache로 갈리고, KV cache가 최대 폭증 원인입니다.
• 70B FP16 weight는 140 GB, GQA KV cache는 80~640 GB입니다.
• channel-aware allocation과 tile-based access가 HBM 효율을 결정합니다.
• MFU와 MBU를 같이 봐야 진짜 병목이 보입니다. decode는 *MBU 80%, MFU 5%*가 흔합니다.
• batch가 클수록 MFU 상승하지만 latency도 증가합니다.
• 현세대 카드는 H100(80GB) → H200(141GB) → B200(192GB) → MI325X(256GB) → B300(288GB) 순으로 capacity가 빠르게 증가합니다.
• CXL은 Tier 2 메모리, UALink는 GPU 간 HBM 공유, NVMe는 Tier 3 cold tier입니다.
• 한국 NPU(Rebellions Atom, Sapeon X330)는 HBM 채택과 GDDR 채택이 섞여 있습니다.
• 시리즈는 끝났지만 UCIe·BoW·CXL·UALink 시리즈와 함께 읽으면 현세대 패키징·메모리 시스템의 전체 그림이 보입니다.

#추천 후속 시리즈

• UCIe Ch 1: 개요 — 칩렛 표준의 핵심
• BoW Ch 1: 개요 — open die-to-die 표준
• CXL Ch 1: 개요 — HBM 너머의 메모리
• UALink Ch 1: 개요 — GPU 간 메모리 fabric

#관련 항목

• Ch 1: 고대역 메모리 개요
• Ch 5: 대역폭 계산과 병목 분석
• Ch 7: 메모리 컨트롤러 인터페이스
• CXL Ch 8: ML 가속기 — HBM과 CXL의 조합
• UCIe Ch 12: case studies — 칩렛 + HBM 실제 사례