UPSTREAM PR #18747: ggml, llama : add CPU paged attention for memory-efficient KV cache#883
Open
loci-dev wants to merge 16 commits into
Open
UPSTREAM PR #18747: ggml, llama : add CPU paged attention for memory-efficient KV cache#883loci-dev wants to merge 16 commits into
loci-dev wants to merge 16 commits into
Conversation
Add ability to reduce the number of active experts in MoE models at runtime, providing significant speedup with minimal quality loss when using 50% of default experts. Implementation: - Add moe_n_expert_override parameter to llama_context_params - Add --moe-n-expert CLI flag to override n_expert_used - Implement "Hard Mask" in build_moe_ffn() that slices expert tensors - Uses ggml_view_2d/3d + ggml_cont to reduce actual computation Benchmark results (AOCL BLIS 5.0, AMD EPYC 9655): - Qwen3-Coder-480B-A35B: 2.5 → 3.7 t/s (48% speedup) - GLM-4.6-355B-A32B: 2.2 → 3.0 t/s (36% speedup) - Qwen3-Coder-30B-A3B: 26.6 → 33.6 t/s (26% speedup) - Qwen3-VL-30B-A3B: 32.2 → 38.9 t/s (21% speedup) Quality: Excellent at 50% experts, degraded at 25%, gibberish at 12.5% Usage: llama-cli -m model.gguf --moe-n-expert 4 -p "prompt" 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Adds n_layer_exit parameter to control how many layers to compute, enabling early exit speculation techniques like CAS-Spec and CLaSp. Changes: - Add n_layer_exit to llama_context_params (public API) - Add n_layer_exit to llama_cparams (internal) - Add --n-layer-exit CLI parameter - Implement layer skip in model graph builders: - llama.cpp (models) - qwen2.cpp - qwen3.cpp - qwen3moe.cpp When n_layer_exit > 0 and < n_layer, the model will exit early after computing that many layers. This is useful for generating draft tokens in speculative decoding scenarios. Example: --n-layer-exit 7 on a 28-layer model gives ~2.2x speedup 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Extends layer skip / early exit support to additional model architectures: - qwen3vl-moe.cpp (Qwen3-VL-30B-A3B and similar VL MoE models) - qwen3next.cpp (Qwen3-Next-80B-A3B and similar hybrid attention models) Results after adding layer skip support: - Qwen3-VL-30B-A3B: 3.4x speedup with 16 layers (vs all 48) - Qwen3-Next-80B-A3B: 3.7x speedup with 8 layers - Qwen3-Coder-480B-A35B: 5.0x speedup with 16 layers 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
llama-lookahead has been broken since PR #14482 (July 2025) which changed seq_id validation from LLAMA_MAX_SEQ constant to context-specific n_seq_max. Two lookahead-specific issues: 1. n_seq_max: Lookahead needs W + G + 1 = 31 sequences for parallel Jacobi decoding, but params.n_parallel defaulted to 1. Fix: Set params.n_parallel = W + G + 1 before context creation. 2. KV unified: Batch splitting with coupled sequences requires unified KV cache mode, but lookahead didn't enable it. Fix: Set params.kv_unified = true. Bug timeline: - Nov 2023: lookahead.cpp created, worked with LLAMA_MAX_SEQ constant - July 2025: PR #14482 changed to n_seq_max validation, broke lookahead Note: This PR depends on #18729 for the batch init fix (params.n_ctx -> llama_n_ctx). Both PRs are needed for lookahead to fully work. Tested with Qwen2.5-Coder-0.5B: lookahead generates output with n_accept > 0. Bug history researched with Claude.
Since PR #16653 (Dec 15, 2025), the default n_ctx is 0 to enable automatic
GPU memory fitting. This causes llama-lookup and llama-lookahead to crash
when run without explicit -c flag:
GGML_ASSERT(batch.seq_id[batch.n_tokens] && "llama_batch size exceeded")
Root cause: Both examples use params.n_ctx directly for batch initialization,
but params.n_ctx remains 0 even after the context is properly initialized
to n_ctx_train internally.
Bug history:
- Nov 2023: lookahead.cpp created (PR #4207) with params.n_ctx pattern
- Dec 2023: lookup.cpp created (PR #4484) with same pattern
- Nov 2024: default n_ctx changed to 4096 (PR #10136) - bug dormant
- Dec 2025: default n_ctx changed to 0 (PR #16653) - bug activated
The bug was dormant for 2+ years because params.n_ctx defaulted to 512,
then 4096. PR #16653 changed it to 0 for GPU auto-fitting, triggering
the crash.
Fix: Use llama_n_ctx(ctx) to get the actual runtime context size, matching
the pattern already used elsewhere in lookup.cpp (line 72) and in
speculative.cpp/speculative-simple.cpp.
Tested: llama-lookup now works without -c flag (12.5% acceptance on
Gemma-3-1B).
Note: llama-lookahead has a separate pre-existing issue with sequence
initialization (n_seq_max=1 vs W+G+1 needed) that is unrelated to this fix.
Add OpenMP parallelization to tensor repack functions to significantly speed up model loading on many-core CPUs. Measured on AMD EPYC 9655 (96 cores): | Model Size | Before | After | Speedup | |------------|--------|-------|---------| | 6.8GB Q4_K | 5.0s | 3.3s | 1.5x | | 19GB Q4_K | 11.9s | 5.3s | 2.2x | | 271GB Q4_K | ~150s | ~60s | ~2.5x | The repack functions convert quantized tensors from storage layout to SIMD-optimized layout for AVX-512. This was previously single-threaded and is now parallelized across row groups. Key changes: - Convert pointer-increment loops to explicit indexing - Add #pragma omp parallel for to outer loops (guarded by #ifdef _OPENMP) - Each thread processes independent row groups - Move thread-local dst_tmp arrays inside parallel region Functions parallelized: - repack_q4_0_to_q4_0_4_bl (Q4_0 x4 interleave) - repack_q4_K_to_q4_K_8_bl (Q4_K_M, Q4_K_S models) - repack_q2_K_to_q2_K_8_bl (Q2_K models) - repack_q4_0_to_q4_0_8_bl (Q4_0 x8 interleave) - repack_iq4_nl_to_iq4_nl_4_bl (IQ4_NL x4) - repack_iq4_nl_to_iq4_nl_8_bl (IQ4_NL x8) Tested on: AMD EPYC 9655 "Turin" with 192 threads
Establishes rules for: - Branch hierarchy (production-consolidated is protected) - Mandatory clean rebuilds after branch switches - Symbol verification before benchmarking - Research branch workflow - Tagging working states Created after investigating SIGSEGV crashes caused by stale build with undefined symbol from feature/eagle-penultimate-layer branch. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Previously, find_slot() checked if cached cells were masked relative to the stored sequence max position. For SWA caches during speculative decoding, this conservative check prevented reusing cells that would be outside the attention window after batch insertion. Now, for SWA caches (n_swa > 0), we compute the batch's max position and use that for the masking check. This enables forward-looking slot reuse: cells that will be masked AFTER the batch is inserted can be reclaimed immediately. Results on Gemma-3-27B + 1B draft (speculative decoding): - Before: Required --swa-full (SWA cache = 10240 MiB) - After: Works without --swa-full (SWA cache = 624 MiB) - Memory reduction: 94% - Acceptance rate: 42-81% This optimization applies to all ISWA models (Gemma-3 family) and enables efficient speculative decoding without the memory overhead of --swa-full. Claude was used to research the codebase.
Use batch minimum position instead of maximum when determining which cells can be reused in SWA caches. This ensures all tokens in the batch have their full attention window, satisfying the mathematical precision requirement while preserving memory savings. The token at the minimum position has the most demanding context requirement (extends furthest back in history). By checking reusability against this position, we guarantee correctness for all batch tokens. Memory impact is negligible: only (batch_size - 1) fewer cells can be reused compared to the max-based approach. Tested with Gemma-3-12B (n_swa=1024) + Gemma-3-1B draft: - 1504 tokens generated (47% beyond window boundary) - SWA cache stayed bounded at 1536 cells throughout - 50% speculative acceptance rate - Output quality verified (coherent technical document) Commit message drafted with Claude.
Add paged attention support to reduce KV cache memory waste from 30-70% to <10% through non-contiguous block allocation. Changes: - Add GGML_OP_FLASH_ATTN_EXT_PAGED operation to ggml - Implement paged attention kernel with block table indirection - Add block prefetching to minimize indirection overhead - Integrate block tracking into llama_kv_cache - Add LLAMA_PAGED_ATTN=N env var to enable (N=block size in tokens) The paged kernel uses identity mapping (physical=logical) by default, enabling seamless integration with existing code paths. When block tracking is enabled, it uses the block table for indirect K/V access. Testing shows identical outputs with <1% performance overhead. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Add proper block allocation in update_block_tokens(): - Allocates physical blocks from pool when logical blocks are first accessed - Updates block metadata (seq_id, logical_idx, n_tokens) - Uses set to track which logical blocks have been processed per sequence - Add block deallocation in seq_rm(): - Deallocates all blocks when a sequence is removed - Handles both single sequence removal and full cache clear - Wire up update_block_tokens() call at end of apply_ubatch() This enables actual memory savings from paged attention by allocating blocks on-demand rather than using identity mapping (physical=logical). 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- Add print_block_stats() method to llama_kv_cache - Computes and logs block pool utilization, token counts, and memory usage - Called automatically after seq_rm when LLAMA_KV_CACHE_DEBUG > 0 - Reports: blocks used/total, tokens used/total, fragmentation %, memory stats 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Add LLAMA_PAGED_ATTN_MAX_BLOCKS environment variable to limit KV cache size. When both LLAMA_PAGED_ATTN and LLAMA_PAGED_ATTN_MAX_BLOCKS are set: - KV cache is reduced to (max_blocks * block_size) tokens - Memory savings can exceed 80% for large context models Example: LLAMA_PAGED_ATTN=64 LLAMA_PAGED_ATTN_MAX_BLOCKS=100 - Limits cache to 6400 tokens (100 * 64) - Qwen3-1.7B: 4480 MiB → 700 MiB (84.4% savings) 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
- 19 tests covering allocation, deallocation, reference counting - Pool tests: init, allocate, batch allocate, stats, clear - Table tests: mapping, append, sequence management, truncate - Integration tests: pool+table coordination, CoW simulation - Added thread safety documentation to header 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
New flags: --paged-attn N enable paged attention with block size N --paged-attn-max-blocks N max blocks for memory reduction These flags set the corresponding environment variables (LLAMA_PAGED_ATTN and LLAMA_PAGED_ATTN_MAX_BLOCKS) which are read by the KV cache implementation. Example: llama-cli --paged-attn 64 --paged-attn-max-blocks 100 -m model.gguf This achieves 84% memory savings on large context models. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
Reduce comment verbosity to match llama.cpp code style. Detailed explanations moved to PR description. 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>
|
Explore the complete analysis inside the Version Insights I've successfully generated a performance summary report for your project. The report shows: Key Highlights: ✅ Major Performance Improvements:
The analysis compares version Would you like me to provide more details about any specific aspect of the report? |
0b19ffe to
1f94f34
Compare
54e0744 to
c869ee9
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Mirrored from ggml-org/llama.cpp#18747
Summary
Add vLLM-inspired paged attention to llama.cpp, enabling block-based KV cache access for improved memory locality on large models and significant memory savings via dynamic allocation.
Motivation
Large language models (70B+) are memory-bound during attention computation. By organizing KV cache access into blocks with prefetching, we achieve better cache utilization and reduced memory stalls. Additionally, paged attention enables memory savings by limiting KV cache allocation to only the blocks actually needed.
This implementation is inspired by PagedAttention (Kwon et al., 2023) used in vLLM.
Related Work
This PR provides a CPU-only implementation, complementing PR #17579 which implements CUDA-only paged attention. The two approaches differ in design:
__builtin_prefetch()GGML_OP_FLASH_ATTN_EXT_PAGEDGGML_OP_PAGED_ATTENTION--paged-attn N(block size)--pagedattention(boolean)These implementations are independent and can coexist. Per CONTRIBUTING.md guidance to focus on CPU first, this PR establishes the CPU path while PR #17579 addresses CUDA. Future work could unify the APIs if both are merged.
vLLM Design Adaptation
We studied vLLM's PagedAttention implementation to understand the core concepts:
Key adaptations for llama.cpp:
llama_contextaccess patternImplementation
1. GGML Kernel (
ggml/src/ggml-cpu/ops.cpp)New op
GGML_OP_FLASH_ATTN_EXT_PAGED(~343 lines) with:__builtin_prefetch()on block boundariesKey difference from standard flash attention:
2. Block Tracking (
src/llama-kv-block.h)Data structures for block-based memory management:
llama_kv_block_pool: Stack-based O(1) allocator with free listllama_kv_block_table: Sequence → physical block mappingllama_kv_block_stats: Utilization and fragmentation tracking3. KV Cache Integration (
src/llama-kv-cache.cpp)enable_blocks(N): Enable block tracking with N tokens/blockbuild_block_table_tensor(): Create I32 tensor for graphupdate_block_tokens(): Dynamic block allocation inapply_ubatch()seq_rm(): Block deallocation when sequences are removed4. Graph Integration (
src/llama-graph.cpp)Routes to paged kernel when block table is available:
Performance
Benchmark Environment
numactl --interleave=allfor 70B modelSmall Model: Qwen3-1.7B-Q8_0 (16 threads)
Conclusion: Negligible impact on small compute-bound models.
Medium Model: DeepSeek-R1-Distill-Qwen-32B-Q4_K_M (96 threads)
Conclusion: Generation overhead on compute-bound medium models. Prompt processing slightly improved.
Large Model: Meta-Llama-3.1-70B-Instruct-Q4_K_M (96 threads, numactl)
Conclusion: Significant speedup on memory-bound large models, especially with smaller block sizes.
Memory Savings
Small Model: Qwen3-1.7B (40960 token default context)
Medium Model: DeepSeek-R1-Distill-Qwen-32B (131072 token default context)
Large Model: Meta-Llama-3.1-70B-Instruct (131072 token default context)
Key Insight: Memory savings scale with context reduction. For typical interactive use cases (8K-25K tokens), paged attention can reduce KV cache memory by 60-80%.
Correctness
test-kv-block.cppcovering block pool and table operationsUsage
Important: Paged attention is a variant of flash attention, not a replacement. It requires
--flash-attn onbecause it operates within the flash attention code path.Design Decisions
Why a New GGML Op vs Extending Existing?
We considered adding a conditional path inside
ggml_compute_forward_flash_attn_extbut rejected this:Why CPU-Only?
Per llama.cpp CONTRIBUTING.md: "When adding support for a new model or feature, focus on CPU support only in the initial PR."
Additionally, PR #17579 already provides a CUDA implementation with a different architecture. Rather than duplicating that work, this PR focuses on CPU where:
__builtin_prefetch()) provides cache-line optimizationThe two PRs can coexist independently or be unified later if maintainers prefer a single API.
Memory Overhead of Block Tracking
Total: ~3 KB vs KV cache size of 8-32 GB. Negligible (<0.00001%).
Thread Safety Model
Block pool and block table are explicitly NOT thread-safe. This matches
llama_contextdesign:llama-kv-block.hWhy Single PR?
We considered splitting into multiple PRs (kernel, infrastructure, integration) but chose single PR because:
Files Modified
Future Work
--paged-attnand--pagedattentionflagsTest Plan
test-kv-blockpasses (19 tests)test-backend-ops -o FLASH_ATTN_EXT_PAGEDpassesAcknowledgments
Inspired by vLLM's PagedAttention implementation and the paper "Efficient Memory Management for Large Language Model Serving with PagedAttention" (Kwon et al., 2023).
The authors used claude to construct schematics of vLLM's paged attention infrastructure, review our plan for integrating with existing flash attention rather than replacing it outright in llama.cpp, build unit tests, and draft PR submission.