# PrismML's llama.cpp fork (Ternary Bonsai 2 PTQ1_0/PQ2_0 kernels) with sudoingX's
# bonsai2 branch on top: the fast PTQ1_0 mat-vec kernel (PrismML-Eng/llama.cpp#218)
# and the MTP fixes. Built for one CUDA architecture and without the web UI so
# it needs neither the full arch matrix nor npm.
{
lib,
cmake,
ninja,
pkg-config,
fetchFromGitHub,
autoAddDriverRunpath,
cudaPackages_12_9,
# RTX 3060 (sm_86) and RTX 4060 (sm_89), so either card runs the same build.
cudaArchitectures ? "86;89",
...
}:
let
cudaPackages = cudaPackages_12_9;
rev = "ff414120c343e6e6cb868013c99f1dde52b27e70";
in
cudaPackages.backendStdenv.mkDerivation (finalAttrs: {
pname = "llama-cpp-bonsai";
version = "prism-b10743-ff41412";
src = fetchFromGitHub {
owner = "sudoingX";
repo = "llama.cpp";
inherit rev;
hash = "sha256-CQOTOY5IdnQA6c6DceD5xYxYddMnJ9SPy8ihbBJhvME=";
};
patches = [
# Batch-invariant mode sends 1 to 4 query tokens over an F16 cache through one tensor-core tile,
# instead of the vector kernel that slows down at long contexts.
./llama-cpp-bonsai-bi-mma-f16.patch
# Time to first token on cached prompts: no scheduler reserve per request with backend sampling,
# no rewrite of a just-restored checkpoint, recycled checkpoint buffers, no tiny batch split
# for a nearby last user message.
./llama-cpp-bonsai-ttft.patch
# Gated delta net prefill: next-token loads issued ahead, column-batched warp reductions with the same
# summation tree, 8 columns per warp. Bit-identical, ~1.9x faster per 1024-token ubatch on the 3060.
./llama-cpp-bonsai-gdn-prefill.patch
# Decode: MTP catch-up rows of a verify batch ride in the next draft batch (one MTP pass per round
# instead of two, rows computed identically), and the conv-state concat takes the transposing kernel.
# Greedy and sampled output identical; ~+6% decode on the 3060.
./llama-cpp-bonsai-mtp-decode.patch
# Flash attention reads a q4_0/q8_0 cache in place on the tensor cores (professorpalmer's
# bonsai-ada-surgery 0009 + 0029), also under batch-invariant mode: a q8_0 cache decodes as fast as F16
# at depth in about half the memory.
./llama-cpp-bonsai-q8-kv.patch
# PTQ1_0 mat-vec for 3 to 8 columns on the tensor cores (mma.m16n8k32 over decoded trits), so a
# speculative verify costs about one column instead of ~1.8x at 4. Same bits as the dp4a kernel.
./llama-cpp-bonsai-ptq1-tc-verify.patch
# The MTP graph publishes h only for output rows, as a masked draft context reads it back. With the
# catch-up rows riding in the draft batch, the second draft step was conditioned on the first catch-up
# row's h: n-max 2 acceptance 56% -> 73%.
./llama-cpp-bonsai-mtp-chain-h.patch
# --cache-disk DIR / --cache-disk-size MiB: prompts the RAM prompt cache evicts, and everything it and the
# slots hold at shutdown, go to files that a later prefix match reads back (an 8K-token conversation
# with its DeltaNet checkpoints is ~1 GB and restores in ~0.6 s instead of a ~14 s prefill).
./llama-cpp-bonsai-prompt-cache-disk.patch
# Prefill fusions, bit-identical: ADD+RMS_NORM+MUL on every NVIDIA GPU (PrismML-Eng/llama.cpp#209 by
# @cklxx, from professorpalmer's 0040), and the SWIGLU feeding a Hadamard-rotated matmul computed inside
# the transform's load. ~+2% prefill.
./llama-cpp-bonsai-prefill-fusions.patch
# PTQ1_0 MMQ: activations double-buffered with cp.async and the next weight tile fetched into registers
# during compute. Bit-identical, +3.3% on the matmuls.
./llama-cpp-bonsai-mmq-pipeline.patch
# PTQ1_0 MMQ quantizes activations with one scale per 128 values and accumulates a whole weight block in
# int32, so the float epilogue runs once per block instead of per 32: +15-18% on the matmuls, ~+11%
# prefill. Not bit-identical: PPL +0.12%, top token unchanged at 99.42% of positions. Decode and
# speculative verify (mat-vec) are untouched.
./llama-cpp-bonsai-mmq-y128.patch
# The Hadamard transform in front of a PTQ1_0 matmul that takes the MMQ path also writes the per-128 q8 tiles
# MMQ reads, so the matmul skips its own quantize pass (and the f32 copy when nothing else reads it).
# Bit-identical; +3% prefill.
./llama-cpp-bonsai-fwht-q8.patch
# PTQ1_0 MMQ main loop staged one 128-value block at a time: one barrier per block, x and y copied with
# cp.async, trits decoded between the MMAs, B fragments via ldmatrix. Bit-identical; +10% prefill.
# GGML_CUDA_NO_PTQ1_0_STAGED=1 restores the old loop.
./llama-cpp-bonsai-mmq-staged.patch
# Prompt batches over a q8_0 cache tile the 6 query heads of a KV head in pairs instead of padding them
# to 8, stage raw q8_0 rows with cp.async and dequantize exactly in half2. Bit-identical; prefill +8% at
# 16K, +14% at 32K, +21% at 64K. Decode and verify keep the 8-wide tile.
./llama-cpp-bonsai-fattn-q8-prefill.patch
# GET /cache (slots and both prompt-cache tiers) and POST /cache/clear?scope=slots|ram|disk|all, run on the
# server loop like the slot actions; a slot that is generating is left alone.
./llama-cpp-bonsai-cache-api.patch
# Prompt-cache retention: per-entry class (skip, ephemeral, default, keep, pin) from the request's "cache" field,
# learned per prompt family from reuse, or an optional classifier; TTLs, eviction by class and reuse, pins.
./llama-cpp-bonsai-cache-policy.patch
# A prompt family is the system block (the user turn carries each call's query or page), or without one
# the opening of the first user turn, so a service's one-shot calls share a family and learn "ephemeral".
./llama-cpp-bonsai-cache-family.patch
# Suffix cache reuse: after an edit in the middle of a cached prompt (compaction, a trimmed tool output, a system
# prompt change), the attention KV of the text after the edit is moved to its new positions (q8_0 keys
# re-rotated) and the DeltaNet layers continue from the pre-edit state, so only the new text is prefilled.
# Approximate (KL ~0.017 nats/token after a compaction); per request "scr", server default --scr.
./llama-cpp-bonsai-scr.patch
# Image conversations in the NVMe tier (.kvc v3: the media chunk map next to the tokens, keyed by image
# sha256), an image-embedding cache (256 MiB RAM + 4 GiB under <cache-disk>/embd), and a fix for SCR
# consuming RAM entries that only shared a few chat-template tokens with a new conversation.
./llama-cpp-bonsai-imgcache.patch
# Request telemetry: one record per request (cache/SCR/prefill split, TTFT, t/s, draft, images) plus minute
# samples, in a ring and capped JSONL under <cache-disk>/telemetry (on by default with a disk tier,
# --no-telemetry turns it off); GET /stats and /stats/requests, shown on the status page.
./llama-cpp-bonsai-telemetry.patch
# PTQ1_0 tensor-core verify mat-vec (MTP verify, 2-8 columns): the 32 residue slots fold in registers in
# warp_reduce_sum's pairing and only log2(warps) steps go through smem, with CTA geometry per shape.
# Bit-identical to the 1-column kernel; MTP decode +8.7% (two slots verifying +15%).
./llama-cpp-bonsai-decode-verify.patch
# Two-slot fairness: while a prompt is pending beside another active request, each server step is either a
# decode step or one prompt chunk (cut at the solo ubatch boundaries), picked by owed GPU time, so chat keeps
# ~22 t/s during a long prefill instead of ~1. No contention runs the old path unchanged.
# --fair-chunk (default n_ubatch, 0 = off), --fair-decode-weight (default 1).
./llama-cpp-bonsai-fairness.patch
# The DeltaNet layers read each seq's recurrent state straight from its cache row instead of gathering
# it first (only while every seq's row is its own cell; otherwise the old gather). Bit-identical; MTP
# decode +1.6-2.3%. GGML_GDN_STATE_GATHER=1 restores the gather.
./llama-cpp-bonsai-gdn-rows.patch
# One CUDA kernel for the DeltaNet conv path (concat -> ssm_conv -> silu -> conv-state writes): the concat
# is never materialized; same arithmetic. Bit-identical; MTP decode +1.5%.
# GGML_CUDA_SSM_CONV_CONCAT_FUSION=0 turns it off.
./llama-cpp-bonsai-gdn-conv.patch
# q8_0 decode/verify flash attention (the 4x8 tile) stages raw K/V with cp.async and starts the next
# tile's K and V a phase early. Bit-identical; decode +2.9% at 64K depth.
./llama-cpp-bonsai-fattn-verify-lead.patch
# MTP drops the rejected rows of a verification from its deferred queue at accept. Without it a fair
# prompt step of the other slot flushed them into the draft context, and every later draft of the
# decoding slot failed (draft positions ahead of the target).
./llama-cpp-bonsai-mtp-fair-deferred.patch
# Relax prefix cache and SCR thresholds for short-prefix one-shot services (web-extract, openwebui):
# - SCR reuse floor 512 -> 128 tokens
# - keep warm slot state on failed cache load if common prefix > 0 instead of wiping slot
# - allow matching prefix cache entries if lcp >= 32 tokens even when f_keep < 0.25
# - route to warm slot on system prompt match (lcp >= 32)
./llama-cpp-bonsai-cache-tuning.patch
# Flush MTP rows from short edited gaps before an SCR run disables the splice guard.
./llama-cpp-bonsai-scr-mtp-flush.patch
# Keep up to 64 matching spans when a summarizer flattens a long transcript.
./llama-cpp-bonsai-scr-spans.patch
# Pin a system prefix once three requests repeat it exactly, apart from the conversations after it.
./llama-cpp-bonsai-prefix-learn.patch
# Reprocess once from a system prefix that cached turns already cover, so agents learn it too.
./llama-cpp-bonsai-prefix-replay.patch
# Learn the start that distinct prompts share, so per-session text late in a system block does not stop it.
./llama-cpp-bonsai-prefix-shared.patch
# Keep a learned prefix's recurrent state as a splice restore point, learned during splices too.
./llama-cpp-bonsai-prefix-restore.patch
# Pin/keep buttons and family verdicts leave learned prefixes learned; never restore an empty state.
./llama-cpp-bonsai-prefix-entry.patch
# Learning outlasts restarts, pins after two prompts, and learned prefixes stay in RAM.
./llama-cpp-bonsai-prefix-persist.patch
# Conversations are stored against their learned prefix: only the KV rows after it and their own state.
./llama-cpp-bonsai-prefix-share.patch
];
nativeBuildInputs = [
cmake
ninja
pkg-config
cudaPackages.cuda_nvcc
autoAddDriverRunpath
];
buildInputs = with cudaPackages; [
cuda_cccl
cuda_cudart
libcublas
];
cmakeFlags = [
(lib.cmakeBool "GGML_NATIVE" false)
(lib.cmakeBool "GGML_CUDA" true)
(lib.cmakeFeature "CMAKE_CUDA_ARCHITECTURES" cudaArchitectures)
(lib.cmakeBool "BUILD_SHARED_LIBS" true)
(lib.cmakeBool "LLAMA_BUILD_EXAMPLES" false)
(lib.cmakeBool "LLAMA_BUILD_SERVER" true)
(lib.cmakeBool "LLAMA_BUILD_TESTS" false)
(lib.cmakeBool "LLAMA_BUILD_UI" false)
(lib.cmakeBool "LLAMA_CURL" false)
(lib.cmakeBool "LLAMA_OPENSSL" false)
(lib.cmakeFeature "LLAMA_BUILD_NUMBER" "10743")
(lib.cmakeFeature "LLAMA_BUILD_COMMIT" (builtins.substring 0 7 rev))
];
meta = {
description = "llama.cpp with the Ternary Bonsai 2 PTQ1_0 kernels and MTP, CUDA only";
homepage = "https://github.com/sudoingX/llama.cpp/tree/bonsai2";
license = lib.licenses.mit;
mainProgram = "llama-server";
platforms = [ "x86_64-linux" ];
};
})