【#3】ds4.c 분석하기

본 시리즈는 DeepSeek V4 Flash / Pro 전용 추론 엔진인 DwarfStar(ds4)의 코드를 분석하는 연재입니다.

제3회는 추론 엔진의 심장부인 Metal 그래프와 압축 KV 캐시(KV Cache)의 구현을 추적합니다. 주요 참조 위치:

ds4.c, ds4_metal.m, metal/dsv4_kv.metal, metal/dsv4_misc.metal, metal/dsv4_hc.metal, metal/flash_attn.metal

연재 「DwarfStar(ds4) 읽기」 전 8회

• 제1회 왜 전용 엔진을 작성하는가

• 제2회 비대칭 2bit 양자화(Quantization)와 imatrix
  제3회 Metal 그래프와 압축 KV (본 기사)

• 제4회 디스크 KV 캐시

• 제5회 서버와 DSML 툴 호출

• 제6회 TCP 파이프라인 분산 추론

• 제7회 네이티브 에이전트

• 제8회 스티어링(Steering)・MTP・평가 기반

• DS4의 어텐션(Attention)은 단순한 전체 KV 캐시가 아니다. raw sliding-window KV, 압축 KV, ratio-4 층용 인덱서(indexer)를 조합한다.

• Flash의 압축률은 고정되어 있지 않다. ds4_expected_layer_compress_ratio()에 따르면, Flash는 0-1층이 비압축이며, 이후 짝수층은 ratio 4, 홀수층은 ratio 128이다. Pro는 0-1층도 ratio 128이며, 그 이후는 동일한 교차 구성이다.

• 컴프레서(compressor)는 단순 평균이 아니다. wkv와 wgate로 KV 후보와 스코어(score)를 만들고, APE를 더한 뒤, softmax 가중 풀링(weighted pooling)을 하여 압축 행을 만든다.

• ratio-4 층에서는 어텐션용 컴프레서와 인덱서용 컴프레서의 2개 레인을 가진다. 인덱서는 압축 행 전체에서 top-k를 선택하며, FlashAttention 측은 raw SWA와 선택된 압축 행을 동시에 어텐드(attend)한다.

• DS4 Flash는 Hyper-Connection n_hc=4를 가지며, Metal 측에서는 Sinkhorn 정규화된 4x4 혼합 행렬을 만드는 전용 커널이 있다.

• RoPE는 헤드 전체가 아니라 말단 64차원에만 적용된다. 압축층에서는 일반적인 RoPE base가 아닌 압축용 base를 사용한다.

지난 회까지의 설명에서는 "4 토큰마다 1 압축 행"이라는 ratio-4를 중심으로 살펴보았습니다. 하지만 코드를 읽어보면 DS4의 압축은 그뿐만이 아닙니다.

ds4.c의 ds4_expected_layer_compress_ratio()는 모델의 변체(variant)별 기대 압축률을 명시하고 있습니다.

static uint32_t ds4_expected_layer_compress_ratio(uint32_t il) {
switch (DS4_MODEL_VARIANT) {
case DS4_VARIANT_FLASH:
...

Flash의 경우:

층	압축
0-1	없음
...

Pro의 경우:

층	압축
0-1	ratio 128
...

이 값은 단순히 코드 측에서 가정하는 것이 아니라, GGUF 메타데이터의 deepseek4.attention.compress_ratios와 대조됩니다. 불일치할 경우 기동 시에 종료됩니다.

validate_compress_ratio_metadata(m);

DS4는 "이 GGUF라면 아마 돌아갈 것이다"가 아니라, "이 형태·이 메타데이터·이 텐서 레이아웃(tensor layout)이 아니면 동작시키지 않겠다"는 설계입니다. 전용 엔진다운 과감한 결정입니다.

CPU 경로의 주석이 DS4의 KV 캐시를 이해하는 최단 경로입니다.

* The CPU path is the correctness reference. It maintains raw SWA KV rows,
* optional compressed KV rows, the indexer mask for ratio-4 layers, and a
* reusable decode scratch arena ...

실체는 ds4_layer_cache

입니다.

typedef struct {
float *raw_kv;
uint32_t n_raw;
...

여기서부터 각 층(layer)이 가지는 상태는 크게 3가지 종류라는 것을 알 수 있습니다.

상태	역할
raw_kv	직근 토큰용 sliding-window KV
attn_comp_kv	어텐션(Attention) 본체가 읽는 압축 KV
index_comp_kv	ratio-4 층에서 압축 행을 선택하기 위한 indexer KV

raw_kv는 전체 컨텍스트(context) 분량을 모두 가지지 않습니다. DS4_N_SWA, 즉 Flash의 형태에서는 128 토큰 분량이 상한입니다.

static uint32_t ds4_default_raw_cap(uint32_t ctx_size) {
uint32_t raw_cap = DS4_N_SWA;
if (raw_cap > ctx_size) raw_cap = ctx_size;
...

이를 통해 직근의 국소 문맥(local context)은 고정밀도로 유지하고, 오래된 문맥은 압축 행으로 접어 넣는(fold into) 구조가 됩니다.

긴 프롬프트를 한꺼번에 Metal 그래프로 흘려보내면, 스크래치(scratch)나 임시 버퍼가 거대해집니다. DS4는 긴 프롬프트의 prefill을 청크(chunk)화합니다.

ds4_default_prefill_cap_for_prompt()는 환경 변수 DS4_METAL_PREFILL_CHUNK가 없다면, 프롬프트가 4096 토큰을 초과할 경우 청크 상한을 4096으로 설정합니다.

const char *env = getenv("DS4_METAL_PREFILL_CHUNK");
...
} else if (prompt_len > 4096) {
...

README의 Benchmarking 절에서도 긴 프롬프트는 기본적으로 4096 토큰 청크로 prefill 된다고 설명하고 있습니다. 청크 크기를 바꾸면 KV 체크포인트 / logits의 경로가 바뀌므로, 비교 시에는 실행 설정으로서 명시해야 합니다.

이 청크화는 단순한 메모리 절약이 아닙니다. compressor/indexer의 절대 위치 경계를 유지하면서, 동일한 범위에 대응하는 layer-major 그래프를 재사용하기 위한 설계입니다.

압축 KV라고 하면, 단순히 4개 토큰을 평균하여 1개 행으로 만드는 것으로 상상하기 쉽습니다. DS4의 compressor는 그렇지 않습니다.

디코드 참조 구현 compressor_decode_one()은 1 토큰마다 다음 처리를 수행합니다.

• wkv로 압축 KV 후보를 생성
• wgate로 스코어(score)를 생성 - APE, 즉 학습된 위치 편향(positional bias)을 스코어에 더함
• 재귀 상태(recursive state)에 kv와 score를 저장
• ratio 경계에 도달하면 softmax 가중 풀링(weighted pooling)으로 1개의 압축 행을 생성
• RMS 정규화(normalization)를 적용
• tail-only RoPE를 적용
• 어텐션 압축 KV라면 FP8 양자화(quantization)를 왕복하고, indexer 압축 KV라면 QAT를 왕복함

스코어 + APE의 저장은 Metal 커널에서도 전용화되어 있습니다.

kernel void kernel_dsv4_compressor_store_one(
...,
device const float * kv,
...

풀링은 softmax 가중 평균입니다. compressor_pool_decode_state()는 차원별로 스코어의 최댓값을 취하고, exp(score - max)를 가중치로 하여 KV를 평균합니다.

const float w = expf(state_score[r * width + j] - max_score);
denom += w;
sum += w * state_kv[r * width + j];
...

Metal 측에도 융합 커널(fused kernel) kernel_dsv4_softmax_pool이 있습니다.

for (int64_t ir = 0; ir < args.ne00; ++ir) {
const float s = ... score ...
const float w = exp(s - max_s);
...

따라서, compressor(압축기)는 "학습된 스코어에 의한 차원별 풀링 (pooling by dimension via learned scores)"입니다. 이는 긴 문맥을 버리는 것이 아니라, 모델이 학습한 압축 표현 (compressed representation)으로 변환하는 것으로 보아야 합니다.

ratio 4 층은 특별합니다. compressor_pool_decode_state()의 주석에 명시된 것처럼, ratio-4 층에서는 어텐션 압축 (attention compression)과 indexer 압축 (indexer compression)이라는 2개의 레인 (lane)을 가집니다.

/* Ratio-4 layers keep two lanes: attention compression and indexer compression. */
const uint32_t coff = compress_ratio == 4 ? 2u : 1u;
const uint32_t width = coff * head_dim;

ratio 4의 경우, 상태 (state)는 4행이 아니라 8행에 상당하는 양을 다룹니다. 이는 전반부와 후반부를 가진 회전하는 frontier (rotating frontier)입니다. 압축된 행을 내보낸 후, Metal 커널 kernel_dsv4_ratio4_shift_f32가 후반부를 전반부로 이동시킵니다.

ratio-4 frontier (8행)의 흐름을 도식화하면 다음과 같습니다.

frontier (ratio-4, 8행) ── 4 토큰마다 1개의 압축행을 emit (내보냄)
┌─────────── 이전 사이클 ───────────┐┌─────────── 현재 사이클 ───────────┐
│ row0 row1 row2 row3 ││ row4 row5 row6 row7 │
...

// Ratio-4 compression keeps two 4-row halves of recurrent state.
kernel void kernel_dsv4_ratio4_shift_f32(
...,
...

이 2개 레인 구조는 어텐션 본체가 읽는 압축 KV (compressed KV)와 indexer가 읽는 압축 KV를 동일한 ratio 경계에서 업데이트하기 위해 필요합니다.

Flash의 형상 (shape)에서는 n_indexer_head = 64, n_indexer_head_dim = 128, n_indexer_top_k = 512입니다. 일반적인 어텐션 헤드 (attention head)의 차원은 512이지만, indexer는 128 너비의 좁은 표현으로 압축행을 스코어링 (scoring)합니다.

ds4.c의 텐서 검증 (tensor validation)에서는 ratio 4인 층에 대해서만 indexer 관련 텐서를 요구합니다.

if (compress_ratio == 4) {
l->indexer_attn_q_b = required_tensorf(...);
l->indexer_proj = required_tensorf(...);
...

Metal 측에는 디코딩 전용인 kernel_dsv4_indexer_score_one_direct가 있습니다. 주석에 따르면, 1개의 압축행을 스레드 그룹 (thread group)에 배치하고, 64개의 indexer 헤드를 4개 헤드 단위의 그룹으로 처리하여 스코어를 생성한다고 되어 있습니다.

indexer의 출력은 "어떤 압축행을 어텐션 본체에 보여줄 것인가"입니다. 모든 압축행에 밀집하게 (densely) 어텐션하는 대신, 선택된 top-k 행만을 봄으로써 1M 컨텍스트에서도 어텐션 비용을 억제합니다.

ratio-4 경로의 어텐션은 kernel_dsv4_indexed_mixed_attention_heads8를 통해 잘 이해할 수 있습니다. 주석에 "DS4 ratio-4 indexed mixed attention"이라고 적혀 있으며, 이는 밀집된 top-k 마스크 경로를 대체하는 커널입니다.

처리는 크게 2단계로 나뉩니다.

• 가장 최근의 raw sliding-window 행에 어텐드 (attend) 한다.
• top-k로 선택된 압축행에 어텐드 (attend) 한다.

raw 측은 윈도우 (window)를 확인하여, qpos에서 볼 수 있는 범위만 처리합니다.

const uint window_first = (args.window != 0u && qpos + 1u > args.window) ?
qpos + 1u - args.window : 0u;
uint first = max(first_raw_pos, window_first);
...

압축 측은 indexer가 생성한 topk를 읽습니다.

for (uint i = 0; i < args.top_k; i++) {
const int32_t idx = row_topk[i];
if (idx < 0) continue;
...

이 커널은 online softmax 형태로 M과 S를 업데이트하면서, raw 행(row)과 압축 행(compressed row)을 동일한 어텐션 출력(attention output)으로 통합합니다. 최신 데이터는 raw, 오래된 문맥은 선택된 압축 행을 사용한다는 설계가 커널 내부에 직접적으로 나타나 있습니다.

DeepSeek V4의 RoPE는 헤드 전체가 아니라, 헤드의 끝부분인 n_rot = 64 차원에만 적용됩니다. Flash의 헤드 차원(head dimension)은 512이므로, 앞쪽 448 차원은 비 RoPE(non-RoPE)이며, 마지막 64 차원만 회전합니다.

rope_tail_ext_inplace()는 그 이름처럼 끝부분(tail)을 가리킵니다.

const uint32_t n_nope = head_dim - n_rot;
float *tail = x + h * head_dim + n_nope;

나아가 DS4는 압축층(compression layer)에서 일반적인 DS4_ROPE_FREQ_BASE가 아닌 DS4_COMPRESS_ROPE_FREQ_BASE를 사용합니다.

static float layer_rope_freq_base(uint32_t il) {
return ds4_layer_compress_ratio(il) != 0 &&
DS4_COMPRESS_ROPE_FREQ_BASE > 0.0f
...

압축층에서 YaRN 확장을 사용하는 조건도 여기에 포함되어 있습니다. 긴 문맥(long context)을 성립시키기 위한 위치 처리(position processing)가 dense 층과 압축층에서 분기된다는 점이 중요합니다.

compressor_decode_one()은 압축 행을 만든 후, 헤드 차원에 따라 다른 연산 과정(round-trip)을 거칩니다.

if (head_dim == DS4_N_HEAD_DIM) {
dsv4_fp8_kv_quantize_row_inplace_cpu(out_comp, head_dim, DS4_N_ROT);
} else if (head_dim == DS4_N_INDEXER_HEAD_DIM) {
...

일반적인 어텐션 압축 KV는 512 너비이며, 비 RoPE 부분에 E4M3 형식의 FP8 연산을 적용합니다. indexer 압축 KV는 128 너비이며, Hadamard 128과 FP4 활성화 양자화(activation quantization)를 거칩니다.

이는 단순히 "저장 형식만 가볍게 하기 위함"이 아니라, 공식 DeepSeek V4 그래프와 일치시키기 위한 과정입니다. ds4.c의 주석에도 indexer Q와 indexer compressor KV에 QAT(Quantization-Aware Training) 연산이 필요하며, 이것이 없으면 top-k 압축 행 선택이 모델의 그래프와 달라지게 된다고 명시되어 있습니다.

Flash의 형상(shape)에서는 n_hc = 4, n_hc_sinkhorn_iter = 20입니다. DS4는 은닉 상태(hidden state)를 단일 스트림으로 갖는 것이 아니라, Hyper-Connection의 4개 스트림을 다룹니다.

Metal 커널 kernel_dsv4_hc_split_sinkhorn은 믹서 행(mixer row)을 다음과 같이 분해합니다.

• pre weights (사전 가중치)
• post gates (사후 게이트)
• HC 간의 결합 행렬 (coupling matrix)

HC=4인 고속 경로에서는 4x4 행렬을 만든 후 Sinkhorn 정규화(Sinkhorn normalization)를 반복합니다.

for (int iter = 1; iter < args.sinkhorn_iters; ++iter) {
r0 *= 1.0f / (row_sum0 + epsv);
...
...

Sinkhorn 정규화는 행 방향과 열 방향의 정규화를 교대로 수행하여 행렬을 이중 확률 행렬(doubly stochastic matrix)에 가깝게 만드는 기법입니다. 여기서는 4개의 잔차 스트림(residual streams)을 어떻게 섞을지를 학습된 믹서(mixer)로부터 생성합니다.

DS4가 "일반적인 Transformer 러너"가 아니라는 점은 이러한 모델 고유의 그래프에서 잘 드러납니다.

ds4_metal.m의 도입부에는 전용 파이프라인 상태(pipeline state)가 대량으로 나열되어 있습니다.

static id<MTLComputePipelineState> g_hc_split_sinkhorn_pipeline;
static id<MTLComputePipelineState> g_dsv4_fp8_kv_quantize_pipeline;
static id<MTLComputePipelineState> g_dsv4_indexer_qat_pipeline;
...

이 목록만 보더라도, DwarfStar가 "범용 텐서 그래프 실행기 (General-purpose Tensor Graph Executor)"가 아니라, DeepSeek V4의 핫패스 (Hot path)를 직접 커널화 (Kernelize)하고 있음을 알 수 있습니다.

예를 들어 compressor store는 범용 그래프라면 APE 복사 (Copy), 스코어 가산 (Score addition), set_rows를 여러 번 디스패치 (Dispatch)해야 할 부분을, kernel_dsv4_compressor_store_one을 통해 단 1회의 디스패치로 처리합니다. ratio-4 shift 또한 범용 복사가 아닌 전용 커널 (Dedicated kernel)입니다.

이러한 종류의 특화는 유지보수 대상 모델이 늘어날수록 부담이 됩니다. 그렇기에 DS4는 제1회에서 보았듯이, "한 번에 하나의 모델"이라는 좁은 도박을 선택하고 있습니다.

제1회에서는 DS4의 압축 어텐션 (Compressed Attention)을 개념적으로 살펴보았습니다. 코드까지 읽어보면 중요한 점은 다음과 같습니다.

• 압축률은 4뿐만 아니라, Flash에서는 0/4/128이 층(Layer)마다 혼재함
• raw KV는 전체 컨텍스트가 아닌 SWA 128행
• compressor는 학습된 스코어 (Learned score) + APE + softmax 풀링 (Pooling)
• ratio-4 층만 indexer 경로를 가지며, top-k 압축 행을 선택함
• 어텐션 커널 (Attention kernel)은 raw 행과 선택된 압축 행을 동일한 online softmax로 통합함
• RoPE, FP8/QAT 왕복, HC Sinkhorn까지 DeepSeek V4 전용 그래프로 구현되어 있음

다음 회차에서는 이 KV 상태를 어떻게 저장하고 어떻게 재사용하는지를 다룹니다. DwarfStar의 README에서 말하는 "KV cache is a first-class disk citizen"은 단순한 철학이 아니라, DSV4 페이로드 (Payload)와 KVC 스토어 (Store)라는 구체적인 파일 형식으로 구현되어 있습니다.

본 기사는 Quick Iterate의 로컬 LLM 연구의 일환으로, 공개 리포지토리(Repository) antirez/ds4의 코드를 분석한 것입니다. 행 번호, 상수, 벤치마크 값은 열람 커밋 ba00a8a (2026-05-30) / README 취득일 2026-06-01 시점의 것입니다. ds4-agent는 alpha, 엔진 본체는 beta 품질로 활발하게 변화하므로, 인용된 부분은 각자 최신 README 및 소스 코드를 통해 다시 확인하시기 바랍니다.

Quick Iterate 주식회사 ― IoT / 전력 모니터링 / AI / 위성·무선 통신 / 시스템 통합 (System Integration)

로컬 LLM·에이전트 기반에 관한 문의는 언제든 환영합니다.