Why 8× smaller can load slower: mixed-precision KV caches inside the kernel8× 작은데 왜 더 느리게 읽히는가: kernel 안의 mixed-precision KV cache
Quantization shrinks the bytes; the kernel still has to unpack, dequantize, and feed
them to the math units. This page walks that path at the instruction level — using only
published kernel-engineering results (MARLIN, FLUTE, QServe, KVQuant, and friends) and
vendor-documented A100 facts — to show exactly when byte savings turn into time savings,
and when they quietly don’t.
Quantization은 byte를 줄이지만, kernel은 여전히 그 byte를 unpack하고 dequantize해서
연산 유닛에 먹여야 한다. 이 페이지는 그 경로를 instruction 수준에서 따라간다 — 공개된
kernel engineering 결과(MARLIN, FLUTE, QServe, KVQuant 등)와 vendor가 문서화한 A100
사실만으로, byte 절약이 언제 시간 절약이 되고 언제 조용히 안 되는지를 정확히 보인다.
Scope note: various published schemes assign different bit-widths per token, head, or
channel (MiKV, QAQ, DiffKV). How widths get assigned is out of scope here —
everything below takes an assignment as given and asks only how it is stored, loaded,
and computed fast.
범위 주석: 공개된 여러 기법이 token·head·channel별로 서로 다른 bit-width를 부여한다
(MiKV, QAQ, DiffKV). width를 어떻게 정하는지는 이 페이지의 범위 밖이다 —
아래의 모든 내용은 부여가 주어졌다고 놓고, 그것을 어떻게 저장하고 읽고 빠르게
계산하는지만 다룬다.
Before you start: this page assumes you know roughly what the KV cache
is and why decode-time inference re-reads it every step — the
KV-cache quantization walkthrough
covers that, including how the low-bit codes are made. It also helps to have seen the
GPU memory hierarchy once (HBM, L2, shared
memory, registers). §2 rebuilds the minimum of both inline, so you can also just read on.
시작 전에: 이 페이지는 KV cache가 무엇이고 decode 추론이 왜 매 step
그것을 다시 읽는지 정도는 안다고 가정한다 —
KV-cache quantization 워크스루가
low-bit code가 만들어지는 과정까지 다룬다.
GPU memory hierarchy(HBM, L2, shared memory,
register)를 한 번 본 적이 있으면 더 좋다. §2에서 둘의 최소한을 다시 쌓으니 그냥
읽어 내려가도 된다.
§1The itch: 8× smaller should mean 8× faster — it often doesn’t문제의식: 8× 작아지면 8× 빨라져야 하는데 — 자주 그렇지 않다
When a language model generates text, most of what the GPU does at each step is not
arithmetic — it is re-reading a growing record of everything generated so far
(the KV cache). That re-reading is the speed limit, so the obvious fix is to make the
record smaller: store each number in 4 or 2 bits instead of 16. The storage claim is a
single bar chart, and it looks unbeatable:
언어 모델이 텍스트를 생성할 때 GPU가 매 step 하는 일의 대부분은 연산이 아니다 —
지금까지 생성한 모든 것의 기록(KV cache)을 다시 읽는 것이다. 그 재독이
속도의 상한이므로, 자연스러운 해법은 기록을 줄이는 것이다: 숫자 하나를 16 bit 대신
4 bit나 2 bit로 저장한다. 저장의 주장은 막대 그래프 하나이고, 이길 수 없어 보인다:
FP16
256 B / row
INT2
32 B / row · 8×
The promise — payload bytes only. The rest of the page is about what those bytes cost to use.약속 — payload byte만 센 것이다. 페이지의 나머지는 그 byte를 쓰는 데 드는 비용 이야기다.
And yet the published record contains this result: on the same A100 GPU, a
straightforwardly-written 4-bit KV kernel ran 1.1–1.2× slower than the 8-bit
one it was supposed to beat (QServe) — half the bytes, more time. Meanwhile another
paper’s 4-bit KV kernels beat full-precision fp16 outright, 1.2–1.7× (KVQuant). Same
silicon, same idea, opposite outcomes. Something between “fewer bytes” and “less time”
is not free — and this page is about naming that something precisely.
그런데 공개된 기록에는 이런 결과가 있다: 같은 A100에서, 평범하게 작성한 4-bit KV
kernel이 이겨야 할 8-bit보다 1.1–1.2× 느렸다 (QServe) — byte는 절반인데
시간은 더 걸렸다. 한편 다른 논문의 4-bit KV kernel은 full-precision fp16을 1.2–1.7×로
완파했다 (KVQuant). 같은 실리콘, 같은 아이디어, 정반대의 결과. “byte가 준다”와
“시간이 준다” 사이의 무언가가 공짜가 아니라는 뜻이고, 이 페이지는 그 무언가를
정확히 이름 붙이는 페이지다.
What does the GPU actually have to do to use a byte it just saved?GPU는 방금 아낀 byte를 쓰기 위해 실제로 무엇을 해야 하는가?
§2Foundations: the four facts everything else stands on기초: 이후의 전부가 딛고 서는 네 가지 사실
1 · Decode is a re-read, and it is memory-bound. Generating one token
multiplies one query against every stored K and V row — a matrix-vector product where
each loaded element is used in roughly one multiply-accumulate, no matter the
batch size (QServe calls this out explicitly). One operation per element loaded means
the math units finish long before the memory system does: the kernel’s time is
bytes ÷ bandwidth. This is exactly why shrinking bytes should work — and it is the
regime all of §3–§7 lives in.
1 · decode는 재독이고, memory-bound다. token 하나를 생성하려면 query
하나를 저장된 모든 K·V row와 곱한다 — batch 크기와 무관하게 load된 element 하나가
대략 multiply-accumulate 한 번에 쓰이는 matrix-vector 곱이다 (QServe가
명시적으로 짚는 사실). element당 연산 1개라는 것은 연산 유닛이 memory system보다
훨씬 먼저 끝난다는 뜻이다: kernel의 시간은 bytes ÷ bandwidth다. byte를 줄이는 것이
통해야 하는 이유가 정확히 이것이고, §3–§7 전체가 이 regime 안의 이야기다.
2 · The machine executes 32 threads in lockstep. A GPU streaming
multiprocessor (SM) issues instructions to warps — groups of 32 threads that
all run the same instruction at once. Bytes reach those threads through a hierarchy:
HBM (the big, slow-per-request DRAM) → L2 cache → shared memory (a small scratchpad
per SM) → registers, where instructions actually operate. Two kinds of execution units
sit at the end: tensor cores, which do dense matrix math at enormous rates, and
ordinary CUDA cores, which do everything else — shifts, masks, type conversions —
about 50× slower per op (QServe’s A100 comparison).
2 · 하드웨어는 thread 32개를 발맞춰 실행한다. GPU의 streaming
multiprocessor(SM)는 warp — 같은 instruction을 동시에 실행하는 thread 32개
묶음 — 단위로 instruction을 발행한다. byte는 계층을 지나 thread에 도달한다: HBM(크지만
요청당 느린 DRAM) → L2 cache → shared memory(SM마다 있는 작은 scratchpad) → register,
여기서 instruction이 실제로 값을 만진다. 끝에는 두 종류의 연산 유닛이 있다: dense
행렬 연산을 엄청난 속도로 처리하는 tensor core, 그리고 나머지 전부 — shift, mask,
type 변환 — 를 맡는 평범한 CUDA core. 후자는 op당 약 50× 느리다 (QServe의 A100 비교).
3 · A quantized cache stores codes, not numbers. Quantization maps
each group of values (say 128 of them) onto small integer codes plus two recovery
constants — a scale s and zero-point z. To use a code q the kernel must compute
x̂ = (q − z) · s. So every low-bit read implies extra work that fp16 never did:
extracting packed codes from a 32-bit word (unpack) and applying the recovery
arithmetic (dequant) — all of it on the slow CUDA cores, because tensor cores
only consume cleanly-typed, cleanly-laid-out operands.
3 · quantized cache는 숫자가 아니라 code를 저장한다. quantization은
값들의 그룹(예: 128개)을 작은 정수 code와 복원 상수 두 개 — scale s, zero-point z —
로 사상한다. code q를 쓰려면 kernel이 x̂ = (q − z) · s를 계산해야 한다. 즉 모든
low-bit 읽기에는 fp16이 하지 않던 추가 작업이 붙는다: 32-bit word에서 packed code를
꺼내는 unpack, 복원 산술을 적용하는 dequant — 전부 느린 CUDA core
몫이다. tensor core는 type과 layout이 깔끔한 피연산자만 먹기 때문이다.
4 · So there is a budget. While the memory system delivers bytes, the
CUDA cores can execute only so many instructions “for free” under that shadow. Divide
CUDA-core throughput by DRAM bandwidth and you get a number: on an A100-80GB,
≈ 9.8 ALU ops per byte moved (QServe; ≈ 12.5 on the 40 GB part — estimate). Stay under
it and dequant hides entirely under DRAM time; cross it and the kernel stops waiting
on memory and starts waiting on its own unpack instructions.
4 · 그래서 예산이 있다. memory system이 byte를 나르는 동안 CUDA
core가 그 그늘 아래에서 “공짜로” 실행할 수 있는 instruction 수에는 한계가 있다.
CUDA-core throughput을 DRAM bandwidth로 나누면 숫자가 나온다: A100-80GB에서 이동
byte당 ≈ 9.8 ALU ops (QServe; 40 GB 모델은 ≈ 12.5 — 추정). 이 밑에 있으면 dequant는
DRAM 시간 밑에 완전히 숨고, 넘는 순간 kernel은 memory가 아니라 자기 unpack
instruction을 기다리기 시작한다.
Takeaway핵심 한 줄The whole page is one accounting identity: a decode kernel is fast when its time equals bytes ÷ bandwidth — and every instruction spent on unpack, dequant, addressing, or gather is a claim against that identity.페이지 전체가 하나의 회계 항등식이다: decode kernel은 시간이 bytes ÷ bandwidth와 같을 때 빠르고, unpack·dequant·주소 계산·gather에 쓰는 모든 instruction은 그 항등식에 대한 청구다.
Where exactly, along a byte’s journey, does that extra work get inserted?그 추가 작업은 byte의 여정 어디에 정확히 끼어드는가?
§3The promise vs the path약속 vs 경로
We now have the budget (§2). What §1’s bar chart hid is a pipeline: every saved byte still has to travel the full memory hierarchy — and then through two stages fp16 never needed.이제 예산(§2)이 손에 있다. §1의 막대가 숨긴 것은 파이프라인이다: 아낀 byte도 여전히 memory hierarchy 전체를 지나야 하고 — 그 다음 fp16에는 없던 두 stage를 지나야 한다.
The performance story, stage by stage: every byte traverses HBM → L2 → shared memory →
registers, and then an unpack + dequant instruction sequence, before any math
happens. Hover each stage.
성능의 이야기를 stage별로: 모든 byte는 HBM → L2 → shared memory → register를 지나고
그 다음 unpack + dequant instruction 시퀀스를 지나야 비로소 연산이 시작된다.
각 stage에 hover해 보라.
Why this can break: decode attention reads the KV cache at an arithmetic intensity of
~1 MAC per element regardless of batch size (QServe) — it is
DRAM-bandwidth-bound, which is exactly why KVQuant argues compression is worth “a more
complex dequantization process.” But dequant runs on CUDA cores, and the A100
CUDA-core roofline is only ≈ 9.8 ops per byte moved (QServe, A100-80GB; ≈ 12.5 on the
40 GB part — estimate). Cross that line and the kernel stops waiting on DRAM and
starts waiting on its own unpack instructions.
이것이 깨질 수 있는 이유: decode attention은 batch 크기와 무관하게 element당
~1 MAC의 arithmetic intensity로 KV cache를 읽는다 (QServe) — DRAM
bandwidth-bound이고, 바로 그래서 KVQuant는 “더 복잡한 dequantization을 감수하고라도”
압축할 가치가 있다고 논증한다. 그런데 dequant는 CUDA core에서 돌고, A100의 CUDA-core
roofline은 이동 byte당 겨우 ≈ 9.8 ops다 (QServe, A100-80GB; 40 GB 모델은 ≈ 12.5 —
추정). 이 선을 넘는 순간 kernel은 DRAM이 아니라 자기 unpack instruction을 기다리기
시작한다.
Takeaway핵심 한 줄Quantization doesn’t remove work — it moves work from the memory system to the CUDA cores: UNPACK and DEQUANT are now first-class stages of the dataflow, running on the pipe with ~1/50th of the tensor-core budget (QServe).quantization은 일을 없애지 않는다 — memory system의 일을 CUDA core로 옮긴다: UNPACK과 DEQUANT는 이제 dataflow의 일급 stage이고, tensor-core 예산의 ~1/50짜리 pipe에서 돈다 (QServe).
How many instructions do those two stages actually cost? Follow one warp and count.그 두 stage는 실제로 instruction 몇 개짜리인가? warp 하나를 따라가며 세어 보자.
§4Inside one warp: load → unpack → dequant → mathwarp 하나의 내부: load → unpack → dequant → 연산
§3 named the added stages; it did not price them. This section walks one warp through one transaction of packed KV, instruction by instruction, and hands you the bill in step 6.§3은 추가된 stage의 이름만 붙였지 가격은 매기지 않았다. 이 섹션은 warp 하나가 packed KV transaction 하나를 끌고 가는 과정을 instruction 단위로 따라가고, step 6에서 청구서를 건넨다.
One warp, one transaction of packed uniform low-bit KV, followed to the FMA. The
instruction counts are estimates anchored to QServe’s published 5 ops/element (naive)
and ~2 ops/element (optimized); the verdict meter in step 6 is the punchline.
warp 하나가 packed uniform low-bit KV transaction 하나를 FMA까지 끌고 가는 과정.
instruction 수는 QServe가 제시한 element당 5 ops (naive)와 ~2 ops (최적화)에 앵커한
추정치이고, step 6의 verdict meter가 punchline이다.
Step 1 / 6
1One transaction, four times the elementstransaction 하나에 element는 4배
2Into registers — where the budget livesregister로 — 예산이 사는 곳
3Unpack: per-element loops vs register-wide bit tricksunpack: element별 루프 vs register 단위 bit trick
4Fetch the group’s scale and zero-point그룹의 scale과 zero-point 가져오기
5Dequant FMA, then the one MAC that mattersdequant FMA, 그리고 진짜 중요한 MAC 하나
6The bill, per 8 elements — and where the bottleneck flips8 element당 청구서 — bottleneck이 뒤집히는 곳
at / over the wall — QServe measured 1.1–1.2× slower than KV8 on A100벽에 걸리거나 넘는다 — QServe가 A100에서 KV8보다 1.1–1.2× 느림을 실측
format포맷
bytes / 8 elem
instr / 8 elem (est.)
ops / byte
relative time / 8 elem (A100-80GB roofline, est.)8 elem당 상대 시간 (A100-80GB roofline, 추정)
FP16 passthrough
16 B
≈ 1
≈ 0.1
DRAMissue
INT4 · fused + bit-tricked
4 B
≈ 10–12
≈ 2.5–3
DRAMissue
INT4 · naive per-element
4 B
≈ 32–40
≈ 8–10
DRAMissue
INT2 · fused + bit-tricked
2 B
≈ 12–16
≈ 6–8
DRAMissue
INT2 · naive per-element
2 B
≈ 32–40
≈ 16–20
DRAMissue
Bars are per-row normalized: the longer bar is the bottleneck. Issue time is ops ÷
the 9.8 ops/byte CUDA-core budget (QServe, A100-80GB) — an estimate, like every
instruction count here; the anchors (5 ops/element naive, ~2 optimized, 1.1–1.2×
measured slowdown for naive 4-bit KV vs KV8) are QServe’s. KVQuant’s fused
kernels are the counterpoint: at batch 1 they beat fp16 outright, 1.2–1.7×.
막대는 행별 정규화다: 더 긴 쪽이 bottleneck이다. issue 시간은 ops ÷ 9.8
ops/byte CUDA-core 예산 (QServe, A100-80GB) — 여기의 모든 instruction 수처럼
추정치이고, anchor(naive element당 5 ops, 최적화 ~2, naive 4-bit KV의 KV8 대비
1.1–1.2× 실측 감속)는 QServe의 것이다. KVQuant의 fused kernel은 반례다: batch
1에서 fp16을 1.2–1.7× 완파한다.
Takeaway핵심 한 줄Per-element work does not shrink with bit-width: naive INT4 lands at ~10 ops/byte against a 9.8 budget — which is why QServe measured it 1.1–1.2× slower than 8-bit — and INT2 halves the bytes without halving the instructions, the arithmetic reason “2-bit is not 2× faster than 4-bit” without heroic kernels.element당 작업은 bit-width와 함께 줄지 않는다: naive INT4는 예산 9.8 대비 ~10 ops/byte에 도달하고 — 그래서 QServe가 8-bit보다 1.1–1.2× 느린 것을 실측했다 — INT2는 instruction은 그대로 둔 채 byte만 반으로 줄인다. heroic kernel 없이는 “2-bit가 4-bit보다 2× 빠르지 않은” 산술적 이유다.
Step 4 quietly fetched a scale and zero-point from somewhere. How many bytes does that second stream really add?step 4가 어디선가 scale과 zero-point를 조용히 가져왔다. 그 두 번째 stream은 실제로 몇 byte를 더하는가?
§5The second stream: what “2-bit” actually stores두 번째 stream: “2-bit”가 실제로 저장하는 것
The walkthrough priced the instructions. But §2’s fact 3 said a quantized cache stores codes plus recovery constants — and those constants are traffic too, uncounted so far.walkthrough는 instruction의 가격을 매겼다. 그런데 §2의 사실 3은 quantized cache가 code에 더해 복원 상수를 저장한다고 했다 — 그 상수들도 트래픽이고, 아직 세지 않았다.
Every group carries an FP16 scale (and usually a zero-point) — a second memory stream
with its own bytes, its own loads, its own addresses. Head dim 128, 4 B of FP16
scale+zero per group; hover each bar for the arithmetic.
모든 그룹은 FP16 scale(과 대개 zero-point)을 달고 다닌다 — 자기 byte, 자기 load, 자기
주소를 가진 두 번째 memory stream이다. head dim 128, 그룹당 FP16 scale+zero 4 B; 각
막대에 hover하면 산술이 보인다.
effective bits per value (payload + metadata; estimates from verified group sizes)값당 effective bits (payload + metadata; 검증된 group size에서의 추정)
payload codespayload codescales + zerosscale + zeroFP16 residual window (KIVI R=128; metadata-precision fp16 = standard practice, not stated in the paper)FP16 residual window (KIVI R=128; metadata 정밀도 fp16은 관행 — 논문에 명시되진 않음)
Takeaway핵심 한 줄Scales and zeros are real traffic: at KIVI’s group size 32 a “2-bit” cache stores 3 effective bits — metadata alone is half the payload — and the FP16 residual window (R=128) lifts a 512-token context to ≥ 6.25 bits/element (estimates from verified settings).scale과 zero는 실제 트래픽이다: KIVI의 group size 32에서 “2-bit” cache는 3 effective bits를 저장하고 — metadata만으로 payload의 절반 — FP16 residual window(R=128)는 512-token context를 ≥ 6.25 bits/element로 끌어올린다 (검증된 설정 위의 추정).
Uniform low-bit adds a second stream and survives. What if the width itself varies token by token?uniform low-bit는 두 번째 stream을 달고도 살아남는다. 그런데 width 자체가 token마다 달라지면?
§6Deep water: where mixed widths break the machine’s assumptions깊은 물: 혼합 width가 하드웨어의 가정을 깨는 곳
This is the hardest section of the page — take it slowly. Everything so far assumed one width for every token; §2’s machine facts (lockstep warps, whole 32-byte sectors, affine addresses) were all quietly load-bearing. Mixed precision removes the uniformity they stand on, and four things break at once.여기가 이 페이지에서 가장 어려운 섹션이다 — 천천히 가자. 지금까지는 모든 token이 같은 width라고 가정했고, §2의 하드웨어 사실들(lockstep warp, 통째로 소비되는 32-byte sector, affine 주소)이 전부 조용히 하중을 받치고 있었다. mixed precision은 그들이 딛고 선 균일성을 치우고, 네 가지가 동시에 부러진다.
Take a generic example alphabet {0, 2, 4, 8, 16} bits — some tokens evicted, some
low-bit, some full precision (how they’re assigned doesn’t matter here). Three GPU
assumptions break at once: affine addressing, one-path-per-warp, and aligned vector
loads. DiffKV states the systems version plainly: heterogeneous precision “introduces
irregular memory usage patterns … posing significant scalability challenges.”
일반적인 예시 alphabet {0, 2, 4, 8, 16} bits를 놓자 — 어떤 token은 evict, 어떤
것은 low-bit, 어떤 것은 full precision (여기서 부여 방식은 중요하지 않다). GPU의 세
가정이 동시에 깨진다: affine 주소, warp당 단일 경로, aligned vector load. DiffKV는
시스템 관점을 그대로 적는다: 이종 precision은 “irregular memory usage pattern을
도입해 … 심각한 scalability 문제를 일으킨다.”
failure 1Addressing becomes data-dependent주소가 데이터 의존적이 된다
failure 2Warp divergence: you pay the sum, not the maxwarp divergence: max가 아니라 합을 낸다
failure 3Coalescing and vector loads break on variable-length records가변 길이 record가 coalescing과 vector load를 깬다
failure 4The 0-bit class: eviction means gather0-bit class: eviction은 곧 gather다
Takeaway핵심 한 줄The four failures are one fact in different units: the GPU’s cheap paths — affine addresses, one instruction per warp, whole 32-byte sectors — assume uniformity. Data-dependent widths buy pointer chases, serialized lanes, wasted sectors, and gathers; SqueezeLLM’s 0.45%-sparse ⇒ ~10%-latency measurement is the cleanest published price tag for that bookkeeping.네 failure는 같은 사실을 다른 단위로 본 것이다: GPU의 싼 경로 — affine 주소, warp당 한 instruction, 통째로 소비되는 32-byte sector — 는 균일성을 가정한다. 데이터 의존적 width는 pointer chase, lane 직렬화, sector 낭비, gather를 사들인다; SqueezeLLM의 0.45% sparse ⇒ latency ~10% 실측이 그 장부 관리의 가장 깔끔한 공개 가격표다.
Every one of these breakages appears in published systems — so how do shipped kernels put the uniformity back?이 붕괴들은 전부 공개된 시스템에 등장한다 — 그렇다면 실제로 배포된 kernel들은 균일성을 어떻게 되찾는가?
§7The published fix patterns공개된 fix 패턴들
§6 ended in wreckage on purpose. The literature’s answer is not one trick but a repertoire — and every entry works the same way: restore, somewhere, the uniformity the hardware was promised.§6은 일부러 잔해로 끝냈다. 문헌의 답은 한 가지 트릭이 아니라 레퍼토리다 — 그리고 모든 항목이 같은 방식으로 작동한다: 하드웨어가 약속받았던 균일성을 어딘가에서 복원하는 것.
Each failure above has a published answer. Flip each card between the broken layout and
the engineered one — every pattern is cited to the paper that documented it.
위의 각 failure에는 공개된 답이 있다. 카드를 깨진 layout과 엔지니어링된 layout 사이로
뒤집어 보라 — 모든 패턴은 그것을 기록한 논문을 인용한다.
F-a · Atom → DiffKV · Any-Precision LLM
Precision-partitioned poolsprecision별 분할 pool
before — one stream, mixed recordsbefore — 한 stream에 혼합 recordafter — per-width contiguous regions + uniform tilesafter — width별 연속 영역 + uniform tile
Atom is the ancestor: it “reorders the scattered outlier channels … to the end of
the matrix” so each region is uniform. DiffKV builds the KV-cache version —
differentiated pages plus a GPU memory manager that “compacts fragmented free
memory … in parallel” (2.7–5.7× compression, 1.9–5.4× throughput). Any-Precision
LLM goes further with bitplanes: width becomes which planes you read, and
lower width “directly translates into proportional speedup.”
조상은 Atom이다: “흩어진 outlier channel을 행렬 끝으로 reorder”해 각 영역을
uniform하게 만든다. DiffKV는 KV-cache 버전을 만든다 — differentiated page와
“파편화된 free memory를 병렬로 compact하는” GPU memory manager (압축 2.7–5.7×,
throughput 1.9–5.4×). Any-Precision LLM은 bitplane으로 더 나간다: width가
읽는 plane의 수가 되고, 낮은 width는 “비례하는 speedup으로 직결된다.”
F-b · KIVI · KVQuant · Atom
Fused dequant-in-registersregister 안에서 fused dequant
sequential pack 0 1 2 3…→unpack→cross-lane shuffle at runtime→mma fragments
before — decoded values land in the wrong lanes; fixing that costs instructions every iterationbefore — decode된 값이 엉뚱한 lane에 떨어지고, 매 iteration마다 instruction으로 고쳐야 한다
pack once, permuted: 6 4 2 0 7 5 3 1→lop3 unpack→already in fragment layout
after — the shuffle happened offline, at pack timeafter — shuffle은 pack 시점에 offline으로 끝났다
MARLIN stores weights “interleaved, according to the pattern 64207531” for parallel
in-register decode; FLUTE reorders offline “such that after dequantization, the
weights are already laid out exactly in the expected format”; QServe’s
compute-aware reorder concatenates each thread’s channels “into a single 128-bit
word” — 128-bit loads with zero runtime address math, and INT4→INT8 unpack in 3
logic ops. KV caveat: K/V are generated at runtime, so “offline” becomes
once, at quantization time — QServe’s KV page layout with inline
scales/zeros is that instance. (Deeper cut: LUT-GEMM precomputes partial sums and
“eliminates the resource-intensive dequantization process” entirely — 2.1× over
the OPTQ kernel at 3-bit.)
MARLIN은 weight를 “64207531 패턴으로 interleave”해 저장해 register 내 병렬
decode를 만든다; FLUTE는 “dequantization 후 weight가 기대 format 그대로
놓이도록” offline으로 reorder한다; QServe의 compute-aware reorder는 각 thread의
channel을 “하나의 128-bit word로 연결” — 주소 계산 없는 128-bit load에 INT4→INT8
unpack이 3 logic ops다. KV 주의: K/V는 runtime에 생성되므로 “offline”은
quantize 시점에 한 번이 된다 — QServe의 inline scale/zero KV page
layout이 그 사례다. (더 깊은 수: LUT-GEMM은 partial sum을 미리 계산해
“dequantization 자체를 제거”한다 — 3-bit에서 OPTQ kernel 대비 2.1×.)
F-d · FLUTE
LUT replication vs bank conflictsLUT 복제 vs bank conflict
before — one naive table, 32 banks, colliding lanesbefore — 순진한 table 하나, bank 32개, 충돌하는 lane들after — pay ~1 KB of SMEM per copy, get parallel lookupsafter — 복사본당 SMEM ~1 KB를 내고 병렬 lookup을 얻는다
FLUTE documents both the problem (“up to 8-way bank conflicts” for a naive
vectorized 4-bit table) and the fix: vectorize the LUT into all 2⁸ index-pairs and
“duplicate the … lookup table multiple times, placing copies in different memory
banks.” Hardware rule: 32 banks, 32 bits per bank per clock, successive words in
successive banks; n-way conflicts serialize n× (CUDA docs). MARLIN’s non-LUT
analogue is an XOR-swizzled SMEM layout (element (i, j) placed via i⊕j).
FLUTE는 문제(“순진한 vectorized 4-bit table은 최대 8-way bank conflict”)와
해법을 모두 기록한다: LUT를 2⁸개 index-pair 전체로 vectorize하고 “복사본을 서로
다른 memory bank에 여러 벌” 둔다. 하드웨어 규칙: bank 32개, bank당 클럭당 32
bit, 연속 word는 연속 bank로; n-way conflict는 n× 직렬화된다 (CUDA docs).
MARLIN의 non-LUT 대응물은 XOR-swizzle SMEM layout이다 (element (i, j)를 i⊕j로
배치).
F-e · MARLIN · (cost side: QServe on Atom)
Pipelining and the occupancy taxpipelining과 occupancy 세금
before — load and compute take turnsbefore — load와 compute가 번갈아 논다after — dequant hides under DRAM time … while registers lastafter — dequant가 DRAM 시간 밑에 숨는다 … register가 버티는 동안은
MARLIN runs a 4-stage cp.async pipeline with an evict_first cache hint, multiple
warps per output tile and a logarithmic parallel reduction. The tax is registers:
QServe observes that Atom’s dual INT32 + FP32 accumulator sets, “consuming a large
number of registers within each warp, limit the number of warps that can be
executed simultaneously” — occupancy loss on a 256 KB/SM register file. Deeper
pipelines and fatter unpack sequences both spend the same budget.
MARLIN은 evict_first cache hint의 4-stage cp.async pipeline, output tile당 다중
warp, logarithmic parallel reduction으로 돈다. 세금은 register다: QServe는 Atom의
INT32 + FP32 이중 accumulator가 “warp마다 다량의 register를 소모해 동시에 실행될
수 있는 warp 수를 제한한다”고 지적한다 — 256 KB/SM register file 위의 occupancy
손실. 더 깊은 pipeline과 더 살찐 unpack 시퀀스는 같은 예산을 나눠 쓴다.
Takeaway핵심 한 줄Every fix restores a uniformity the hardware wants: pools make each warp’s tile single-width (Atom → DiffKV), bitplanes turn width into a read-count instead of a branch (Any-Precision), offline permutation moves shuffles to pack time (MARLIN, QServe, FLUTE), LUT replication de-serializes shared memory (FLUTE), and pipelining hides the rest under DRAM time — until registers run out.모든 fix는 하드웨어가 원하는 균일성을 복원한다: pool은 warp의 tile을 단일 width로 만들고 (Atom → DiffKV), bitplane은 width를 분기가 아닌 읽기 횟수로 바꾸고 (Any-Precision), offline permutation은 shuffle을 pack 시점으로 옮기고 (MARLIN, QServe, FLUTE), LUT 복제는 shared memory의 직렬화를 풀고 (FLUTE), pipelining은 나머지를 DRAM 시간 밑에 숨긴다 — register가 바닥나기 전까지.
With the full repertoire applied, does low-bit KV always win — or only in a regime?레퍼토리를 전부 적용하면 low-bit KV는 항상 이기는가 — 아니면 특정 regime에서만인가?
§8When does it actually win?그래서 언제 실제로 이기는가?
The fixes buy back the instructions; they do not repeal §2’s roofline. Whether byte savings convert to time still depends on which side of the budget the kernel sits — and on whether the kernel is memory-bound at all.fix들은 instruction을 되사 오지만 §2의 roofline을 폐지하진 못한다. byte 절약이 시간으로 환전되는지는 여전히 kernel이 예산의 어느 쪽에 앉아 있는지 — 그리고 애초에 memory-bound인지 — 에 달려 있다.
Two dials decide the regime: how many instructions each byte costs (bit-width × kernel
quality), and whether the kernel is memory-bound at all (batch). Drag both.
두 개의 다이얼이 regime을 정한다: byte 하나가 몇 개의 instruction을 요구하는가
(bit-width × kernel 품질), 그리고 kernel이 애초에 memory-bound인가 (batch). 둘 다
움직여 보라.
bytes saved vs instructions added아낀 byte vs 추가된 instruction
bitsINT4
bytes / 8 elem4 B
instr / 8 elem≈ 10–12
INT4, fused: 4× fewer bytes for ~10–12 instructions per 8 elements — stays memory-bound; the savings are real (KVQuant: 1.2–1.7× faster than fp16 at batch 1).INT4, fused: byte 4× 감소에 8 element당 instruction ~10–12개 — memory-bound가 유지되고 절약은 진짜다 (KVQuant: batch 1에서 fp16보다 1.2–1.7× 빠름).
speedup vs batch — MARLIN’s measured decay (A10, 4-bit weights)speedup vs batch — MARLIN의 실측 감쇠 (A10, 4-bit weight)
batchb = 32 · ≈ 3.9×
Memory-bound regime: MARLIN holds close to the maximum 4× for batches up to 16–32; FLUTE’s LUT gains likewise concentrate below batch 32.Memory-bound regime: MARLIN은 batch 16–32까지 최대 4×에 근접하게 유지한다; FLUTE의 LUT 이득도 batch 32 아래에 몰린다.
The honest summary, regime by regime: at batch 1 with fused kernels the naive
byte-counting story largely holds — KVQuant’s low-bit KV kernels beat fp16 outright.
Naive per-element dequant loses even there (QServe: 4-bit KV 1.1–1.2× slower than
8-bit on A100). And as batch grows the baseline roofline rises to meet you: MARLIN’s
fully-engineered kernel decays from ~3.9× to ~1.5× by batch 128 (end-to-end in vLLM:
2.8× at batch 16 → ~1.2× at 128). Every added op per byte moves that crossover
earlier.
regime별 정직한 요약: batch 1 + fused kernel에서는 byte 계산 그대로의 이야기가 대체로
성립한다 — KVQuant의 low-bit KV kernel은 fp16을 완전히 이긴다. naive element별
dequant는 거기서도 진다 (QServe: A100에서 4-bit KV가 8-bit보다 1.1–1.2× 느림). 그리고
batch가 커지면 baseline의 roofline이 올라와 마중 나온다: 완전히 엔지니어링된 MARLIN
kernel도 batch 128에서는 ~3.9× → ~1.5×로 감쇠한다 (vLLM end-to-end: batch 16에서
2.8× → 128에서 ~1.2×). byte당 op가 하나 늘 때마다 그 crossover는 앞당겨진다.
Takeaway핵심 한 줄No kernel trick moves the regime boundary itself: past the roofline crossover the baseline is no longer waiting on the bytes you saved — even MARLIN’s fully-engineered kernel decays from ~3.9× to ~1.5× by batch 128.어떤 kernel 트릭도 regime의 경계 자체는 옮기지 못한다: roofline crossover를 지나면 baseline은 더 이상 당신이 아낀 byte를 기다리지 않는다 — 완전히 엔지니어링된 MARLIN kernel조차 batch 128에서 ~3.9× → ~1.5×로 감쇠한다.
If the mechanics are this well mapped, what is the field still fighting over?메커니즘이 이렇게까지 정리돼 있다면, 이 분야는 아직 무엇을 두고 싸우고 있는가?
Why this matters왜 중요한가
“Bytes ÷ bandwidth” is the most seductive model in inference acceleration, and it is
right just often enough to be dangerous. The published record is unambiguous on both
sides: fused low-bit KV kernels genuinely beat fp16 at batch 1 (KVQuant), and a
naive 4-bit KV kernel genuinely lost to 8-bit on the same silicon (QServe) — the
difference is nothing but the instruction stream between the load and the FMA.
Reading compression papers with this page’s meter in mind turns “N× smaller” claims
into checkable questions: how many ops per byte, on which pipe, against which
baseline, at what batch.
“bytes ÷ bandwidth”는 inference 가속에서 가장 유혹적인 모델이고, 위험할 만큼만
자주 맞는다. 공개된 기록은 양쪽 모두 명확하다: fused low-bit KV kernel은 batch
1에서 fp16을 실제로 이기고 (KVQuant), naive 4-bit KV kernel은 같은 실리콘에서
8-bit에 실제로 졌다 (QServe) — 차이는 load와 FMA 사이의 instruction stream뿐이다.
이 페이지의 meter를 염두에 두고 압축 논문을 읽으면 “N× 작다”는 주장이 검증
가능한 질문으로 바뀐다: byte당 몇 ops를, 어느 pipe에서, 어떤 baseline 상대로,
어떤 batch에서.
The mechanics above are settled science by systems standards; the open questions are about what to do with them. Three directions are visibly active in the literature.위의 메커니즘은 시스템 기준으로는 정리된 과학이다; 열린 질문은 그것으로 무엇을 하느냐다. 문헌에서 눈에 띄게 활발한 방향이 셋 있다.
Below 4 bits with better codes. Uniform integer grids waste precision
on well-behaved distributions, so newer quantizers spend math to earn bits — non-uniform
LUTs (KVQuant), rotations that spread outliers before
coding (TurboQuant,
LiftQuant). Every new code re-opens this page’s
question: what does its decode cost per byte, and on which pipe?
Choosing the widths. This page took the assignment as given; deciding
it — which tokens, heads, or layers deserve bits — is its own literature, from eviction
scores (H2O) to head-level retrieval-vs-streaming splits
(DuoAttention), and each policy’s output lands on
§6’s systems bill. The budget itself moves. Every ops-per-byte number
here is A100-specific; new accelerator generations shift the CUDA-core-to-bandwidth
ratio and add native low-bit datapaths, so the crossover lines in §8 must be re-derived
per part — the analysis pattern transfers, the constants do not.
더 좋은 code로 4 bit 아래로. uniform 정수 grid는 얌전한 분포에서
정밀도를 낭비하므로, 최신 quantizer들은 연산을 써서 bit를 번다 — non-uniform LUT
(KVQuant), coding 전에 outlier를 퍼뜨리는 rotation
(TurboQuant,
LiftQuant). 새 code가 나올 때마다 이 페이지의 질문이
다시 열린다: 그 decode는 byte당 얼마이고, 어느 pipe에서 도는가?
width를 정하는 문제. 이 페이지는 부여를 주어진 것으로 놓았지만,
그것을 정하는 일 — 어떤 token·head·layer가 bit를 받을 자격이 있는가 — 은 그 자체로
하나의 문헌이다. eviction score(H2O)부터 head 단위
retrieval-vs-streaming 분할(DuoAttention)까지
있고, 모든 정책의 출력은 §6의 시스템 청구서 위에 떨어진다. 예산 자체가
움직인다. 여기의 모든 ops-per-byte 숫자는 A100 전용이다; 새 가속기 세대는
CUDA-core 대 bandwidth 비율을 바꾸고 native low-bit datapath를 더하므로, §8의
crossover 선은 part마다 다시 유도해야 한다 — 분석 패턴은 이식되지만 상수는 아니다.