Study · GPU kernels

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을 기다리기 시작한다.

time per byte moved = max(DRAM time, instruction-issue time) — the budget, sketched ALU ops issued per byte moved → time DRAM time — flat: the bytes cost what they cost instruction-issue time — rises with every op per byte budget ≈ 9.8–12.5 ops/byte (A100, QServe / est.) left: memory-bound — byte savings pay right: ALU-bound — they don’t
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해 보라.

HBM 1,555 GB/s L2 40 MB SMEM ≤164 KB/SM Registers 256 KB/SM UNPACK shf · prmt · lop3 DEQUANT (q − z) · s MMA / FMA the actual math the path every KV byte takes at decode time the stages quantization adds — CUDA-core instructions, not tensor-core math QServe: on A100, one CUDA-core op ≈ 50 INT4 tensor-core ops

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배

one warp 32 threads cp.async 16 B / thread aligned · contiguous one up-to-128 B transaction = 4 × 32 B sectors 32 B32 B32 B32 B what 128 B carries: FP16 · 64 elements INT4 · 256 elements · INT2 · 512 elements SMEM staged tiles 32 banks × 32 b/clk Peak-bandwidth pattern: aligned, contiguous, 16 B per thread (128-bit) — anything else wastes fetched sectors. Packing already paid off here: same 128 B, 4–8× the elements.

2Into registers — where the budget livesregister로 — 예산이 사는 곳

SMEM tile ld.shared 128 bits / thread 4 × 32-bit registers per thread — R0 zoomed: 8 INT4 codes R0R1R2R3 q7 q6 q5 q4 q3 q2 q1 q0 Register file: 256 KB/SM = 64 K × 32-bit (GA100). Issue: 4 warp schedulers × 1 instruction/clock — INT32 and FP32 pipes run concurrently at full rate, but every unpack instruction still eats a shared issue slot.

3Unpack: per-element loops vs register-wide bit tricksunpack: element별 루프 vs register 단위 bit trick

naive — per element shr and cvt fma ≈ 5 ALU ops × every element (QServe) → ≈ 32–40 instr / 8 elements (est.) engineered — per register lop3 prmt shf whole registers at once → ≈ 10–12 instr / 8 elements incl. dequant (est.) MARLIN’s trick, one register at a time R: …q3 q2 q1 q0 (packed INT4) lop3(R, mask, fp16-exponent magic) fp16(q1) fp16(q0) two FP16 values per instruction — SIMD within the register QServe: with offline interleave (w0 w16 w1 w17 …), a register’s worth of INT4→INT8 unpacks in 3 logic ops bit-width drops → bytes shrink, this work doesn’t

4Fetch the group’s scale and zero-point그룹의 scale과 zero-point 가져오기

payload stream — packed codes g0 · 64 Bg1 · 64 Bg2 · 64 Bg3 · 64 B metadata stream — FP16 scale + zero per group s·zs·zs·zs·z MARLIN reorganizes scales like the weights and reloads them per sub-tile so the compiler’s pipelined instruction order survives; QServe stores FP16 scales/zeros inline in each KV-cache page, right after the quantized features. §5 prices this stream.

5Dequant FMA, then the one MAC that mattersdequant FMA, 그리고 진짜 중요한 MAC 하나

dequant, two elements per instruction x̂ = (q − z) × s HSUB2 · HFMA2 QServe moves zero-point subtraction into the epilogue algebraically — off the per-element critical path. the actual math softmax(q·Kᵀ)·V — batched GEMV ~1 MAC per element, regardless of batch (QServe) tensor-core mma needs ldmatrix — which “will not work when storage and compute data types differ” (QServe) everything before this box was overhead

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× 느림을 실측
ALU ops per byte moved (per 8 elements; estimates on verified anchors) 9.8 · A100-80GB (QServe) ≈12.5 · A100-40GB (est.) DRAM-bound: byte savings pay ALU-bound: they don’t FP16 passthrough INT4 fused INT4 naive INT2 fused INT2 naive
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 DRAM issue
INT4 · fused + bit-tricked 4 B ≈ 10–12 ≈ 2.5–3 DRAM issue
INT4 · naive per-element 4 B ≈ 32–40 ≈ 8–10 DRAM issue
INT2 · fused + bit-tricked 2 B ≈ 12–16 ≈ 6–8 DRAM issue
INT2 · naive per-element 2 B ≈ 32–40 ≈ 16–20 DRAM issue

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하면 산술이 보인다.

uniform low-bit: two streams codes · 64 B / group (INT4, g=128) s·z 4 B/g mixed precision adds a third: tags + offsets variable-length records (§6) s·z width tags · ~0.25 B/token/head offsets · ~4 B/token → dependent load KVQuant’s full inventory: per-channel (K) / per-token (V) scales + 4-bit LUT indices + ~1% outliers in CSR/CSC. QServe: pointer arithmetic for strided layouts runs on CUDA cores — “32× lower throughput than the INT8 tensor core” — and blocks packed 128-bit loads.
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 1 Addressing becomes data-dependent주소가 데이터 의존적이 된다

uniform INT4 — affine: addr(i) = base + 64·i, every address known instantly t0 t1 t2 t3 t4 t5 t6 t7 base+64i mixed {0,2,4,8,16} — variable records: offset(i) = Σ widths before it (a prefix sum) 4b 2b 16b · 8b 4b 2b 8b either compute the prefix sum at runtime, or store an offset table (~4 B/token) — a dependent load before every payload load. Both defeat static addressing and 128-bit vector loads. (estimate)

failure 2 Warp divergence: you pay the sum, not the maxwarp divergence: max가 아니라 합을 낸다

CUDA Programming Guide (SIMT): on a data-dependent branch “the warp executes each branch path taken, disabling threads that are not on that path.” lane 0 lane 1 lane 2 lane 3 lane 4 lane 5 lane 6 lane 7 width 2-bit path 4-bit path 8-bit path total time = 2-bit path + 4-bit path + 8-bit path, serialized — a direct consequence of the SIMT rule above

failure 3 Coalescing and vector loads break on variable-length records가변 길이 record가 coalescing과 vector load를 깬다

uniform: records tile the 32-byte sectors exactly — every fetched byte is used 32 B · used 32/3232 B · used 32/3232 B · used 32/3232 B · used 32/32 mixed: records start mid-sector — loads straddle boundaries, sectors arrive partly wasted record straddles 2 sectors ← extra sector fetched FLUTE: threads “should access the quantized weight in granularity of 128 bits (or at least in powers of 2)” — for 3-bit it must split the weight into 1-bit + 2-bit slices and issue two separate async copies. Po2 widths {2,4,8,16} keep records byte-aligned; 3/5-bit fields straddle 32-bit words — the case shf funnel shifts exist for (PTX ISA); several kernels simply refuse odd widths (bitsandbytes/BitBLAS: no 3-bit — FLUTE).

failure 4 The 0-bit class: eviction means gather0-bit class: eviction은 곧 gather다

evicted tokens leave holes; the survivors must be gathered through a compacted index t0 t1 · t3 t4 · t6 · t8 t9 · t11 compacted row (gather via index list) Measured proxies for what this bookkeeping costs: SqueezeLLM moved just 0.45% of weights to a CSR side-band and paid ~10% kernel latency (2.4×→2.2×), needing “balanced hybrid kernels” with extra synchronization; KVQuant pays a CSR/CSC gather for its ~1% outliers; DiffKV makes parallel compaction of fragmented free memory the paper’s headline mechanism.
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

4b 2b 16b 8b 4b 2b 8b 4b interleaved widths: every warp hits every failure in §6
before — one stream, mixed recordsbefore — 한 stream에 혼합 record
4-bit pool · uniform tiles 8-bit pool tags within a pool: affine addresses, one code path, aligned loads
after — 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

HBM · packed→ dequant kernel→ HBM · fp16 buffer→ GEMV
before — dequantize to a DRAM buffer: the fp16 bytes you deleted get written and re-readbefore — DRAM buffer로 dequantize: 지웠던 fp16 byte를 다시 쓰고 다시 읽는다
HBM · packed→ registers: unpack + dequant + FMA→ output
after — only low-bit values ever cross DRAMafter — DRAM을 건너는 것은 low-bit 값뿐

KIVI: “We have fused the dequantization process with matrix multiplication.” KVQuant fuses 4-bit LUT dequant and on-the-fly RoPE into the matvec — Key matvec 33.3 → 25.6 µs at 2K context — 1.2–1.6× (Key) / 1.3–1.7× (Value) faster than fp16. Atom “fuses dequantization with FlashInfer … so that only low-bit values from KV-cache are loaded.” KIVI: “dequantization을 matrix multiplication과 fuse했다.” KVQuant는 4-bit LUT dequant와 on-the-fly RoPE를 matvec에 fuse한다 — 2K context에서 Key matvec 33.3 → 25.6 µs — fp16보다 1.2–1.6× (Key) / 1.3–1.7× (Value) 빠르다. Atom은 “dequantization을 FlashInfer에 fuse해 KV-cache에서 low-bit 값만 load되게” 한다.

F-c · MARLIN · QServe · FLUTE

Offline interleaved packingoffline interleaved packing

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

bank0 lanes pile onto one bank → up to 8-way conflict (4-bit LUT) → 8× serialized
before — one naive table, 32 banks, colliding lanesbefore — 순진한 table 하나, bank 32개, 충돌하는 lane들
vectorized LUT (all 2⁸ index pairs ≈ 1 KB), copies in different banks — conflict-free
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 세금

load tile 0 unpack+math 0 load tile 1 unpack+math 1 load tile 2 unpack+math 2 serial: the math units idle during every load, the memory system idles during every unpack
before — load and compute take turnsbefore — load와 compute가 번갈아 논다
cp.async tile 0 cp.async tile 1 cp.async tile 2 cp.async tile 3 unpack+math 0 unpack+math 1 unpack+math 2 P=4 stages in flight; prefetch P−1 tiles ahead; “all shared memory addressing completely static” (MARLIN)
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

bits INT4
bytes / 8 elem 4 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)

1× 2× 3× 4× crossover b ≈ 50–64 16326496128 batch size kernel speedup over FP16 GEMM (curve sketched through MARLIN’s reported values — labeled points)
batch b = 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에서.

~1 MAC/element, any batch (QServe) CUDA-core budget ≈ 9.8–12.5 ops/B naive INT4 = 5 ops/elem → 1.1–1.2× slower fused KV4 @ b1 = 1.2–1.7× faster (KVQuant) MARLIN: 3.9× @ b16–32 → 1.5× @ b128

§9Where this goes다음은 어디로 가는가

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마다 다시 유도해야 한다 — 분석 패턴은 이식되지만 상수는 아니다.

Going deeper더 깊이

Primary sources1차 자료

arXiv 2408.11743

MARLIN — Frantar et al., 2024

The reference INT4×FP16 GEMM kernel: interleaved packing, lop3 dequant, cp.async pipelines — and the honest speedup-vs-batch decay curve.INT4×FP16 GEMM kernel의 표준: interleaved packing, lop3 dequant, cp.async pipeline — 그리고 정직한 speedup-vs-batch 감쇠 곡선.

the kernel-engineering playbookkernel engineering 플레이북
arXiv 2407.10960

FLUTE — Guo et al., 2024

LUT-quantized GEMM: 128-bit access granularity, bit-slicing odd widths, and vectorized + duplicated lookup tables against bank conflicts.LUT-quantized GEMM: 128-bit access granularity, 홀수 width의 bit-slicing, bank conflict에 맞서는 vectorized + duplicated lookup table.

odd widths & LUTs홀수 width와 LUT
arXiv 2405.04532

QServe — Lin et al., 2024

W4A8KV4 serving: the 5-ops-per-element dequant accounting, compute-aware weight reordering, and why naive 4-bit KV lost to 8-bit.W4A8KV4 serving: element당 5 ops의 dequant 회계, compute-aware weight reordering, 그리고 naive 4-bit KV가 8-bit에 진 이유.

the ops-per-byte accountingops-per-byte 회계
arXiv 2401.18079

KVQuant — Hooper et al., 2024

Fused non-uniform KV dequant that is faster than fp16 at batch 1 — the existence proof that low-bit KV loading can win outright.batch 1에서 fp16보다 빠른 fused non-uniform KV dequant — low-bit KV loading이 완승할 수 있다는 존재 증명.

the batch-1 counterpointbatch-1 반례
arXiv 2402.10517

Any-Precision LLM — Park et al., 2024

Bitplane storage: a lower width means reading fewer planes, not taking a different branch — width variation without divergence.bitplane 저장: 낮은 width = 다른 branch가 아니라 더 적은 plane 읽기 — divergence 없는 width 가변.

bitplanesbitplane
arXiv 2310.19102

Atom — Zhao et al., 2023

Mixed-precision made regular by reordering: scattered outlier channels moved to a contiguous tail so each region is uniform.reordering으로 규칙화한 mixed-precision: 흩어진 outlier channel을 연속된 꼬리로 옮겨 각 영역을 uniform하게.

the pool ancestorpool 패턴의 조상
arXiv 2412.03131

DiffKV — Zhang et al., 2024

Heterogeneous KV precision as a memory-management problem: differentiated pages plus parallel compaction of fragmented free memory.이종 KV precision을 memory management 문제로: differentiated page와 파편화된 free memory의 parallel compaction.

mixed-precision KV systemsmixed-precision KV 시스템
arXiv 2306.07629

SqueezeLLM — Kim et al., 2023

The price of a sparse side-band, measured: 0.45% of weights in CSR costs ~10% kernel latency and needs balanced hybrid kernels.sparse side-band의 가격 실측: weight의 0.45%를 CSR로 빼는 데 kernel latency ~10%, balanced hybrid kernel 필요.

what gather costsgather의 비용

Back to all study notes모든 study 노트로 돌아가기