FLOP classes (165)
One record per operation class, each with a bound_by and an arithmetic-intensity derivation. This is the "why is my inference slow, and which hardware property would fix it" layer — it sits between a benchmark and a part.
Speculative decoding losslessness dispute: lossy verification rewrites the decoding distribution
memory boundThe theoretical guarantee of speculative decoding (Leviathan et al. 2023, Chen et al. 2023) is that modified rejection sampling preserves the target distribution EXACTLY - this is the 'lossless' property. EAGLE-2 claims to be a 'lossless acceleration algorithm' that 'ensures that…flops/sd-lossless-verification-dispute
contestedmedium · 0.70
Vision encoder forward pass cost (ViT/SigLIP/CLIP)
compute boundViT forward pass is compute-bound at typical resolutions. For ViT-L/16 at 224x224: 190.7 GFLOPs total, with attention FLOPs ~2*L^2*d_h*n_q = 2*256^2*1024*16 = 2.1 GFLOPs (L=256 patches) and MLP FLOPs ~2*L*d_h*4*d_h = 2*256*1024*4096 = 2.1 GFLOPs per layer, 24 layers = ~103 GFLOPs…flops/vision-encoder-forward-pass-cost-vit-siglip-clip
contestedhigh · 0.85
Vision encoder vs LLM prefill cost comparison
compute boundThe vision encoder forward pass and the LLM prefill are both compute-bound, but they scale differently. Vision encoder cost scales with image resolution (quadratically in patches for attention, linearly for MLP). LLM prefill cost scales with text length (linearly). At 224x224, th…flops/vision-encoder-vs-llm-prefill-cost-comparison
contestedhigh · 0.85
Architecture-conforming engines remove the decode GEMM cost model entirely: Groq, Cerebras, TPU
compute boundTHE MODEL DEFAULT IS WRONG FOR THESE ENGINES IN A WAY THAT IS NOT A CONSTANT ERROR. The standard decode cost model (fct-batch1-step-cost-formula) is T = max(2*P_t/BW, 2*B*P_t/PEAK): the memory-bound term 2*P_t/BW comes from STREAMING THE WEIGHTS once per step, and the compute-bou…flops/eng-profile-arch-conforming-no-decode-gemm
draftmedium · 0.66
Diffusion LM practical serving: denoising steps, per-step cost, and throughput competitiveness vs AR decoding
compute boundThe serving picture for diffusion LMs is fundamentally different from AR. An AR model does 1 forward pass per token: flops = 2*N_active, constant per token. A diffusion LM does S denoising steps, each a full forward pass over the sequence: flops per step = 2*N_active + attention …flops/diffusion-lm-practical-serving-picture
draftlow · 0.60
Engine choice changes prefill cost PER REQUEST: prefix caching and radix prefix reuse are not the same accounting
compute boundTHE PER-TOKEN ARITHMETIC IS UNCHANGED BY THE ENGINE - prefill attention flops are 2*n_q*d_h*L^2 and the projection GEMMs are 2*L*d^2 regardless of who schedules them. What the engine changes is the DENOMINATOR: how many of those L^2 terms are executed at all for a given REQUEST. …flops/eng-profile-prefix-cache-prefill-cost-per-request
draftmedium · 0.65
Engine choice changes the overhead ledger, not just the constant: vLLM vs TensorRT-LLM
launch_overhead boundINHERITS the ledger from fct-overhead-flop-ledger unchanged in form - the arithmetic intensities of RMSNorm, residual add, SwiGLU, RoPE and softmax are properties of the MATHEMATICS, identical on both engines. What differs is how many of those terms ever reach HBM. The ledger's u…flops/eng-profile-vllm-vs-trtllm-overhead-ledger
draftlow · 0.62
Engine choice moves the speculative-decoding break-even, because the engine sets which term dominates it
both boundTHE BREAK-EVEN FORMULA IS THE SAME ONE IN fct-spec-decoding-cost-model-break-even - A_be = T_spec/(k*T_base) memory-bound, A_be = 1 + P_d/(B*P_t) compute-bound - and it is a property of the MODEL and the HARDWARE, not of the engine. What the engine changes is WHERE a given (k, B,…flops/eng-profile-spec-decoding-engine-break-even-shift
draftlow · 0.60
Engines split the FLOP map differently for MoE and SSM: routing, expert-parallel dispatch, and the recurrent-scan mismatch
both boundFOR MoE, THE ENGINE DECIDES WHETHER THE ROUTING TERMS ARE A ROUNDING ERROR OR A DOMINANT CLASS, and the mechanism is specific. A dense model has no routing cost at all. An MoE model has top-k selection over E experts plus the all-to-all dispatch and combine that moves each token'…flops/eng-profile-moe-and-ssm-engine-split
draftmedium · 0.68
Hybrid attention-SSM KV cache sizing: how many SSM layers vs attention layers and what the KV-per-token becomes
memory boundIn a hybrid model with A attention layers and S SSM layers, the KV cache per token is A * num_kv_heads * head_dim * 2 * 2 bytes (only attention layers contribute). The SSM layers contribute a fixed per-sequence state of S * d_inner * N * 2 bytes, independent of context. The KV-pe…flops/hybrid-attn-ssm-kv-cache-sizing
draftmedium · 0.70
KV-tiering and disaggregation engines move cost from flops to bytes: the hit path is a bandwidth problem wearing a FLOP record's clothes
memory boundTHIS IS THE CLEAREST CASE IN THE REPO WHERE A FLOP CLASS IS ACTUALLY A BYTES CLASS, AND IT IS WORTH NAMING EXPLICITLY SO A READER DOESN'T PRICE IT WRONG. Prefill-decode disaggregation and KV tiering (lmcache, mooncake, nixl, and the vLLM/SGLang disaggregated paths) do not reduce …flops/eng-profile-kv-tiering-moves-flops-to-bytes
draftlow · 0.64
MoE at the inference level: routing overhead, load balancing, expert parallelism, and the memory-bandwidth story
memory boundMoE changes the memory-bandwidth story in a way that is the opposite of what the flop count suggests. A dense model with N params has flops = 2*N per token and bytes = 2*N (the weight stream), so AI = 1 flop/byte. An MoE with E experts, top-k, has flops = 2*N*(k/E) per token but …flops/moe-inference-serving-level
draftmedium · 0.70
SSM decode cost: O(1) per token vs O(n) for attention — the constant-state advantage and its limits
compute boundThe per-token decode cost of an SSM layer is O(1) in sequence length, while attention is O(n). For an SSM layer with state size N and inner dimension d_inner, each decode step does: (1) input projection: 2*d_inner*hidden flops, (2) state update: 2*d_inner*N flops (the B_t@C_t pro…flops/ssm-decode-cost-o1-vs-attention-on
draftmedium · 0.70
2:4 structured sparsity (mma.sp): halving bytes by dropping exactly half of every four
memory boundThe only mechanism in this repo that genuinely reduces work rather than reducing overhead: the tensor core skips the zero elements. From the PTX ISA, mma.sp has a dense-equivalent shape list and a sparse counterpart at doubled K - half-precision .m16n8k16 dense becomes .m16n8k32 …flops/structured-2-4-sparsity-mma-sp
verifiedmedium · 0.75
AITER (AI Tensor Engine for ROCm) kernel library and measured speedups
compute boundflops/amd-rocm-aiter-kernels
verifiedhigh · 0.90
Arc Pro B60/B70 llama.cpp GGUF Quantization Optimization
memory boundflops/xpu-arc-pro-llamacpp-quantization
verifiedhigh · 0.90
Attention kernel regime map: what wins at each context length and batch size
both boundThe optimal attention kernel depends on context length, batch size, and hardware. SHORT CONTEXT (L < 4K), SMALL BATCH (B < 8): FA-2 or FA-3 with standard tiling. The kernel is compute-bound; tiling efficiency is high. FA-2 achieves 50-73% peak on A100. FA-3 achieves 75% on H100. …flops/attn-regime-map-context-batch
verifiedhigh · 0.85
Attention over the KV cache in the decode step
memory boundStructurally ~1 flop/byte for MHA at EVERY context length and EVERY batch size - and CONFIRMED by KIVI's abstract, which states the two-part failure independently: as batch size and context length grow the KV cache 'becomes the new bottleneck in speed and memory usage', AND separ…flops/decode-attention
verifiedhigh · 0.90
Attention over the prompt (prefill / training-style causal attention)
compute boundPrefill attention AI grows linearly with sequence length, and with the GQA group ratio on top of that. Derivation for causal attention with n_q query heads, n_kv KV heads, head_dim d_h, sequence length L, bf16 KV: flops = 2*L^2*d_h*n_q (QK^T, halved by the causal mask) + 2*L^2*d_…flops/prefill-attention
verifiedhigh · 0.90
Audio input token count explodes relative to text output
memory boundflops/aud-flop-audio-token-explosion
verifiedhigh · 0.90
AWS Neuron: dtype and configuration tradeoffs for inference
both boundInferentia2 supports FP32, TF32, BF16, FP16, UINT8, and cFP8 (configurable FP8). Trainium2 adds FP8 at 1300 TFLOPS per chip. The auto-cast path takes FP32/FP16 models and casts down to BF16/FP8/cFP8 at compile time. Stochastic rounding is available to protect accuracy. The dtype …flops/alt-neuron-dtype-tradeoffs
verifiedhigh · 0.88
AWS Neuron: neuroncc compilation, NEFF format, and auto-casting
both boundThe Neuron Graph Compiler (neuronx-cc) transforms ML models from TensorFlow/MXNet/PyTorch/XLA HLO into NEFF (Neuron Executable File Format). It auto-casts FP32 operations to BF16 by default for optimal performance. Supports optimization levels 1-3 (compilation time vs runtime per…flops/alt-neuroncc-compilation
verifiedhigh · 0.90
AWS Neuron: NRT memory model, collectives, and multi-chip inference
interconnect boundThe NeuronX Runtime (NRT) loads NEFF files and executes them on NeuronCores. It handles device allocation, memory management, and collective communications. For multi-chip inference, NeuronLink provides chip-to-chip bandwidth (192 GB/s for Inf2, 1024 GB/s for Trn2). The runtime s…flops/alt-nrt-memory-model
verifiedhigh · 0.90
Blackwell-specific failure modes, numerical issues, and kernel availability gaps
compute boundSeveral Blackwell-specific failure modes and gaps are documented. (1) SM100 vs SM120 incompatibility: CUTLASS README explicitly warns that sm_100a-compiled kernels are NOT compatible with RTX 50 series (SM120) because SM120 has a different compute capability from datacenter SM100…flops/bw-blackwell-failure-modes-kernel-gaps
verifiedhigh · 0.85
Block-scaled FP8/FP4 tensor cores: MXFP and NVFP formats inside the MMA instruction
compute boundThis is the mechanism by which precision reduction becomes FLOPs rather than a dequantization tax. From the PTX ISA shape tables: warp-level mma has .e4m3/.e5m2 at .m16n8k32 (PTX 8.4) and .m16n8k16 (8.7); .e3m2/.e2m3/.e2m1 at .m16n8k32 (8.7); and BLOCK-SCALED forms where the mult…flops/block-scaled-tensor-core-fp4-fp8
verifiedhigh · 0.85
CDNA3 Matrix Core MFMA instruction precision support (native vs emulated)
compute boundflops/amd-cdna3-mfma-precision-support
verifiedhigh · 0.90
CDNA4 Infinity Fabric 20% faster links (38.4 Gbps, 76.8 GB/s per direction)
interconnect boundflops/amd-cdna4-infinity-fabric-change
verifiedhigh · 0.95
CDNA4 Matrix Core MFMA instruction precision support (MXFP8, MXFP6, MXFP4)
compute boundflops/amd-cdna4-mfma-precision-support
verifiedhigh · 0.90
Cerebras WSE-3: wafer-scale SRAM bandwidth advantage for inference
compute boundThe WSE-3 packs 44 GB of on-chip SRAM at 21 PB/s bandwidth (CS-3) or 43.2 PB/s (CS-4 Turbo). This is ~1000x the bandwidth of GPU HBM. Weights stay on-chip — there is no weight streaming from HBM. Tokens flow uninterrupted through model layers pipelined across wafers. The bandwidt…flops/alt-cerebras-wafer-sram
verifiedhigh · 0.85
Cerebras WSE-3: what does NOT work well — capacity, flexibility, and ecosystem
memory boundCerebras WSE-3 limitations: (1) capacity — 44 GB SRAM per wafer is small vs GPU HBM (80-192 GB), so models >44 GB require multi-wafer pipelining; (2) no dynamic batching — the wafer-scale architecture is optimized for single-stream low-latency decode, not high-throughput batched …flops/alt-cerebras-limitations
verifiedmedium · 0.80
Chunked prefill and hybrid batching: tail-latency elimination vs raw throughput and TTFT tax
both boundMonolithic prefill executes all prompt tokens L in a single large-M GEMM, achieving high compute bound arithmetic intensity (AI ~ L). However, at long context (16K-128K), monolithic prefill monopolizes GPU streaming multiprocessors for hundreds to thousands of milliseconds, causi…flops/chunked-prefill-ttft-throughput-tradeoff
verifiedhigh · 0.90
Chunked prefill and prefill/decode interleaving: deciding when to interrupt a prefill for a decode step
both boundNot an arithmetic mechanism; it is the knob that decides the arithmetic-intensity mix of a step. A full prefill step has AI ~ prompt_length (large-M GEMM, compute-bound); a decode step has AI = batch (memory-bound). Chunked prefill splits a long prompt into chunks of `max_num_bat…flops/chunked-prefill-interleaving-scheduling
verifiedmedium · 0.80
Claimed vs effective context length per model (RULER): the advertised number and the measured one, side by side
compute boundThis record is not an arithmetic-intensity claim; the cost side of long context lives in [[ctx-kv-cache-gib-at-claimed-context]] and [[ctx-128k-prefill-compute-cost]]. What it establishes is the input side: how much of the advertised context a model actually uses. RULER defines e…flops/ctx-claimed-vs-effective-context-table
verifiedhigh · 0.85
Consumer Blackwell (SM120) vs datacenter Blackwell (SM100a): one architecture name, two incompatible kernel targets
compute boundA pure compatibility fact with an immediate performance consequence, and the clearest example in this slice of a kernel-library support matrix changing the arithmetic. CUTLASS's README states it outright: 'The NVIDIA Blackwell SM100 architecture used in the datacenter products ha…flops/sm120-vs-sm100a-kernel-binary-incompatibility
verifiedhigh · 0.85
Cross-encoder reranking: N independent bidirectional forwards, compute-bound at batch 1
compute boundReranking N passages is N INDEPENDENT forward passes with no KV-cache reuse, so it is prefill-shaped and compute-bound even at batch 1 - the exact opposite of the generative decode regime. Arithmetic: for a bidirectional encoder of P parameters run over one sequence of t tokens, …flops/cross-encoder-reranking
verifiedhigh · 0.88
CUTLASS collective GEMM: the arch-tagged Tensor Core mainloop that most production GEMM kernels wrap
both boundA GEMM at batch B is memory-bound below the ridge point and compute-bound above it; a tile-based mainloop does not change that crossover, it changes how close a real kernel gets to it. What CUTLASS's collective mainloop actually reduces is shared-memory traffic per FLOP: a tile o…flops/cutlass-collective-gemm-tensorop-tiles
verifiedhigh · 0.85
Decode-phase GEMM: memory-bandwidth bound at batch size 1
memory boundAt batch size 1, the GEMM is memory-bandwidth bound: each weight is loaded from HBM and used for only one multiply-accumulate. The arithmetic intensity is ~1 flop/byte, far below the GPU's ridge point.flops/qhw-decode-gemm-memory-bound
verifiedhigh · 0.90
DeepEP expert-parallel dispatch/combine: an all-to-all that is itself a kernel and overlaps compute
interconnect boundNot FLOPs - bytes, and specifically bytes that cross NVLink or RDMA rather than HBM. Dispatch moves each routed token's activation to the rank owning its expert; combine moves expert outputs back. The kernel-level content is the overlap: num_sms bounds how many SMs the communicat…flops/deepep-expert-parallel-dispatch-combine
verifiedhigh · 0.85
DeepGEMM FP8/FP4 MQA scoring kernels for the lightning indexer
both boundThe indexer computes, per token i and each KV token j in its span, a compressed scalar logit: reconstruct kv_j by applying the per-token KV scale (kv[0][j,:] * kv[1][j]), take q[i] . kv_j, apply ReLU, multiply by a per-head weight, and sum over heads. The FLOPs are a small GEMM p…flops/deepgemm-mqa-lightning-indexer-scoring
verifiedmedium · 0.80
DeepGEMM grouped/masked GEMM and Mega MoE: shaping MoE expert GEMMs for decode and for NVLink overlap
both boundGrouped GEMM changes the tiling, not the math: instead of B separate GEMMs each with a token-thin M, the tokens routed to one expert are concatenated so the expert's GEMM has a usable M dimension, which is what keeps tensor cores busy. DeepGEMM groups ONLY the M axis and requires…flops/deepgemm-grouped-masked-moe-gemm
verifiedhigh · 0.85
Dense weight GEMMs in the decode step (the per-token matmul tax)
memory boundAI = B flop/byte, exactly, at batch B in bf16 - this is the definitional result and it explains most decode benchmarks. Derivation: every weight in the model is read from HBM exactly once per step and reused across all B rows, so bytes = 2*N_params (2 bytes per weight) while flop…flops/decode-gemm
verifiedhigh · 0.92
Diffusion / masked-diffusion LMs: N denoising passes over the whole sequence instead of one token per forward pass
compute boundThis inverts the decode cost model, so derive it explicitly. An AR model does 1 forward pass per generated token: flops = 2*N_active, and the per-token cost is flat because each pass sees only one new query position. A masked diffusion LM instead starts from a fully masked sequen…flops/diffusion-lm-denoising-steps
verifiedmedium · 0.80
Documented ROCm failure modes: Vulkan PPL explosion, HIP gfx1151 corrupt output, HIP ubatch wrong logits
compute boundflops/amd-rocm-failure-modes
verifiedhigh · 0.90
Draft model memory cost: competition with KV cache for HBM budget
memory boundThe draft model's weights occupy HBM alongside the target model weights and the KV cache. For a target model of size N_target and draft model of size N_d, the total weight footprint is N_target + N_d + KV_cache. The KV cache grows as batch_size * context_length * num_layers * num…flops/sd-draft-model-memory-cost
verifiedmedium · 0.80
Draft-model/speculative decoding support on Blackwell
memory boundSpeculative decoding on Blackwell follows the same bandwidth-trade principle as on Hopper: the draft pass reads the draft model's weights for k tokens at once, giving AI ~k flop/byte, and is memory-bound exactly like decode_gemm at effective batch k. Blackwell's advantages for sp…flops/bw-speculative-decoding-blackwell
verifiedmedium · 0.75
Dynamic resolution and any-resolution vision token cost
compute boundDynamic resolution approaches (NaViT, Pixtral, LLaVA-NeXT/AnyRes) allow the vision encoder to process images at their native resolution and aspect ratio, rather than resizing to a fixed grid. The token count is determined by the image dimensions: tokens = (H/patch_size) * (W/patc…flops/dynamic-resolution-and-any-resolution-vision-token-cost
verifiedhigh · 0.85
Embedding / batch encoding: BERT-style bi-encoder over documents, compute-bound at batch 1
compute boundBatch encoding a corpus for retrieval is a bi-encoder forward per document with no KV cache, and like cross-encoder reranking it is compute-bound well before batch 1 - but unlike reranking its corpus cost is LINEAR, because each document is encoded once and its vector cached. The…flops/embedding-batch-encoding
verifiedhigh · 0.88
Embedding lookup and output projection over the vocabulary
memory boundInput embedding is a pure gather: ~0 flops against hidden_width*2 bytes read per token, so AI is effectively zero and it is strictly a memory-access pattern problem. The output side (LM head, untied) is a real GEMM: flops = 2*hidden_width*V against bytes = 2*hidden_width*V + 2*V,…flops/embedding
verifiedmedium · 0.70
Encoder-decoder vs decoder-only multimodal serving
both boundEncoder-decoder architectures (Whisper, Flamingo) process the input modality in a separate encoder, then decode the output autoregressively. The encoder runs once per input, the decoder runs once per output token. Decoder-only with visual tokens (LLaVA, Pixtral, Qwen2-VL) process…flops/encoder-decoder-vs-decoder-only-multimodal-serving
verifiedhigh · 0.85
Encoder-serving engines invert the roofline: at batch 1 an embedding or rerank workload wants FLOPS, the inverse of LLM decode on the same GPU
compute boundTHE INVERSION IS EXACT AND IT IS THE MOST COMMON MISMATCH IN THIS REPO'S DOMAIN. An LLM decode step streams each weight once per generated token, so arithmetic intensity is ~1 flop per byte at batch 1 and HBM bandwidth alone decides tok/s (the qhw-decode-gemm-memory-bound regime)…flops/eng-profile-encoder-serving-inverts-the-roofline
verifiedmedium · 0.75
Expert parallelism for MoE: when EP beats TP and the dispatch/combine overhead
interconnect boundEP shards experts across GPUs rather than sharding each expert's weights. For a MoE model with E experts, EP assigns E/N experts to each GPU. The arithmetic intensity of each expert GEMM is the same as a dense GEMM at the same batch size, but the batch is split across experts by …flops/sched-expert-parallelism
verifiedmedium · 0.80
FlashAttention 1-4 generation comparison: what each changed mechanically and measured performance
both boundEach FA generation attacks a different bottleneck. FA-1 (2022): IO-aware tiling eliminates S x S HBM round trip, 2-4x wallclock vs PyTorch. FA-2 (2023): work partitioning fix, ~2x vs FA-1, 50-73% of peak FLOPs/s on A100 (FA-1 was 25-40%). FA-3 (2024): Hopper warp specialization +…flops/attn-flashattention-generation-comparison
verifiedhigh · 0.90
FlashAttention IO-aware tiling (FA1): eliminating the S x S HBM round trip
compute boundDoes not raise AI arithmetically - the FLOP count of causal attention is unchanged. It removes bytes: the naive path writes the L x L score matrix to HBM and reads it back, adding Theta(L^2) HBM traffic on top of the Theta(L*d) KV traffic, which drops the effective AI from Theta(…flops/flashattention-io-aware-tiling
verifiedhigh · 0.90
FlashAttention-2 work partitioning: fixing warp-level underutilization
compute boundNo change to bytes or FLOPs relative to FA-1. FA-2 is a pure scheduling fix: it removes idle tensor-core cycles caused by poor thread-block/warp partitioning and redundant shared-memory traffic, i.e. it converts achieved FLOPs/s toward the FLOPs the tiles already imply. The metri…flops/flashattention-2-work-partitioning
verifiedhigh · 0.85
FlashAttention-3 (Hopper): warp specialization, TMA, and FP8 to overlap matmul with softmax
both boundThree different reductions, worth keeping separate. (1) Bytes: TMA (cp.async.bulk.tensor) does address generation and out-of-bounds predication in hardware, freeing registers, which lets the tile be larger - larger tiles amortize each HBM->SMEM byte over more FLOPs. (2) Occupancy…flops/flashattention-3-hopper-warp-specialization
verifiedhigh · 0.85
FlashAttention-3 vs FA-4: register pressure vs shared-memory bandwidth bottleneck distinction
both boundFA-3 and FA-4 attack different bottlenecks. FA-3 (Hopper): the bottleneck is SFU/matmul overlap. On H100 SXM5, tensor cores do 989 TFLOPS FP16 but only 3.9 TFLOPS of special functions (16 ops/SM/clock * 132 SM * 1.83 GHz). At head_dim 128, there are ~512x more matmul than exponen…flops/attn-fa3-vs-fa4-distinction
verifiedhigh · 0.90
FlashAttention-4 (Blackwell): tensor memory, 2-CTA MMA, and emulated exponential
both boundFA-4 is the first attention kernel whose bottleneck moves off tensor cores. Feed-and-speeds per B200 SM at M=N=D=128: tensor cores 8192 ops/cycle, exponential unit 16 ops/cycle, shared memory 128 bytes/cycle. Forward per tile: 1024 tensor-core cycles vs 1024 exp cycles vs 768 SME…flops/flashattention-4-blackwell-tmem
verifiedhigh · 0.85
FlashDecoding / split-KV attention: turning a single-token decode into a bandwidth-saturating GEMM
memory boundThe problem it solves is parallelism, not arithmetic. At batch 1 the decode step's attention has L query rows total, so a non-split kernel launches far fewer CTAs than the GPU has SMs and runs at a small fraction of memory bandwidth while the KV cache streams from HBM. Split-KV d…flops/flashdecoding-split-kv-decode
verifiedmedium · 0.80
FlashInfer: a kernel-selection layer over FlashAttention, cuDNN, CUTLASS and TensorRT-LLM
memory boundFlashInfer is not itself an arithmetic mechanism; it is the layer where backend choice decides which arithmetic is done. Recorded as custom_kernel because the engine-level effect is a discrete change in the kernels executing the same math. Its kernels change the bytes/FLOPs trade…flops/flashinfer-backend-selection
verifiedmedium · 0.80
FlashMLA decode: FP8/FP4 KV cache with scales inlined per token (DeepSeek MLA)
memory boundThe mechanism is a change in bytes moved, achieved by changing the KV cache element size, not by changing the attention math. MLA decode reads one KV head per token instead of a full K and V pair per head, so the cache is already much smaller than a GQA cache; FlashMLA pushes tha…flops/flashmla-decode-fp8-fp4-kv
verifiedhigh · 0.85
FlashMLA fused norm-RoPE-attention-RoPE-cast kernel: absorbing per-token scalar work into the attention kernel
both boundThis kernel changes nothing about the attention arithmetic; it removes kernels and bytes at the boundaries. It fuses Q-norm (V4 only, not V4.1), Q-RoPE, core attention, the conjugate O-RoPE and the cast of the output to FP8 into a single kernel. Each of those steps on its own is …flops/flashmla-fused-norm-rope-attn-rope-cast
verifiedmedium · 0.80
FP8 fine-grained scaling on Hopper vs Blackwell: per-tensor vs per-block scales and accuracy implications
compute boundFP8 (E4M3/E5M2) was introduced on Hopper with per-tensor scaling: a single FP32 scale factor for the entire tensor, applied before the FP8 GEMM and after in the epilogue. This coarse scaling is problematic for tensors with high dynamic range or outliers, as a single scale cannot …flops/bw-fp8-fine-grained-scaling-hopper-vs-blackwell
verifiedmedium · 0.80
FP8 KV cache: 2x concurrency at 128K, and why it fails differently at long context than at short
memory boundKV-cache dtype is the cleanest 2x in the whole stack: bytes per cached token halve from bf16's 2 bytes/element to fp8_e4m3's 1, so decode attention - which the repo already derives as memory-bound at AI = g (the GQA group ratio) flop/byte, independent of both batch and context le…flops/ctx-fp8-kv-cache-long-context-accuracy
verifiedmedium · 0.82
Gaudi Fusion Pass: Operator Fusion on SynapseAI
memory boundflops/xpu-gaudi-fusion-pass
verifiedmedium · 0.80
Gaudi HPU Graphs: Static Shape Capture and Replay
launch_overhead boundflops/xpu-gaudi-hpu-graphs
verifiedhigh · 0.85
Gaudi routed-expert gather: beat the dense MoE pass only while few experts are live
memory boundflops/xpu-gaudi-moe-routed-expert-gather
verifiedmedium · 0.72
Gaudi Static Shape Requirement for Inference Optimization
launch_overhead boundflops/xpu-gaudi-static-shapes
verifiedhigh · 0.85
Grammar-mask evaluation per token in constrained / structured decoding
memory boundGrammar-mask evaluation is a per-token cost incurred AFTER the logits exist, and it scales with vocabulary size rather than with model size. A naive CFG-constrained decoder, as the XGrammar paper's abstract describes it, 'requires going through several stack states over all token…flops/constrained-decoding-grammar-mask
verifiedhigh · 0.86
Groq LPU: compiler and kernel constraints from static scheduling
launch_overhead boundThe Groq compiler is model-independent — it does not require model-specific kernel coding (unlike GPUs). It accepts workloads from multiple frameworks and maps/schedules them across one or more LPUs. The compiler outputs a program encompassing all data movement information. Const…flops/alt-groq-compiler-constraints
verifiedhigh · 0.88
Groq LPU: deterministic single-core-per-token architecture and SRAM-to-DRAM ratio
memory boundThe Groq LPU uses a deterministic assembly-line architecture where each token is processed by a single core (SIMD function unit) with no dynamic scheduling, no cache misses, and no variance. All data movement is statically scheduled at compile time. The SRAM-to-DRAM ratio is effe…flops/alt-groq-deterministic-architecture
verifiedhigh · 0.90
HIP graph capture and replay for ROCm kernel launch overhead reduction
launch_overhead boundflops/amd-rocm-hip-graph-capture
verifiedmedium · 0.80
Host-side sampling loop: full-logits D2H transfer plus the CPU sampler chain
launch_overhead boundThis is the same V-wide tensor the sampling record describes, but measured in HOST time rather than device bytes, and that changes the answer by four orders of magnitude. On the device the logit read is negligible - 4*128256 = 501.0 KiB per token at fp32 against a Llama-3-8B batc…flops/host-side-sampling-loop
verifiedhigh · 0.85
Image tiling and visual token KV cache cost
memory boundImage tiling multiplies visual token count by the number of tiles. For a 4-tile image (2x2 grid), the LLM sees 4x the visual tokens. KV cache cost scales linearly with token count: KV_bytes = num_layers * num_kv_heads * head_dim * 2 * bytes_per_element. For a 7B model with 28 lay…flops/image-tiling-and-visual-token-kv-cache-cost
verifiedhigh · 0.85
Iteration-level scheduling: the CPU overhead that CUDA graphs cannot remove
launch_overhead boundNot a flop class in isolation - it is the class that decides which flop classes a step contains. Iteration-level scheduling (Orca, OSDI 2022) admits a new request as soon as any sequence in the running batch finishes, so the batch never sits at one sequence and the decode GEMM's …flops/sched-iteration-level-scheduling
verifiedhigh · 0.85
Iteration-level scheduling: the overlap of prefill and decode in one batch
both boundNot a flop class in isolation - it is the class that decides which flop classes a step contains. Orca's iteration-level scheduling admits a new request as soon as any sequence in the running batch finishes, so the batch never sits at one sequence and the decode GEMM's AI (= B) ac…flops/continuous-batching-scheduling
verifiedmedium · 0.80
KV cache at the claimed context length, and how many 128K-context users one H100 or H200 can actually serve
memory boundKV bytes per token is exact arithmetic with no measurement in it: B_per_token = n_attn_layers * n_kv_heads * head_dim * 2 (K and V) * bytes_per_element, with bytes_per_element = 2 for bf16, 1 for fp8_e4m3 or int8, 0.5 for fp4. Two cases need care. MLA does not use that formula: D…flops/ctx-kv-cache-gib-at-claimed-context
verifiedmedium · 0.80
KV cache movement: block copy, prefix sharing, prefill-to-decode handoff, and offload
interconnect boundPure data movement, AI = 0 - no flops at all, so every definition of this class is a bandwidth or interconnect problem. Volume per token is 2 (K and V) x num_kv_heads x head_dim x layers x kv_dtype_bytes; for a 32-layer, hidden-4096, GQA-8-head, bf16 model that is on the order of…flops/kv-transfer
verifiedhigh · 0.85
KV cache quantization at 128K: FP8 scaling granularity, attention drift, and two-level accumulation
memory boundKV cache quantization directly halves memory bandwidth consumption during decode attention. For a model in BF16, each KV token requires 2 * num_layers * num_kv_heads * head_dim * 2 bytes. For Llama-3.1-70B (80 layers, 8 KV heads, head_dim 128), BF16 KV cache consumes 327,680 byte…flops/kv-cache-quantization-fp8-scale-granularity
verifiedhigh · 0.90
KV cache-aware scheduling: LMCache/offloading to CPU and SSD, and the measured hit-rate effect on TTFT
memory boundNot an arithmetic mechanism at all, which is why it matters as a kernel-level decision input. A KV tier between the GPU and host/remote storage changes how many bytes the decode attention kernel pulls from HBM, but it does not change the attention FLOPs at all. The strongest vers…flops/sched-kv-cache-aware-scheduling
verifiedmedium · 0.82
KV-cache offload to a CPU/DRAM tier: bytes over PCIe, why it is bandwidth-bound, and when it beats recompute
memory boundAI = 0 in the strict sense - offload is a DMA copy with no arithmetic - so cost = KV bytes / PCIe bandwidth, and the useful question is never 'how fast is the copy' but 'is the copy cheaper than the recompute it replaces'. That comparison is a flop-versus-byte trade between two d…flops/kv-cache-offload-host-dram-tier
verifiedhigh · 0.85
KV-cache tiering (LMCache): deciding, outside the kernel, how much KV the attention kernel must read from HBM
memory boundNot an arithmetic mechanism at all, which is why it matters as a kernel-level decision input. A KV tier between the GPU and host/remote storage changes how many bytes the decode attention kernel pulls from HBM, but it does not change the attention FLOPs at all. The strongest vers…flops/lmcache-kv-tier-outside-attention
verifiedmedium · 0.70
llama.cpp MMQ: INT8 activation quantization with per-arch tile configs, no tensor cores
memory boundA different answer to the same problem, and the reason llama.cpp's quantized numbers are not comparable to Marlin's or Machete's. MMQ quantizes the ACTIVATIONS to q8_1 blocks of 32 elements and performs the dot product in INT8 accumulating to INT32, using DP4A where available. So…flops/llamacpp-mmq-int8-dp4a-gemm
verifiedmedium · 0.80
Logit processing and token sampling
memory boundAlso memory-bound and batch-invariant in the worst way. Reading the final logits for a vocab V from HBM costs 2*V bytes per token at bf16 (a single fp32 row is 4*V), and the softmax/top-k/renormalize arithmetic is a few multiples of V - so AI is on the order of 1-2 flop/byte and,…flops/sampling
verifiedmedium · 0.80
Logit processors in the decode loop: penalties, min_p, bias, min_tokens, and the host work they force
launch_overhead boundLogit processors are not one cost, they are two, and conflating them is the whole reason this class has been invisible. (a) DEVICE half: each active processor is one or more full-width passes over the [B, V] logits tensor. vLLM's shipped MinPLogitsProcessor.apply does softmax ove…flops/logits-processors-cpu-pipeline
verifiedmedium · 0.82
Long-context RAG ingest: a 32k document prefilled into a KV cache (attention-dominated, plus the KV write)
compute boundLong-context ingest is the most compute-dense step in the entire serving stack, and the arithmetic says so unambiguously. Take Llama-3-8B from its own config.json (d=4096, 14336 intermediate, 32 layers, 32 query / 8 KV heads, head_dim 128, vocab 128256): P_layer = d*n_q*dh + 2*d*…flops/long-context-rag-prefill
verifiedhigh · 0.90
Low-entropy / near-greedy decode: when the distribution collapses, which flops can be skipped
memory boundSampling's cost is fixed by vocabulary size, not by how peaked the distribution is - and that is the central fact about low-entropy decode. Per token the sampler must read the logits out of HBM (2*V bytes at bf16, 4*V at fp32) before it can discover that the answer was obvious th…flops/low-entropy-greedy-decode
verifiedhigh · 0.85
Low-rank key projection and KV compression: research-only vs deployed
memory boundLow-rank key projection and KV compression aim to reduce the KV cache by projecting to a lower-dimensional space. MLA (Multi-head Latent Attention) is the deployed extreme of this: it compresses K and V JOINTLY through one shared down-projection W_DKV, so every head's K and V liv…flops/attn-low-rank-kv-compression-status
verifiedhigh · 0.85
Machete weight-only 4-bit GEMM: CUTLASS 3.x tiles plus an offline prepacked int4 B operand, Hopper only
compute boundSame class of claim as Marlin - no FLOP change, dequant removed from the runtime path - but achieved differently: Machete keeps CUTLASS's Hopper WGMMA tensor-core mainloop and changes the B OPERAND LAYOUT instead. The kernel source is explicit that LayoutB is ColumnMajor only 'no…flops/machete-weight-only-hopper-gemm
verifiedhigh · 0.90
Marlin weight-only INT4 GEMM: repacking so dequant is free and tensor cores stay saturated
compute boundMarlin does not change the FLOPs; it removes the dequantization tax by moving it offline and restructuring the kernel so dequant and MMA are ordered so neither pipeline stalls. The author's framing is a ratio: 'as long as we perform less than 25-50 (tensor core) multiply-accumula…flops/marlin-weight-only-int4-gemm
verifiedhigh · 0.85
MI300X 192GB HBM3 memory pooling and KV cache capacity advantage
memory boundflops/amd-mi300x-memory-pooling
verifiedhigh · 0.90
MI300X Infinity Fabric 7-link topology and tensor-parallel scaling efficiency
interconnect boundflops/amd-mi300x-infinity-fabric-topology
verifiedhigh · 0.90
MLA kernel support matrix: which architectures and engines support Multi-head Latent Attention
memory boundMLA kernel support is arch-specific and engine-specific. FlashMLA (DeepSeek's kernel library) requires NVIDIA SM100/SM103 (Blackwell) with CUDA >= 13.1. The 2026.09.30 release REMOVED Hopper support and earlier models (V3/V3.2/V4.0). Decode supports only FP8 and FP4 KV cache - bf…flops/attn-mla-kernel-support-matrix
verifiedhigh · 0.85
MoE expert GEMMs (the sparse feed-forward compute)
memory boundMoE AI is worse than dense by exactly the activation sparsity factor, and that is the whole serving story. With E experts, top-k, total expert params N and batch B: each token touches k/E of the parameters, so flops = 2*B*N*(k/E) while bytes = 2*N (the entire expert weight set is…flops/moe-experts
verifiedhigh · 0.88
MoE gating: router GEMM, softmax/top-k, and the token scatter-gather around it
interconnect boundThe router projection itself is tiny and bandwidth-bound: for t tokens, flops = 2*t*d_router*E while bytes = 2*d_router*E (weights, read once) + 2*t*d (activations), so AI is a small constant that grows only weakly with token count and is dominated by reading the router weights. …flops/moe-routing
verifiedhigh · 0.87
Multi-GPU tensor-parallel inference: the per-layer all-reduce and the batch at which it matters
interconnect boundPure data movement, AI = 0 - like kv_transfer, this class performs no arithmetic, so every question about it is a bandwidth question. Tensor parallelism shards each layer's weight matrices across N devices and inserts two communication ops per layer; Megatron-LM's contribution is…flops/tensor-parallel-allreduce
verifiedhigh · 0.88
Multi-head Latent Attention: KV cache compressed to a single low-rank latent vector per token
memory boundMLA changes the BYTES, not the flops, and it changes them by a large published factor. DeepSeek-V2 Table 1 states KV cache per token in elements: MHA 2*n_h*d_h*l, GQA 2*n_g*d_h*l, MQA 2*d_h*l, MLA (d_c+d_rh)*l. With the DeepSeek-V2 setting d_c = 4*d_h and d_rh = d_h/2, MLA = 4.5*…flops/mla-latent-attention
verifiedhigh · 0.90
Multi-turn chat with prefix caching: two flop profiles from one workload (cache hit vs miss)
both boundA single workload with two genuinely different arithmetic-intensity profiles, and conflating them is the single most common error in prefix-cache benchmarks. MISS path = full prefill: AI = g*L/2 flop/byte for the attention term (Llama-3-8B GQA-4 at L = 2,048 gives 4*2048/2 = 4,09…flops/prefix-cache-hit-miss
verifiedhigh · 0.85
MXFP4 native hardware support on CDNA4 (MI350X/MI355X)
compute boundflops/amd-fp4-mxfp4-support
verifiedhigh · 0.90
NUMA-aware process binding and NPS (NUMA Partitioning Services) on MI300X
interconnect boundflops/amd-numa-nps-binding
verifiedmedium · 0.80
NVFP4 and MXFP4 block-scaled formats on Blackwell: microscaling block size, per-block scales, and kernel consumption
compute boundBlackwell supports three 4-bit floating-point formats with different block-scaling strategies. (1) FP4 (E2M1): 4 bits per value (1 sign, 2 exponent, 1 mantissa), no hardware scaling - requires software-managed scales. (2) MXFP4: E2M1 + 1 shared power-of-two scale (UE8M0) per 32-v…flops/bw-nvfp4-mxfp4-block-scaled-formats
verifiedhigh · 0.85
NVLink 5 and GB200/GB300 NVL72 domain topology: 72-GPU domain and crossing it
interconnect boundNVLink 5 provides 1.8 TB/s bidirectional bandwidth per GPU (vs 900 GB/s on Hopper NVLink 4), which is 14x PCIe Gen5. The GB200 NVL72 connects 72 Blackwell GPUs in a single NVLink domain using 9 NVLink switch trays, each with 144 NVLink ports at 100 GB/s, delivering 130 TB/s total…flops/bw-nvlink5-nvl72-domain-topology
verifiedhigh · 0.85
OCP-FP8 (E4M3/E5M2) vs MXFP8 precision differences on CDNA3 and CDNA4
compute boundflops/amd-fp8-ocp-vs-mxfp8-precision
verifiedhigh · 0.85
On PCIe-attached gfx906, the cross-GPU all-reduce in decode is interconnect-bound and makes decode throughput decline as GPUs are added
interconnect boundExtremely low at batch 1 decode: each added GPU adds a full activation all-reduce over PCIe while contributing only its share of a weight-streaming GEMV. MI50 moves ~1 TB/s of weights per card, so per-GPU weight traffic is the term that shrinks with added GPUs while the all-reduc…flops/comm-flop-gfx906-pcie-allreduce-decode-critical-path
verifiedlow · 0.60
oneAPI/Level Zero/SYCL Programming Model: Intel-Specific Kernel Optimization
compute boundflops/xpu-oneapi-sycl-programming-model
verifiedhigh · 0.85
oneDNN Graph API and IPEX Fusion for Transformer Inference
memory boundflops/xpu-onednn-graph-fusion
verifiedmedium · 0.80
Paged attention block-size constraints: kernels that hard-code block size and the compression ratio trap
memory boundPagedAttention (vLLM) manages KV cache in fixed-size blocks. The block size is a kernel constraint, not just a storage choice. FlashAttention's paged path requires page_block_size to be a multiple of 256. Some kernels hard-code the block size entirely, so a different compression …flops/attn-paged-attention-block-size-hardcoded
verifiedhigh · 0.85
PagedAttention kernels: block-table indirection for the decode KV cache
memory boundPaging changes no arithmetic. It changes the memory layout so that decode attention reads a list of contiguous pages via a per-request block table instead of one dense per-request cache. FLOPs are identical and total bytes are identical; what changes is coalescing efficiency and …flops/paged-attention-block-table-kernel
verifiedmedium · 0.80
Per-kernel launch and dispatch overhead at batch 1
launch_overhead boundNot a byte/flop quantity - it is a fixed CPU-side cost per kernel launch, measured in microseconds, paid before any bytes move. At batch 1 a decode step runs on the order of a few hundred to a few thousand kernels (attention, norms, activations, GEMMs, residuals, sampling), so la…flops/custom-kernel-launch-overhead
verifiedhigh · 0.85
Per-token detokenisation, EOS/stop-string detection, and the host work they put on the critical path
launch_overhead boundThis class has no bytes-per-flop at all: it is host-side text processing that runs once per generated token, on the API server's critical path, outside the model's arithmetic. Two costs, and they scale differently. (1) PER-TOKEN DETOKENISATION: the new token id is turned into tex…flops/detokenization-stop-string-scan
verifiedmedium · 0.80
Pipeline parallelism for LLM inference: when PP beats TP and the bubble problem
interconnect boundPP splits a model layer-wise, where each GPU is responsible for a subset of layers. The arithmetic intensity of each stage is the same as a dense model at the same batch size, but the batch is split into micro-batches that flow through the pipeline. The key advantage over TP: PP …flops/sched-pipeline-parallelism
verifiedmedium · 0.82
Prefill deflection scheduling: deflecting prefill to idle decode nodes instead of queueing
both boundNot an arithmetic mechanism - it is a scheduling decision that changes where prefill executes. The key insight from Kairos (arXiv 2607.02043) is that under bursty, heavy-tailed workloads, prefill nodes saturate while decode nodes have compute underutilized. On a 2P2D A100 cluster…flops/sched-prefill-deflection
verifiedhigh · 0.85
Prefill FLOP crossover: the context length at which attention equals the MLP, computed per architecture family from each model's own config
compute boundTHE CROSSOVER IS A CONFIG RATIO, NOT A CONTEXT NUMBER. Within one prefill of L tokens, per layer: attention flops = 2*n_q*head_dim*L^2 (causal: the two L x L score/output matmuls each contribute 2*n_q*head_dim*L^2 before halving, and the mask halves both, so the factor 2 already …flops/fct-prefill-attention-gemm-crossover-by-family
verifiedhigh · 0.85
Prefill-decode disaggregation: measured TTFT improvements and the KV transfer cost that dominates
interconnect boundAI = 0 in the strict sense - disaggregation is a data movement and scheduling decision, not arithmetic. The KV cache transfer is a DMA copy with no arithmetic, so cost = KV bytes / interconnect bandwidth. The useful question is never 'how fast is the copy' but 'is the copy cheape…flops/sched-prefill-decode-disaggregation
verifiedhigh · 0.88
Prefill/decode disaggregation: the KV handoff tax that separates a compute-bound phase from a memory-bound phase
interconnect boundThe handoff itself is AI = 0 - it is a copy, not arithmetic - so its cost is KV bytes / link bandwidth, and that ratio is the entire design constraint. Volume per request = (prompt_tokens + 1) x 2 (K and V) x num_kv_heads x head_dim x num_layers x kv_dtype_bytes, plus any MLA lat…flops/prefill-decode-disaggregation-kv-handoff
verifiedhigh · 0.85
Prefix caching and PagedAttention block sharing: granularity of sharing and the 128K cache miss penalty
both boundIn standard PagedAttention (SOSP'23, arXiv:2309.06180, [[sched-vllm-pagedattention-sosp23]]), KV cache memory is allocated in non-contiguous physical blocks of fixed size B (typically 16 tokens). Within a single request using parallel sampling or beam search, all prompt tokens an…flops/prefix-caching-paged-attention-sharing-and-miss-penalty
verifiedhigh · 0.90
Prefix/radix-cache hit path: the compute-to-bandwidth inversion and the provenance requirement
memory boundThis record covers the cost-inversion and the provenance question on the hit path; the arithmetic itself (the AI = t_new GEMM identity and the 2*n_q*dh*t_new*L attention term, versus the miss path's 2*n_q*dh*L^2) is derived in [[prefix-cache-hit-miss]]. What belongs here is the s…flops/prefix-cache-hit-path-gemm-only
verifiedmedium · 0.80
Production speculative decoding deployments: engines, settings, and measured throughput
memory boundProduction deployments of speculative decoding are limited and settings vary by engine. (1) vLLM: supports n-gram speculation (prompt lookup) and EAGLE/MTP. Default is often disabled for throughput workloads. CUDA graph downgrades for spec-decode have been reported. (2) SGLang: s…flops/sd-production-deployments
verifiedmedium · 0.80
Quantization runtime cost: dequant, scale unpacking, and batch crossover
both boundQuantization changes the bytes-per-flop ratio of every weight GEMM by the compression factor: a 4-bit weight format cuts weight bytes 4x versus bf16, so decode GEMM AI goes from ~B to ~4B flop/byte and the effective ridge crossing batch drops by roughly the compression factor. Th…flops/quantization-overhead
verifiedhigh · 0.85
Quantized decode at high batch: where the dequantization tax comes back
both boundQuantization raises decode arithmetic intensity by the compression factor, which moves the roofline crossover batch DOWN by the same factor - and that is precisely why dequantization overhead reappears at high batch. For 4-bit weights, bytes fall 4x while flops are unchanged, so …flops/quantized-decode-high-batch
verifiedhigh · 0.90
Rebellions Rebel100: 4-chiplet UCIe mesh and predictive DMA for inference
memory boundRebel100 uses 4 chiplets connected via UCIe-Advanced mesh at 1 TB/s per bidirectional channel at 11 ns latency. The four chiplets present as one virtual die with no software-visible boundary. 144 GB HBM3E at 4.8 TB/s. The predictive DMA engine enables 2.7 TB/s effective bandwidth…flops/alt-rebellions-chiplet-mesh
verifiedhigh · 0.85
RMSNorm, LayerNorm, and softmax / activation elementwise work
memory boundThe most memory-bound class in the model - AI well below 1 flop/byte, and it does not improve with batch or context. RMSNorm on a t x d activation: flops ~ 4*t*d (square, sum, scale) against bytes = 2*t*d*2 = 4*t*d read plus 2*t*d*2 = 4*t*d written, so AI ~ 4*t*d/8*t*d = 0.5 flop…flops/norm-softmax
verifiedmedium · 0.75
rocBLAS/hipBLASLt tuning for MI300X/MI355X GEMM performance
compute boundflops/amd-rocm-rocblas-tuning
verifiedmedium · 0.80
SageAttention is research-only: absent from vLLM and SGLang LLM serving paths
compute boundSageAttention is a quantization-based attention optimization that is NOT deployed in production engines. It is absent from the vLLM source tree and in SGLang it appears only under multimodal_gen, not the LLM serving path. SageAttention quantizes attention computations to int8/fp8…flops/attn-sageattention-not-deployed
verifiedhigh · 0.90
SGLang ROCm-specific configuration and AITER integration for AMD GPUs
compute boundflops/amd-sglang-rocm-flags
verifiedmedium · 0.80
Sliding-window and interleaved local/global attention: local layers are O(L*w) in prefill and O(1) in cache
compute boundA local layer replaces the L x L causal attention matrix with L x w, so prefill flops fall by the factor L/w while bytes for K and V fall by the same factor only for the retained window - the cache per token per LOCAL layer is bounded at w and does not grow with context at all. G…flops/sliding-window-hybrid-attention
verifiedhigh · 0.85
Sliding-window and SSM hybrid decode: the fixed per-sequence state cost invisible in KV-per-token calculations
memory boundHybrid attention-SSM models have a per-sequence cost that does NOT show up in a KV-per-token calculation. For SSM layers (Mamba, etc.), the recurrent state is fixed-size per sequence: state = d_inner * N elements, where d_inner = expand * hidden and N = ssm_state_size. This state…flops/attn-sliding-window-ssm-fixed-state-cost
verifiedhigh · 0.90
Sparse and structured attention variants: NSA, DSA, and their honest deployment status
both boundSparse attention reduces the WORK SET but adds a full-context index pass and non-contiguous gather. NSA (Native Sparse Attention): three paths per query - compressed coarse tokens, top-k selected fine tokens, sliding window. Per-query decode work falls from L to (compressed_count…flops/attn-sparse-structured-attention-status
verifiedhigh · 0.85
Sparse attention mechanisms for long context: hardware-aligned block selection in NSA and MoBA
both boundSparse attention replaces the full O(L^2) attention computation with selective subsets of tokens, but token-granular selection collapses arithmetic intensity by turning contiguous tensor loads into non-contiguous gathers. Hardware-aligned sparse attention enforces block-level reg…flops/sparse-attention-hardware-aligned-nsa-moba
verifiedhigh · 0.88
Sparsity-aware expert cache: request-level activation trace driving prefetch and eviction under host-RAM expert offload
interconnect boundThe expert GEMM is unchanged. What changes is the byte term. With experts resident in host DRAM, a decode step moves k experts per layer over PCIe instead of reading all E experts from HBM, so the step is interconnect-bound at PCIe bandwidth rather than HBM-bound at the full expe…flops/moe-infinity-expert-cache-prefetch
verifiedhigh · 0.85
Speculative decoding at scale: draft versus verify cost, and why the win shrinks as batch grows
memory boundBoth passes are memory-bound in the same way a decode GEMM is, and that is the entire reason the scheme has a ceiling. Draft: AI ~ k (reads the draft weights once for k positions). Verify: AI ~ k (reads the target weights once for k positions). A plain decode step has AI = 1. So …flops/speculative-decoding-bandwidth-trade-at-batch
verifiedmedium · 0.75
Speculative decoding batch dilution: why speedup collapses at production batch sizes
memory boundThe core problem: speculative decoding raises arithmetic intensity from 1 to k on both passes, but a real decode step at batch B already has AI = B. So a k=4 speculative step at batch B has AI = B + k - the increment shrinks as a fraction of the total. At batch 1, speculative dec…flops/sd-batch-dilution-problem
verifiedhigh · 0.90
Speculative decoding cost model: the break-even acceptance rate a(B, k, draft size) actually requires, and the draft model's HBM cost against KV cache
both boundTHE COST MODEL, with each term defined. Baseline decode step at batch B: T_base(B) = max(2*P_t/BW, 2*B*P_t/PEAK) for the target's weight stream alone, with P_t the target parameter count and the same terms added for KV. One speculative iteration that drafts k positions and expect…flops/fct-spec-decoding-cost-model-break-even
verifiedmedium · 0.75
Speculative decoding drafter bottlenecks: draft LM-head projection and quantized drafter scaling
memory boundIn modern auxiliary-head speculative drafters (EAGLE-3, Medusa, MTP), the drafter backbone is compact (1-2 layers, hidden size d), but its LM-head must project hidden states to the full target vocabulary V (e.g. V = 128k in Llama-3, V = 152k in Qwen3). Consequently, computing log…flops/speculative-quantized-drafter-lm-head
verifiedhigh · 0.88
Speculative decoding hardware sensitivity: GPU generations, FP8/FP4-native methods
memory boundAcceptance rate is primarily a function of draft-model quality and task, NOT hardware - the same draft model achieves the same acceptance rate on any GPU. What changes across hardware is the OVERHEAD PROFILE: (1) On GPUs with higher memory bandwidth (H100 → B200 → GB300), the ver…flops/sd-hardware-sensitivity
verifiedmedium · 0.75
Speculative decoding method families: draft-model vs n-gram/lookahead vs tree variants
memory boundThree distinct families with fundamentally different cost profiles. (1) Draft-model (EAGLE-1/2/3, Medusa, Hydra): AI ~ k for the draft pass (reads draft weights once for k positions), AI ~ k for verify (reads target weights once for k positions). Draft model is a separate neural …flops/sd-speculative-method-families
verifiedhigh · 0.85
Speculative decoding rejection sampling: exact distribution math and acceptance-rate mechanics
memory boundVerification pass computes target model logits for all k draft tokens in a single batched forward pass. For each draft token x_i proposed by draft distribution q(x_i | x_<i) and evaluated by target distribution p(x_i | x_<i), modified rejection sampling accepts x_i with probabili…flops/speculative-decoding-rejection-sampling-math
verifiedhigh · 0.90
Speculative decoding: draft-model proposal pass
memory boundThe draft pass reads the draft model's weights for k tokens' worth of positions at once. With a draft of size N_d, flops = 2*k*N_d against bytes = 2*N_d (weights) + small activations, giving AI ~ k flop/byte - the batch-k of a batched decode GEMM. So the draft pass is memory-boun…flops/speculative-draft
verifiedmedium · 0.80
Speculative decoding: target-model verification pass
memory boundThe verify pass runs the target model over k draft positions as a batch-k forward pass. flops = 2*k*N_target, bytes = 2*N_target (weights read ONCE for all k positions) plus k times the KV-cache attention and norm work. Weights dominate: AI ~ k flop/byte. So verify is the same me…flops/speculative-verify
verifiedhigh · 0.85
SSM / linear-RNN decode: the constant-size recurrent state update that replaces the KV cache entirely
memory boundThis class has no KV cache at all, which inverts the entire decode cost model. Per layer the state is S = d_inner * N elements (d_inner = expand*hidden, N = d_state) and each decode step does ONE state update: for Mamba-style, flops = 2*(d_inner*N*2 + d_inner*N) for the B_t@C_t p…flops/ssm-recurrent-state-update
verifiedhigh · 0.90
tcgen05 5th-gen Tensor Core: native FP4/FP6 support and emulation paths
compute boundThe tcgen05.mma instruction on Blackwell (sm_100a) supports multiple .kind values that determine which precision formats are natively executed. From the PTX ISA 9.4 Table 48: .kind::f16 supports FP16/BF16 at K=16 (dense) or K=32 (sparse); .kind::f8f6f4 supports FP8 (E4M3/E5M2) an…flops/bw-tcgen05-fp4-fp6-native-support
verifiedhigh · 0.85
Tensor Memory (TMEM): a third accumulator store, between registers and shared memory
both boundTMEM changes the capacity budget of a GEMM or attention mainloop rather than the roofline. It is a dedicated on-chip scratchpad wired into the tensor cores, with structure and rules that come straight from the PTX ISA: on sm_100a/sm_100f it is a two-dimensional matrix of 512 colu…flops/tensor-memory-tmem-allocation
verifiedhigh · 0.85
Tensor parallel crossover batch B*: the batch above which the all-reduce dominates
interconnect boundPure data movement, AI = 0 - like kv_transfer, this class performs no arithmetic, so every question about it is a bandwidth question. Tensor parallelism shards each layer's weight matrices across N devices and inserts two communication ops per layer. The op that matters in infere…flops/sched-tensor-parallel-crossover
verifiedhigh · 0.88
Tensor-core instruction generations: wmma/mma (Ampere) -> wgmma (Hopper) -> tcgen05 (Blackwell)
compute boundThis is the layer that determines what a kernel can possibly do. Three generations exist and they change how a kernel must be STRUCTURED, not just how fast it runs. (1) Ampere and earlier: warp-level wmma (implicit load/store of fragments, the compiler handles element distributio…flops/tensor-core-instruction-generations-wgmma-tcgen05
verifiedhigh · 0.90
Tenstorrent Tensix: small-SRAM, modest-bandwidth, open-source architecture
memory boundTenstorrent Wormhole (n150) has 108 MB SRAM, 12 GB GDDR6 at 288 GB/s, 74 TFLOPS FP16. Blackhole (p150) has 180 MB SRAM, 32 GB GDDR6 at 512 GB/s, 664 TFLOPS BLOCKFP8. The architecture uses Tensix cores with dedicated RISC-V control cores and a network-on-chip. SRAM acts as a L2-li…flops/alt-tenstorrent-tensix
verifiedhigh · 0.88
The cost of one forward pass with no batching: the batch-1 decode-step formula, and the batch at which it stops being memory-bound
memory boundBATCH 1 IS A REGIME, NOT A SCALE. At batch 1 a decode step streams the whole weight set once and does one multiply-accumulate per weight, so AI = 2*P_flops/(w_b*P_bytes) = 2/w_b, where w_b is bytes per stored weight: exactly 1.0 flop/byte for bf16 (w_b=2), 2.0 for fp8/int8 (w_b=1…flops/fct-batch1-step-cost-formula
verifiedmedium · 0.80
The energy form: deriving joules per token from watts and throughput, and from bytes and bandwidth when you have neither, with a stated error bound
both boundTWO ROUTES, AND THE CHOICE BETWEEN THEM IS THE WHOLE SUBJECT. ROUTE 1 (measured watts) has ZERO model error and should always be preferred: E/token [J] = P_mean [W] / throughput [tok/s]. The only requirements are that P_mean is the MEAN over the same window that produced the toke…flops/fct-energy-form-joules-per-token-derivation
verifiedmedium · 0.75
The MLA up-projection that absorption removes: recomputing K and V from the cached latent
memory boundThis is the cost MLA hides, and it only appears if you decline the absorption trick. DeepSeek-V2 states that during inference 'since W_UK can be absorbed into W_Q, and W_UV can be absorbed into W_O, we even do not need to compute keys and values out for attention'. Do the naive t…flops/mla-absorbed-upprojection
verifiedhigh · 0.85
The overhead FLOP ledger: norms, RoPE, activations, softmax, residuals, KV-cache writes and sampling priced per decode step, fused versus unfused
memory boundTHE HEADLINE IS A BYTE RESULT, NOT A FLOP RESULT, AND THAT IS THE POINT. The elementwise classes carry real flops but almost no arithmetic intensity, so they are invisible in a FLOP count and dominant in a bandwidth count. Per decode step, per token-slot, per layer, with d = hidd…flops/fct-overhead-flop-ledger
verifiedmedium · 0.75
TileLang / TileKernels: a tile DSL as the middle layer between hand-written CUDA and Triton
both boundThe point of a tile DSL is that the author writes tiles and the compiler decides the instruction selection, register/shared-memory allocation and async-copy pipelining. That means the same source can target NVIDIA SM90/SM100 and Huawei Ascend 950 with identical Python APIs, selec…flops/tilelang-tile-dsl-moe-routing-quant
verifiedmedium · 0.75
TMA (cp.async.bulk.tensor): hardware address generation and predication for global->shared copies
memory boundTMA reduces register pressure and address-computation work in the copy path, which is what allows larger tiles and therefore more FLOPs per byte moved from HBM. The mechanism from the PTX ISA: cp.async.bulk.tensor is a non-blocking instruction that initiates an asynchronous copy …flops/tma-cp-async-bulk-tensor
verifiedhigh · 0.85
TMA + tcgen05 MMA with async copy and warp-specialization in CUTLASS Blackwell kernels
both boundBlackwell kernels use TMA (cp.async.bulk.tensor) for global->shared memory copies and tcgen05.mma for tensor-core operations, with warp specialization to overlap data movement with compute. The producer-consumer pattern: producer warps issue TMA copies (cp.async.bulk.tensor) to m…flops/bw-tma-tcgen05-async-copy-warp-specialization
verifiedhigh · 0.85
TPU bucketing and padding: batch size rules for MXU efficiency
memory boundThe TPU runtime splits a batch across all 8 cores of a TPU device. For v5e/v5p/v6e, each core uses 8x128 vector registers. Batch size should be evenly divisible by 8 or 128. The XLA compiler rounds up tensor sizes in HBM for efficiency. Two rules: (1) batch multiple of 64 with fe…flops/alt-tpu-bucketing-padding
verifiedhigh · 0.95
TPU MFU: performance guide rules for achieving good model flops utilization
both boundMFU (Model FLOPs Utilization) on TPU is governed by: (1) input-bound models — ensure steady data stream with parallel reads, (2) small batch from sharding — use largest batch that fits, (3) memory management — TPU_PREMAPPED_BUFFER_SIZE tuning, pinned_host vs unpinned_host, (4) ne…flops/alt-tpu-mfu-performance-rules
verifiedhigh · 0.95
TPU SPMD partitioning: how layout constrains model architecture
interconnect boundTPUs use SPMD (Single Program, Multiple Data) where the same program runs on all cores. The XLA compiler automatically partitions tensors across cores using layout constraints. This means the model architecture must be amenable to automatic sharding — manual control is limited co…flops/alt-tpu-spmd-partitioning
verifiedhigh · 0.90
TPU XLA compilation: static shapes, padding, and fusion
both boundXLA compiles the entire HLO graph to a TPU executable at compile time. All tensor shapes must be known statically — there is no dynamic shape support. The compiler pads tensors to 128x8 chunks (the MXU tile size); padding wastes compute and memory. Batch size must be a multiple o…flops/alt-tpu-xla-static-shapes
verifiedhigh · 0.95
TT-Metal Flash-Decode: split the KV cache across Tensix cores, reduce over the NoC
memory boundDecode attention reads the whole KV cache for one query token, so bytes moved grow with sequence length while the matmul is a single query row against that cache. The prefill kernel's own report finds the fused op is not fully bandwidth-bound — BFP8 is only moderately faster than…flops/tt-metal-flash-decode-kv-split
verifiedmedium · 0.75
Video inference: frame sampling and per-frame token cost
both boundVideo inference cost = frame_sampling_rate * tokens_per_frame * num_frames. For a 1-minute video at 1 fps = 60 frames. At 448x448 per frame with 2x2 spatial merge = 256 tokens per frame. Total visual tokens = 60*256 = 15,360. LLM prefill cost = 2*7e9*15,360 = 215 TFLOPs. KV cache…flops/video-inference-frame-sampling-and-per-frame-token-cost
verifiedmedium · 0.80
Vision-language connector cost (MLP projector / Perceiver resampler / Q-Former)
compute boundThe connector is a small transformer or MLP that maps vision encoder output to LLM embedding space. For a 2-layer MLP projector (LLaVA style): input dim 1024 (ViT-L hidden), output dim 4096 (LLM hidden), 256 tokens. FLOPs = 2*1024*4096*256*2 = 4.3 GFLOPs. For a Perceiver resample…flops/vision-language-connector-cost-mlp-projector-perceiver-resampler
verifiedhigh · 0.85
vLLM ROCm-specific flags, environment variables, and known-good configuration
launch_overhead boundflops/amd-vllm-rocm-flags
verifiedhigh · 0.90
What a 128K prefill actually costs: flops, wall-clock, energy and dollars
compute boundPrefill flops are 2*N_active*L for the weight GEMMs plus 2*n_q*d_h*L^2 for causal attention, and at 128K the FIRST term dominates for every current dense model - which is the counter-intuitive result and the one worth internalizing. Using active (not total) parameters so MoE mode…flops/ctx-128k-prefill-compute-cost
verifiedmedium · 0.78
Which attention kernels can run long-context attention on which hardware: the SM-level gate on whether long context is servable at all
compute boundKernel availability at long context is a gate, not a tax: it decides whether a context length is servable on a given card, before any efficiency question is asked. The gate is the streaming-multiprocessor generation, because each FlashAttention generation needs a specific tensor-…flops/ctx-long-context-attention-kernel-matrix
verifiedmedium · 0.82
Which scheduling optimizations survive at production scale versus which are benchmark artifacts
both boundNot an arithmetic mechanism - it is a meta-observation about which scheduling techniques work in production versus which only work in benchmarks. The key distinction: production workloads have heavy-tailed prompt lengths, bursty arrivals, and mixed prefill/decode traffic, while b…flops/sched-production-vs-benchmark
verifiedmedium · 0.78
Why SSM and linear-RNN FLOPs do not convert into GPU FLOPs: the sequential scan bypasses the tensor cores
both boundThe arithmetic is small and easy; the hardware is what fails. A selective-SSM step is h_t = a_t*h_{t-1} + b_t*x_t with a per-channel, input-dependent decay a_t - a rank-1-scaled recurrence whose dependence between time steps is a MULTIPLICATIVE chain. Tensor cores compute sums of…flops/ssm-sequential-scan-tensor-core-mismatch
verifiedhigh · 0.90
Why the gap exists: which RoPE/context-extension method each long-context model actually used, and trained-at-shorter-than-licensed
compute boundPosition-embedding scaling does not change attention FLOPs or KV bytes at all - it is a pure index remap on the rotary phase. AI is unchanged; what changes is whether the model's learned function of relative distance still matches the distances it is now being asked about. That i…flops/ctx-rope-scaling-method-per-model
verifiedhigh · 0.85
YOCO, NSA, DSA, RetNet, lightning attention, HGRN2, Kimi Linear: attention replacements that change the complexity class
both boundGroup the architectures by what they do to the flops, because the four groups behave completely differently on hardware. (1) SINGLE-CACHE DECODER-DECODER (YOCO): a self-decoder computes the global KV cache once and a cross-decoder consumes it, so the cache is stored ONCE instead …flops/modern-attention-variants
verifiedhigh · 0.85
Most referenced
The records other records point at. Useful as an entry point: a heavily referenced accelerator or engine is one many decisions in this repo assume.