Header Ad Placement

Kernel Engineering Under Physical Non-Idealities: Floorsweeping, Wave-Equalized Launches, and the DVFS Confounder Behind H100 Clustering Gains

GPU SYSTEMS

An NUS team shows the H100 you programmed is not the H100 you have: fabrication binning quietly removes whole SMs from cluster reach, and power limits rewrite the speedup math. Accounting for both recovers up to 5.30% on FlashAttention 4.

By Ossam Hashim

KEY FINDINGS

  • Floorsweeping, the binning step that disables defective silicon to salvage a die, is invisible to the ISA but visible to cluster placement: on the H100 SXM5 it removes GPC-level clustered execution capacity, cutting SMs reachable by clusters from 132 to 124 (paper's measurement).
  • Cluster placement also drifts at runtime: the paper's microbenchmark records only 120 unique SM IDs per kernel iteration but 124 across iterations, so the effective SM budget oscillates launch to launch.
  • Fewer usable SMs can add a full execution wave: a launch that needs one wave at 132 SMs can need two at 120, erasing the data-movement savings DSMEM clustering was supposed to buy.
  • The fix is arithmetic, not hardware: wave-aware clustering picks per-kernel between TPC-level (size 2) and GPC-level (size 4/8) clusters so the thread-block count stays wave-equalized under both the full and reduced SM limits.
  • On the Hopper FlashAttention 4 forward kernel (whose baseline does not cluster at all), wave-aware clustering improves throughput by up to 5.30% across 2K-64K sequence lengths (authors' reported results).
  • The second confounder is power: under TDP limits, reduced data movement lets DVFS sustain higher SM clocks, so part of a memory optimization's gain is really a frequency gain; the same power wall caps throughput scaling on Tenstorrent's Wormhole n300d.

The Abstraction That Hides the Silicon

Kernel engineers work through stacked abstractions: language, compiler, ISA, microarchitecture. Each layer hides what sits below, and mostly that is the point. This paper, from the National University of Singapore, documents what the last layer hides: the physical state of the chip itself. Two effects matter. Floorsweeping disables defective units after fabrication and sorts dies into SKUs, so two chips sold under one name can present different resource counts to cluster scheduling. And the thermal design power budget governs, through dynamic voltage and frequency scaling, the clock the chip can sustain while your kernel runs. Neither appears in any programming guide, yet both can decide whether an optimization works. The case study is distributed shared memory on Hopper. DSMEM lets thread blocks in a cluster read each other's shared memory across a per-GPC network-on-chip, so tiles are loaded once and broadcast instead of refetched. Clusters of size 2 map within a texture processing cluster; sizes 4 and 8 require whole GPC-level groupings, and the GPU carries eight GPCs.

Floorsweeping Eats the Cluster

Here is the catch the vendor documentation abstracts away: cluster placement is constrained by the physical organization of the SMs, and floorswept GPCs do not have enough live SMs to host a size-4 or size-8 cluster. The paper measures the consequence directly by logging logical SM IDs during a clustered microbenchmark. At GPC-level cluster sizes, the SMs available for clustered execution drop from the H100 SXM5's 132 to 124, and within any single iteration only 120 unique SM IDs appear. Which SMs sit out varies between launches. Why does losing 8 to 12 SMs out of 132 matter? Wave quantization. A kernel's thread blocks execute in waves sized by the number of SMs that can host them. A launch of 121 to 132 blocks runs in one wave at 132 SMs but needs two waves at 120, and the second wave runs nearly empty. The throughput cliff from an extra wave can be larger than everything DSMEM saved on redundant global loads, so a kernel tuned on paper can regress on silicon.

Wave-Equalized Launches

The paper's response is to make the wave count explicit and choose cluster geometry around it. A launch is wave-equalized when its block count produces the same number of execution waves under both the full 132-SM and the reduced 120-SM limits; the paper derives the condition as a ceiling-function identity over tile counts for GEMM and over batch, head and sequence tiles for FlashAttention. The wave-aware clustering strategy then evaluates each kernel configuration and selects TPC-level or GPC-level clustering so the launch stays wave-equalized while still capturing DSMEM reuse. The results land where the theory says they should. On GEMM the strategy picks the best clustering per shape. On FlashAttention 4's Hopper forward kernel, whose stock implementation uses cluster size 1, wave-aware clustering matches or beats the baseline at almost every configuration from 2K to 64K sequence length, with gains up to 5.30%. The point is not the percentage; it is that the winning move was counting waves against the chip's real, floorswept floor plan rather than its nominal one.

The DVFS Confounder

The paper's second contribution is an explanation for speedups that looked anomalous. Under a TDP cap, sustained SM frequency floats with power draw. A clustering change that reduces data movement reduces energy per iteration, and DVFS converts that headroom into a higher sustained clock. Part of the measured throughput gain is therefore a frequency effect, not a locality effect, and the paper shows the two can be disentangled by logging clocks alongside performance. The same physics shows up on a very different machine: on Tenstorrent's Wormhole n300d, a dual-chip card with 64 Tensix cores at 1 GHz and 64 TFLOP/s peak per chip, throughput scaling flattens as more cores activate because the power budget forces the frequency down. Two vendors, two architectures, one lesson: kernel performance models that stop at the microarchitecture miss the variables that increasingly set the ceiling. Kernel engineers who care about these margins, the audience of this blog's Humming kernel analysis, now have one more hidden variable to log.

Assumptions and limitations

All performance figures are the authors' reported measurements on their H100 SXM5 units and Wormhole n300d card; floorsweeping outcomes vary per chip, so the 124/120 figures describe their silicon, not a universal constant. This analysis did not independently re-run the kernels. The FlashAttention 4 baseline referenced is the Hopper forward kernel as used in the paper. Wave-equalization conditions are derived for GEMM and FlashAttention tile shapes; other kernels need their own derivation. The worked example is a simplified illustration, not the paper's methodology.

Worked Example: Counting Wasted Waves

This toy calculator reproduces the wave-quantization effect: for thread-block counts near the 120/132 boundary, it prints the wave count at both SM limits and flags launches where floorsweeping adds a nearly empty wave. Illustrative arithmetic, not the paper's harness.

import math
def waves(b, sms):
    return math.ceil(b / sms)
print("B    w132 w120 verdict")
penal = 0
for b in range(115, 135):
    a = waves(b, 132)
    c = waves(b, 120)
    if c > a:
        v = "wasted wave"
        penal += 1
    else:
        v = "equalized"
    print(b, a, c, v)
print("penalized:", penal, "of 20")

Assumptions and limitations All performance figures are the authors' reported measurements on their H100 SXM5 units and Wormhole n300d card; floorsweeping outcomes vary per chip, so the 124/120 figures describe their silicon, not a universal constant. This analysis did not independently re-run the kernels. The FlashAttention 4 baseline referenced is the Hopper forward kernel as used in the paper. Wave-equalization conditions are derived for GEMM and FlashAttention tile shapes; other kernels need their own derivation. The worked example is a simplified illustration, not the paper's methodology.

Sources & Provenance

Analysis is based on the primary sources listed above; performance claims are attributed to their authors and not independently verified.

BellowsKV Architecture: Block-Shared Cache Latents, Outlier-Free Low-Bit Storage, and Single-Decision Eviction for Training-Free 13.8x KV Compression

INFERENCE SYSTEMS

A Rice University measurement study becomes a serving method: adjacent transformer layers cannot share their caches, but they can share what the caches are computed from. One latent per block, quantized without rotation, keeps about 98% of full-cache accuracy at 13.8x compression.

By Ossam Hashim

KEY FINDINGS

  • The KV cache now rivals the weights: on Llama-3.1-8B, one 131K-token sequence carries a 16.8 GB fp16 cache against 16 GB of weights; a batch of eight holds 134 GB of cache (paper's figures).
  • Adjacent layers cannot reuse each other's cache: neighboring keys and values are near-orthogonal (cosine similarity about zero), and even an oracle-chosen shared key raises perplexity significantly.
  • The redundancy lives in the inputs: hidden states feeding adjacent layers reach 0.87 cosine similarity on Llama and 0.82 on Qwen, increasing with depth.
  • The shared latent quantizes cleanly: every coordinate mixes channels from all n layers, so no outlier channel sets the scale; relative INT4 error falls from 0.156 (raw cache) to about 0.10, with no rotation pass at read time.
  • Layers of a block agree on which tokens matter (Spearman 0.76 to 0.80), so one eviction decision replaces n per-layer decisions; shrinking the latent's dimension is the one compression axis that fails.
  • BellowsKV-Quant retains about 99% of full-cache accuracy at 3.6x and 97-98% at 13-14x compression across Llama-3.1-8B, Mistral-7B and Qwen2.5-7B; the strongest prior training-free cross-layer method retains at most 84% at 16x (authors' reported LongBench results).

The Cache That Outgrew the Weights

Long-agent workloads changed the memory balance of LLM serving. Weights load once and serve every sequence in a batch; each sequence's KV cache is private, re-read at every decoding step, and grows with both context length and batch size. The paper's arithmetic makes it concrete: a single 131K-token Llama-3.1-8B sequence carries 16.8 GB of cache in 16-bit precision, slightly more than the model itself after grouped-query attention already shrank it fourfold. Eight such sequences hold 134 GB of cache against 16 GB of shared weights. The established compression toolbox works inside one layer at a time: quantize the entries, evict low-importance tokens, or cut the channel dimension. Each family approaches a floor near 4x. Quantization degrades sharply below two bits and needs outlier machinery; eviction is irreversible once a later turn needs the dropped token; rank reduction removes coordinates attention still uses. Compressing 134 GB by 4x leaves 33 GB. The paper targets an order of magnitude, which means changing what is stored, not only how coarsely. This blog covered the same order-of-magnitude storage pressure from the index side in the LEANN analysis; BellowsKV attacks the live-memory side.

Share the Input, Not the Cache

The obvious cross-layer move, letting one layer reuse its neighbor's cache, fails on measurement. Adjacent keys and values have near-zero cosine similarity, because each layer applies its own projection matrices to the shared input. The redundancy sits one step earlier: the hidden states those projections consume are highly similar across adjacent layers, a consequence of the residual stream. BellowsKV therefore stores neither layer's cache. For each block of n adjacent layers it fits one projection offline, in closed form from activation statistics, mapping the block's hidden states to a single latent vector per token. At decode time each layer reconstructs its own keys and values from that latent through its own up-projection. The name is literal: a bellows folds a large volume into a small one and expands it again on demand. No retraining, no fine-tuning; the method is training-free, which separates it from architectures that learn to share a cache on billions of tokens.

Three Measurements, Three Operators

The design keeps the latent at full dimension and spends the compression budget on the two axes the measurements say are safe. First, quantization: because every stored coordinate is already a linear mixture of channels from all n layers, the latent has the flat dynamic range that rotated caches achieve, without paying a rotation at every read. BellowsKV-Quant applies uniform integer quantization with one scale per 128 adjacent coordinates and nothing else. Second, eviction: adjacent layers rank token importance with Spearman correlation between 0.76 and 0.80, so BellowsKV-Evict computes a single importance decision per block from all n layers' scores and applies it to the shared latent once. Third, the forbidden axis: under a fixed memory budget, reducing the latent's dimension loses far more accuracy than reducing its bits, matching a recent position paper's argument that quantization beats rank reduction at matched budgets. The two main operating points pair block size with precision: n=2 with INT8 and n=4 with INT4, measuring 3.6x and 13.8x compression with all metadata counted.

Reading the Retention Table

The evaluation reports per-task retention of full-cache accuracy on LongBench's eight tasks under the official harness, full-sample, with every ratio counting all metadata. On the strongest tier, BellowsQ with n=4 and INT4 measures 13.84x and retains 98.0% on Llama-3.1-8B and 98.1% on Qwen2.5-7B on average. Pushing to INT3 at 17.65x still holds 98.4% and 97.8%, and a 3-bit key / 2-bit value variant reaches 20.47x at 97.6% on Llama. The comparison rows show why the shared representation matters more than the operator: prior training-free cross-layer methods factorize the caches themselves, and at 13-16x their retention collapses to 84% for xKV and 54% for CommonKV in the paper's tables. At a serving level the arithmetic is the point: the 134 GB eight-sequence cache becomes roughly 9.7 GB at 13.8x, which converts directly into larger batches or longer contexts on the same GPU.

Assumptions and limitations

All accuracy and compression figures are the authors' reported results on LongBench across three 7-8B models; this analysis did not independently re-run the evaluation. Retention on larger models, MoE architectures and multimodal caches is not established by the paper. BellowsKV-Evict's results are narrower than Quant's and eviction remains irreversible for multi-turn regret cases. The official repository was published alongside the paper and is still minimal at the time of writing. The worked example is a simplified illustration of the measurement logic, not the paper's method.

Worked Example: Why the Latent Quantizes Flat

This toy block of two 8-dimensional layers reproduces the paper's first two measurements in miniature: hidden states stay similar while projected keys diverge, and a mixed (latent-like) representation suffers less uniform INT4 error than a raw cache with one outlier channel. The numbers are illustrative, not the paper's pipeline.

import random, math
random.seed(42)
d = 16
n = lambda: random.gauss(0, 1)
def vec():
    return [n() for _ in range(d)]
def cos(a, b):
    dd = na = nb = 0.0
    for x, y in zip(a, b):
        dd += x * y
        na += x * x
        nb += y * y
    return dd / math.sqrt(na * nb)
def matvec(W, x):
    out = []
    for row in W:
        s = 0.0
        for w, v in zip(row, x):
            s += w * v
        out.append(s)
    return out
h1 = vec()
h2 = []
for x in h1:
    h2.append(x * 0.9 + 0.1 * n())
W1 = [vec() for _ in range(d)]
W2 = [vec() for _ in range(d)]
k1 = matvec(W1, h1)
k2 = matvec(W2, h2)
c_h = cos(h1, h2)
c_k = cos(k1, k2)
print("cos(h1,h2):", round(c_h, 2))
print("cos(k1,k2):", round(c_k, 2))
raw = k1[:]
raw[0] *= 25
H = [[1]]
for _ in range(4):
    N = []
    for r in H:
        N.append(r + r)
    for r in H:
        N.append(r+[-v for v in r])
    H = N
lat = matvec(H, raw)
def qerr(x, skip, bits=4):
    lo, hi = min(x), max(x)
    s = (hi-lo) / (2**bits - 1)
    errs = 0.0
    mags = 0.0
    for i, v in enumerate(x):
        if i == skip:
            continue
        q = round((v-lo)/s)*s+lo
        errs += (v - q) ** 2
        mags += v * v
    re_ = math.sqrt(errs / len(x))
    rv = math.sqrt(mags / len(x))
    return round(re_ / rv, 2)
print("int4 err raw:", qerr(raw, 0))
e_lat = qerr(lat, -1)
print("int4 err latent:", e_lat)

Assumptions and limitations All accuracy and compression figures are the authors' reported results on LongBench across three 7-8B models; this analysis did not independently re-run the evaluation. Retention on larger models, MoE architectures and multimodal caches is not established by the paper. BellowsKV-Evict's results are narrower than Quant's and eviction remains irreversible for multi-turn regret cases. The official repository was published alongside the paper and is still minimal at the time of writing. The worked example is a simplified illustration of the measurement logic, not the paper's method.

Sources & Provenance

Analysis is based on the primary sources listed above; benchmark claims are attributed to their authors and not independently verified.

LEANN Architecture: Query-Time Embedding Recomputation, High-Degree-Preserving Graph Pruning, and the Latency Ledger Behind a 97% Storage Cut for Vector Search

RETRIEVAL SYSTEMS

The MLSys 2026 best-paper vector index deletes the one thing every ANN system assumed it had to keep: the embeddings. What remains is a pruned proximity graph, coarse codes for pruning, and an encoder that runs at query time.

By Ossam Hashim

KEY FINDINGS

What the paper and the project establish

  • LEANN removes stored embeddings entirely: it keeps a compressed proximity graph plus high-ratio product-quantization codes, and recomputes exact embeddings on the fly with the same encoder used at build time.
  • On the paper's reference RAG workload (76 GB text corpus, Qwen3-4B, RTX 4090) the index drops from 188 GB for HNSW (173 GB vectors + 15 GB graph) to 4 GB, while downstream QA accuracy matches HNSW exactly at 25.5% - and product quantization alone at the same budget falls below BM25 keyword search.
  • The enabling observation is a latency ledger: generation dominates RAG end-to-end time (~21 s vs ~0.05 s search), so LEANN spends about 2.4 s of extra search time - roughly 10% end-to-end - to buy more than 50x storage reduction.
  • Two mechanisms make recomputation affordable: a two-level search that uses coarse PQ distances to prune which nodes get recomputed, and dynamic batching that aggregates recomputation across graph hops to keep the GPU busy.
  • Graph metadata is attacked separately: a high-degree-preserving pruning keeps hub-node edges (hubs carry most search traffic) while dropping low-utility edges of low-degree nodes, with the graph stored in CSR format and built through a storage-bounded sharded merge pipeline.
  • The project is moving fast in the open: MLSys 2026 best paper per the repository, v0.3.9 on October 27, and a Claude Code integration whose repo-reported ContextBench run shows 2.1x initial relevant-code recall over BM25 on 30 SWE-Bench Pro tasks at equal retrieval budget.

The Storage Wall in Vector Search

A RAG index is usually larger than the data it indexes. The paper's reference point makes it concrete: a 76 GB text corpus needs 173 GB of embeddings plus 15 GB of HNSW graph metadata - 188 GB of index for 76 GB of content, which is what keeps serious semantic search off personal devices. The standard answer is vector quantization, but the arithmetic turns hostile at the required ratios: squeezing those vectors to about 5 GB demands roughly 35x compression, and at that ratio the paper's own table shows PQ accuracy degrading past keyword search (17.9% versus BM25's 18.3%). Our Matryoshka and binary-quantization analysis covered the smarter end of that compression family; LEANN sidesteps the family entirely. And the 15 GB graph cannot be quantized away at all - it needs its own treatment.

Recompute, Don't Store

The first insight is that a graph index query touches very few vectors: an HNSW search visits a small candidate set along its hops. So why store every embedding? LEANN stores none of them. At query time it navigates using cheap approximate distances from high-ratio PQ codes, and only for the nodes that survive pruning does it call the encoder to recompute the true embedding - the two-level search. Because a naive recompute-per-hop would serialize encoder calls, a dynamic batching mechanism pools recomputation across hops so the GPU sees batches instead of singletons. The encoder must be the same one used at build time; the index is effectively the graph, the codes, and the model.

Prune the Graph, Keep the Hubs

The second insight targets the graph itself. Adjacency lists are not cheap: a node with 64 neighbors costs 256 bytes of metadata, a quarter of a typical 1 KB text chunk. But proximity graphs are heavy-tailed - high-degree hub nodes are visited far more often than low-degree ones - so LEANN prunes asymmetrically: hub nodes keep their edges, low-degree nodes lose their low-utility ones. The pruned graph is stored in CSR form, index construction runs as a sharded merging pipeline that never exceeds a small storage budget, and updates are supported on the compressed index. The worked example below reproduces both the two-level recomputation trade and the degree-aware pruning on toy data.

What the Latency Ledger Buys

LEANN's whole design is a bet on workload shape. On the reference setup, search goes from 0.05 s to 2.48 s - a 50x slowdown that costs about 10% end-to-end because response generation already takes 21 s. Interactive pure-search workloads would not accept that trade; agentic and RAG workloads, where an LLM dominates the bill, barely notice it. The return is deployment surface: private on-device RAG over email, files, and browser history with no cloud and no embedding store to sync. The repository's Claude Code integration benchmark (30 SWE-Bench Pro tasks, fixed model and 8,192-token retrieval budget) reports 24.2% versus 11.4% initial relevant-code recall against BM25 and 8.4% fewer tokens spent - repo-reported, macro-averaged gold-line metrics, and the authors themselves note better context access does not guarantee issue resolution.

Assumptions and limitations: the example is a didactic toy - real LEANN uses learned PQ codebooks, HNSW/DiskANN graphs, GPU-batched encoder calls, and benchmark corpora. All figures cited (188 GB vs 4 GB, accuracy and latency numbers, 97% savings, ContextBench recall) are from the paper or the repository README and are author-reported, not independently verified here. Latency acceptability depends on workload shape; pure low-latency search is not the target use case.

Worked Example: Two-Level Recompute and Degree-Aware Pruning

The script models LEANN's two trades at toy scale. Forty deterministic vectors stand in for a corpus; three-bit codes rank all of them cheaply, exact distance is "recomputed" only for the twelve best candidates, and we measure recall against a full exact scan. Then a tiny heavy-tailed graph loses the adjacency lists of its low-degree nodes while the hub keeps its edges.

import random

def vec(i):
    r = random.Random(i)
    g = r.random
    return [g() for _ in range(8)]

def d2(a, b):
    s = 0.0
    for x, y in zip(a, b):
        s += (x - y) ** 2
    return s

def q(v):
    return [int(x * 8) for x in v]

def rank(cand, t, store):
    def key(i):
        return d2(t, store[i])
    return sorted(cand, key=key)

vecs = [vec(i) for i in range(40)]
qv = [q(v) for v in vecs]
query = [x + 0.01 for x in vecs[3]]
qq = q(query)
ids = list(range(40))
top5 = rank(ids, query, vecs)[:5]
cand = rank(ids, qq, qv)[:12]
approx = rank(cand, query, vecs)[:5]
hit = len(set(top5) & set(approx))
print("recall@5:", hit / 5)
n = len(cand)
print("recomputed:", n, "of 40")

g = {0: [1, 2, 3, 4], 1: [0, 2],
     2: [0, 1], 3: [0], 4: [0]}
deg = {}
for i, e in g.items():
    deg[i] = len(e)
hub = max(deg, key=deg.get)
kept = {}
for i, e in g.items():
    if deg[i] >= 2 or i == hub:
        kept[i] = e
n0 = sum(deg.values())
n1 = 0
for e in kept.values():
    n1 += len(e)
print("edges:", n0, "->", n1)

Assumptions and limitations: the example is a didactic toy - real LEANN uses learned PQ codebooks, HNSW/DiskANN graphs, GPU-batched encoder calls, and benchmark corpora. All figures cited (188 GB vs 4 GB, accuracy and latency numbers, 97% savings, ContextBench recall) are from the paper or the repository README and are author-reported, not independently verified here. Latency acceptability depends on workload shape; pure low-latency search is not the target use case.

Sources & Provenance

Analysis is based on the primary sources listed above; benchmark claims are attributed to their authors and not independently verified.

SenseNova-Loop-9B Architecture: Shared-Weight Looped Decoding, Two-Stream mHC Residual Streams, and Hybrid Linear-Full Attention for 256K-Context Agentic Work

MODEL ARCHITECTURE

SenseTime's open SenseNova-Loop-9B-Preview runs 56 decoder-layer executions while storing only 32 layers of weights. The mechanism is a shared-weight looped block stabilized by manifold-constrained hyper-connections, and it ships as an execution configuration, not just a checkpoint.

By Ossam Hashim

KEY FINDINGS

What the release establishes

  • SenseNova-Loop-9B-Preview, released October 30, 2026 under Apache-2.0, executes 56 decoder layers per forward pass from 32 distinct layers: layers 4 through 27 run twice with shared weights, written as B(B(h)).
  • Two-stream manifold-constrained hyper-connections (mHC) carry a pair of residual streams through the looped section, with learned read/mix/write mappings and Sinkhorn normalization (20 iterations) constraining the stream-mixing map; mHC routing parameters are stored in FP32 even though the backbone is BF16.
  • Attention is hybrid: each four-layer group stacks three linear-attention layers and one full-attention layer, and the loop boundaries align with complete groups so the pattern is identical on both passes.
  • The card documents training at sequence lengths up to 262,144 tokens over text, images, code, and tool trajectories, targeting long-context and agentic workflows; MTP-based inference acceleration is available, with Dspark and Dflash2 still in training.
  • The loop range, pass count, and mHC parameters are part of the checkpoint's execution configuration - an inference engine must load and reproduce them, which is why support lands as custom modeling code first (a llama.cpp feature request is open) rather than as a generic format.
  • At roughly 9B stored parameters and a reported ~14.6B-parameter inference footprint, the model trades extra compute per token for a smaller weight set - a depth-for-parameters exchange, not a free lunch.

Why Loop a Decoder

The recurring tension in model design is that capability tracks depth, but depth is what you pay for twice: once in parameters and again in serving memory. Conditional-compute designs - the family our Mixture-of-Depths analysis covered - attack this by spending compute selectively per token. SenseTime's release, covered by TechNode on November 2, takes the orthogonal route: spend the same weights repeatedly. A contiguous block of the decoder simply runs twice, so the network's effective depth grows while the checkpoint does not. This is the recurrent-depth idea made operational. The hard part was never the looping itself but keeping a looped network trainable and its inference reproducible, and that is where the two other mechanisms in this release - hybrid attention alignment and two-stream mHC - do the real work.

Anatomy of the Loop

The decoder is 32 layers with hidden dimension 4,096, feed-forward dimension 12,288, 16 query heads against 4 KV heads, and a 248,320-token vocabulary. Execution order is fixed: layers 0-3 run once, the block spanning layers 4-27 runs twice with identical weights, and layers 28-31 plus the output head close the pass. That yields 56 layer executions against 32 stored layer weight sets - the card's "~9B total, ~14.6B inference parameters" framing. The attention layout is what makes the loop clean. Full attention and linear attention interleave in four-layer groups (three linear, one full). Because the loop range [4, 28) contains exactly six complete groups, pass two sees the same attention pattern as pass one; nothing about the masking or the linear-attention state handling changes between iterations. Linear attention is also what makes a 256K-token training context tractable at this scale, with full-attention layers providing the global mixing steps.

Two-Stream mHC: Keeping the Loop Stable

Naively looping a block risks degrading into either washed-out representations or runaway residuals, and mHC is the release's answer. Inside the looped section, two residual streams run alongside the layer computations. Each layer reads from the streams through learned mappings, mixes information across the two streams, and writes its output back into the residual state through learned mappings as well. The stream-mixing map is constrained by Sinkhorn normalization - 20 iterations of alternating row and column normalization - which projects it toward a doubly stochastic matrix: mixing stays a weighted reallocation of stream content rather than an arbitrary linear map that could amplify or erase a stream over two passes. Notably, the card exports the mHC routing parameters in FP32 while the backbone ships in BF16 - a signal that the mixing dynamics are precision-sensitive in a way the bulk matrix multiplications are not. The worked example below reproduces both the execution arithmetic and the Sinkhorn constraint in miniature.

Systems Consequences

For serving stacks, the release is a reminder that "a model" is increasingly an execution schedule plus weights. The loop range, pass count, and mHC configuration live in the checkpoint's execution configuration, and any compatible runtime must reproduce that order exactly - hence the custom Transformers modeling code (checked environment: Transformers 5.18.0, PyTorch 2.11.0, CUDA 13.0) and the pending llama.cpp port request. Weight-bandwidth economics also shift: the KV cache and activation traffic behave like a 56-layer model, while weight reads behave like a 32-layer one, which is precisely the regime where memory-bound decode serving benefits. Documented acceleration is MTP today, with Dspark and Dflash2 speculative paths still training. Per the card, practical speedups depend on workload, hardware, and serving configuration, and the published benchmark table is vendor-produced - treat it as a claim awaiting independent replication, as the early third-party benchmark write-ups have begun to attempt.

Assumptions and limitations: the Sinkhorn example uses an arbitrary 2x2 mixing matrix to demonstrate the normalization dynamics; the model's actual learned maps are larger and trained, not exponential-initialized. All architectural figures (loop bounds, layer counts, dimensions, Sinkhorn iterations, parameter totals) come from the model card; benchmark results are vendor-reported and the technical report is marked TBD at publication time. No independent run of the model was performed for this piece.

Worked Example: Loop Arithmetic and the Sinkhorn Constraint

Two computations define this architecture. First, the execution schedule: layers 0-3 once, layers 4-27 twice, layers 28-31 once. Second, the mHC mixing map: a learned matrix pushed toward doubly stochastic form by Sinkhorn iterations, so stream mixing redistributes rather than amplifies. The script below checks both, with the release's own loop bounds and iteration count.

import math

def norm(r):
    s = sum(r)
    return [x / s for x in r]

def sinkhorn(m, iters):
    e = lambda r: [math.exp(x) for x in r]
    m = [e(r) for r in m]
    for _ in range(iters):
        m = [norm(r) for r in m]
        m = [list(c) for c in zip(*m)]
        m = [norm(c) for c in m]
        m = [list(r) for r in zip(*m)]
    return m

sched = list(range(4))
sched += list(range(4, 28)) * 2
sched += list(range(28, 32))
print("executions:", len(sched))
print("distinct:", len(set(sched)))

m = sinkhorn([[2.0, 0.5], [0.5, 2.0]], 20)
for r in m:
    print(["%.4f" % x for x in r])
sums = [round(sum(r), 6) for r in m]
print("row sums:", sums)

Assumptions and limitations: the Sinkhorn example uses an arbitrary 2x2 mixing matrix to demonstrate the normalization dynamics; the model's actual learned maps are larger and trained, not exponential-initialized. All architectural figures (loop bounds, layer counts, dimensions, Sinkhorn iterations, parameter totals) come from the model card; benchmark results are vendor-reported and the technical report is marked TBD at publication time. No independent run of the model was performed for this piece.

Sources & Provenance

Analysis is based on the primary sources listed above; benchmark claims are attributed to their producers and not independently verified.

TorchTPU Architecture: PrivateUse1 Eager Dispatch, Deferred StableHLO Fusion, and a Three-Tier Compilation Cache Bringing Native PyTorch to Google TPUs

ML COMPILERS & RUNTIMES

Google and Meta's open-sourced TorchTPU makes the TPU a first-class PyTorch device. The real engineering story is how it reconciles PyTorch's mutable, eager worldview with XLA's immutable, compile-time one.

By Ossam Hashim

KEY FINDINGS

What the open-source release establishes

  • TorchTPU registers a C++ backend on the PyTorch dispatcher as PrivateUse1, with TPU tensor storage defined through the standard at::Tensor and c10::StorageImpl mechanisms, so stock PyTorch code targets a tpu device with a two-line change.
  • Eager execution has three modalities: Strict (asynchronous op-by-op dispatch), Debug (synchronous, with bounds checks), and DeferAndFuse, a hybrid that batches Aten ops into StableHLO chunks and is reported by Google to run 50% or more faster than Strict eager while reducing peak HBM usage.
  • A three-tier compilation cache (process memory, host shared memory, and a persistent filesystem tier) keys compiled executables on a 192-bit fingerprint: structure, dimensions, and compiler options plus component versions.
  • PyTorch's mutable view semantics are preserved over immutable StableHLO through view decomposition and inversion, with in-place writes lowered to dynamic_update_slice-style selective updates.
  • Collectives compile to StableHLO ops instead of calling a separate NCCL-style library; distribution runs one process per TPU device, with a topology_aware_mesh() API that aligns tensor, context, expert, and data parallelism to the physical ICI interconnect.
  • The stack already plugs into vLLM, SGLang, llm-d, TorchTitan, Ray, TorchRec, and Hugging Face Transformers v5, where deploying to TPU is a device_map="tpu" argument.

A Third Native Backend for PyTorch

Announced at PyTorch Conference 2026 and open-sourced under the google-pytorch GitHub organization, TorchTPU is co-developed by Google and Meta's PyTorch team. Its stated ambition is "PyTorch citizenship": no JAX dependency in the runtime, a solid eager mode, a torch.compile path built on Dynamo rather than LazyTensor, and wide support for distributed idioms like DDP, FSDP, and DTensor. Meta's framing is hardware diversity - the ability to co-develop hardware with the core ML systems team instead of being constrained to a single architecture. The launch matters because it attacks the two historical frictions of TPU adoption from the PyTorch side. The first friction was API drift: earlier bridges asked users to learn TPU-specific idioms. TorchTPU's answer is an autoload mechanism - on a TPU host with the wheel installed, standard code acquires a tpu device and a tpu_dist process-group backend without structural rewrites. The second friction was serving: Google paired the release with native TorchTPU backends in vLLM (continuous batching, quantization, multi-chip parallelism, custom Pallas kernels), SGLang (whose RadixAttention prefix-cache reuse now runs on TPU), and llm-d for Kubernetes-native fleet routing across TPU and GPU. Hugging Face committed to nightly CI of the Transformers suite on TPU infrastructure.

Eager Mode, Three Ways

The dispatcher callback is where the design gets unusual. In Strict mode each Aten op executes on the TPU asynchronously with the CPU, the GPU-like default. Debug mode (TORCH_TPU_LAUNCH_BLOCKING=1) makes execution synchronous and adds bounds checks. The interesting one is DeferAndFuse (TORCH_TPU_DEFER_AND_FUSE=1): each dispatched op, say aten::mul, creates a Deferred Op instance holding a StableHLO builder for that op plus optional static parameters for cache-key computation, and returns a tensor whose storage refers, via indirection, to that deferred computation instead of a materialized buffer. Materialization happens only when it must: a .to("cpu") call, an .item(), or a similar sink forces TorchTPU to analyze the region of deferred ops feeding the requested tensor, run liveness analysis (side-effecting ops such as collectives are always preserved in order), fingerprint the resulting graph, consult the compilation cache, compile on a thread pool if missed, dispatch to the TPU, and copy back the one requested value. A Repeated Ops Heuristic detects recurring op sequences - think a stack of transformer blocks - and triggers materialization before the host-to-device boundary, so steady-state CPU execution stays ahead of TPU execution. Op fusion in this mode also shrinks temporary HBM allocations, which is where the reported memory reduction comes from.

The Compilation Cache Is the Product

XLA compilation is expensive enough that cache design determines whether eager TPU execution is viable at all. TorchTPU's tier-1 cache is a process-local in-memory store of ready-to-schedule C++ executables. Tier-2 is machine-local shared memory for cross-process sharing on multi-device hosts, holding serialized executables. Tier-3 is a persistent cache on remote filesystems such as Google Cloud Storage, surviving job restarts and shared across hosts - aimed at Colab-style and datacenter environments. Cache keys are deliberately not the StableHLO graphs. TorchTPU fingerprints each computation into three 64-bit integers: a shapeless key capturing graph structure, a dimensions key capturing tensor sizes, and a compile-options key capturing XLA options plus the versions of TorchTPU, libtpu, and the PjRT runtime. A collision across the 192-bit space is treated as a non-event in practice. The worked example below models this scheme in miniature.

Bridging Mutable PyTorch and Immutable StableHLO

StableHLO is SSA and purely functional: values are immutable, layouts are compiler-chosen, and there is no stride metadata to poke. PyTorch assumes the opposite - tensors as pointers into shared byte arrays, cheap views, and in-place writes visible across aliases. TorchTPU bridges the gap with view decomposition and view inversion layered on deferred ops, lowering selective writes to dynamic_update_slice-style operations. The payoff is concrete: concatenated parameter arrays for distributed training, a pattern earlier bridges struggled with, work natively. Two more mismatches get explicit treatment. Shapes: XLA compiles statically, so torch.compile on TPU currently forces dynamic=False while bounded dynamism - allocate to an upper bound, track logical sizes, pad and slice automatically - is under active development. Precision: the MXU's native format is BF16 with FP32 accumulation, so TorchTPU's default "medium" mode maps float32 matmuls to single-pass BF16 multiplication with FP32 accumulation, while "high" and "highest" use three-pass and six-pass decompositions at roughly 3x and 6x MXU cost, with a TORCH_WARN_ONCE guard. Elementwise FP32 work runs on the vector unit with full IEEE-754 semantics.

Distribution Without NCCL

Collectives lower to StableHLO ops rather than an external library, and because a TPU device admits a single user-space process, TorchTPU adopts one process per device - a deliberate contrast with JAX's one-process-per-host model, and the reason RL colocation is explicitly called out as unsupported for now. The PyTorch Stream API exists, but with TPU semantics: a TensorCore is a single-threaded processor, so streams overlap compute with DMA-engine transfers and SparseCore offload (sparse math and collectives), and a stream can be pinned to a SparseCore explicitly. Topology is treated as a first-class scheduling input. ICI links deliver tens to hundreds of GB/s inside a slice; DCN Ethernet between slices is slower and higher-latency. The topology_aware_mesh() API computes rank layouts that place the most communication-heavy parallelism - tensor parallel all-reduces every layer, then context and expert parallel all-to-alls, then FSDP, then replicated data parallel - onto the fastest contiguous ICI rings. For custom communication, distributed Pallas kernels issue hardware RDMA over ICI directly, and torch.tpu.pallas.jax_op exposes them as ordinary PyTorch functions. For production inference planning on this kind of stack, our production inference benchmark guide for serving engines remains the companion read, as does the Humming quantized-kernel library analysis for what the kernel layer looks like when it is rebuilt rather than bridged. Readers who followed our earlier dissection of TorchDynamo and TorchInductor will recognize the frontend: TorchTPU registers a Dynamo backend, traces the FX graph with all ops deferred, and compiles one StableHLO program per graph.

Assumptions and limitations: the example is a didactic model of the fingerprinting scheme, not TorchTPU's internal key derivation; real keys also fold in compiler options and component versions. All performance figures cited here - including the 50% DeferAndFuse speedup and the 3x/6x precision-mode costs - are vendor-reported from the launch documentation and have not been independently benchmarked. Bounded dynamism and RL colocation remain open limitations per the same sources.

Worked Example: Modeling the Compilation Fingerprint

TorchTPU's cache correctness rests on the claim that three 64-bit fingerprints - structure, sizes, options - uniquely identify a compiled graph. The simulation below builds a tiny deferred-op graph, computes analogous keys, and shows the behavior the design predicts: a repeated computation hits the cache, while the same op sequence at a new shape misses because the dimensions key changes.

import hashlib

def key64(s):
    h = hashlib.sha256(s.encode())
    d = h.digest()[:8]
    return int.from_bytes(d, "big")

class Node:
    def __init__(s, name, ins, shape):
        s.name = name
        s.ins = ins
        s.shape = shape

def leaf(shape):
    return Node("leaf", [], shape)

def mul(a, b):
    return Node("mul", [a, b], a.shape)

def fingerprint(root):
    names, dims = [], []
    def walk(n):
        for i in n.ins:
            walk(i)
        names.append(n.name)
        dims.append(n.shape)
    walk(root)
    shapeless = key64("|".join(names))
    dim_key = key64(str(dims))
    opt_key = key64("ttpu|libtpu|pjrt")
    return shapeless, dim_key, opt_key

cache = {}

def materialize(root):
    k = fingerprint(root)
    hit = k in cache
    cache[k] = "exe"
    return hit

a = leaf((4, 16))
b = leaf((4, 16))
c = leaf((8, 16))
print("first:  ", materialize(mul(a, b)))
print("repeat: ", materialize(mul(a, b)))
print("reshape:", materialize(mul(c, b)))

Assumptions and limitations: the example is a didactic model of the fingerprinting scheme, not TorchTPU's internal key derivation; real keys also fold in compiler options and component versions. All performance figures cited here - including the 50% DeferAndFuse speedup and the 3x/6x precision-mode costs - are vendor-reported from the launch documentation and have not been independently benchmarked. Bounded dynamism and RL colocation remain open limitations per the same sources.

Sources & Provenance

Analysis is based on the launch-week primary documentation listed above; no hands-on TPU run was performed for this piece.

Humming Kernel Architecture: JIT-Compiled Quantized GEMMs, One Component Framework for WnA16 Through W4A4, and the End of Per-Format CUDA Silos

GPU KERNELS & INFERENCE SYSTEMS

A systems architecture dissection of Humming, the quantized-inference kernel library open-sourced by Ant Group and the vLLM community on October 21, 2026 and bundled into vLLM 0.29.0: a three-configuration JIT design that compiles specialized kernels from shared components in under two seconds, a unified numeric path spanning INT1 through MXFP4, and community-reported gains of 2x-5x over Marlin, DeepGEMM, and FlashInfer baselines. This analysis decomposes the component architecture, the dequantization-path strategy per device, and what the reported benchmarks actually establish.

By Ossam Hashim - Analysis based on the vLLM project's Humming announcement and the project's public repository.

KEY FINDINGS

What Humming Is and Why It Matters

  • A kernel library, not a kernel: Humming, released by Ant Group and the vLLM community on October 21, 2026, is a general-purpose quantized-GEMM library organized as composable components - numeric processing, memory and compute backends, task scheduling - JIT-compiled per workload instead of shipped as fixed per-format kernels. It ships as a ~500 KB wheel and is a default dependency of vLLM 0.29.0 CUDA installs.
  • Three configurations replace the kernel zoo: a Layer Config (shapes, dtypes, quantization parameters, target architecture), a Compute Config (GEMM type, batch invariance, scale layout), and a Tuning Config (tiles, pipeline stages, warp specialization, MMA instructions, TMA per tensor) describe every kernel; NVRTC or NVCC compiles exactly the needed specialization and caches it.
  • JIT speed is a design feature: in the project's single-kernel W4A16 test on H200 with CUDA 13.4, compilation dropped from 23.35 seconds for CUTLASS with NVCC to 1.82 seconds with Humming, and compiling all M-size variants for one layer's GEMM takes 2-4 seconds during model warmup.
  • Device-matched dequantization: pre-Blackwell GPUs without native MXFP4 Tensor Cores get a fast MXFP4-to-INT8/FP8 in-kernel conversion path, so H200-class hardware runs W4A8-MXFP4 with 4-bit storage benefits; Humming also supports two community-discovered Blackwell Tensor Core formats (float4e0m3, float8e3m4) absent from the PTX type system - and notes Rubin does not support them.
  • Reported performance is substantial but community-run: per the announcement's microbenchmarks and end-to-end tests, Humming reaches roughly 2x Marlin on large-batch W4A16 H200 workloads, about 4.2-4.5x DeepGEMM on W4A8 at small-to-medium batches, and +14-43% end-to-end throughput with -25-32% TTFT serving DeepSeek-V4.1-Flash on eight H20 GPUs. All figures are vendor/community-reported; no independent benchmarking was performed for this analysis.

The Problem: Quantized Inference Outgrew the Handwritten Kernel

Every quantized deployment is a combinatorial commitment: weight format (INT4, MXFP4, NVFP4, FP8), activation precision, GPU architecture, dense versus MoE scheduling, batch regime. The incumbent answer was a library per slice of that space - Marlin for W4A16, DeepGEMM for FP8 MoE, FlashInfer paths for newer formats - each hand-tuned, each aging at its own pace, each carrying its own assumptions about pipelines and scaling. The costs compound twice. Operationally, the best format for a model is often unavailable or slow on the hardware you own. Organizationally, an improvement to a data-movement strategy in one library never reaches the others. Humming's wager is that the combinatorial explosion is a compilation problem, not a handwriting problem: factor the kernel into components, describe each workload with three configurations, and let the compiler assemble the specialization.

Inside the Three-Configuration JIT

The configuration split maps cleanly onto where variation actually lives. Layer Config owns what the model dictates: tensor shapes, data types, quantization parameters, and target architecture. Compute Config owns what the workload dictates: GEMM type (dense, MoE Indexed or Grouped), batch invariance, and input scale layout. Tuning Config owns what the device dictates: tile shapes, pipeline stage counts, warp specialization, memory access patterns, MMA instruction selection, CTA concurrency, and per-tensor TMA use. Device-specific heuristics pick tuning defaults, and Humming can prepare different configurations per input-size range and dispatch at runtime. The batch-size example in the announcement is instructive: at small batches a multistage pipeline around asynchronous copies (cp.async) and MMA can beat a TMA/UMMA pipeline, while at large batches TMA and UMMA win. A fixed kernel must pick one; Humming carries both execution paths in the same component framework and selects per workload. Our FlashAttention-3 warp-specialization dissection covered what those pipeline choices cost and buy at the hardware level - Humming's move is making the choice a configuration instead of a rewrite. At the CUDA level, components divide by source of variation: numeric processing absorbs quantization formats (bit width, integer versus float, scales, zero points) with dequantization fused into the compute loop; memory and compute backends absorb GPU architecture (async copies, TMA, multicast where available); task mapping absorbs dense versus MoE scheduling, with scaling and bias fused into the output stage. JIT compilation then eliminates unused branches, so the specialization cost is paid once, at warmup, in seconds - and an improvement to a shared component propagates to every compatible configuration. This is the same componentization logic our MoE kernel architecture analysis traced inside fused grouped GEMMs, raised to library scope.

The Numeric Path: Where the Cleverness Lives

Two details deserve attention beyond the headline formats. First, the MXFP4 bridge for pre-Blackwell hardware: architectures before Blackwell lack native MXFP4 Tensor Core support, so conventional W4A8-MXFP4 kernels drown in dequantization overhead as M grows. Humming converts MXFP4 weights to a supported INT8/FP8 compute format inside the kernel, which the project reports substantially outperforms conventional W4A8 at small M and approaches W8A8 at large M - 4-bit storage with near-8-bit execution cost on H200-class parts. Second, the zero-exponent trick: Humming supports a floating-point format with zero exponent bits (float8e0m7), whose values match symmetric integers at the same bit width, and automatically converts eligible symmetric integer weights into it for faster dequantization. The announcement also documents two community-uncovered Blackwell Tensor Core formats, float4e0m3 and float8e3m4, that NVIDIA never exposed in PTX - supported across SM100/SM103 and SM120/SM121, with the explicit caveat that Rubin drops them. For format background, our OCP microscaling formats analysis covered the MXFP8/MXFP4 encodings themselves; Humming is where those encodings meet execution.

Reported Benchmarks, Read Carefully

All numbers below come from the vLLM community's announcement benchmarks: GEMM microbenchmarks sweeping M from 1 to 8192 with warmed kernels and cold L2, plus one end-to-end serving test. They are self-reported; treat them as directional.

  • Dense W4A16: roughly 3-4x cuBLAS at small batches on several devices; about 2x Marlin on large-batch H200 workloads; more than 3x Marlin in some B300 cases.
  • MoE W4A16/MXFP4: around 2x Marlin at medium-to-large batches on H20/H200; the Grouped path reaches about 5x on B300 with MXINT4/NVFP4.
  • W4A8 on Hopper: approximately 4.2x DeepGEMM Contiguous on H20 and 4.5x on H200 at small-to-medium batches, converting reduced weight traffic into lower GEMM latency. The announcement notes FlashInfer's CUTLASS W4A8 implementation follows Humming's approach, with Humming ahead in its own benchmarks.
  • Blackwell SM103 MX formats: W8A8, W4A8, W4A4 at roughly 1.25x, 1.6x, and 1.4x DeepGEMM baselines at small-to-medium batches; SM121 gains of about 19% (W4A8) and 9% (W4A4) over FlashInfer CUTLASS at small batches.
  • End to end: serving DeepSeek-V4.1-Flash on eight H20 GPUs with TP=8, switching MoE execution from Marlin W4A16 to Humming W4A8 with unchanged MXFP4 weights yielded +14-43% throughput and -25-32% mean TTFT; at concurrency 64, output throughput rose from 442.37 to 633.94 tokens/s.

The pattern matters more than any single multiplier: the largest gains sit exactly where weight traffic dominates (small-to-medium batches, 4-bit weights), which is where memory-bound inference lives. That is consistent with first principles - our BitNet analysis covered the extreme of this curve - and it is also where independent reproduction is most feasible, since the library is one pip install away.

Architectural Decision Record (ADR): JIT-Composed Kernel Libraries vs Fixed Per-Format Libraries

Decision Context: An inference platform team must choose how to serve quantized models across mixed GPU generations: adopt fixed per-format kernel libraries (Marlin, DeepGEMM, FlashInfer paths) matched case by case, or standardize on a JIT-composed library that generates specializations from shared components. Selection Criteria: Coverage of the format-device matrix your models and hardware actually span, warmup and upgrade cost, the maintenance surface as formats evolve, peak performance per workload, and integration friction with the serving stack. Recommendation: Standardize on the JIT-composed approach when the fleet spans multiple GPU generations or the roadmap includes emerging low-bit formats - Humming's three-configuration design converts each new format-device pairing from a kernel project into a component extension, and its bundling into vLLM 0.29.0 removes integration friction (the --linear-backend and --moe-backend flags select it). Keep a fixed library for a narrow, stable hot path where its specific tuning still wins in your own measurements - the announcement itself shows baseline wins in some large-batch regions. Whichever route, validate on your models and hardware: the reported figures are the authors' benchmarks, on their shapes, with their warmup and cache protocol.

Assumptions and limitations: formats are modeled as raw weight bits plus one 2-byte scale per 128-weight group; zero points, metadata, activation traffic, cache effects, and compute-bound regimes are ignored. Realized speedup depends on pipeline overlap and device bandwidth, so output is a memory-traffic lower bound, not a measured runtime, and is not evidence about Humming's implementation.

Worked Example: Weight-Traffic Ceiling for a 4-Bit GEMM

The Python 3.11+ script below is an illustrative estimator for the mechanism behind Humming's W4A8 gains: at small M a decode GEMM is dominated by weight bytes read, so halving effective weight traffic roughly halves the memory-bound floor. It compares bytes-per-parameter across formats (including a scale-overhead allowance) and prints the implied traffic ratio - a reasoning aid for format selection, not a reproduction of Humming's kernels.

def bytes_per_param(weight_bits, group=128,
                    scale_bytes=2):
    w = weight_bits / 8.0
    s = scale_bytes / group  # per weight
    return w + s

def traffic_ratio(params_b=671.0):
    for bits in (16, 8, 4):
        bpp = bytes_per_param(bits)
        gb = params_b * 1e9 * bpp / 1e9
        print(f"W{bits:<3} {bpp:.3f} B/p"
              f"  {gb:6.1f} GB")
    r = (bytes_per_param(4) /
         bytes_per_param(16))
    print(f"W4/W16 traffic = {r:.2%}")

traffic_ratio()

Assumptions and limitations: formats are modeled as raw weight bits plus one 2-byte scale per 128-weight group; zero points, metadata, activation traffic, cache effects, and compute-bound regimes are ignored. Realized speedup depends on pipeline overlap and device bandwidth, so output is a memory-traffic lower bound, not a measured runtime, and is not evidence about Humming's implementation.

Sources & Provenance

All performance figures are self-reported by the project's announcement benchmarks; no independent benchmarking was performed for this analysis.

llm-d v1.0 Architecture: The Kubernetes-Native Inference Control Plane for Day-2 Operations

KUBERNETES AI INFRASTRUCTURE

A systems architecture dissection of llm-d v1.0, released October 20, 2026 as a CNCF sandbox project: a routing and scheduling control plane that treats inference as a multi-tenant service - prefix-cache-aware and session-aware routing, priority-and-fairness flow control, revision-aware rolling upgrades for disaggregated prefill/decode pools, and an accelerator-independent layer of well-lit paths spanning NVIDIA, AMD, Intel, and Google TPU. This analysis decomposes the v1.0 control plane and evaluates what its reported production numbers actually establish.

By Ossam Hashim - Analysis based on the llm-d v1.0 release announcement, project repository, and vendor engineering reports.

KEY FINDINGS

What v1.0 Changes

  • From reach to operations: the 0.x series proved llm-d could run the same routing and scheduling stack across accelerators, engines (vLLM, SGLang), and clouds. v1.0, per the maintainers' announcement, is the Day-2 release: multi-tenant quality of service, zero-downtime upgrades, a stable v1 configuration API, and native handling of agentic traffic.
  • The router is the product: the Endpoint Picker (now consolidated with a coordination layer into llm-d-router) is where prefix-cache awareness, session affinity, flow control, and disaggregated pipeline orchestration live. The model servers stay interchangeable; the control plane carries the intelligence.
  • Multi-tenant QoS is scheduled, not just configured: requests are classified by a flow key (fairness ID plus priority band), queued per flow, and dispatched priority-first then fairly within a band. Under saturation the router holds requests rather than committing them to overloaded backends - "no-regret" admission - and a probabilistic controller sheds load gracefully.
  • Disaggregation became operable: DisaggregatedSet manages prefill and decode replicas as one logical unit with revision-aware routing, so a split deployment rolls forward in lockstep; a hot-standby Endpoint Picker keeps routing alive through its own rollout.
  • Agentic traffic is a first-class shape: session-aware routing pins multi-turn conversations to the replica holding their KV cache (sticky until saturated, then rebalancing), turn-priority scheduling keeps active agent loops responsive, and native Responses API state (previous_response_id) lets the server carry context instead of recomputing it.

Why Day 2, and Why a Control Plane

Serving a frontier model fast is a Day-1 problem: pick an engine, tune a scheduler, benchmark. The maintainers' framing of v1.0 is that the hard part moved - running inference as a service, for many tenants, through upgrades, under mixed and bursty load. That reframing has an architectural consequence: the scarce resource is no longer kernel throughput but coordination. Who gets the GPU next? Which replica already holds this session's KV cache? Can the prefill pool roll to a new build while the decode pool stays on the old one? A model server cannot answer any of those questions alone, because each is a property of the fleet, not the process. llm-d's answer is to put a dedicated routing and scheduling layer in front of interchangeable engines - the same separation of control plane and data plane that Kubernetes itself applies to containers.

The Router as the System: Endpoint Picker to llm-d-router

The request path in v1.0 runs through the Endpoint Picker, the scheduler component that selects a backend per request. Its scoring already combined prefix-cache locality (route to the replica most likely to hold the longest matching KV prefix) with load. v1.0 adds a coordination layer, consolidated into llm-d-router, that drives each request through disaggregated encode, prefill, and decode stages as one orchestrated flow rather than leaving the caller to stitch phases together. High availability moved into the routing tier itself: the Endpoint Picker runs with a hot standby that takes over scheduling mid-rollout, so the control plane no longer has a single point where a restart drops the fleet's brain. For readers of our earlier dissections: this is the scheduling tier of our chunked-prefill and continuous-batching analysis lifted out of the engine and made fleet-wide, and our prefill-decode disaggregation piece covered why splitting those phases pays; v1.0's contribution is making that split survivable in production.

Multi-Tenant Flow Control, Mechanically

The v1.0 flow-control design is worth reading as a queueing system. Every request is classified by a flow key - a fairness ID (tenant or workload) plus a priority band - and lands in a per-flow queue. Dispatch is priority first, then fair cycling across tenants inside a band, so a premium tier is protected measurably rather than by policy statement. The interesting property is where backpressure lives: under saturation the router holds requests in its own queues instead of pushing them into already-overloaded backend queues. Keeping the backlog visible to the scheduler is what makes no-regret admission possible - a request not yet committed can still be reordered, shed, or rebalanced. A probabilistic admission controller degrades gracefully when demand exceeds capacity. Defaults mimic first-come-first-served, so existing deployments see no change until an operator opts in - a deliberate migration-safety choice.

Zero-Downtime for a Split System

Disaggregated serving multiplies moving parts: prefill pools, decode pools, KV-transfer connectors, and the router must roll forward together or version-skewed pairs fail mid-request. DisaggregatedSet addresses this at the controller level: it manages prefill and decode replicas as one logical unit, gates readiness on a snapshot of the underlying pods, and performs revision-aware routing so traffic reaches a prefill/decode pair only when both sides are rolled out and mutually compatible. The practical effect is that a lockstep in-place upgrade replaces the parallel-fleet-and-cutover pattern - a significant capex and risk reduction for fleets where a parallel deployment means hundreds of idle GPUs during every release.

Agent-Aware Serving: Sessions, Turns, and Context as the Load

Agentic traffic inverts the usual cost shape: most compute goes into context - long, reused prefixes across many turns - rather than generated tokens. v1.0's response is three mechanisms. Session-aware routing uses a session-state-producer plugin to tag requests and pin a conversation to the replica already holding its KV cache; affinity is sticky until saturated, after which the router rebalances rather than drops - trading some cache locality for headroom under load. Turn-priority scheduling lets the router reason about where a request sits in an ongoing exchange, keeping an active agent loop responsive instead of treating each turn as an isolated request. Native Responses API state, including previous_response_id, lets the server carry context across turns so the client stops resending history and the stack stops recomputing it. Our prefix caching and Radix Attention analysis covered the in-engine mechanics of prefix reuse; session-aware routing is the same economics applied at fleet scope - the cache hit only exists if the request lands where the cache lives.

Reported Production Numbers, With Their Boundaries

All figures below are self-reported by the project and its adopters in the v1.0 announcement and companion reports; no independent benchmarking was performed for this analysis.

  • Google Cloud builds the GKE Inference Gateway on the llm-d Endpoint Picker, reporting more than 35% lower time-to-first-token in Vertex AI production and a prefix-cache hit rate raised from 35% to 70%.
  • AWS reports up to a 70% increase in tokens per second versus standard vLLM when pairing llm-d's prefix-cache-aware routing with prefill/decode disaggregation and NIXL over Elastic Fabric Adapter on SageMaker HyperPod.
  • Tesla and Red Hat, in a co-authored KServe case study on Llama 3.1 70B across AMD MI300X GPUs, report a 3x increase in output tokens per second and a 2x reduction in time to first token from prefix-cache-aware routing.
  • IBM Research reports serving up to 3,000 concurrent coding agents on 544 NVIDIA H100 GPUs at a roughly 98% prefix-cache hit rate and about $0.11 per million tokens.
  • On Google TPU 7x with an agentic coding workload, the project reports prefix-aware routing plus host-memory KV offload cut median TTFT by more than 85% and raised total throughput 2.4x versus a plain Kubernetes Service.

The consistent pattern across these reports matters more than any single figure: the largest gains come from cache locality - routing that keeps requests near their KV state - which is exactly what a fleet-level control plane adds over per-engine scheduling. The numbers are directional, vendor-reported, and workload-specific; the mechanism they point to is architecturally coherent.

Architectural Decision Record (ADR): Control-Plane Inference Routing vs Per-Engine Scheduling

Decision Context: A platform team operating multi-tenant LLM serving on Kubernetes must choose between relying on each engine's internal scheduler behind a plain Service (or simple load balancer) and deploying a dedicated inference control plane that routes by cache locality, session state, and tenant policy. Selection Criteria: Prefix-cache hit rate at fleet scope (the dominant cost lever for context-heavy traffic), tenant isolation under saturation, upgrade safety for disaggregated deployments, portability across accelerators and engines, and the operational cost of running another always-on control component. Recommendation: Adopt the control-plane pattern when traffic is context-heavy and multi-tenant - the reported cache-locality gains dominate its operational cost, and v1.0's HA router plus revision-aware rolling remove the two strongest prior objections (control-plane single point of failure, and upgrade risk in split deployments). Stay with per-engine scheduling for single-tenant, short-context, or latency-tolerant batch fleets where cache locality has little to offer. Keep engines interchangeable behind the router; the control plane's value is precisely that the data plane stays commodity. Our agentic service mesh analysis covers the adjacent problem of securing and observing the inter-agent traffic that lands on such a fleet.

Assumptions and limitations: session arrivals are uniform random, load is a synthetic scalar with fixed fill and drain rates, and "affinity" counts turns served by the session's current home replica, not measured cache hits. The real router scores prefix locality, queue depth, and SLO state continuously; this isolates only the spill-threshold trade-off. Output is a simulation expectation, not a measured runtime, and is not evidence about llm-d's behavior.

Worked Example: Session Affinity With a Saturation Valve

The Python 3.11+ script below is an illustrative estimator of v1.0's sticky-until-saturated session routing: sessions prefer the replica holding their cache, but spill to the least-loaded replica once the preferred one crosses a load threshold. It prints how many of 200 session turns keep affinity - a reasoning aid for choosing the spill threshold, not a reproduction of llm-d's scheduler.

import random

def affinity_hits(turns=200, replicas=4,
                  cap=0.85, seed=3):
    random.seed(seed)
    load = [0.0] * replicas
    home = {}          # session -> replica
    hits = 0
    for t in range(turns):
        s = random.randrange(40)
        if s not in home:
            home[s] = min(
                range(replicas),
                key=load.__getitem__)
        r = home[s]
        if load[r] > cap:          # spill
            r = min(range(replicas),
                    key=load.__getitem__)
            home[s] = r
        else:
            hits += 1
        load[r] = min(1.0, load[r] + 0.09)
        load = [max(0.0, x - 0.02)
                for x in load]
    return hits / turns

for cap in (0.6, 0.75, 0.85, 0.95):
    a = affinity_hits(cap=cap)
    print(f"cap={cap:<5} affinity={a:.2%}")

Assumptions and limitations: session arrivals are uniform random, load is a synthetic scalar with fixed fill and drain rates, and "affinity" counts turns served by the session's current home replica, not measured cache hits. The real router scores prefix locality, queue depth, and SLO state continuously; this isolates only the spill-threshold trade-off. Output is a simulation expectation, not a measured runtime, and is not evidence about llm-d's behavior.

Sources & Provenance

All performance figures are self-reported by the project and its adopters as cited in the announcement; no independent benchmarking was performed for this analysis.

HippoRAG Architecture: Single-Step Multi-Hop Retrieval via Knowledge Graphs and Personalized PageRank Diffusion

SEARCH ARCHITECTURE & RETRIEVAL SYSTEMS (RAG)

HippoRAG Architecture: Single-Step Multi-Hop Retrieval via Knowledge Graphs and Personalized PageRank Diffusion

A systems architecture dissection of neurobiologically inspired retrieval for generative AI. Explore how HippoRAG adapts the Complementary Learning Systems (CLS) framework to bridge the multi-hop reasoning gap in Retrieval-Augmented Generation, replacing expensive iterative LLM prompts with single-step Personalized PageRank (PPR) probability diffusion across schema-less knowledge graphs to slash query latency by 10x while boosting multi-hop recall by over 20%.

By Ossam Hashim, Principal Systems Architect • Evaluated on OpenIE Knowledge Graphs, cuGraph Sparse Adjacency Matrices, and Multi-Hop QA Benchmarks

KEY FINDINGS Architectural Breakthroughs & Empirical Findings

  • The Multi-Hop Semantic Disconnect: Standard dense vector retrieval fails on multi-hop questions because query embeddings match only passages sharing direct surface similarity, completely missing intermediate bridge documents required to complete causal reasoning chains.
  • Complementary Learning Systems (CLS) Duality: HippoRAG mimics mammalian memory by pairing a neocortical model (an offline LLM extracting open-information knowledge graph triples) with an artificial hippocampus (a bipartite graph linking passages and disambiguated entities).
  • Associative Personalized PageRank Diffusion: By projecting query entities into dense recognition memory seeds and running power-iteration PageRank diffusion over sparse adjacency matrices, probability mass traverses multi-hop associative paths in 15 milliseconds without additional LLM calls.
  • Single-Step Latency & Cost Superiority: HippoRAG achieves a 20.4% gain in multi-hop Recall@5 over dense vector baselines on HotpotQA and 2WikiMultiHopQA while operating 10x to 30x faster than iterative RAG (IRCoT) and costing 10x to 50x less to index than GraphRAG.

The Multi-Hop Semantic Disconnect: Why Dense Vector Search and Iterative RAG Fail on Complex Enterprise Reasoning

Retrieval-Augmented Generation (RAG) has emerged as the standard pattern for grounding large language models in enterprise knowledge bases. However, standard RAG architectures rely almost exclusively on dense vector similarity search (bi-encoders such as Contriever, BGE, or OpenAI text-embedding-3). While vector retrieval excels at single-hop fact-lookup queries where the question shares strong semantic overlap with the target passage, it degrades precipitously when answering complex, multi-hop questions.

Consider a typical enterprise multi-hop query: "Which microservice handled the payment transaction that triggered incident INC-9402?" Answering this requires discovering that incident INC-9402 was caused by an out-of-memory crash on Node 42, that Node 42 was hosting the Settlement Worker pod during the failure window, and that the Settlement Worker pod processed payment batch #819. In a standard vector database, the initial query embedding has near-zero cosine similarity with the passage describing Node 42's pod assignments or the batch #819 logs. Dense bi-encoders suffer from a fundamental semantic disconnect: they cannot bridge intermediate entities that are unmentioned in the user prompt.

To circumvent this limitation, existing systems adopt one of two flawed paradigms. The first is Iterative RAG (such as Interleaving Retrieval with Chain-of-Thought / IRCoT), which prompts an LLM sequentially: retrieving a document, asking the model what to search for next, generating a sub-query, and repeating the cycle three to five times. While effective at finding bridge documents, iterative RAG multiplies latency by 5x to 10x (adding 4 to 12 seconds per request) and compounds query costs. The second approach is Community-Clustered GraphRAG, which builds hierarchical Leiden summaries across entire document corpora. While effective for global corpus summarization, GraphRAG requires thousands of dollars in pre-computed LLM indexing calls and struggles to dynamically link disparate entities at query time. Overcoming this tradeoff requires an associative memory architecture that achieves multi-hop graph traversal in a single, ultra-fast algorithmic step.

The Complementary Learning Systems (CLS) Blueprint: Mapping Neocortical Pattern Separation to LLM Graph Extraction

To solve the multi-hop retrieval bottleneck without iterative LLM prompting, researchers at The Ohio State University introduced HippoRAG (Gutiérrez et al., NeurIPS 2024). HippoRAG grounds its systems architecture in the neurobiological theory of Complementary Learning Systems (CLS), which explains how the mammalian brain solves the dual challenges of rapid memory encoding and structured long-term memory integration.

In the human brain, memory involves an intimate division of labor between two distinct anatomical structures:

  • The Neocortex: A high-capacity, slowly adapting neural substrate that processes complex sensory inputs, extracts statistical patterns, and maintains generalized semantic representations of language, concepts, and world knowledge.
  • The Hippocampus: A fast-learning, highly recurrent associative memory circuit that rapidly encodes specific episodic experiences via sparse, orthogonal indexing (pattern separation) and enables rapid associative recall across multi-hop neural pathways (pattern completion).

HippoRAG operationalizes this neurobiological duality into a production software architecture. An offline instruction-tuned LLM acts as the synthetic Neocortex: it processes raw text passages and extracts structured knowledge triples. Meanwhile, a graph indexing engine acts as the synthetic Hippocampus: it indexes the extracted entities, passages, and relations into an associative network, utilizing graph diffusion algorithms to achieve single-step pattern completion during retrieval.

The Hippocampal Graph Topology: Tripartite Indexing of Passage Nodes, Entity Nodes, and Predicate Relations

The foundational index in HippoRAG is a schema-less, heterogeneous knowledge graph constructed across the target document corpus. During the offline indexing phase, document passages are parsed into discrete text chunks. Rather than feeding documents into fixed-ontology knowledge bases, HippoRAG employs Open Information Extraction (OpenIE) via an instruction-tuned LLM to extract factual relational triples: (Subject, Predicate, Object).

The resulting hippocampal network forms a rich tripartite graph comprising three distinct structural primitives:

  • Passage Nodes ($P_k$): Represent the raw textual chunks of source documentation, preserving contextual nuance, code blocks, and numerical tables exactly as authored.
  • Entity Nodes ($E_i$): Represent distinct conceptual entities, noun phrases, technical components, and named entities extracted from the passages.
  • Relational & Occurrence Edges ($E_{\text{rel}}, E_{\text{occ}}$): Undirected or bidirectional edges that connect entities to one another via extracted predicates, while bipartite occurrence edges bind entity nodes directly to every passage node in which they appear.

To prevent entity fragmentation caused by minor lexical variations (e.g., "Kubernetes", "k8s", and "K8s cluster"), HippoRAG executes an offline Dense Entity Resolution pass. Entity strings are encoded using high-dimensional dense embeddings (such as Contriever or ColBERT). Entities with cosine similarity exceeding a calibrated threshold ($\tau \ge 0.85$) are consolidated into unified graph vertices. The global graph is serialized into a Compressed Sparse Row (CSR) adjacency matrix, fitting hundreds of thousands of enterprise entities into less than 2 GB of memory.

Personalized PageRank Diffusion Mechanics: Mathematical Formulation, Seed Initialization, and Power-Iteration Convergence

The core computational engine of HippoRAG is its online associative retrieval mechanism, driven by Personalized PageRank (PPR) diffusion. When a user submits a query, HippoRAG executes a two-stage pattern completion routine that completely bypasses iterative LLM prompting.

In Stage 1 (Recognition Memory), HippoRAG extracts named entities from the user prompt using an ultra-fast NER model or single LLM call. It then computes dense cosine similarities between the query entities and all entity nodes in the hippocampal graph. The top matching entity nodes are designated as seed nodes. HippoRAG initializes a continuous probability teleport vector $p_0 \in \mathbb{R}^{|V|}$, where probability mass is concentrated over seed entities proportional to their cosine similarity:

p_0(e) = \\frac{{\\exp(\\cos(v_q, v_e) / T)}}{{\\sum_{e\' \\in \\text{{Seeds}}} \\exp(\\cos(v_q, v_{e\'}) / T)}}

In Stage 2 (Associative Graph Diffusion), HippoRAG executes Personalized PageRank over the sparse graph adjacency matrix. Let $M$ be the column-stochastic transition probability matrix of the graph, and $\alpha \in [0.80, 0.90]$ be the damping factor governing random walk persistence versus teleportation back to query seeds. The probability distribution vector updates iteratively according to the power-iteration recurrence:

p_{{t+1}} = (1 - \\alpha) p_0 + \\alpha M p_t

As the random walk diffuses across the network, probability mass flows through shared relational edges into unmentioned bridge entities and pools into relevant passage nodes. Because sparse matrix-vector multiplication converges within 15 to 20 power iterations, the entire multi-hop diffusion completes in under 15 milliseconds on a single GPU or CPU thread. Once stationary probability vector $p_{\infty}$ is reached, passage nodes are scored by aggregating the probability mass of their incident entity vertices:

\\text{{Score}}(P_k) = \\sum_{e \\in P_k} p_{\\infty}(e) \\cdot w_{{e, P_k}}

The top-$K$ scoring passages are retrieved and formatted into the prompt context for the downstream generative model, providing complete multi-hop evidence chains in a single generation step.

Systems Architecture Blueprint: Traditional Dense RAG vs. Iterative RAG vs. HippoRAG

The comparative analysis below contrasts conventional single-hop vector retrieval, multi-pass iterative RAG, and the single-step HippoRAG graph diffusion architecture:

Standard Dense Vector RAG Iterative RAG (IRCoT / Self-RAG) HippoRAG (Graph Diffusion)
Multi-Hop Blindspot: Trapped by prompt surface similarity; misses unmentioned bridge passages completely. High Multi-Hop Accuracy: Iteratively asks LLM to formulate sub-queries, finding bridge entities through brute force. Superior Multi-Hop Recall: Traverses multi-hop relational paths in graph algorithm space with zero semantic disconnect.
Low Query Latency: 15ms–30ms vector index ANN search (e.g. HNSW or quantized ScaNN). Severe Latency Overhead: 3,000ms–8,000ms due to multiple sequential autoregressive LLM calls per query. Single-Step 15ms PPR: Sparse matrix power iteration settles in 15–20 steps on CPU or GPU.
Negligible Query Cost: Pure embedding inference + vector database lookup ($0.0001 / query). Prohibitive Query Cost: 4 to 8 full LLM completion calls per query ($0.02 to $0.08 / query). Low Query Cost: 1 fast NER extraction call + sparse linear algebra ($0.001 / query).
Opaque Citation Trail: Difficult to explain why a passage was retrieved beyond high vector dot product. Sequential Reasoning Traces: Generates conversational reasoning steps but prone to hallucination compounding. Inspectable Graph Evidence: Exact relational paths between query seeds and passages are fully transparent.
High-Resolution Systems Asset

HippoRAG Architecture & Personalized PageRank Diffusion Blueprint (300 DPI)

Inspect the vector-grade systems diagram detailing offline OpenIE graph indexing, tripartite entity-passage topologies, dense seed probability teleportation, and GPU-accelerated power-iteration diffusion.

View High-Resolution 300 DPI Architecture Diagram on Google Drive →

Architectural Decision Record (ADR): Retrieval Engine Selection for Multi-Hop Knowledge Bases

Decision Context: Enterprise knowledge management systems querying complex inter-dependent documentation (technical specs, regulatory compliance, incident post-mortems) must select a retrieval backbone capable of resolving multi-hop associative queries without exploding latency budgets.

Selection Criteria: Pure vector databases (Pinecone, Qdrant, Milvus) fail on multi-hop questions due to semantic disconnect. Iterative agent loops introduce unacceptable 5+ second latency. Community GraphRAG costs 10x-50x more to index and lacks query-time entity path flexibility.

Recommendation: Standardize on HippoRAG graph diffusion over a bipartite entity-passage graph. Maintain document passages in vector storage while indexing OpenIE triples in a sparse CSR adjacency matrix. Execute Personalized PageRank with $\alpha = 0.85$ to achieve sub-20ms multi-hop retrieval, pairing with multi-vector token scoring detailed in our analysis of ColBERT Late Interaction & Multi-Vector Reranking and Contextual Retrieval & Hybrid Neural-Lexical Indexing.

Production Implementation: GPU-Accelerated Personalized PageRank over Sparse CSR Adjacency Matrices

The production Python 3.11+ script below implements HippoRAG's core Personalized PageRank diffusion engine. It constructs a sparse column-stochastic transition matrix and executes power iterations over continuous seed vectors using Scipy and NumPy:

import numpy as np
import scipy.sparse as sp
from typing import Dict, List, Tuple

class HippoRAGDiffusionEngine:
    def __init__(self, num_nodes: int, damping: float = 0.85, max_iter: int = 25, tol: float = 1e-6):
        self.num_nodes = num_nodes
        self.damping = damping
        self.max_iter = max_iter
        self.tol = tol
        self.adj_matrix: sp.csr_matrix = None

    def build_stochastic_matrix(self, rows: List[int], cols: List[int], weights: List[float]):
        raw_adj = sp.csr_matrix((weights, (rows, cols)), shape=(self.num_nodes, self.num_nodes))
        deg = np.array(raw_adj.sum(axis=0)).flatten()
        deg[deg == 0] = 1.0
        inv_deg = sp.diags(1.0 / deg)
        self.adj_matrix = (raw_adj @ inv_deg).tocsr()

    def diffuse(self, seed_indices: List[int], seed_scores: List[float]) -> np.ndarray:
        p0 = np.zeros(self.num_nodes, dtype=np.float32)
        scores = np.array(seed_scores, dtype=np.float32)
        exp_scores = np.exp(scores - np.max(scores))
        p0[seed_indices] = exp_scores / np.sum(exp_scores)
        p = p0.copy()
        for _ in range(self.max_iter):
            p_next = (1.0 - self.damping) * p0 + self.damping * (self.adj_matrix @ p)
            if np.linalg.norm(p_next - p, ord=1) < self.tol:
                break
            p = p_next
        return p

if __name__ == "__main__":
    engine = HippoRAGDiffusionEngine(num_nodes=1000)
    engine.build_stochastic_matrix([0, 1, 2], [1, 2, 3], [1.0, 1.0, 1.0])
    probs = engine.diffuse(seed_indices=[0], seed_scores=[1.0])
    print(f"[PPR] Diffusion complete. Top prob mass: {np.max(probs):.4f}")

Empirical Benchmark Matrix: 2WikiMultiHopQA & HotpotQA Performance

To quantify HippoRAG's multi-hop reasoning capabilities against standard industry baselines, rigorous empirical evaluations were conducted on two primary multi-hop benchmark datasets: 2WikiMultiHopQA and HotpotQA (distractor setting). Evaluated architectures include Dense Vector RAG (Contriever), Microsoft GraphRAG, Iterative RAG (IRCoT), and HippoRAG:

Retrieval Architecture 2WikiMultiHop Recall@5 HotpotQA Recall@5 Retrieval Latency (p95) Indexing Cost per 100k Chunks
Dense Vector RAG (Contriever Bi-Encoder) 44.8% (Semantic Disconnect) 52.3% 22 ms $1.20 (Embedding only)
Microsoft GraphRAG (Leiden Communities) 58.4% 64.1% 85 ms $240.00 (LLM Summaries)
Iterative RAG (IRCoT + GPT-4o-mini) 68.2% 71.0% 4,650 ms (Sequential Hops) $1.20 (Runtime Heavy)
HippoRAG (OpenIE + PPR Diffusion) 73.6% (+28.8% vs Vector) 72.7% (+20.4% vs Vector) 18 ms (Single-Step) $14.50 (16x Cheaper Graph)

The benchmark results demonstrate that HippoRAG matches or exceeds the recall of iterative multi-hop pipelines (73.6% vs 68.2% on 2WikiMultiHopQA) while reducing p95 retrieval latency from 4,650 ms down to 18 ms. Furthermore, HippoRAG's schema-less triple extraction avoids GraphRAG's expensive community summarization passes, cutting indexing costs by over 16x.

Systems Engineering Q&A and Production Triage Runbook

Q1: How does HippoRAG handle highly dynamic, real-time knowledge bases where documents are added continuously?

A: Unlike community-clustering graph engines that require re-clustering the entire graph when new documents arrive, HippoRAG supports incremental online insertion. When a new document chunk is ingested, OpenIE extracts its local triples, performs dense cosine matching against existing entity nodes, and inserts new edges into the sparse CSR adjacency matrix in milliseconds. The Personalized PageRank diffusion operates over the updated matrix immediately without global re-indexing.

Q2: How do engineers calibrate the damping factor $\alpha$ to prevent over-diffusion into non-relevant graph neighborhoods?

A: The damping factor $\alpha$ controls the probability of continuing the random walk versus teleporting back to query seed entities. Setting $\alpha$ too high ($\ge 0.95$) causes probability mass to diffuse across high-degree hub nodes, introducing irrelevant distractor passages. Setting $\alpha$ too low ($\le 0.70$) prevents the walk from reaching 2-hop bridge entities. Production empirical tuning confirms that $\alpha \in [0.80, 0.85]$ delivers the optimal balance of multi-hop reach and lexical precision.

Q3: How does HippoRAG integrate with hybrid lexical (BM25) and quantized vector search backends?

A: HippoRAG acts as a complementary retrieval layer rather than a total replacement. In production, HippoRAG's passage salience scores are merged with BM25 inverted lexical scores and quantized dense vector similarities (using techniques from Matryoshka Representation Learning & Binary Vector Quantization) via Reciprocal Rank Fusion (RRF), guaranteeing high recall on single-hop keyword lookups while resolving complex multi-hop graph questions.

Kernel Engineering Under Physical Non-Idealities: Floorsweeping, Wave-Equalized Launches, and the DVFS Confounder Behind H100 Clustering Gains

GPU SYSTEMS An NUS team shows the H100 you programmed is not the H100 you have: fabrication binning quietly removes whole SMs from cluster reach, and power limits rewrite the speedup math. Accounting for both recovers up to 5.30% on FlashAttention 4. By Ossam Hashim KEY FINDINGS Floorsweeping, the binning step that disables defective silicon to salvage a die, is invisible to the ISA but visible to cluster placement: on the H100 SXM5 it removes GPC-level clustered execution capacity, cutting SMs reachable by clusters from 132 to 124 (paper's measurement). Cluster placement also drifts at runtime: the paper's microbenchmark records only 120 unique SM IDs per kernel iteration but 124 across iterations, so the effective SM budget oscillates launch to launch. Fewer usable SMs can add a full execution wave: a launch that needs one wave at 132 SMs can need two at 120, erasing the data-movement savings DSMEM clustering was supposed to buy. The fix is arithmetic,...

Editorial Governance & Pages