* wip: port MTP architecture
Ports the Multi-Token Prediction (MTP) architecture to the older `llama.cpp` codebase used by `ikllama`.
Changes include:
- Updating `llama_batch` to support `mtp_params`.
- Modifying `llama_decode_internal` (and `encode`) to handle MTP operations (Warmup, Update, Draft).
- Adding public APIs for MTP state management (`llama_set_draft_input_hidden_state`).
- Adapting the embedding extraction logic to skip MTP update passes.
* Refactors `server_slot` to support generic speculative decoding (MTP or Draft Model).
* core: enable hybrid outputs (logits + embeddings) for MTP support
* fix(mtp): correct KV-cache slot finding for updates
* fix(mtp): persist hidden states to prevent context corruption during drafting
* refactor(mtp): clean unused code
* fix(mtp): update server to new functions name
* fix(mtp): fix graph and save hidden state
* mtp: refactor integration, context params and kv cache search
* mtp: fix hidden state extraction and speculative acceptance flow
* server: fix MTP warmup for long prompts and reset token buffer
* llama: refactor MTP operation state to context parameters
* server: fix n_past calculation in MTP acceptance
* llama: fix mtp enable flags
* speculative: refactor MTP to use common_speculative interface
* context: remove unused signatures
* clip: fix deprecated enum-enum conversion warning
* common: fix format string crash in help message
* context: fix mtp activation logic
* llamat: always use the extracted embedding
* llama: get all embeddings to kv cache
* llama: revert logit to not run mtp for not supported arch
* llama: allocate all the n_outputs for MTP
* wip
* server-context: get only the last embedding for hidden state
* ggml-backend: fix array of bounds in debug build
* server-context: run mt kv update to each prompt batch
* revert segmentation fault fixes
* glm-mtp(feat): optimize graph embedding and recursive drafting
* Update nix flake: sync with upstream, fix for newer nixpkgs
- Sync package.nix with upstream: add rocmGpuTargets/useRpc params,
fix env optionals→optionalAttrs, remove unused inputs
- Rewrite devshells.nix to use mkShell directly instead of passthru
- Fix CUDA license check in nixpkgs-instances.nix for list-type licenses
- Update flake inputs (nixpkgs, flake-parts) to latest
- Update references from ggerganov/llama.cpp to ikawrakow/ik_llama.cpp
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
* Apply nix review suggestions from baileylu121 (#1371)
- Use legacyPackages instead of import to avoid extra nixpkgs instances
- Use inherit (pkgs) stdenv idiom
- Remove unnecessary lib.pipe wrapper in devshells.nix
- Simplify license normalization with lib.toList in nixpkgs-instances.nix
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
* Update ik_llama-cuda.Containerfile
- DGGML_NATIVE=ON, as most users are building images locally
- Latest llama-swap version
* Update ik_llama-cpu.Containerfile
- DGGML_NATIVE=ON, as most users are building images locally
- Latest llama-swap version
* Update README.md
- DGGML_NATIVE=ON is now default
* Update ik_llama-cuda-swap.config.yaml
- Adapt to the latest version
* Update ik_llama-cpu-swap.config.yaml
- Adapt to the latest version
* Update ik_llama-cuda-swap.config.yaml
- Add model with aliases
* Update ik_llama-cpu-swap.config.yaml
- Add model with aliases
Add a separate HAVE_VNNI256 code path using _mm256_dpwssd_epi32 and
_mm256_dpbusd_epi32 for the Q3_K R4 kernel. The existing HAVE_FANCY_SIMD
(AVX-512 VNNI) path is preserved unchanged.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Add a separate HAVE_VNNI256 code path using _mm256_dpwssd_epi32 and
_mm256_dpbusd_epi32 for the Q6_K R4 kernel. The existing HAVE_FANCY_SIMD
(AVX-512 VNNI) path is preserved unchanged.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
Purpose:
Add --minilog flag to llama-sweep-bench that filters log output to show only essential GPU/layer distribution information while suppressing verbose model metadata and per-layer device assignment messages.
Changes:
- Add llama_selective_log_callback with blacklist approach (sweep-bench.cpp)
Blacklisted patterns (hidden):
- Per-layer device assignments ('Setting default device in layer')
- KV metadata dump header and entries
- Tensor type counts
- Model validation messages
- EOG/special token cache info
- Metadata printout (llm_load_print_meta, print_info)
- Layer sizes table
- Tensor loading info (llm_load_tensors)
- Separator lines
- Most common cases of incomplete/continuation lines are also hidden
All other log output is shown, including:
- GPU VRAM info
- Split/buffer distribution per device
- Graph split estimates
- Final benchmark table and timings
* fix: FA vec kernels for D=256 quantized KV cache
Three bugs prevented Flash Attention from working with quantized
(q8_0) KV cache on models with head dimension 256, such as
Qwen3-Coder-Next:
1. need_f16_K/V used dimension-based logic (Dk != 128) that forced
q8_0→f16 conversion at D=256. The kernel then read f16 data as
q8_0 blocks. Changed to type-based logic (type_K == GGML_TYPE_F16)
matching mainline.
2. quantize_q8_1_to_shared output pointers were not advanced between
loop iterations, so the second iteration overwrote the first half
of Q data in shared memory. Added i0 offset to output pointers.
3. Q_i32 register array in vec_f32 kernel was sized to 1 instead of 2
for D=256 due to an errant comparison (Dk >= expr evaluates to
boolean 1, not expr). Removed the comparison.
Also adds extern template declarations for (256, Q8_0, Q8_0) so the
kernel is compiled and linked.
Tested on Pascal (GTX 1080, CC 6.1) with Qwen3-Coder-Next (D=256,
n_head_kv=2) using --cache-type-k q8_0 --cache-type-v q8_0.
* fix: use dimension-based need_f16 for D=256 vec kernels
Use the sign trick with dpbusd instead of maddubs+madd+add,
replacing 3 AVX2 instructions with 1 fused VNNI instruction.
Removes dead HAVE_FANCY_SIMD code left over from the R16 split.
Co-authored-by: Adam Caldwell <accaldwell@users.noreply.github.com>
hksdpc255
Kawrakow
Samuel Oliveira Alves
Alec Koumjian
Adam Caldwell
Aliez Ren
mcm007
Nexes the Elder
firecoperana
gapeleon
Juk Armstrong
Max Homilius