Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
48 commits
Select commit Hold shift + click to select a range
64a6d4f
Optimized ARM NEON q1_0 dot (#33)
pl752 Jun 8, 2026
720e06b
Q1_0 repack kernels for Arm NEON+DP (#34)
pl752 Jun 8, 2026
ac71e9f
Add release-prism workflow
khosravipasha Apr 19, 2026
984bf97
Add Q2_0 quantization: type definition and CPU backend
khosravipasha Apr 18, 2026
f51b5aa
Add Q2_0 Metal backend
khosravipasha Apr 18, 2026
9c0edea
cuda: Q2_0
khosravipasha Apr 18, 2026
5f33471
release-prism: install spirv-headers for ubuntu-arm64 vulkan build
khosravipasha Apr 19, 2026
34dc581
vulkan: Q2_0
Vort3xed May 7, 2026
98e900b
fix: fp32_to_fp16 saturates finite overflow to inf
Vort3xed May 7, 2026
0e0353d
ggml-cpu: AVX-512-VNNI dot-products for Q1_0/Q2_0 (#37)
bri-prism Jun 4, 2026
4d88cd4
cpu: move Q2_0 AVX-512-VNNI dot into arch/x86, restore generic scalar
bri-prism Jun 4, 2026
21eff18
cuda: opt-in Hopper wgmma path for Q1_0 prefill (sm_90a)
bri-prism Jun 10, 2026
6f25c34
cuda: extend Hopper wgmma path to Q2_0
bri-prism Jun 10, 2026
02860ff
cuda: hopper path review follow-up — multi-GPU correctness + hardening
bri-prism Jun 10, 2026
fdb74a0
cuda: clang-format the hopper-q1 path
bri-prism Jun 11, 2026
3ff0bfc
cuda: reject Blackwell in the hopper wgmma gate (#47)
bri-prism Jun 12, 2026
6700b53
ci(release): resolve OpenMP redist DLL by glob instead of pinned MSVC…
khosravipasha Jul 6, 2026
0ad1dab
ci(release): disable LLAMA_BUILD_APP in ios-xcode-build
khosravipasha Jul 7, 2026
afc74b7
kv-cache: optional per-channel K-cache mean-centering for Q4_0 (#51)
bri-prism Jul 9, 2026
a18e55e
kv-mean-center: add make-calib-corpus.sh self-generated corpus helper
bri-prism Jul 10, 2026
f28050e
make-calib-corpus: address review feedback
bri-prism Jul 10, 2026
3277aaa
Merge pull request #52 from PrismML-Eng/feat/kv-mean-center-selfgen-c…
bri-prism Jul 10, 2026
a5527fc
kv-cache: mean-centering on hybrid models + calibration-basis guard (…
bri-prism Jul 10, 2026
422901c
tests: add Q2_0 error thresholds to test-quantize-fns
bri-prism Jul 12, 2026
ceb0790
Merge pull request #56 from PrismML-Eng/fix/test-quantize-fns-q2_0-th…
bri-prism Jul 12, 2026
ba62b70
speculative: dspark block-diffusion drafter + CUDA Markov resample; Q…
bri-prism Jul 13, 2026
b28d513
dspark: add log-SNR embed support to the GGUF converter (#57)
bri-prism Jul 13, 2026
3560f10
speculative: Metal DSpark Markov resample + quantized markov heads (#59)
bri-prism Jul 13, 2026
80570cb
metal: fix M5 device creation + add Q2_0 multi-column mul_mv kernels …
bri-prism Jul 13, 2026
46755d6
metal: GDN rows-indexed state read + snapshot write-fold (ring decode…
bri-prism Jul 14, 2026
157e758
metal gdn: review follow-ups that missed #61 (CPU workspace overflow,…
bri-prism Jul 14, 2026
972086d
dspark: independent unmasked capture path, avoid full-vocab lm_head o…
bri-prism Jul 14, 2026
da9c580
dspark: address review feedback on unmasked capture path
bri-prism Jul 14, 2026
5d2aa86
Merge pull request #63 from PrismML-Eng/fix/dspark-unmasked-capture
bri-prism Jul 14, 2026
887f007
metal: default-on nr1=2 multi-column Q2_0 matvec (spec-decode verify)…
bri-prism Jul 14, 2026
c024aa2
server, speculative-simple: wire dspark tap capture (#67)
khosravipasha Jul 14, 2026
c0616d5
clip: fix wrong-type format arguments flagged by CodeQL (#68)
bri-prism Jul 14, 2026
4b2f05a
test-arg-parser: skip download tests when the network is unreachable …
bri-prism Jul 14, 2026
62061f9
ci(release): build examples so llama-speculative-simple ships in preb…
khosravipasha Jul 14, 2026
e373b73
test: fix -Werror=format in test-arg-parser download test (#80)
bri-prism Jul 17, 2026
41e362d
cuda: speed up Q1_0 extraction with byte permutes (#73)
bri-prism Jul 17, 2026
38c66ad
ci(release): full self-contained Windows Vulkan/HIP bundles + README …
khosravipasha Jul 17, 2026
79697f2
perf(ggml-cpu): enable Q2_0 fast path on AVX-VNNI CPUs (#72)
thtro Jul 17, 2026
9fcaed7
ggml-cpu: x86 AVX512-VNNI repack GEMV/GEMM for Q1_0 and Q2_0 (#86)
bri-prism Jul 18, 2026
7529fda
force qwen35 to treat IGPU like a GPU so DSpark does not fall back to…
thadreber-web Jul 20, 2026
d2ad88a
Fix legacy SSE transport detection and transport selection
joydolma Jul 21, 2026
227e15c
docs(ui): update createTransport documentation
joydolma Jul 23, 2026
04cb388
Respect configured MCP transport and add regression tests
joydolma Jul 23, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
752 changes: 752 additions & 0 deletions .github/workflows/release-prism.yml

Large diffs are not rendered by default.

12 changes: 12 additions & 0 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -118,6 +118,18 @@ option(LLAMA_TESTS_INSTALL "llama: install tests" ON)
# 3rd party libs
option(LLAMA_OPENSSL "llama: use openssl to support HTTPS" ON)
option(LLAMA_LLGUIDANCE "llama-common: include LLGuidance library for structured output in common utils" OFF)
option(LLAMA_DSPARK_MARKOV_BLAS "llama-common: use BLAS for dspark Markov resample" OFF)

# Device-side dspark Markov resample. Compiled by default whenever the CUDA
# backend is built and -- when the drafter actually carries a Markov head -- it
# is the DEFAULT resample path at runtime. Set LLAMA_DSPARK_MARKOV_CUDA=0 in the
# environment to fall back to the BLAS/scalar host paths.
if (GGML_CUDA)
set(LLAMA_DSPARK_MARKOV_CUDA_DEFAULT ON)
else()
set(LLAMA_DSPARK_MARKOV_CUDA_DEFAULT OFF)
endif()
option(LLAMA_DSPARK_MARKOV_CUDA "llama-common: use CUDA for dspark Markov resample" ${LLAMA_DSPARK_MARKOV_CUDA_DEFAULT})


# Required for relocatable CMake package
Expand Down
17 changes: 17 additions & 0 deletions README.md
Original file line number Diff line number Diff line change
@@ -1,5 +1,22 @@
# llama.cpp

> [!IMPORTANT]
> **This is the PrismML fork of llama.cpp.** It adds the `Q2_0` 2-bit quantization used by the [Bonsai](https://huggingface.co/collections/prism-ml/bonsai) models.
>
> **New here? Start with the [Bonsai-demo](https://github.com/PrismML-Eng/Bonsai-demo) repo.** It downloads the right models and the correct prebuilt binaries for your hardware/backend automatically.
>
> Ternary (`Q2_0`) support is migrating into mainline llama.cpp backend-by-backend, so which build + model file to use depends on where you run:
>
> - `*-Q2_0.gguf` (group size 128): the format **this fork** uses. Run it with this fork's builds / [releases](https://github.com/PrismML-Eng/llama.cpp/releases). Does not load on mainline llama.cpp.
> - `*-Q2_0_g64.gguf` (group size 64): the **official mainline** llama.cpp format (currently CPU and Metal). Use a recent `ggml-org/llama.cpp` build for these, not this fork.
> - `*-PQ2_0.gguf`: planned future fork format, **not supported anywhere yet**.
>
> Use a complete matching build. Do NOT drop this fork's `ggml-*` libraries into a stock llama.cpp build (ABI/format mismatch, models fail to load).
>
> **For the latest backend-by-backend migration status, see [Upstream Status for Ternary](https://github.com/PrismML-Eng/Bonsai-demo#upstream-status-for-ternary) in the Bonsai-demo README.**

---

![llama](https://user-images.githubusercontent.com/1991296/230134379-7181e485-c521-4d23-a0d6-f7b3b61ba524.png)

[![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT)
Expand Down
58 changes: 58 additions & 0 deletions common/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -83,6 +83,8 @@ add_library(${TARGET}
json-partial.cpp
json-partial.h
json-schema-to-grammar.cpp
kv-mean-center.cpp
kv-mean-center.h
llguidance.cpp
log.cpp
log.h
Expand Down Expand Up @@ -171,4 +173,60 @@ if (LLAMA_LLGUIDANCE)
endif()
endif()

if (LLAMA_DSPARK_MARKOV_BLAS)
find_package(BLAS REQUIRED)

if ("${BLAS_INCLUDE_DIRS}" STREQUAL "")
find_path(LLAMA_DSPARK_MARKOV_CBLAS_INCLUDE_DIR
NAMES cblas.h
HINTS
/usr/include
/usr/local/include
/usr/include/openblas
/opt/homebrew/opt/openblas/include
/usr/local/opt/openblas/include
/usr/include/x86_64-linux-gnu/openblas/include
)
set(BLAS_INCLUDE_DIRS ${LLAMA_DSPARK_MARKOV_CBLAS_INCLUDE_DIR})
endif()

if ("${BLAS_INCLUDE_DIRS}" STREQUAL "")
message(FATAL_ERROR "LLAMA_DSPARK_MARKOV_BLAS was enabled but cblas.h could not be found")
endif()

target_compile_definitions(${TARGET} PRIVATE LLAMA_DSPARK_MARKOV_BLAS)
target_include_directories(${TARGET} SYSTEM PRIVATE ${BLAS_INCLUDE_DIRS})
target_link_libraries(${TARGET} PUBLIC ${BLAS_LIBRARIES})
endif()

if (LLAMA_DSPARK_MARKOV_CUDA)
# The CUDA language and CMAKE_CUDA_ARCHITECTURES are already established by
# ggml's CUDA backend build; reuse them for the self-contained resample TU.
enable_language(CUDA)
find_package(CUDAToolkit REQUIRED)

target_sources(${TARGET} PRIVATE dspark-markov.cu)
target_compile_definitions(${TARGET} PRIVATE LLAMA_DSPARK_MARKOV_CUDA)
target_link_libraries(${TARGET} PRIVATE CUDA::cudart)

if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
# CMAKE_CUDA_ARCHITECTURES is resolved inside ggml's CUDA backend
# subdirectory (auto-detect / GGML_CUDA_ARCHITECTURES) with a plain set(),
# so it does not propagate up to this scope. Pinning "80" would build the
# resample TU for SM80 only -- forcing a PTX JIT on Hopper/H100 (sm_90) and
# failing outright on the pre-Ampere GPUs the rest of ggml still supports.
# Inherit the exact architecture list the ggml CUDA backend was built for.
if (TARGET ggml-cuda)
get_target_property(_dspark_markov_cuda_archs ggml-cuda CUDA_ARCHITECTURES)
endif()
if (_dspark_markov_cuda_archs)
set_target_properties(${TARGET} PROPERTIES CUDA_ARCHITECTURES "${_dspark_markov_cuda_archs}")
else()
# ggml-cuda target/property unavailable: detect the build machine's GPU
# (e.g. sm_90 on H100) rather than defaulting to SM80 only.
set_target_properties(${TARGET} PROPERTIES CUDA_ARCHITECTURES "native")
endif()
endif()
endif()

target_link_libraries(${TARGET} PUBLIC llama Threads::Threads)
13 changes: 11 additions & 2 deletions common/arg.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1379,7 +1379,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
[](common_params & params, int value) {
params.n_chunks = value;
}
).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_PERPLEXITY, LLAMA_EXAMPLE_RETRIEVAL}));
).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_PERPLEXITY, LLAMA_EXAMPLE_RETRIEVAL, LLAMA_EXAMPLE_KV_MEAN_CENTER}));
add_opt(common_arg({ "-fa", "--flash-attn" }, "[on|off|auto]",
string_format("set Flash Attention use ('on', 'off', or 'auto', default: '%s')",
llama_flash_attn_type_name(params.flash_attn_type)),
Expand Down Expand Up @@ -2074,6 +2074,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.cache_type_v = kv_cache_type_from_str(value);
}
).set_env("LLAMA_ARG_CACHE_TYPE_V"));
add_opt(common_arg(
{"--kv-mean-center"}, "FNAME",
"path to a K-cache mean-centering bias file (GGUF), generated with tools/kv-mean-center\n"
"subtracts a fixed per-(head,channel) bias from K before it is quantized into the cache;\n"
"requires --cache-type-k q4_0 (see docs/kv-mean-center.md)",
[](common_params & params, const std::string & value) {
params.kv_mean_center_path = value;
}
).set_env("LLAMA_ARG_KV_MEAN_CENTER"));
add_opt(common_arg(
{"--hellaswag"},
"compute HellaSwag score over random tasks from datafile supplied with -f",
Expand Down Expand Up @@ -2706,7 +2715,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.out_file = value;
}
).set_examples({LLAMA_EXAMPLE_IMATRIX, LLAMA_EXAMPLE_CVECTOR_GENERATOR, LLAMA_EXAMPLE_EXPORT_LORA, LLAMA_EXAMPLE_TTS, LLAMA_EXAMPLE_FINETUNE,
LLAMA_EXAMPLE_RESULTS, LLAMA_EXAMPLE_EXPORT_GRAPH_OPS}));
LLAMA_EXAMPLE_RESULTS, LLAMA_EXAMPLE_EXPORT_GRAPH_OPS, LLAMA_EXAMPLE_KV_MEAN_CENTER}));
add_opt(common_arg(
{"-ofreq", "--output-frequency"}, "N",
string_format("output the imatrix every N iterations (default: %d)", params.n_out_freq),
Expand Down
4 changes: 4 additions & 0 deletions common/common.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1589,6 +1589,10 @@ struct llama_context_params common_context_params_to_llama(const common_params &
cparams.type_k = params.cache_type_k;
cparams.type_v = params.cache_type_v;

// note: params (and therefore params.kv_mean_center_path) is kept alive by the caller for
// at least as long as it takes to call llama_init_from_model() with the returned cparams
cparams.path_kv_mean_center = params.kv_mean_center_path.empty() ? nullptr : params.kv_mean_center_path.c_str();

return cparams;
}

Expand Down
20 changes: 19 additions & 1 deletion common/common.h
Original file line number Diff line number Diff line change
Expand Up @@ -96,6 +96,7 @@ enum llama_example {
LLAMA_EXAMPLE_FIT_PARAMS,
LLAMA_EXAMPLE_RESULTS,
LLAMA_EXAMPLE_EXPORT_GRAPH_OPS,
LLAMA_EXAMPLE_KV_MEAN_CENTER,

LLAMA_EXAMPLE_COUNT,
};
Expand Down Expand Up @@ -161,6 +162,7 @@ enum common_speculative_type {
COMMON_SPECULATIVE_TYPE_DRAFT_SIMPLE, // standalone draft model speculative decoding
COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3, // Eagle3 speculative decoding
COMMON_SPECULATIVE_TYPE_DRAFT_MTP, // Multi-token prediction
COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK, // dspark: EAGLE-style block-diffusion drafter
COMMON_SPECULATIVE_TYPE_NGRAM_SIMPLE, // simple self-speculative decoding based on n-grams
COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K, // self-speculative decoding with n-gram keys only
COMMON_SPECULATIVE_TYPE_NGRAM_MAP_K4V, // self-speculative decoding with n-gram keys and 4 m-gram values
Expand Down Expand Up @@ -362,8 +364,20 @@ struct common_params_speculative {
}

uint32_t need_n_rs_seq() const {
// Both MTP and dspark verify a whole draft block against the target in
// one llama_decode(), then crop the target's cache back to the accepted
// length with a PARTIAL llama_memory_seq_rm(). On a hybrid GDN/attention
// target (e.g. QWEN35/QWEN35MOE) that partial removal only succeeds if
// the recurrent-state rollback ring (n_rs_seq) was sized up front --
// see llama_memory_recurrent::seq_rm()'s "partial rollback via
// per-token snapshot index" path and llm_arch_supports_rs_rollback().
// Omitting a block-verify draft type here silently leaves n_rs_seq=0,
// so the post-verify crop on ctx_tgt no-ops instead of failing loudly
// (llama_memory_hybrid::seq_rm short-circuits to `return false` without
// mutating either sub-cache) -- the target's GDN state then keeps
// absorbing every future round's rejected draft tail.
bool needs_rs_seq = std::any_of(types.begin(), types.end(), [&](auto t) {
return t == COMMON_SPECULATIVE_TYPE_DRAFT_MTP;
return t == COMMON_SPECULATIVE_TYPE_DRAFT_MTP || t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK;
});

return needs_rs_seq ? draft.n_max : 0u;
Expand Down Expand Up @@ -565,6 +579,10 @@ struct common_params {
ggml_type cache_type_k = GGML_TYPE_F16; // KV cache data type for the K
ggml_type cache_type_v = GGML_TYPE_F16; // KV cache data type for the V

// path to a K-cache mean-centering bias file (GGUF), or empty to disable.
// only takes effect when cache_type_k == GGML_TYPE_Q4_0; see docs/kv-mean-center.md
std::string kv_mean_center_path = "";

common_conversation_mode conversation_mode = COMMON_CONVERSATION_MODE_AUTO;

// multimodal models (see tools/mtmd)
Expand Down
Loading