mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-08-28 11:07:40 +02:00
Compare commits
2
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Xuan Son Nguyen
|
4e2f54f8f3 | ||
|
|
4a6ad487a6 |
@@ -714,10 +714,10 @@ jobs:
|
||||
with:
|
||||
key: release-windows-2025-vs2026-${{ matrix.arch }}-cpu
|
||||
|
||||
# note: builds only the ggml-hip backend - llama-server is injected from the
|
||||
# windows-cpu zip during the release "Merge artifacts" step
|
||||
# TODO: build only the ggml-hip backend like the other windows backend jobs
|
||||
# (windows-cuda, windows-sycl), then drop the ui-build dependency
|
||||
windows-rocm:
|
||||
needs: [check-release]
|
||||
needs: [check-release, ui-build]
|
||||
if: ${{ needs.check-release.outputs.should_release == 'true' }}
|
||||
|
||||
runs-on: windows-2022
|
||||
@@ -736,9 +736,11 @@ jobs:
|
||||
with:
|
||||
fetch-depth: 0
|
||||
|
||||
- name: Install Ninja
|
||||
run: |
|
||||
choco install ninja
|
||||
- name: Download UI build
|
||||
uses: actions/download-artifact@v7
|
||||
with:
|
||||
name: llama-ui.zip
|
||||
path: tools/ui/dist
|
||||
|
||||
- name: ccache
|
||||
uses: ggml-org/ccache-action@v1.2.21
|
||||
@@ -793,28 +795,33 @@ jobs:
|
||||
|
||||
- name: Build
|
||||
run: |
|
||||
cmake -S . -B build `
|
||||
-G "Ninja Multi-Config" `
|
||||
mkdir build
|
||||
cd build
|
||||
cmake .. `
|
||||
-G "Unix Makefiles" `
|
||||
-DCMAKE_PREFIX_PATH="${env:HIP_PATH}" `
|
||||
-DCMAKE_BUILD_TYPE=Release `
|
||||
-DGGML_BACKEND_DL=ON `
|
||||
-DGGML_NATIVE=OFF `
|
||||
-DGGML_CPU=OFF `
|
||||
-DGGML_CPU=ON `
|
||||
-DGGML_CPU_ALL_VARIANTS=ON `
|
||||
-DGGML_HIP=ON `
|
||||
-DCMAKE_C_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
|
||||
-DCMAKE_CXX_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang++.exe" `
|
||||
-DCMAKE_C_FLAGS="-Wno-error=incompatible-pointer-types" `
|
||||
-DCMAKE_HIP_COMPILER="${env:HIP_PATH}\lib\llvm\bin\clang.exe" `
|
||||
-DHIP_PATH="${env:HIP_PATH}" `
|
||||
-DGGML_HIP_ROCWMMA_FATTN=ON `
|
||||
-DAMDGPU_TARGETS="${{ matrix.gpu_targets }}"
|
||||
cmake --build build --config Release --parallel ${env:NUMBER_OF_PROCESSORS} --target ggml-hip
|
||||
cmake --build . --config Release --parallel ${env:NUMBER_OF_PROCESSORS}
|
||||
|
||||
- name: Verify HIP backend was built
|
||||
run: |
|
||||
$hipDll = Get-ChildItem -Path build\bin\Release -Filter "ggml-hip*.dll" -ErrorAction SilentlyContinue
|
||||
$hipDll = Get-ChildItem -Path build\bin -Filter "ggml-hip*.dll" -ErrorAction SilentlyContinue
|
||||
if (-not $hipDll) {
|
||||
Write-Host "##[error]ggml-hip*.dll was NOT produced. The HIP backend silently failed to build."
|
||||
Write-Host "Contents of build\bin\Release:"
|
||||
Get-ChildItem build\bin\Release | Format-Table -AutoSize
|
||||
Write-Host "Contents of build\bin:"
|
||||
Get-ChildItem build\bin | Format-Table -AutoSize
|
||||
exit 1
|
||||
}
|
||||
Write-Host "HIP backend artifact found:"
|
||||
@@ -829,40 +836,10 @@ jobs:
|
||||
$rocmVersionShort = ('${{ matrix.ROCM_VERSION }}'.Split('.')[0..1] -join '.')
|
||||
echo "ROCM_VERSION_SHORT=$rocmVersionShort" >> $env:GITHUB_ENV
|
||||
|
||||
- name: Bundle HIP runtime DLLs (amdhip64_7.dll, rocm_kpack.dll, amd_comgr.dll)
|
||||
run: |
|
||||
$ErrorActionPreference = "Stop"
|
||||
# See issue https://github.com/ggml-org/llama.cpp/issues/26929.
|
||||
# ggml-hip.dll loads amdhip64_7.dll at run time. The Adrenalin driver
|
||||
# ships an amdhip64_7.dll in System32, which the loader searches before PATH,
|
||||
# so a matching DLL from PATH cannot win. Copy amdhip64 next to the
|
||||
# binaries (exe directory is searched before System32) so the correct
|
||||
# runtime is used. rocm_kpack.dll is amdhip64_7's direct dependency, so
|
||||
# copy the matching version too. amd_comgr is copied as well to keep it
|
||||
# in sync with the bundled amdhip64, avoiding a version mismatch with a
|
||||
# amd_comgr from System32.
|
||||
# rocblas/hipblaslt kernels resolve fine via PATH and are not copied.
|
||||
$binPath = (rocm-sdk path --bin).Trim()
|
||||
if (-not $binPath) { throw "rocm-sdk path --bin returned empty" }
|
||||
write-host "ROCm bin path: $binPath"
|
||||
|
||||
$patterns = @("amdhip64_7.dll", "rocm_kpack.dll", "amd_comgr.dll")
|
||||
foreach ($pattern in $patterns) {
|
||||
$files = Get-ChildItem -Path $binPath -Filter $pattern -ErrorAction SilentlyContinue
|
||||
if (-not $files) { throw "no match for $pattern in $binPath" }
|
||||
foreach ($f in $files) {
|
||||
Copy-Item $f.FullName -Destination build\bin\Release -Force
|
||||
write-host " copied $($f.Name)"
|
||||
}
|
||||
}
|
||||
|
||||
- name: Pack artifacts
|
||||
run: |
|
||||
7z a -snl llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip `
|
||||
.\build\bin\Release\ggml-hip.dll `
|
||||
.\build\bin\Release\amdhip64_7.dll `
|
||||
.\build\bin\Release\rocm_kpack.dll `
|
||||
.\build\bin\Release\amd_comgr.dll
|
||||
cp "LICENSE" "build\bin\"
|
||||
7z a -snl llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip .\build\bin\*
|
||||
|
||||
- name: Upload artifacts
|
||||
uses: actions/upload-artifact@v6
|
||||
|
||||
@@ -1643,14 +1643,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
|
||||
}
|
||||
}
|
||||
).set_env("LLAMA_ARG_CTX_SIZE"));
|
||||
add_opt(common_arg(
|
||||
{ "--kv-unified-per-slot" }, "N",
|
||||
"context limit per parallel slot (default: unset, behavior unchanged).\n"
|
||||
"when set without -c/--ctx-size, the shared KV pool is sized to n_parallel*N",
|
||||
[](common_params & params, int value) {
|
||||
params.kv_unified_per_slot = value;
|
||||
}
|
||||
).set_env("LLAMA_ARG_KV_UNIFIED_PER_SLOT").set_examples({ LLAMA_EXAMPLE_SERVER }));
|
||||
add_opt(common_arg(
|
||||
{"-n", "--predict", "--n-predict"}, "N",
|
||||
string_format(
|
||||
|
||||
@@ -627,7 +627,6 @@ struct common_params {
|
||||
bool cache_prompt = true; // whether to enable prompt caching
|
||||
bool cache_idle_slots = true; // save and clear idle slots upon starting a new task
|
||||
int32_t n_ctx_checkpoints = 32; // max number of context checkpoints per slot
|
||||
int32_t kv_unified_per_slot = 0; // max context per parallel slot; 0 = unset
|
||||
int32_t checkpoint_min_step = 8192; // minimum spacing between context checkpoints
|
||||
int32_t cache_ram_mib = 8192; // -1 = no limit, 0 - disable, 1 = 1 MiB, etc.
|
||||
|
||||
|
||||
+1
-17
@@ -935,9 +935,6 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
|
||||
// dspark speculators
|
||||
bool sample_from_anchor = true;
|
||||
|
||||
// block-internal attention
|
||||
bool causal_attn = false;
|
||||
|
||||
const int32_t * target_layer_ids = nullptr; // model_dft's extract layer indices
|
||||
uint32_t target_layer_ids_n = 0;
|
||||
|
||||
@@ -975,25 +972,12 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
|
||||
if (llama_model_meta_val_str(model_dft, "dflash.sample_from_anchor", buf, sizeof(buf)) >= 0) {
|
||||
sample_from_anchor = std::strcmp(buf, "true") == 0;
|
||||
}
|
||||
if (llama_model_meta_val_str(model_dft, "dflash.attention.causal", buf, sizeof(buf)) >= 0) {
|
||||
causal_attn = std::strcmp(buf, "true") == 0;
|
||||
}
|
||||
}
|
||||
|
||||
selector_top_k = llama_model_dflash_selector_top_k(model_dft);
|
||||
is_dflash2 = selector_top_k > 0;
|
||||
mask_token_id = llama_vocab_mask(llama_model_get_vocab(model_dft));
|
||||
|
||||
if (is_dspark && this->params.p_min > 0.0f) {
|
||||
char buf[16] = {};
|
||||
const bool has_conf =
|
||||
llama_model_meta_val_str(model_dft, "dflash.has_confidence_head", buf, sizeof(buf)) < 0 ||
|
||||
std::strcmp(buf, "true") == 0;
|
||||
if (!has_conf) {
|
||||
throw std::runtime_error("DSpark draft has no confidence head: please set --spec-draft-p-min 0");
|
||||
}
|
||||
}
|
||||
|
||||
LOG_INF("%s: adding speculative implementation '%s'\n", __func__, common_speculative_type_to_str(type).c_str());
|
||||
LOG_INF("%s: - n_max=%d, n_min=%d, p_min=%.2f\n", __func__, this->params.n_max, this->params.n_min, this->params.p_min);
|
||||
LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u, sample_from_anchor=%s\n", __func__,
|
||||
@@ -1052,7 +1036,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
|
||||
|
||||
// DFlash2 reads its selector lattice from h_nextn and never consumes raw logits.
|
||||
llama_set_embeddings_nextn(ctx_dft, true, /*masked*/ !is_dflash2);
|
||||
llama_set_causal_attn(ctx_dft, causal_attn); // DFlash needs non-causal attention unless the model says otherwise
|
||||
llama_set_causal_attn(ctx_dft, false); // DFlash needs non-causal attention
|
||||
}
|
||||
|
||||
~common_speculative_impl_draft_dflash() override {
|
||||
|
||||
@@ -236,8 +236,6 @@ TEXT_MODEL_MAP: dict[str, str] = {
|
||||
"Qwen3_5ForConditionalGeneration": "qwen",
|
||||
"Qwen3_5MoeForCausalLM": "qwen",
|
||||
"Qwen3_5MoeForConditionalGeneration": "qwen",
|
||||
"Qwen4ExpForCausalLM": "qwen4exp",
|
||||
"Qwen4ExpForConditionalGeneration": "qwen4exp",
|
||||
"RND1": "qwen",
|
||||
"RWForCausalLM": "falcon",
|
||||
"RWKV6Qwen2ForCausalLM": "rwkv",
|
||||
@@ -335,7 +333,6 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
|
||||
"Qwen3VLMoeForConditionalGeneration": "qwen3vl",
|
||||
"Qwen3_5ForConditionalGeneration": "qwen3vl",
|
||||
"Qwen3_5MoeForConditionalGeneration": "qwen3vl",
|
||||
"Qwen4ExpForConditionalGeneration": "qwen4exp",
|
||||
"RADIOModel": "nemotron",
|
||||
"Sarashina2VisionForCausalLM": "sarashina2",
|
||||
"SmolVLMForConditionalGeneration": "smolvlm",
|
||||
|
||||
+2
-6
@@ -1006,16 +1006,12 @@ class ModelBase:
|
||||
else:
|
||||
raise ValueError(f"Unknown file type: {self.ftype.name}")
|
||||
|
||||
# a chunked tensor quantizes as one chunk at a time, while it is written
|
||||
quantize = data.quantize if isinstance(data, gguf.LazyChunkedTensor) else (
|
||||
lambda qtype, d=data: gguf.quants.quantize(d, qtype))
|
||||
|
||||
try:
|
||||
data = quantize(data_qtype)
|
||||
data = gguf.quants.quantize(data, data_qtype)
|
||||
except gguf.QuantError as e:
|
||||
logger.warning("%s, %s", e, "falling back to F16")
|
||||
data_qtype = gguf.GGMLQuantizationType.F16
|
||||
data = quantize(data_qtype)
|
||||
data = gguf.quants.quantize(data, data_qtype)
|
||||
|
||||
shape = gguf.quant_shape_from_byte_shape(data.shape, data_qtype) if data.dtype == np.uint8 else data.shape
|
||||
|
||||
|
||||
+3
-15
@@ -709,20 +709,14 @@ class DFlashModel(Qwen3Model):
|
||||
extract_layer_ids = [i + 1 for i in target_layer_ids]
|
||||
self.gguf_writer.add_target_layers(extract_layer_ids)
|
||||
|
||||
use_sliding_window = self.hparams.get("use_sliding_window", False) or dflash_config.get("use_swa", False)
|
||||
sliding_window = dflash_config.get("swa_window_size") or self.hparams.get("sliding_window")
|
||||
use_sliding_window = self.hparams.get("use_sliding_window", False)
|
||||
sliding_window = self.hparams.get("sliding_window")
|
||||
layer_types = self.hparams.get("layer_types")
|
||||
if use_sliding_window and sliding_window and layer_types:
|
||||
is_swa = [lt == "sliding_attention" for lt in layer_types]
|
||||
self.gguf_writer.add_sliding_window(sliding_window)
|
||||
self.gguf_writer.add_sliding_window_pattern(is_swa)
|
||||
|
||||
causal = self.hparams.get("is_causal")
|
||||
if causal is None:
|
||||
causal = dflash_config.get("causal")
|
||||
if causal is not None:
|
||||
self.gguf_writer.add_causal_attention(bool(causal))
|
||||
|
||||
# M-RoPE target: the draft ropes on the temporal dim only, so write
|
||||
# degenerate sections [n_rot/2, 0, 0, 0]
|
||||
if self._target_uses_mrope():
|
||||
@@ -743,8 +737,6 @@ class DFlashModel(Qwen3Model):
|
||||
name, gen = item
|
||||
if not name.startswith("model."):
|
||||
name = "model." + name
|
||||
if "sink" in name and not name.endswith(".weight"):
|
||||
name += ".weight"
|
||||
return super().filter_tensors((name, gen))
|
||||
|
||||
_ROPE_PERMUTE_SUFFIXES = (
|
||||
@@ -823,10 +815,6 @@ class DSparkModel(DFlashModel):
|
||||
super().set_gguf_parameters()
|
||||
self.gguf_writer.add_sample_from_anchor(self._sample_from_anchor)
|
||||
|
||||
# confidence head is optional: vanilla-markov exports ship without it
|
||||
has_conf = any("confidence_head.proj" in name for name in self.model_tensors)
|
||||
self.gguf_writer.add_has_confidence_head(has_conf)
|
||||
|
||||
@classmethod
|
||||
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
|
||||
if item[0] == "t2d": # not used at runtime
|
||||
@@ -845,7 +833,7 @@ class DSparkModel(DFlashModel):
|
||||
self._d2t = data_torch
|
||||
return
|
||||
|
||||
if self._n_vocab_draft == self.hparams["vocab_size"] and name.endswith("lm_head.weight"):
|
||||
if self._n_vocab_draft == self.hparams["vocab_size"] and name.endswith(("embed_tokens.weight", "lm_head.weight")):
|
||||
return
|
||||
|
||||
# interleaved-rope checkpoints (rope_is_neox_style = false) -> NeoX layout: per head, even dims first then odd
|
||||
|
||||
@@ -1,195 +0,0 @@
|
||||
from __future__ import annotations
|
||||
|
||||
from typing import Iterable, cast
|
||||
|
||||
import torch
|
||||
from torch import Tensor
|
||||
|
||||
import gguf
|
||||
import numpy as np
|
||||
|
||||
from .base import ModelBase
|
||||
from .qwen import _LinearAttentionVReorderBase, _Qwen35MRopeMixin
|
||||
from .qwen3vl import Qwen3VLVisionModel
|
||||
|
||||
|
||||
@ModelBase.register("Qwen4ExpForConditionalGeneration", "Qwen4ExpForCausalLM")
|
||||
@ModelBase.example("Qwen/Qwen3.8-Flash-Next")
|
||||
class Qwen4ExpTextModel(_Qwen35MRopeMixin, _LinearAttentionVReorderBase):
|
||||
"""Qwen3.8-Flash-Next.
|
||||
|
||||
Shares the Qwen3.5 gated delta net and interleaved mrope, and adds three things:
|
||||
hyper-connections in place of every layer norm, QSA sparse attention on the full
|
||||
attention layers, and PLE n-gram hash embeddings on a single layer.
|
||||
"""
|
||||
|
||||
model_arch = gguf.MODEL_ARCH.QWEN4EXP
|
||||
|
||||
# the MTP block is a separate draft head; vLLM drops it too
|
||||
supports_mtp_export = False
|
||||
no_mtp = True
|
||||
|
||||
def __init__(self, *args, **kwargs):
|
||||
super().__init__(*args, **kwargs)
|
||||
# only the shard names, so the table itself is never held
|
||||
self._ple_shards: dict[int, str] = {}
|
||||
self._ple_row_dim: int | None = None
|
||||
|
||||
def _read_hash_constants(self, suffix: str) -> list[int]:
|
||||
"""Read an int64 PLE constant straight from the checkpoint.
|
||||
|
||||
prepare_tensors() casts every non-float dtype to float32 before
|
||||
modify_tensors() sees it (base.py), which would silently round these
|
||||
45-bit multipliers. Reading the lazy tensor here bypasses that.
|
||||
"""
|
||||
for name, gen in self.model_tensors.items():
|
||||
if name.endswith(suffix):
|
||||
t = gen()
|
||||
if t.dtype != torch.int64:
|
||||
t = t.to(torch.int64)
|
||||
return [int(x) for x in t.tolist()]
|
||||
raise ValueError(f"PLE constant {suffix!r} missing from the checkpoint")
|
||||
|
||||
def set_gguf_parameters(self):
|
||||
super().set_gguf_parameters()
|
||||
hp = self.hparams
|
||||
|
||||
self.gguf_writer.add_hyper_connection_count(hp["hc_count"])
|
||||
self.gguf_writer.add_hyper_connection_low_rank(hp["hc_lowrank"])
|
||||
|
||||
n_layer = hp["num_hidden_layers"]
|
||||
self.gguf_writer.add_indexer_head_count(hp["indexer_n_heads"])
|
||||
self.gguf_writer.add_indexer_key_length(hp["indexer_head_dim"])
|
||||
self.gguf_writer.add_indexer_top_k(hp["indexer_budget"])
|
||||
ratio = hp["indexer_compress_ratio"]
|
||||
layer_types = hp["layer_types"]
|
||||
self.gguf_writer.add_attention_compress_ratios(
|
||||
[ratio if layer_types[i] == "full_attention" else 0 for i in range(n_layer)]
|
||||
)
|
||||
|
||||
# ple_layer_ids is 1-based in the HF config; empty means no n-gram table,
|
||||
# so emit no PLE keys rather than optional ones
|
||||
ple_layers = [i - 1 for i in hp["ple_layer_ids"]]
|
||||
if not ple_layers:
|
||||
return
|
||||
self.gguf_writer.add_ple_layers(ple_layers)
|
||||
self.gguf_writer.add_ple_ngram_size(hp["ngram_size"])
|
||||
self.gguf_writer.add_ple_heads_per_ngram(hp["heads_per_ngram"])
|
||||
self.gguf_writer.add_ple_conv_kernel(hp["ple_conv_kernel_size"])
|
||||
self.gguf_writer.add_ple_eos_token_id(self._eos_token_id())
|
||||
# an image is decoded as an embeddings-only batch, so the graph has no placeholder
|
||||
# ids to hash; carry the id and let it stand in for those positions
|
||||
_img = self._image_token_id()
|
||||
if _img is not None:
|
||||
self.gguf_writer.add_ple_image_token_id(int(_img))
|
||||
if self._ple_row_dim is not None:
|
||||
self.gguf_writer.add_embedding_length_per_layer_input(self._ple_row_dim)
|
||||
|
||||
self.gguf_writer.add_ple_layer_multipliers(
|
||||
self._read_hash_constants("ple_embedding.layer_multipliers"))
|
||||
self.gguf_writer.add_ple_head_offsets(
|
||||
self._read_hash_constants("ple_embedding.ngram_heads_offsets"))
|
||||
self.gguf_writer.add_ple_head_vocab_sizes(
|
||||
self._read_hash_constants("ple_embedding.ngram_heads_vocab_sizes"))
|
||||
|
||||
def _image_token_id(self) -> int | None:
|
||||
img = self.hparams.get("image_token_id")
|
||||
return None if img is None else int(img)
|
||||
|
||||
def _eos_token_id(self) -> int:
|
||||
eos = self.hparams.get("eos_token_id")
|
||||
if isinstance(eos, list):
|
||||
# the PLE hash resets n-grams on the primary EOS
|
||||
return int(eos[-1])
|
||||
if eos is None:
|
||||
raise ValueError("eos_token_id is required: the PLE hash resets its n-grams on it")
|
||||
return int(eos)
|
||||
|
||||
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
|
||||
# int64 hash constants must stay exact; 1-D tensors force F32, so use KV
|
||||
if name.endswith("ple_embedding.layer_multipliers"):
|
||||
self._ple_multipliers = [int(x) for x in data_torch.tolist()]
|
||||
return []
|
||||
if name.endswith("ple_embedding.ngram_heads_offsets"):
|
||||
self._ple_head_offsets = [int(x) for x in data_torch.tolist()]
|
||||
return []
|
||||
if name.endswith("ple_embedding.ngram_heads_vocab_sizes"):
|
||||
self._ple_head_vocab_sizes = [int(x) for x in data_torch.tolist()]
|
||||
return []
|
||||
|
||||
if ".ngram_embedding.shard_" in name:
|
||||
return self._place_ple_shard(data_torch, name)
|
||||
|
||||
# one projection feeds indexer q and k; split it, as minimax-m3 does
|
||||
if ".indexer.index_qk_proj.weight" in name:
|
||||
n_q = self.hparams["indexer_n_heads"] * self.hparams["indexer_head_dim"]
|
||||
q = data_torch[:n_q]
|
||||
k = data_torch[n_q:]
|
||||
return [
|
||||
(self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_Q_PROJ, bid, ".weight"), q),
|
||||
(self.format_tensor_name(gguf.MODEL_TENSOR.INDEXER_K_PROJ, bid, ".weight"), k),
|
||||
]
|
||||
|
||||
# Gemma zero-centred gammas the inherited norm.weight rule misses
|
||||
if name.endswith((".ple.norm_key.weight", ".ple.norm_query.weight", ".ple.norm_conv.weight",
|
||||
".indexer.q_layernorm.weight", ".indexer.k_layernorm.weight")):
|
||||
return [(self.map_tensor_name(name), data_torch + 1)]
|
||||
|
||||
if name.endswith(".ple.conv1d.weight"):
|
||||
return [(self.map_tensor_name(name), data_torch.squeeze())]
|
||||
|
||||
return super().modify_tensors(data_torch, name, bid)
|
||||
|
||||
# the shards concatenate into a tensor of well over 100 GB
|
||||
# use LazyChunkedTensor here, a single shard resident at a time
|
||||
def _place_ple_shard(self, data_torch: Tensor, name: str) -> Iterable[tuple[str, Tensor]]:
|
||||
|
||||
idx = int(name.rpartition(".shard_")[2].partition(".")[0])
|
||||
n_parts = self.hparams["split_ngram_parts"]
|
||||
|
||||
self._ple_shards[idx] = name
|
||||
self._ple_row_dim = int(data_torch.shape[-1])
|
||||
|
||||
if len(self._ple_shards) < n_parts:
|
||||
return []
|
||||
|
||||
# the checkpoint may yield the shards in any order, the row order is by index
|
||||
shards = [self._ple_shards[i] for i in sorted(self._ple_shards)]
|
||||
rows = 0
|
||||
for shard in shards:
|
||||
shape = self.model_tensors[shard]().shape
|
||||
if int(shape[-1]) != self._ple_row_dim:
|
||||
raise ValueError(
|
||||
f"PLE shard {shard} has row dim {int(shape[-1])}, expected {self._ple_row_dim}")
|
||||
rows += int(shape[0])
|
||||
|
||||
table = gguf.LazyChunkedTensor(
|
||||
[self._load_ple_shard(shard) for shard in shards],
|
||||
shape=(rows, self._ple_row_dim),
|
||||
dtype=np.float32,
|
||||
)
|
||||
gguf_name = gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.PER_LAYER_TOKEN_EMBD]
|
||||
return [(gguf_name + ".weight", cast(Tensor, table))]
|
||||
|
||||
def _load_ple_shard(self, name: str):
|
||||
def load() -> np.ndarray:
|
||||
from .base import LazyTorchTensor
|
||||
|
||||
# a fresh lazy tensor every call, or to_eager() memoizes every shard
|
||||
eager = LazyTorchTensor.to_eager(self.model_tensors[name]())
|
||||
return eager.to(torch.float32).contiguous().numpy()
|
||||
return load
|
||||
|
||||
def prepare_tensors(self):
|
||||
super().prepare_tensors()
|
||||
n_parts = self.hparams.get("split_ngram_parts", 0)
|
||||
if self._ple_shards and len(self._ple_shards) != n_parts:
|
||||
raise ValueError(
|
||||
f"got {len(self._ple_shards)} PLE embedding shards, expected {n_parts}"
|
||||
)
|
||||
|
||||
|
||||
@ModelBase.register("Qwen4ExpForConditionalGeneration")
|
||||
@ModelBase.example("Qwen/Qwen3.8-Flash-Next")
|
||||
class Qwen4ExpVisionModel(Qwen3VLVisionModel):
|
||||
"""The vision tower is an unmodified Qwen3-VL ViT."""
|
||||
@@ -4643,7 +4643,6 @@ static htp_op_code op_remap_to_htp(const ggml_tensor * t) {
|
||||
case GGML_OP_CLAMP: return HTP_OP_CLAMP;
|
||||
case GGML_OP_SQR: return HTP_OP_SQR;
|
||||
case GGML_OP_SQRT: return HTP_OP_SQRT;
|
||||
case GGML_OP_LOG: return HTP_OP_UNARY_LOG;
|
||||
case GGML_OP_SOFT_MAX: return HTP_OP_SOFTMAX;
|
||||
case GGML_OP_SSM_CONV: return HTP_OP_SSM_CONV;
|
||||
case GGML_OP_GATED_DELTA_NET: return HTP_OP_GATED_DELTA_NET;
|
||||
@@ -4667,7 +4666,6 @@ static htp_op_code op_remap_to_htp(const ggml_tensor * t) {
|
||||
case GGML_UNARY_OP_EXP: return HTP_OP_UNARY_EXP;
|
||||
case GGML_UNARY_OP_SOFTPLUS: return HTP_OP_UNARY_SOFTPLUS;
|
||||
case GGML_UNARY_OP_TANH: return HTP_OP_UNARY_TANH;
|
||||
case GGML_UNARY_OP_ABS: return HTP_OP_UNARY_ABS;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
@@ -5465,7 +5463,6 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
|
||||
|
||||
case GGML_OP_SQR:
|
||||
case GGML_OP_SQRT:
|
||||
case GGML_OP_LOG:
|
||||
supp = ggml_hexagon_supported_unary(sess, op);
|
||||
break;
|
||||
|
||||
@@ -5484,7 +5481,6 @@ static bool ggml_backend_hexagon_device_supports_op(ggml_backend_dev_t dev, cons
|
||||
case GGML_UNARY_OP_SIGMOID:
|
||||
case GGML_UNARY_OP_SOFTPLUS:
|
||||
case GGML_UNARY_OP_TANH:
|
||||
case GGML_UNARY_OP_ABS:
|
||||
case GGML_UNARY_OP_SILU:
|
||||
case GGML_UNARY_OP_GELU:
|
||||
case GGML_UNARY_OP_GELU_QUICK:
|
||||
|
||||
@@ -62,8 +62,6 @@ enum htp_op_code {
|
||||
HTP_OP_UNARY_NEG,
|
||||
HTP_OP_UNARY_SOFTPLUS,
|
||||
HTP_OP_UNARY_TANH,
|
||||
HTP_OP_UNARY_ABS,
|
||||
HTP_OP_UNARY_LOG,
|
||||
HTP_OP_GLU_SWIGLU,
|
||||
HTP_OP_GLU_SWIGLU_OAI,
|
||||
HTP_OP_GLU_GEGLU,
|
||||
|
||||
@@ -358,34 +358,6 @@ static inline void hvx_clamp_scalar_f32(uint8_t * restrict dst, const uint8_t *
|
||||
}
|
||||
}
|
||||
|
||||
//
|
||||
// Abs
|
||||
//
|
||||
|
||||
static inline void hvx_abs_f32_aa(uint8_t * restrict dst, const uint8_t * restrict src, uint32_t n) {
|
||||
assert((unsigned long) dst % 128 == 0);
|
||||
assert((unsigned long) src % 128 == 0);
|
||||
|
||||
HVX_Vector * restrict vdst = (HVX_Vector *) dst;
|
||||
HVX_Vector * restrict vsrc = (HVX_Vector *) src;
|
||||
|
||||
const uint32_t elem_size = sizeof(float);
|
||||
const uint32_t epv = 128 / elem_size;
|
||||
const uint32_t nvec = n / epv;
|
||||
const uint32_t nloe = n % epv;
|
||||
|
||||
uint32_t i = 0;
|
||||
|
||||
_Pragma("unroll(4)")
|
||||
for (; i < nvec; i++) {
|
||||
vdst[i] = hvx_vec_abs_f32(vsrc[i]);
|
||||
}
|
||||
if (nloe) {
|
||||
HVX_Vector v = hvx_vec_abs_f32(vsrc[i]);
|
||||
hvx_vec_store_a((void *) &vdst[i], nloe * elem_size, v);
|
||||
}
|
||||
}
|
||||
|
||||
//
|
||||
// Square
|
||||
//
|
||||
|
||||
@@ -62,28 +62,4 @@ static inline HVX_Vector hvx_vec_log_f32(HVX_Vector x) {
|
||||
return hvx_vec_add_f32_f32(term_e, res);
|
||||
}
|
||||
|
||||
static inline void hvx_log_f32_aa(uint8_t * restrict dst, const uint8_t * restrict src, uint32_t n) {
|
||||
assert((unsigned long) dst % 128 == 0);
|
||||
assert((unsigned long) src % 128 == 0);
|
||||
|
||||
HVX_Vector * restrict vdst = (HVX_Vector *) dst;
|
||||
HVX_Vector * restrict vsrc = (HVX_Vector *) src;
|
||||
|
||||
const uint32_t elem_size = sizeof(float);
|
||||
const uint32_t epv = 128 / elem_size;
|
||||
const uint32_t nvec = n / epv;
|
||||
const uint32_t nloe = n % epv;
|
||||
|
||||
uint32_t i = 0;
|
||||
|
||||
_Pragma("unroll(4)")
|
||||
for (; i < nvec; i++) {
|
||||
vdst[i] = hvx_vec_log_f32(vsrc[i]);
|
||||
}
|
||||
if (nloe) {
|
||||
HVX_Vector v = hvx_vec_log_f32(vsrc[i]);
|
||||
hvx_vec_store_a((void *) &vdst[i], nloe * elem_size, v);
|
||||
}
|
||||
}
|
||||
|
||||
#endif /* HVX_LOG_H */
|
||||
|
||||
@@ -777,8 +777,6 @@ static int execute_op(struct htp_ops_context * octx) {
|
||||
case HTP_OP_UNARY_NEG:
|
||||
case HTP_OP_UNARY_EXP:
|
||||
case HTP_OP_UNARY_TANH:
|
||||
case HTP_OP_UNARY_ABS:
|
||||
case HTP_OP_UNARY_LOG:
|
||||
case HTP_OP_L2_NORM:
|
||||
return op_unary(octx);
|
||||
|
||||
|
||||
@@ -443,34 +443,6 @@ static void tanh_f32(const float * restrict src,
|
||||
}
|
||||
}
|
||||
|
||||
static void abs_f32(const float * restrict src,
|
||||
float * restrict dst,
|
||||
const uint32_t num_rows,
|
||||
const struct htp_unary_context * uctx) {
|
||||
htp_unary_op_preamble;
|
||||
|
||||
for (uint32_t ir = 0; ir < num_rows; ir++) {
|
||||
const uint8_t * restrict src_local = (const uint8_t *)src + (ir * src0_row_size_aligned);
|
||||
uint8_t * restrict dst_local = (uint8_t *)dst + (ir * dst_row_size_aligned);
|
||||
|
||||
hvx_abs_f32_aa(dst_local, src_local, ne0);
|
||||
}
|
||||
}
|
||||
|
||||
static void log_f32(const float * restrict src,
|
||||
float * restrict dst,
|
||||
const uint32_t num_rows,
|
||||
const struct htp_unary_context * uctx) {
|
||||
htp_unary_op_preamble;
|
||||
|
||||
for (uint32_t ir = 0; ir < num_rows; ir++) {
|
||||
const uint8_t * restrict src_local = (const uint8_t *)src + (ir * src0_row_size_aligned);
|
||||
uint8_t * restrict dst_local = (uint8_t *)dst + (ir * dst_row_size_aligned);
|
||||
|
||||
hvx_log_f32_aa(dst_local, src_local, ne0);
|
||||
}
|
||||
}
|
||||
|
||||
#define DEFINE_UNARY_TASK(NAME, IS_RMS_NORM_MUL, IS_TRI, CORE_EXPR) \
|
||||
static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * data) { \
|
||||
const struct htp_unary_context * uctx = (const struct htp_unary_context *) data; \
|
||||
@@ -506,9 +478,6 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
const uint32_t nb11 = src1 ? src1->nb[1] : 0; \
|
||||
const uint32_t nb12 = src1 ? src1->nb[2] : 0; \
|
||||
const uint32_t nb13 = src1 ? src1->nb[3] : 0; \
|
||||
const uint32_t nb11_bc = (src1 && src1->ne[1] > 1) ? nb11 : 0; \
|
||||
const uint32_t nb12_bc = (src1 && src1->ne[2] > 1) ? nb12 : 0; \
|
||||
const uint32_t nb13_bc = (src1 && src1->ne[3] > 1) ? nb13 : 0; \
|
||||
const bool src1_contig = src1 ? ((nb12 == (size_t)ne01 * nb11) && (nb13 == (size_t)ne02 * nb12)) : false; \
|
||||
\
|
||||
uint8_t * src0_vtcm_data = uctx->vtcm_src0 + (ith * uctx->vtcm_src0_size_per_thread); \
|
||||
@@ -528,12 +497,8 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
const struct fastdiv_values * div_ne02 = &uctx->kparams->div_ne02; \
|
||||
const struct fastdiv_values * div_ne012 = &uctx->kparams->div_ne012; \
|
||||
\
|
||||
const bool src1_needs_row_clip = (IS_RMS_NORM_MUL) && !uctx->broadcast_weight && !src1_contig; \
|
||||
const bool block_src0_contig = src0_contig && !src1_needs_row_clip; \
|
||||
const bool block_dst_contig = dst_contig && !src1_needs_row_clip; \
|
||||
\
|
||||
const uint32_t src0_max_block = block_src0_contig ? uctx->block : MIN((uint32_t)uctx->block, ne01); \
|
||||
const uint32_t dst_max_block = block_dst_contig ? uctx->block : MIN((uint32_t)uctx->block, ne1); \
|
||||
const uint32_t src0_max_block = src0_contig ? uctx->block : MIN((uint32_t)uctx->block, ne01); \
|
||||
const uint32_t dst_max_block = dst_contig ? uctx->block : MIN((uint32_t)uctx->block, ne1); \
|
||||
const uint32_t BLOCK = MIN(src0_max_block, dst_max_block); \
|
||||
if (BLOCK == 0) { \
|
||||
FARF(ERROR, "unary-f32 : current VTCM reservation %zu is too small, needed at least %zu\n", \
|
||||
@@ -550,8 +515,8 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
} \
|
||||
\
|
||||
for (uint32_t ir = src0_start_row, vtcm_idx = 0; ir < src0_end_row && vtcm_idx < 2; vtcm_idx++) { \
|
||||
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, block_src0_contig, block_dst_contig, \
|
||||
ne01, div_ne01); \
|
||||
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, \
|
||||
div_ne01); \
|
||||
\
|
||||
dma_queue_push(dma_queue, \
|
||||
dma_make_ptr(data_dst, dst_vtcm_data + (vtcm_idx * dst_vtcm_half_size)), \
|
||||
@@ -565,7 +530,7 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
\
|
||||
if ((IS_RMS_NORM_MUL) && !uctx->broadcast_weight) { \
|
||||
const size_t src1_off = src1_contig ? (ir * nb11) : \
|
||||
unary_row_offset(ir, ne01, ne02, div_ne01, div_ne02, div_ne012, nb11_bc, nb12_bc, nb13_bc); \
|
||||
unary_row_offset(ir, ne01, ne02, div_ne01, div_ne02, div_ne012, nb11, nb12, nb13); \
|
||||
dma_queue_push(dma_queue, \
|
||||
dma_make_ptr(src1_vtcm_data + (vtcm_idx * src1_vtcm_half_size), data_src1 + src1_off), \
|
||||
uctx->src1_row_size_aligned, nb11, uctx->src1_data_row_size, block_size); \
|
||||
@@ -575,8 +540,8 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
} \
|
||||
\
|
||||
for (uint32_t ir = src0_start_row; ir < src0_end_row; ) { \
|
||||
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, block_src0_contig, block_dst_contig, \
|
||||
ne01, div_ne01); \
|
||||
const uint32_t block_size = unary_block_size(ir, src0_end_row, BLOCK, src0_contig, dst_contig, ne01, \
|
||||
div_ne01); \
|
||||
\
|
||||
float * dst_vtcm = (float *) dma_queue_pop(dma_queue).src; \
|
||||
float * src0_vtcm = (float *) dma_queue_pop(dma_queue).dst; \
|
||||
@@ -597,12 +562,12 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
\
|
||||
const uint32_t next_ir = ir + block_size; \
|
||||
if (next_ir < src0_end_row) { \
|
||||
const uint32_t next_block_size = unary_block_size(next_ir, src0_end_row, BLOCK, block_src0_contig, \
|
||||
block_dst_contig, ne01, div_ne01); \
|
||||
const uint32_t next_block_size = unary_block_size(next_ir, src0_end_row, BLOCK, src0_contig, dst_contig,\
|
||||
ne01, div_ne01); \
|
||||
const uint32_t pref_ir = next_ir + next_block_size; \
|
||||
if (pref_ir < src0_end_row) { \
|
||||
const uint32_t pref_block_size = unary_block_size(pref_ir, src0_end_row, BLOCK, block_src0_contig, \
|
||||
block_dst_contig, ne01, div_ne01); \
|
||||
const uint32_t pref_block_size = unary_block_size(pref_ir, src0_end_row, BLOCK, src0_contig, \
|
||||
dst_contig, ne01, div_ne01); \
|
||||
const size_t src0_pref_off = src0_contig ? (pref_ir * nb01) : \
|
||||
unary_row_offset(pref_ir, ne01, ne02, div_ne01, div_ne02, div_ne012, nb01, nb02, nb03); \
|
||||
dma_queue_push(dma_queue, \
|
||||
@@ -611,8 +576,7 @@ static void unary_task_f32_##NAME(unsigned int nth, unsigned int ith, void * dat
|
||||
\
|
||||
if ((IS_RMS_NORM_MUL) && !uctx->broadcast_weight) { \
|
||||
const size_t src1_pref_off = src1_contig ? (pref_ir * nb11) : \
|
||||
unary_row_offset(pref_ir, ne01, ne02, div_ne01, div_ne02, div_ne012, nb11_bc, nb12_bc, \
|
||||
nb13_bc); \
|
||||
unary_row_offset(pref_ir, ne01, ne02, div_ne01, div_ne02, div_ne012, nb11, nb12, nb13); \
|
||||
dma_queue_push(dma_queue, \
|
||||
dma_make_ptr(src1_vtcm, data_src1 + src1_pref_off), \
|
||||
uctx->src1_row_size_aligned, nb11, uctx->src1_data_row_size, pref_block_size); \
|
||||
@@ -639,8 +603,6 @@ DEFINE_UNARY_TASK(unary_silu, false, false, silu_f32(src0_vtcm, dst_vtcm, bl
|
||||
DEFINE_UNARY_TASK(unary_gelu, false, false, gelu_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(unary_softplus, false, false, softplus_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(unary_tanh, false, false, tanh_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(unary_abs, false, false, abs_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(unary_log, false, false, log_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(l2_norm, false, false, l2_norm_f32(src0_vtcm, dst_vtcm, block_size, uctx))
|
||||
DEFINE_UNARY_TASK(tri, false, true, tri_f32(src0_vtcm, dst_vtcm, block_size, ir, uctx))
|
||||
|
||||
@@ -888,8 +850,6 @@ DEFINE_UNARY_TILED_TASK(unary_silu, false, tile_silu_f32(dst_vtcm, src_vtcm,
|
||||
DEFINE_UNARY_TILED_TASK(unary_gelu, false, tile_gelu_f32(dst_vtcm, src_vtcm, tw))
|
||||
DEFINE_UNARY_TILED_TASK(unary_softplus, false, tile_unary_softplus_f32(dst_vtcm, src_vtcm, tw))
|
||||
DEFINE_UNARY_TILED_TASK(unary_tanh, false, hvx_tanh_f32_aa(dst_vtcm, src_vtcm, tw))
|
||||
DEFINE_UNARY_TILED_TASK(unary_abs, false, hvx_abs_f32_aa(dst_vtcm, src_vtcm, tw))
|
||||
DEFINE_UNARY_TILED_TASK(unary_log, false, hvx_log_f32_aa(dst_vtcm, src_vtcm, tw))
|
||||
DEFINE_UNARY_TILED_TASK(tri, true, tri_apply_tile_f32(src_vtcm, dst_vtcm, tw, col, i01, ne0, tri_ttype))
|
||||
|
||||
static int execute_op_unary_f32(struct htp_ops_context * octx) {
|
||||
@@ -915,8 +875,6 @@ static int execute_op_unary_f32(struct htp_ops_context * octx) {
|
||||
case HTP_OP_UNARY_GELU: op_type = "gelu-f32"; break;
|
||||
case HTP_OP_UNARY_SOFTPLUS: op_type = "softplus-f32"; break;
|
||||
case HTP_OP_UNARY_TANH: op_type = "tanh-f32"; break;
|
||||
case HTP_OP_UNARY_ABS: op_type = "abs-f32"; break;
|
||||
case HTP_OP_UNARY_LOG: op_type = "log-f32"; break;
|
||||
case HTP_OP_L2_NORM: op_type = "l2norm-f32"; break;
|
||||
case HTP_OP_TRI: op_type = "tri-f32"; break;
|
||||
|
||||
@@ -1015,8 +973,6 @@ static int execute_op_unary_f32(struct htp_ops_context * octx) {
|
||||
case HTP_OP_UNARY_GELU: task_func = unary_task_f32_tiled_unary_gelu; break;
|
||||
case HTP_OP_UNARY_SOFTPLUS: task_func = unary_task_f32_tiled_unary_softplus; break;
|
||||
case HTP_OP_UNARY_TANH: task_func = unary_task_f32_tiled_unary_tanh; break;
|
||||
case HTP_OP_UNARY_ABS: task_func = unary_task_f32_tiled_unary_abs; break;
|
||||
case HTP_OP_UNARY_LOG: task_func = unary_task_f32_tiled_unary_log; break;
|
||||
case HTP_OP_TRI: task_func = unary_task_f32_tiled_tri; break;
|
||||
default: break;
|
||||
}
|
||||
@@ -1036,8 +992,6 @@ static int execute_op_unary_f32(struct htp_ops_context * octx) {
|
||||
case HTP_OP_UNARY_GELU: task_func = unary_task_f32_unary_gelu; break;
|
||||
case HTP_OP_UNARY_SOFTPLUS: task_func = unary_task_f32_unary_softplus; break;
|
||||
case HTP_OP_UNARY_TANH: task_func = unary_task_f32_unary_tanh; break;
|
||||
case HTP_OP_UNARY_ABS: task_func = unary_task_f32_unary_abs; break;
|
||||
case HTP_OP_UNARY_LOG: task_func = unary_task_f32_unary_log; break;
|
||||
case HTP_OP_L2_NORM: task_func = unary_task_f32_l2_norm; break;
|
||||
case HTP_OP_TRI: task_func = unary_task_f32_tri; break;
|
||||
default: break;
|
||||
|
||||
@@ -55,8 +55,6 @@ static inline bool htp_op_is_unary(uint32_t opcode) {
|
||||
case HTP_OP_UNARY_GELU:
|
||||
case HTP_OP_UNARY_SOFTPLUS:
|
||||
case HTP_OP_UNARY_TANH:
|
||||
case HTP_OP_UNARY_ABS:
|
||||
case HTP_OP_UNARY_LOG:
|
||||
case HTP_OP_L2_NORM:
|
||||
case HTP_OP_TRI:
|
||||
return true;
|
||||
|
||||
@@ -66,7 +66,6 @@ fa_vec_cfg_t fa_vec_baseline_cfg(int dk, int dv) {
|
||||
// One row per kept bucket, plus per-(dtype,dk,dv) ne11-collapsed domain defaults
|
||||
// (ne11_b = FA_VEC_NE11_DEFAULT, ne01_b = domain). To retune or add a device, re-run the
|
||||
// sweep and paste its output. See ggml-metal-tuning.h for the row/lookup semantics.
|
||||
// ref: https://github.com/ggml-org/llama.cpp/pull/27824
|
||||
constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
|
||||
{ { GGML_METAL_DEVICE_M1_PRO, GGML_TYPE_F16, 32, 32, 3, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M1_PRO, GGML_TYPE_F16, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
@@ -450,159 +449,6 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
|
||||
{ { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 1, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 32, 32, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, 3, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, 3, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, 3, 3 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 96, 96, 3, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 128, 128, 2, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 192, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 128, 1, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 128, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 192, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 256, 256, 2, 3 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 320, 256, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 320, 256, 3, 0 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 2, 0 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 3, 0 }, { 4, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 1, 3 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 2, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 2, 3 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 3, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 512, 512, 3, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 576, 512, 2, 0 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 576, 512, 2, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 576, 512, 2, 3 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_F16, 576, 512, 3, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 64, 64, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 96, 96, 3, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 128, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 192, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 192, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3_MAX, GGML_TYPE_Q8_0, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 32, 32, 1, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 32, 32, 2, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 32, 32, 3, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 64, 64, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 96, 96, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 96, 96, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 128, 128, 1, 2 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 192, 128, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 256, 256, 3, 0 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 320, 256, 3, 0 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_F16, 512, 512, 3, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 128, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 128, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 192, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, 1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, 1, 3 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M4_PRO, GGML_TYPE_Q8_0, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M4_MAX, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_MAX, GGML_TYPE_F16, 32, 32, 2, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_MAX, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
@@ -794,215 +640,7 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
|
||||
{ { GGML_METAL_DEVICE_M4_MAX, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M4_MAX, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 1, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 32, 32, 3, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 64, 64, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 96, 96, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 96, 96, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 3, 0 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 1, 2 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 3, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 128, 128, 3, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 192, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 3, 2 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 192, 128, 3, 3 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 256, 256, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 256, 256, 1, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 320, 256, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 320, 256, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 320, 256, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 320, 256, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 2, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 2, 3 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 3, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 3, 2 }, { 4, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 3, 3 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 512, 512, 3, 4 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 576, 512, 2, 0 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 576, 512, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 576, 512, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_F16, 576, 512, 2, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 32, 32, 3, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, 1, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, 1, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, 2, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 64, 64, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 96, 96, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 96, 96, 3, 4 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 192, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 192, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 192, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 128, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 192, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 320, 256, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 512, 512, 3, 0 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 512, 512, 2, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 512, 512, 3, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 32, 32, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 64, 64, 1, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 64, 64, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 64, 64, 3, 4 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 128, 128, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 128, 128, 1, 4 }, { 1, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 192, 128, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 192, 128, 3, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 256, 256, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 320, 256, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 320, 256, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 512, 512, 2, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 512, 512, 2, 3 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 512, 512, 3, 3 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, 1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, 1, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, 1, 3 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_F16, 576, 512, 2, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 32, 32, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 64, 64, 1, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 96, 96, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 128, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, 1, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, 2, 3 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 192, 3, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 192, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 512, 512, 1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 512, 512, -1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 512, 512, 2, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q4_0, 512, 512, 2, 3 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 1, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 1, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 2, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 32, 32, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 64, 64, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 128, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 192, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 192, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 192, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 192, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
|
||||
{ { GGML_METAL_DEVICE_M5_MAX, GGML_TYPE_F16, 32, 32, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_MAX, GGML_TYPE_F16, 32, 32, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_MAX, GGML_TYPE_F16, 32, 32, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_MAX, GGML_TYPE_F16, 32, 32, 1, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M5_MAX, GGML_TYPE_F16, 32, 32, 2, 2 }, { 4, 4 } },
|
||||
|
||||
@@ -1,4 +1,3 @@
|
||||
#include <array>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstring>
|
||||
@@ -151,8 +150,7 @@ struct sdpa_partition {
|
||||
|
||||
// Build + compile the contiguous-input GQA SDPA graph (MatMul->Divide->Add->SoftMax->MatMul), f32 out.
|
||||
// Mirrors the hardware-verified scratch/onednn_sdpa_probe.cpp build_gqa (partitions=1, sdp_primitive_kernel_t).
|
||||
static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int seq, int d,
|
||||
const std::array<int64_t, 5> & k_str, const std::array<int64_t, 5> & v_str) try {
|
||||
static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int seq, int d) {
|
||||
using ltype = logical_tensor::layout_type;
|
||||
using dt = logical_tensor::data_type;
|
||||
using ldims = logical_tensor::dims;
|
||||
@@ -160,12 +158,11 @@ static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int
|
||||
const int rep = H / Hkv;
|
||||
const ldims q_sz = {1, Hkv, rep, q, d}, kv_sz = {1, Hkv, 1, seq, d}, s_sz = {1, Hkv, rep, q, seq},
|
||||
sc = {1, 1, 1, 1, 1}, msk = {1, 1, 1, q, seq}, o_sz = {1, Hkv, rep, q, d};
|
||||
const ldims k_st(k_str.begin(), k_str.end()), v_st(v_str.begin(), v_str.end());
|
||||
int64_t id = 0;
|
||||
sdpa_partition E;
|
||||
|
||||
auto query = logical_tensor(id++, t, q_sz, ltype::strided);
|
||||
auto key = logical_tensor(id++, t, kv_sz, k_st);
|
||||
auto key = logical_tensor(id++, t, kv_sz, ltype::strided);
|
||||
auto score = logical_tensor(id++, fi, s_sz, ltype::strided);
|
||||
auto bmm1 = op(id++, op::kind::MatMul, "bmm1");
|
||||
bmm1.set_attr<bool>(op::attr::transpose_b, true); // key is [.., seq, d]
|
||||
@@ -187,7 +184,7 @@ static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int
|
||||
smax.set_attr<std::string>(op::attr::mode, "inf_as_zero");
|
||||
smax.add_inputs({masked}); smax.add_outputs({probs});
|
||||
|
||||
auto value = logical_tensor(id++, t, kv_sz, v_st);
|
||||
auto value = logical_tensor(id++, t, kv_sz, ltype::strided);
|
||||
// f16 output is REQUIRED to hit sdp_primitive_kernel_t (the systolic micro-kernel); an f32 output
|
||||
// falls to larger_partition_kernel_t which materializes N^2 (confirmed: scratch/onednn_sdpa_kernel_probe.cpp).
|
||||
// converted to the f32 ggml dst in the permute below.
|
||||
@@ -201,7 +198,6 @@ static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int
|
||||
|
||||
auto parts = g.get_partitions();
|
||||
if (parts.size() != 1 || !parts[0].is_supported()) {
|
||||
GGML_LOG_WARN("%s: oneDNN did not fuse the SDPA graph; falling back to TILE kernel\n", __func__);
|
||||
return E; // ok stays false -> caller falls back to TILE
|
||||
}
|
||||
E.ins = parts[0].get_input_ports();
|
||||
@@ -213,12 +209,6 @@ static sdpa_partition build_sdpa(const engine & eng, int H, int Hkv, int q, int
|
||||
E.ok = true;
|
||||
return E;
|
||||
}
|
||||
catch (const std::exception & e) {
|
||||
// compile() can reject a stride set the partitioner never inspects; memoise the failure so the
|
||||
// fallback costs one build rather than one per call.
|
||||
GGML_LOG_WARN("%s: oneDNN SDPA partition build failed (%s); falling back to TILE kernel\n", __func__, e.what());
|
||||
return {};
|
||||
}
|
||||
|
||||
void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tensor * dst) try {
|
||||
const ggml_tensor * Q = dst->src[0];
|
||||
@@ -244,34 +234,13 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
|
||||
ggml_sycl_pool_alloc<sycl::half> Qf(ctx.pool(), (size_t) H * q * d);
|
||||
cont_to_f16_sycl<float>((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream);
|
||||
|
||||
// K/V: bind the f16 cache in place. llama.cpp permutes it to [token][head][dim], so its head
|
||||
// plane is strided rather than dense, which is what an explicit stride vector expresses.
|
||||
// Quantized and f32 KV still stage a dense copy -- the layout the k_str/v_str defaults describe.
|
||||
// K/V: use pool-alloc for both F16 and dequant paths.
|
||||
sycl::half * K_ptr = nullptr;
|
||||
sycl::half * V_ptr = nullptr;
|
||||
std::array<int64_t, 5> k_str{ Hkv * seq * d, seq * d, seq * d, d, 1 };
|
||||
std::array<int64_t, 5> v_str = k_str;
|
||||
std::optional<ggml_sycl_pool_alloc<sycl::half>> Kf_pool;
|
||||
std::optional<ggml_sycl_pool_alloc<sycl::half>> Vf_pool;
|
||||
|
||||
auto bindable = [](const ggml_tensor * t) {
|
||||
return t->nb[0] == sizeof(sycl::half) && t->nb[1] % sizeof(sycl::half) == 0 &&
|
||||
t->nb[2] % sizeof(sycl::half) == 0 && t->nb[3] % sizeof(sycl::half) == 0;
|
||||
};
|
||||
auto elem_strides = [](const ggml_tensor * t) {
|
||||
const int64_t s1 = (int64_t) (t->nb[1] / t->nb[0]);
|
||||
const int64_t s2 = (int64_t) (t->nb[2] / t->nb[0]);
|
||||
const int64_t s3 = (int64_t) (t->nb[3] / t->nb[0]);
|
||||
// dims are {mb=1, Hkv, rep=1, seq, d}; the size-1 dims at 0 and 2 never advance an address.
|
||||
return std::array<int64_t, 5>{ s3, s2, s2, s1, 1 };
|
||||
};
|
||||
|
||||
if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16 && bindable(K) && bindable(V)) {
|
||||
K_ptr = (sycl::half *) K->data;
|
||||
V_ptr = (sycl::half *) V->data;
|
||||
k_str = elem_strides(K);
|
||||
v_str = elem_strides(V);
|
||||
} else if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) {
|
||||
if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) {
|
||||
Kf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d);
|
||||
Vf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d);
|
||||
cont_to_f16_sycl<sycl::half>((const char *) K->data, Kf_pool->get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream);
|
||||
@@ -372,24 +341,19 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
|
||||
|
||||
ggml_sycl_pool_alloc<sycl::half> outf(ctx.pool(), (size_t) H * q * d); // f16 contiguous SDPA out [mb,H,q,d]
|
||||
|
||||
// compile once per (device, shape, KV strides), reuse across layers/calls. Stride 2 always
|
||||
// repeats stride 1 and stride 4 is always 1, so the key covers every entry that can differ.
|
||||
// compile once per (device, shape), reuse across layers/calls.
|
||||
static std::unordered_map<std::string, sdpa_partition> cache;
|
||||
char keyb[256];
|
||||
snprintf(keyb, sizeof(keyb), "%d:%lld:%lld:%lld:%lld:%lld:%lld:%lld:%lld:%lld:%lld:%lld", ggml_sycl_get_device(),
|
||||
(long long) H, (long long) Hkv, (long long) q, (long long) seq, (long long) d,
|
||||
(long long) k_str[0], (long long) k_str[1], (long long) k_str[3],
|
||||
(long long) v_str[0], (long long) v_str[1], (long long) v_str[3]);
|
||||
char keyb[96];
|
||||
snprintf(keyb, sizeof(keyb), "%d:%lld:%lld:%lld:%lld:%lld", ggml_sycl_get_device(),
|
||||
(long long) H, (long long) Hkv, (long long) q, (long long) seq, (long long) d);
|
||||
auto it = cache.find(keyb);
|
||||
if (it == cache.end()) {
|
||||
it = cache.emplace(keyb, build_sdpa(eng, (int) H, (int) Hkv, (int) q, (int) seq, (int) d, k_str, v_str)).first;
|
||||
it = cache.emplace(keyb, build_sdpa(eng, (int) H, (int) Hkv, (int) q, (int) seq, (int) d)).first;
|
||||
}
|
||||
sdpa_partition & E = it->second;
|
||||
if (!E.ok) {
|
||||
// oneDNN can decline a shape or a stride set that _supported() never sees; build_sdpa warns per key.
|
||||
ggml_sycl_flash_attn_ext_tile(ctx, dst);
|
||||
return;
|
||||
}
|
||||
// _supported() is authoritative: if it accepted this op the partition must build.
|
||||
// A failure here is a gap in _supported() -- surface it, don't mask it with a fallback.
|
||||
GGML_ASSERT(E.ok && "oneDNN SDPA partition failed to build for a _supported() shape");
|
||||
|
||||
auto id2ptr = [&](size_t r) -> void * {
|
||||
if (r == E.id_q) return Qf.get();
|
||||
|
||||
@@ -104,6 +104,7 @@ enum best_fattn_kernel {
|
||||
|
||||
|
||||
static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const ggml_tensor * dst) {
|
||||
GGML_UNUSED(device);
|
||||
#ifndef SYCL_FLASH_ATTN
|
||||
GGML_UNUSED(dst);
|
||||
return BEST_FATTN_KERNEL_NONE;
|
||||
@@ -262,11 +263,6 @@ static best_fattn_kernel ggml_sycl_get_best_fattn_kernel(const int device, const
|
||||
}
|
||||
} else {
|
||||
if (Q->ne[1] <= 2) {
|
||||
// TILE is faster for quantized KV decode on Xe2 (BMG); keep VEC on untested archs
|
||||
const gpu_arch arch = ggml_sycl_info().devices[device].hw_info.arch;
|
||||
if (arch == gpu_arch::intel_gpu_bmg_g21 || arch == gpu_arch::intel_gpu_bmg_g31) {
|
||||
return BEST_FATTN_KERNEL_TILE;
|
||||
}
|
||||
return BEST_FATTN_KERNEL_VEC;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -167,7 +167,6 @@ class Keys:
|
||||
SELECTOR_RANK = "{arch}.selector_rank"
|
||||
SELECTOR_TOP_K = "{arch}.selector_top_k"
|
||||
SAMPLE_FROM_ANCHOR = "{arch}.sample_from_anchor"
|
||||
HAS_CONFIDENCE_HEAD = "{arch}.has_confidence_head"
|
||||
NORM_BEFORE_RESIDUAL = "{arch}.norm_before_residual"
|
||||
NORM_BEFORE_FC = "{arch}.norm_before_fc"
|
||||
|
||||
@@ -230,19 +229,6 @@ class Keys:
|
||||
COUNT = "{arch}.hyper_connection.count"
|
||||
SINKHORN_ITERATIONS = "{arch}.hyper_connection.sinkhorn_iterations"
|
||||
EPSILON = "{arch}.hyper_connection.epsilon"
|
||||
# absent means the mix projection is full rank (DeepSeek-V4 behaviour)
|
||||
LOW_RANK = "{arch}.hyper_connection.low_rank"
|
||||
|
||||
class PerLayerEmbedding:
|
||||
LAYERS = "{arch}.ple.layers"
|
||||
NGRAM_SIZE = "{arch}.ple.ngram_size"
|
||||
HEADS_PER_NGRAM = "{arch}.ple.heads_per_ngram"
|
||||
CONV_KERNEL = "{arch}.ple.conv_kernel"
|
||||
LAYER_MULTIPLIERS = "{arch}.ple.layer_multipliers"
|
||||
HEAD_OFFSETS = "{arch}.ple.head_offsets"
|
||||
HEAD_VOCAB_SIZES = "{arch}.ple.head_vocab_sizes"
|
||||
EOS_TOKEN_ID = "{arch}.ple.eos_token_id"
|
||||
IMAGE_TOKEN_ID = "{arch}.ple.image_token_id"
|
||||
|
||||
class Rope:
|
||||
DIMENSION_COUNT = "{arch}.rope.dimension_count"
|
||||
@@ -512,7 +498,6 @@ class MODEL_ARCH(IntEnum):
|
||||
QWEN3VLMOE = auto()
|
||||
QWEN35 = auto()
|
||||
QWEN35MOE = auto()
|
||||
QWEN4EXP = auto()
|
||||
PHI2 = auto()
|
||||
PHI3 = auto()
|
||||
PHIMOE = auto()
|
||||
@@ -655,9 +640,6 @@ class MODEL_TENSOR(IntEnum):
|
||||
HC_HEAD_FN = auto()
|
||||
HC_HEAD_BASE = auto()
|
||||
HC_HEAD_SCALE = auto()
|
||||
HC_HEAD_NORM = auto() # qwen4exp
|
||||
HC_HEAD_DOWN = auto() # qwen4exp
|
||||
HC_HEAD_UP = auto() # qwen4exp
|
||||
ROPE_FREQS = auto()
|
||||
ROPE_FACTORS_LONG = auto()
|
||||
ROPE_FACTORS_SHORT = auto()
|
||||
@@ -802,20 +784,6 @@ class MODEL_TENSOR(IntEnum):
|
||||
HC_FFN_FN = auto()
|
||||
HC_FFN_BASE = auto()
|
||||
HC_FFN_SCALE = auto()
|
||||
HC_ATTN_NORM = auto() # qwen4exp
|
||||
HC_ATTN_DOWN = auto() # qwen4exp
|
||||
HC_ATTN_UP = auto() # qwen4exp
|
||||
HC_ATTN_INJECT = auto() # qwen4exp
|
||||
HC_FFN_NORM = auto() # qwen4exp
|
||||
HC_FFN_DOWN = auto() # qwen4exp
|
||||
HC_FFN_UP = auto() # qwen4exp
|
||||
HC_FFN_INJECT = auto() # qwen4exp
|
||||
PLE_KEY = auto() # qwen4exp
|
||||
PLE_VALUE = auto() # qwen4exp
|
||||
PLE_NORM_KEY = auto() # qwen4exp
|
||||
PLE_NORM_QUERY = auto() # qwen4exp
|
||||
PLE_NORM_CONV = auto() # qwen4exp
|
||||
PLE_CONV1D = auto() # qwen4exp
|
||||
ATTN_COMPRESSOR_WKV = auto()
|
||||
ATTN_COMPRESSOR_WGATE = auto()
|
||||
ATTN_COMPRESSOR_APE = auto()
|
||||
@@ -1260,7 +1228,6 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
|
||||
MODEL_ARCH.QWEN3VLMOE: "qwen3vlmoe",
|
||||
MODEL_ARCH.QWEN35: "qwen35",
|
||||
MODEL_ARCH.QWEN35MOE: "qwen35moe",
|
||||
MODEL_ARCH.QWEN4EXP: "qwen4exp",
|
||||
MODEL_ARCH.PHI2: "phi2",
|
||||
MODEL_ARCH.PHI3: "phi3",
|
||||
MODEL_ARCH.PHIMOE: "phimoe",
|
||||
@@ -1402,9 +1369,6 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
|
||||
MODEL_TENSOR.HC_HEAD_FN: "output_hc_fn",
|
||||
MODEL_TENSOR.HC_HEAD_BASE: "output_hc_base",
|
||||
MODEL_TENSOR.HC_HEAD_SCALE: "output_hc_scale",
|
||||
MODEL_TENSOR.HC_HEAD_NORM: "output_hc_norm", # qwen4exp
|
||||
MODEL_TENSOR.HC_HEAD_DOWN: "output_hc_down", # qwen4exp
|
||||
MODEL_TENSOR.HC_HEAD_UP: "output_hc_up", # qwen4exp
|
||||
MODEL_TENSOR.ROPE_FREQS: "rope_freqs",
|
||||
MODEL_TENSOR.ROPE_FACTORS_LONG: "rope_factors_long",
|
||||
MODEL_TENSOR.ROPE_FACTORS_SHORT: "rope_factors_short",
|
||||
@@ -1549,20 +1513,6 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
|
||||
MODEL_TENSOR.HC_FFN_FN: "blk.{bid}.hc_ffn_fn",
|
||||
MODEL_TENSOR.HC_FFN_BASE: "blk.{bid}.hc_ffn_base",
|
||||
MODEL_TENSOR.HC_FFN_SCALE: "blk.{bid}.hc_ffn_scale",
|
||||
MODEL_TENSOR.HC_ATTN_NORM: "blk.{bid}.hc_attn_norm", # qwen4exp
|
||||
MODEL_TENSOR.HC_ATTN_DOWN: "blk.{bid}.hc_attn_down", # qwen4exp
|
||||
MODEL_TENSOR.HC_ATTN_UP: "blk.{bid}.hc_attn_up", # qwen4exp
|
||||
MODEL_TENSOR.HC_ATTN_INJECT: "blk.{bid}.hc_attn_inject", # qwen4exp
|
||||
MODEL_TENSOR.HC_FFN_NORM: "blk.{bid}.hc_ffn_norm", # qwen4exp
|
||||
MODEL_TENSOR.HC_FFN_DOWN: "blk.{bid}.hc_ffn_down", # qwen4exp
|
||||
MODEL_TENSOR.HC_FFN_UP: "blk.{bid}.hc_ffn_up", # qwen4exp
|
||||
MODEL_TENSOR.HC_FFN_INJECT: "blk.{bid}.hc_ffn_inject", # qwen4exp
|
||||
MODEL_TENSOR.PLE_KEY: "blk.{bid}.ple_key", # qwen4exp
|
||||
MODEL_TENSOR.PLE_VALUE: "blk.{bid}.ple_value", # qwen4exp
|
||||
MODEL_TENSOR.PLE_NORM_KEY: "blk.{bid}.ple_norm_key", # qwen4exp
|
||||
MODEL_TENSOR.PLE_NORM_QUERY: "blk.{bid}.ple_norm_query", # qwen4exp
|
||||
MODEL_TENSOR.PLE_NORM_CONV: "blk.{bid}.ple_norm_conv", # qwen4exp
|
||||
MODEL_TENSOR.PLE_CONV1D: "blk.{bid}.ple_conv1d", # qwen4exp
|
||||
MODEL_TENSOR.ATTN_COMPRESSOR_WKV: "blk.{bid}.attn_compressor_kv",
|
||||
MODEL_TENSOR.ATTN_COMPRESSOR_WGATE: "blk.{bid}.attn_compressor_gate",
|
||||
MODEL_TENSOR.ATTN_COMPRESSOR_APE: "blk.{bid}.attn_compressor_ape",
|
||||
@@ -2863,58 +2813,6 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
|
||||
MODEL_TENSOR.NEXTN_SHARED_HEAD_HEAD,
|
||||
MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM,
|
||||
],
|
||||
MODEL_ARCH.QWEN4EXP: [
|
||||
MODEL_TENSOR.TOKEN_EMBD,
|
||||
MODEL_TENSOR.OUTPUT,
|
||||
# no OUTPUT_NORM / ATTN_NORM / ATTN_POST_NORM: hyper-connections replace every layer norm
|
||||
MODEL_TENSOR.HC_HEAD_NORM,
|
||||
MODEL_TENSOR.HC_HEAD_DOWN,
|
||||
MODEL_TENSOR.HC_HEAD_UP,
|
||||
MODEL_TENSOR.HC_ATTN_NORM,
|
||||
MODEL_TENSOR.HC_ATTN_DOWN,
|
||||
MODEL_TENSOR.HC_ATTN_UP,
|
||||
MODEL_TENSOR.HC_ATTN_INJECT,
|
||||
MODEL_TENSOR.HC_FFN_NORM,
|
||||
MODEL_TENSOR.HC_FFN_DOWN,
|
||||
MODEL_TENSOR.HC_FFN_UP,
|
||||
MODEL_TENSOR.HC_FFN_INJECT,
|
||||
# full attention layers: ATTN_Q holds [q|gate] interleaved per head
|
||||
MODEL_TENSOR.ATTN_Q,
|
||||
MODEL_TENSOR.ATTN_Q_NORM,
|
||||
MODEL_TENSOR.ATTN_K,
|
||||
MODEL_TENSOR.ATTN_K_NORM,
|
||||
MODEL_TENSOR.ATTN_V,
|
||||
MODEL_TENSOR.ATTN_OUT,
|
||||
MODEL_TENSOR.INDEXER_Q_PROJ,
|
||||
MODEL_TENSOR.INDEXER_K_PROJ,
|
||||
MODEL_TENSOR.INDEXER_Q_NORM,
|
||||
MODEL_TENSOR.INDEXER_K_NORM,
|
||||
MODEL_TENSOR.ATTN_QKV,
|
||||
MODEL_TENSOR.ATTN_GATE,
|
||||
MODEL_TENSOR.SSM_A,
|
||||
MODEL_TENSOR.SSM_CONV1D,
|
||||
MODEL_TENSOR.SSM_DT,
|
||||
MODEL_TENSOR.SSM_NORM,
|
||||
MODEL_TENSOR.SSM_BETA,
|
||||
MODEL_TENSOR.SSM_ALPHA,
|
||||
MODEL_TENSOR.SSM_OUT,
|
||||
MODEL_TENSOR.FFN_GATE_INP,
|
||||
MODEL_TENSOR.FFN_GATE_INP_SHEXP,
|
||||
MODEL_TENSOR.FFN_UP_SHEXP,
|
||||
MODEL_TENSOR.FFN_DOWN_SHEXP,
|
||||
MODEL_TENSOR.FFN_GATE_SHEXP,
|
||||
MODEL_TENSOR.FFN_DOWN_EXP,
|
||||
MODEL_TENSOR.FFN_UP_EXP,
|
||||
MODEL_TENSOR.FFN_GATE_EXP,
|
||||
MODEL_TENSOR.FFN_GATE_UP_EXP,
|
||||
MODEL_TENSOR.PER_LAYER_TOKEN_EMBD,
|
||||
MODEL_TENSOR.PLE_KEY,
|
||||
MODEL_TENSOR.PLE_VALUE,
|
||||
MODEL_TENSOR.PLE_NORM_KEY,
|
||||
MODEL_TENSOR.PLE_NORM_QUERY,
|
||||
MODEL_TENSOR.PLE_NORM_CONV,
|
||||
MODEL_TENSOR.PLE_CONV1D,
|
||||
],
|
||||
MODEL_ARCH.PLAMO: [
|
||||
MODEL_TENSOR.TOKEN_EMBD,
|
||||
MODEL_TENSOR.OUTPUT_NORM,
|
||||
|
||||
@@ -1008,9 +1008,6 @@ class GGUFWriter:
|
||||
def add_sample_from_anchor(self, value: bool) -> None:
|
||||
self.add_bool(Keys.LLM.SAMPLE_FROM_ANCHOR.format(arch=self.arch), value)
|
||||
|
||||
def add_has_confidence_head(self, value: bool) -> None:
|
||||
self.add_bool(Keys.LLM.HAS_CONFIDENCE_HEAD.format(arch=self.arch), value)
|
||||
|
||||
def add_target_layers(self, value: Sequence[int]) -> None:
|
||||
self.add_array(Keys.LLM.TARGET_LAYERS.format(arch=self.arch), value)
|
||||
|
||||
@@ -1044,40 +1041,6 @@ class GGUFWriter:
|
||||
def add_hyper_connection_epsilon(self, value: float) -> None:
|
||||
self.add_float32(Keys.HyperConnection.EPSILON.format(arch=self.arch), value)
|
||||
|
||||
def add_hyper_connection_low_rank(self, value: int) -> None:
|
||||
self.add_uint32(Keys.HyperConnection.LOW_RANK.format(arch=self.arch), value)
|
||||
|
||||
def add_ple_layers(self, values: Sequence[int]) -> None:
|
||||
self.add_array(Keys.PerLayerEmbedding.LAYERS.format(arch=self.arch), values)
|
||||
|
||||
def add_ple_ngram_size(self, value: int) -> None:
|
||||
self.add_uint32(Keys.PerLayerEmbedding.NGRAM_SIZE.format(arch=self.arch), value)
|
||||
|
||||
def add_ple_heads_per_ngram(self, value: int) -> None:
|
||||
self.add_uint32(Keys.PerLayerEmbedding.HEADS_PER_NGRAM.format(arch=self.arch), value)
|
||||
|
||||
def add_ple_conv_kernel(self, value: int) -> None:
|
||||
self.add_uint32(Keys.PerLayerEmbedding.CONV_KERNEL.format(arch=self.arch), value)
|
||||
|
||||
# multipliers reach ~2.4e13; default INT32 inference would truncate them
|
||||
def _add_u64_array(self, key: str, values: Sequence[int]) -> None:
|
||||
self.add_key_value(key, list(values), GGUFValueType.ARRAY, GGUFValueType.UINT64)
|
||||
|
||||
def add_ple_layer_multipliers(self, values: Sequence[int]) -> None:
|
||||
self._add_u64_array(Keys.PerLayerEmbedding.LAYER_MULTIPLIERS.format(arch=self.arch), values)
|
||||
|
||||
def add_ple_head_offsets(self, values: Sequence[int]) -> None:
|
||||
self._add_u64_array(Keys.PerLayerEmbedding.HEAD_OFFSETS.format(arch=self.arch), values)
|
||||
|
||||
def add_ple_head_vocab_sizes(self, values: Sequence[int]) -> None:
|
||||
self._add_u64_array(Keys.PerLayerEmbedding.HEAD_VOCAB_SIZES.format(arch=self.arch), values)
|
||||
|
||||
def add_ple_eos_token_id(self, value: int) -> None:
|
||||
self.add_uint32(Keys.PerLayerEmbedding.EOS_TOKEN_ID.format(arch=self.arch), value)
|
||||
|
||||
def add_ple_image_token_id(self, value: int) -> None:
|
||||
self.add_uint32(Keys.PerLayerEmbedding.IMAGE_TOKEN_ID.format(arch=self.arch), value)
|
||||
|
||||
def add_attention_scale(self, value: float) -> None:
|
||||
self.add_float32(Keys.Attention.SCALE.format(arch=self.arch), value)
|
||||
|
||||
|
||||
@@ -226,64 +226,3 @@ class LazyNumpyTensor(LazyBase):
|
||||
return eager.tofile(*args, **kwargs)
|
||||
|
||||
# TODO: __array_function__
|
||||
|
||||
|
||||
# Tensor written to file one row-chunk at a time
|
||||
class LazyChunkedTensor:
|
||||
|
||||
def __init__(
|
||||
self, chunks: list[Callable[[], np.ndarray]], shape: tuple[int, ...], dtype: DTypeLike,
|
||||
qtype: Any = None, byteswap: bool = False,
|
||||
):
|
||||
self._chunks = chunks
|
||||
self._qtype = qtype
|
||||
self._byteswap = byteswap
|
||||
self.shape = tuple(shape)
|
||||
self.dtype = np.dtype(dtype)
|
||||
|
||||
@property
|
||||
def nbytes(self) -> int:
|
||||
n = self.dtype.itemsize
|
||||
for d in self.shape:
|
||||
n *= d
|
||||
return n
|
||||
|
||||
def numpy(self) -> LazyChunkedTensor:
|
||||
return self
|
||||
|
||||
def quantize(self, qtype: Any) -> LazyChunkedTensor:
|
||||
from .constants import GGMLQuantizationType
|
||||
from .quants import QuantError, quant_shape_to_byte_shape
|
||||
|
||||
if qtype == GGMLQuantizationType.F32:
|
||||
shape, dtype = self.shape, np.dtype(np.float32)
|
||||
elif qtype == GGMLQuantizationType.F16:
|
||||
shape, dtype = self.shape, np.dtype(np.float16)
|
||||
else:
|
||||
try:
|
||||
shape, dtype = quant_shape_to_byte_shape(self.shape, qtype), np.dtype(np.uint8)
|
||||
except ValueError as e:
|
||||
# raised here and not per chunk, so callers can still fall back to F16
|
||||
raise QuantError(str(e)) from e
|
||||
return LazyChunkedTensor(self._chunks, shape, dtype, qtype, self._byteswap)
|
||||
|
||||
def byteswap(self, inplace: bool = False) -> LazyChunkedTensor:
|
||||
if inplace:
|
||||
raise NotImplementedError("a chunked tensor cannot be byteswapped in place")
|
||||
return LazyChunkedTensor(self._chunks, self.shape, self.dtype, self._qtype, not self._byteswap)
|
||||
|
||||
def tofile(self, *args, **kwargs) -> None:
|
||||
from .quants import quantize
|
||||
|
||||
written = 0
|
||||
for load_chunk in self._chunks:
|
||||
chunk = load_chunk()
|
||||
if self._qtype is not None:
|
||||
# exact only because chunks split on rows, and blocks never cross one
|
||||
chunk = quantize(chunk, self._qtype)
|
||||
if self._byteswap:
|
||||
chunk = chunk.byteswap(inplace=False)
|
||||
chunk.tofile(*args, **kwargs)
|
||||
written += chunk.nbytes
|
||||
del chunk
|
||||
assert written == self.nbytes, f"chunked tensor wrote {written} bytes, expected {self.nbytes}"
|
||||
|
||||
@@ -2708,65 +2708,6 @@ class TensorNameMap:
|
||||
"model.layers.{bid}.post_attention_layernorm",
|
||||
),
|
||||
},
|
||||
MODEL_ARCH.QWEN4EXP: {
|
||||
MODEL_TENSOR.HC_ATTN_NORM: (
|
||||
"model.layers.{bid}.attn_hyper_connection.hc_norm",
|
||||
),
|
||||
MODEL_TENSOR.HC_ATTN_DOWN: (
|
||||
"model.layers.{bid}.attn_hyper_connection.input_mix_weight_down",
|
||||
),
|
||||
MODEL_TENSOR.HC_ATTN_UP: (
|
||||
"model.layers.{bid}.attn_hyper_connection.input_mix_weight_up",
|
||||
),
|
||||
MODEL_TENSOR.HC_ATTN_INJECT: (
|
||||
"model.layers.{bid}.attn_hyper_connection.block_inject_weight",
|
||||
),
|
||||
MODEL_TENSOR.HC_FFN_NORM: (
|
||||
"model.layers.{bid}.mlp_hyper_connection.hc_norm",
|
||||
),
|
||||
MODEL_TENSOR.HC_FFN_DOWN: (
|
||||
"model.layers.{bid}.mlp_hyper_connection.input_mix_weight_down",
|
||||
),
|
||||
MODEL_TENSOR.HC_FFN_UP: (
|
||||
"model.layers.{bid}.mlp_hyper_connection.input_mix_weight_up",
|
||||
),
|
||||
MODEL_TENSOR.HC_FFN_INJECT: (
|
||||
"model.layers.{bid}.mlp_hyper_connection.block_inject_weight",
|
||||
),
|
||||
MODEL_TENSOR.HC_HEAD_NORM: (
|
||||
"model.hyper_connection_mixer.hc_norm",
|
||||
),
|
||||
MODEL_TENSOR.HC_HEAD_DOWN: (
|
||||
"model.hyper_connection_mixer.input_mix_weight_down",
|
||||
),
|
||||
MODEL_TENSOR.HC_HEAD_UP: (
|
||||
"model.hyper_connection_mixer.input_mix_weight_up",
|
||||
),
|
||||
MODEL_TENSOR.INDEXER_Q_NORM: (
|
||||
"model.layers.{bid}.self_attn.indexer.q_layernorm",
|
||||
),
|
||||
MODEL_TENSOR.INDEXER_K_NORM: (
|
||||
"model.layers.{bid}.self_attn.indexer.k_layernorm",
|
||||
),
|
||||
MODEL_TENSOR.PLE_KEY: (
|
||||
"model.layers.{bid}.ple.key_proj",
|
||||
),
|
||||
MODEL_TENSOR.PLE_VALUE: (
|
||||
"model.layers.{bid}.ple.value_proj",
|
||||
),
|
||||
MODEL_TENSOR.PLE_NORM_KEY: (
|
||||
"model.layers.{bid}.ple.norm_key",
|
||||
),
|
||||
MODEL_TENSOR.PLE_NORM_QUERY: (
|
||||
"model.layers.{bid}.ple.norm_query",
|
||||
),
|
||||
MODEL_TENSOR.PLE_NORM_CONV: (
|
||||
"model.layers.{bid}.ple.norm_conv",
|
||||
),
|
||||
MODEL_TENSOR.PLE_CONV1D: (
|
||||
"model.layers.{bid}.ple.conv1d",
|
||||
),
|
||||
},
|
||||
}
|
||||
|
||||
mapping: dict[str, tuple[MODEL_TENSOR, str]]
|
||||
|
||||
@@ -31,7 +31,6 @@ add_library(llama
|
||||
llama-memory.cpp
|
||||
llama-memory-hybrid.cpp
|
||||
llama-memory-hybrid-iswa.cpp
|
||||
llama-memory-hybrid-idx.cpp
|
||||
llama-memory-recurrent.cpp
|
||||
llama-mmap.cpp
|
||||
llama-model-loader.cpp
|
||||
|
||||
@@ -40,7 +40,6 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
|
||||
{ LLM_ARCH_QWEN3VLMOE, "qwen3vlmoe" },
|
||||
{ LLM_ARCH_QWEN35, "qwen35" },
|
||||
{ LLM_ARCH_QWEN35MOE, "qwen35moe" },
|
||||
{ LLM_ARCH_QWEN4EXP, "qwen4exp" },
|
||||
{ LLM_ARCH_PHI2, "phi2" },
|
||||
{ LLM_ARCH_PHI3, "phi3" },
|
||||
{ LLM_ARCH_PHIMOE, "phimoe" },
|
||||
@@ -294,17 +293,6 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
|
||||
{ LLM_KV_HYPER_CONNECTION_COUNT, "%s.hyper_connection.count" },
|
||||
{ LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS, "%s.hyper_connection.sinkhorn_iterations" },
|
||||
{ LLM_KV_HYPER_CONNECTION_EPSILON, "%s.hyper_connection.epsilon" },
|
||||
{ LLM_KV_HYPER_CONNECTION_LOW_RANK, "%s.hyper_connection.low_rank" },
|
||||
|
||||
{ LLM_KV_PLE_LAYERS, "%s.ple.layers" },
|
||||
{ LLM_KV_PLE_NGRAM_SIZE, "%s.ple.ngram_size" },
|
||||
{ LLM_KV_PLE_HEADS_PER_NGRAM, "%s.ple.heads_per_ngram" },
|
||||
{ LLM_KV_PLE_CONV_KERNEL, "%s.ple.conv_kernel" },
|
||||
{ LLM_KV_PLE_LAYER_MULTIPLIERS, "%s.ple.layer_multipliers" },
|
||||
{ LLM_KV_PLE_HEAD_OFFSETS, "%s.ple.head_offsets" },
|
||||
{ LLM_KV_PLE_HEAD_VOCAB_SIZES, "%s.ple.head_vocab_sizes" },
|
||||
{ LLM_KV_PLE_EOS_TOKEN_ID, "%s.ple.eos_token_id" },
|
||||
{ LLM_KV_PLE_IMAGE_TOKEN_ID, "%s.ple.image_token_id" },
|
||||
|
||||
{ LLM_KV_HASH_LAYER_COUNT, "%s.hash_layer_count" },
|
||||
|
||||
@@ -518,29 +506,12 @@ static const std::map<llm_tensor, const char *> LLM_TENSOR_NAMES = {
|
||||
{ LLM_TENSOR_HC_HEAD_FN, "output_hc_fn" },
|
||||
{ LLM_TENSOR_HC_HEAD_BASE, "output_hc_base" },
|
||||
{ LLM_TENSOR_HC_HEAD_SCALE, "output_hc_scale" },
|
||||
{ LLM_TENSOR_HC_HEAD_NORM, "output_hc_norm" },
|
||||
{ LLM_TENSOR_HC_HEAD_DOWN, "output_hc_down" },
|
||||
{ LLM_TENSOR_HC_HEAD_UP, "output_hc_up" },
|
||||
{ LLM_TENSOR_HC_ATTN_FN, "blk.%d.hc_attn_fn" },
|
||||
{ LLM_TENSOR_HC_ATTN_BASE, "blk.%d.hc_attn_base" },
|
||||
{ LLM_TENSOR_HC_ATTN_SCALE, "blk.%d.hc_attn_scale" },
|
||||
{ LLM_TENSOR_HC_FFN_FN, "blk.%d.hc_ffn_fn" },
|
||||
{ LLM_TENSOR_HC_FFN_BASE, "blk.%d.hc_ffn_base" },
|
||||
{ LLM_TENSOR_HC_FFN_SCALE, "blk.%d.hc_ffn_scale" },
|
||||
{ LLM_TENSOR_HC_ATTN_NORM, "blk.%d.hc_attn_norm" },
|
||||
{ LLM_TENSOR_HC_ATTN_DOWN, "blk.%d.hc_attn_down" },
|
||||
{ LLM_TENSOR_HC_ATTN_UP, "blk.%d.hc_attn_up" },
|
||||
{ LLM_TENSOR_HC_ATTN_INJECT, "blk.%d.hc_attn_inject" },
|
||||
{ LLM_TENSOR_HC_FFN_NORM, "blk.%d.hc_ffn_norm" },
|
||||
{ LLM_TENSOR_HC_FFN_DOWN, "blk.%d.hc_ffn_down" },
|
||||
{ LLM_TENSOR_HC_FFN_UP, "blk.%d.hc_ffn_up" },
|
||||
{ LLM_TENSOR_HC_FFN_INJECT, "blk.%d.hc_ffn_inject" },
|
||||
{ LLM_TENSOR_PLE_KEY, "blk.%d.ple_key" },
|
||||
{ LLM_TENSOR_PLE_VALUE, "blk.%d.ple_value" },
|
||||
{ LLM_TENSOR_PLE_NORM_KEY, "blk.%d.ple_norm_key" },
|
||||
{ LLM_TENSOR_PLE_NORM_QUERY, "blk.%d.ple_norm_query" },
|
||||
{ LLM_TENSOR_PLE_NORM_CONV, "blk.%d.ple_norm_conv" },
|
||||
{ LLM_TENSOR_PLE_CONV1D, "blk.%d.ple_conv1d" },
|
||||
{ LLM_TENSOR_ATTN_COMPRESSOR_WKV, "blk.%d.attn_compressor_kv" },
|
||||
{ LLM_TENSOR_ATTN_COMPRESSOR_WGATE, "blk.%d.attn_compressor_gate" },
|
||||
{ LLM_TENSOR_ATTN_COMPRESSOR_APE, "blk.%d.attn_compressor_ape" },
|
||||
@@ -746,29 +717,12 @@ static const std::map<llm_tensor, llm_tensor_info> LLM_TENSOR_INFOS = {
|
||||
{LLM_TENSOR_HC_HEAD_FN, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_HEAD_BASE, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_ADD}},
|
||||
{LLM_TENSOR_HC_HEAD_SCALE, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_HEAD_NORM, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_HEAD_DOWN, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_HEAD_UP, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_ATTN_FN, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_ATTN_BASE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_ADD}},
|
||||
{LLM_TENSOR_HC_ATTN_SCALE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_FFN_FN, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_FFN_BASE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_ADD}},
|
||||
{LLM_TENSOR_HC_FFN_SCALE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_ATTN_NORM, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_ATTN_DOWN, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_ATTN_UP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_ATTN_INJECT, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_FFN_NORM, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_HC_FFN_DOWN, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_FFN_UP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_HC_FFN_INJECT, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_PLE_KEY, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_PLE_VALUE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_PLE_NORM_KEY, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_PLE_NORM_QUERY, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_PLE_NORM_CONV, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL}},
|
||||
{LLM_TENSOR_PLE_CONV1D, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_SSM_CONV}},
|
||||
{LLM_TENSOR_ATTN_COMPRESSOR_WKV, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_ATTN_COMPRESSOR_WGATE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
|
||||
{LLM_TENSOR_ATTN_COMPRESSOR_APE, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_GET_ROWS}},
|
||||
@@ -1075,7 +1029,6 @@ bool llm_arch_is_hybrid(const llm_arch & arch) {
|
||||
case LLM_ARCH_KIMI_K3:
|
||||
case LLM_ARCH_QWEN35:
|
||||
case LLM_ARCH_QWEN35MOE:
|
||||
case LLM_ARCH_QWEN4EXP:
|
||||
case LLM_ARCH_DEEPSEEK4:
|
||||
case LLM_ARCH_MINIMAX_01:
|
||||
return true;
|
||||
|
||||
@@ -45,7 +45,6 @@ enum llm_arch {
|
||||
LLM_ARCH_QWEN3VLMOE,
|
||||
LLM_ARCH_QWEN35,
|
||||
LLM_ARCH_QWEN35MOE,
|
||||
LLM_ARCH_QWEN4EXP,
|
||||
LLM_ARCH_PHI2,
|
||||
LLM_ARCH_PHI3,
|
||||
LLM_ARCH_PHIMOE,
|
||||
@@ -299,17 +298,6 @@ enum llm_kv {
|
||||
LLM_KV_HYPER_CONNECTION_COUNT,
|
||||
LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS,
|
||||
LLM_KV_HYPER_CONNECTION_EPSILON,
|
||||
LLM_KV_HYPER_CONNECTION_LOW_RANK,
|
||||
|
||||
LLM_KV_PLE_LAYERS,
|
||||
LLM_KV_PLE_NGRAM_SIZE,
|
||||
LLM_KV_PLE_HEADS_PER_NGRAM,
|
||||
LLM_KV_PLE_CONV_KERNEL,
|
||||
LLM_KV_PLE_LAYER_MULTIPLIERS,
|
||||
LLM_KV_PLE_HEAD_OFFSETS,
|
||||
LLM_KV_PLE_HEAD_VOCAB_SIZES,
|
||||
LLM_KV_PLE_EOS_TOKEN_ID,
|
||||
LLM_KV_PLE_IMAGE_TOKEN_ID,
|
||||
|
||||
LLM_KV_HASH_LAYER_COUNT,
|
||||
|
||||
@@ -582,29 +570,12 @@ enum llm_tensor {
|
||||
LLM_TENSOR_HC_HEAD_FN,
|
||||
LLM_TENSOR_HC_HEAD_BASE,
|
||||
LLM_TENSOR_HC_HEAD_SCALE,
|
||||
LLM_TENSOR_HC_HEAD_NORM, // qwen4exp
|
||||
LLM_TENSOR_HC_HEAD_DOWN, // qwen4exp
|
||||
LLM_TENSOR_HC_HEAD_UP, // qwen4exp
|
||||
LLM_TENSOR_HC_ATTN_FN,
|
||||
LLM_TENSOR_HC_ATTN_BASE,
|
||||
LLM_TENSOR_HC_ATTN_SCALE,
|
||||
LLM_TENSOR_HC_FFN_FN,
|
||||
LLM_TENSOR_HC_FFN_BASE,
|
||||
LLM_TENSOR_HC_FFN_SCALE,
|
||||
LLM_TENSOR_HC_ATTN_NORM, // qwen4exp
|
||||
LLM_TENSOR_HC_ATTN_DOWN, // qwen4exp
|
||||
LLM_TENSOR_HC_ATTN_UP, // qwen4exp
|
||||
LLM_TENSOR_HC_ATTN_INJECT, // qwen4exp
|
||||
LLM_TENSOR_HC_FFN_NORM, // qwen4exp
|
||||
LLM_TENSOR_HC_FFN_DOWN, // qwen4exp
|
||||
LLM_TENSOR_HC_FFN_UP, // qwen4exp
|
||||
LLM_TENSOR_HC_FFN_INJECT, // qwen4exp
|
||||
LLM_TENSOR_PLE_KEY, // qwen4exp
|
||||
LLM_TENSOR_PLE_VALUE, // qwen4exp
|
||||
LLM_TENSOR_PLE_NORM_KEY, // qwen4exp
|
||||
LLM_TENSOR_PLE_NORM_QUERY, // qwen4exp
|
||||
LLM_TENSOR_PLE_NORM_CONV, // qwen4exp
|
||||
LLM_TENSOR_PLE_CONV1D, // qwen4exp
|
||||
LLM_TENSOR_ATTN_COMPRESSOR_WKV,
|
||||
LLM_TENSOR_ATTN_COMPRESSOR_WGATE,
|
||||
LLM_TENSOR_ATTN_COMPRESSOR_APE,
|
||||
|
||||
+8
-87
@@ -661,19 +661,11 @@ void llama_context::sched_reserve() {
|
||||
|
||||
// reserve again with pp graph to avoid ggml-alloc reallocations during inference
|
||||
{
|
||||
// TODO: the worst case graph is not always reached for `n_seqs > 1`
|
||||
// need to implement a more robust mechanism that tries a few different inputs and analyzes the results
|
||||
ggml_cgraph * gf = nullptr;
|
||||
switch (model.arch) {
|
||||
case LLM_ARCH_MINIMAX_01:
|
||||
// the `inp_diag_decay` tensor size scales with `n_seq_tokens^2` which
|
||||
// makes `n_seqs == 1` use more memory for the compute graph compared to `n_seqs > 1`
|
||||
gf = graph_reserve(n_tokens, 1, n_outputs_pp, mctx.get(), model.hparams.no_alloc);
|
||||
break;
|
||||
default:
|
||||
gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get(), model.hparams.no_alloc);
|
||||
};
|
||||
|
||||
// TODO: not sure if the following graph would be worst case for multi-stream KV caches:
|
||||
//
|
||||
// auto * gf = graph_reserve(n_tokens, 1, n_tokens, mctx.get());
|
||||
//
|
||||
auto * gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get(), model.hparams.no_alloc);
|
||||
if (!gf) {
|
||||
throw std::runtime_error("failed to allocate compute pp buffers");
|
||||
}
|
||||
@@ -2309,7 +2301,6 @@ uint32_t llama_context::graph_max_nodes(uint32_t n_tokens) const {
|
||||
model.arch == LLM_ARCH_BAILINGMOE3 ||
|
||||
model.arch == LLM_ARCH_QWEN35 ||
|
||||
model.arch == LLM_ARCH_QWEN35MOE ||
|
||||
model.arch == LLM_ARCH_QWEN4EXP ||
|
||||
model.arch == LLM_ARCH_DEEPSEEK4 ||
|
||||
(model.arch == LLM_ARCH_DFLASH && model.hparams.dsv4_hc_mult > 0) ||
|
||||
model.arch == LLM_ARCH_NANBEIGE ||
|
||||
@@ -2900,83 +2891,13 @@ public:
|
||||
for (auto & [buft, mbuf] : mbufs_new) {
|
||||
const auto & mbuf_cur = mbufs.at(buft);
|
||||
|
||||
if (!mbuf_cur.buf || mbuf_cur.total_size != mbuf.total_size) {
|
||||
if (!mbuf_cur.buf || mbuf_cur.n_tensors != mbuf.n_tensors || mbuf_cur.total_size != mbuf.total_size) {
|
||||
GGML_ABORT("%s: memory buffer mismatch\n", __func__);
|
||||
}
|
||||
|
||||
if (mbuf_cur.n_tensors == mbuf.n_tensors) {
|
||||
// same chunking: copy 1:1 by index
|
||||
for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
|
||||
GGML_ASSERT(ggml_nbytes(mbuf_cur.cpy[i]) == ggml_nbytes(mbuf.org[i]));
|
||||
ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf.org[i]);
|
||||
}
|
||||
continue;
|
||||
for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
|
||||
ggml_backend_tensor_copy(mbuf_cur.cpy[i], mbuf.org[i]);
|
||||
}
|
||||
|
||||
// different chunking: copy the write-side data (mbuf_cur.cpy) into the read-side targets (mbuf.org)
|
||||
// with a byte cursor. Write and read enumerate the same logical data in the same order but may chunk
|
||||
// it differently, so copy across tensor boundaries rather than 1:1 by index.
|
||||
const size_t total = mbuf_cur.total_size;
|
||||
|
||||
ggml_init_params params_scratch = {
|
||||
/*.mem_size =*/ 2*(mbuf_cur.cpy.size() + mbuf.org.size())*ggml_tensor_overhead(),
|
||||
/*.mem_buffer =*/ NULL,
|
||||
/*.no_alloc =*/ true,
|
||||
};
|
||||
ggml_context * ctx_scratch = ggml_init(params_scratch);
|
||||
|
||||
size_t src_pos = 0;
|
||||
size_t dst_pos = 0;
|
||||
size_t src_j = 0;
|
||||
size_t dst_i = 0;
|
||||
size_t src_base = 0;
|
||||
size_t dst_base = 0;
|
||||
|
||||
while (src_pos < total) {
|
||||
const auto & src_t = mbuf_cur.cpy[src_j];
|
||||
const auto & dst_t = mbuf.org[dst_i];
|
||||
|
||||
const size_t src_size = ggml_nbytes(src_t);
|
||||
const size_t dst_size = ggml_nbytes(dst_t);
|
||||
|
||||
const size_t src_off = src_pos - src_base;
|
||||
const size_t dst_off = dst_pos - dst_base;
|
||||
|
||||
const size_t n_copy = std::min(src_size - src_off, dst_size - dst_off);
|
||||
|
||||
const size_t el = ggml_element_size(src_t);
|
||||
const int64_t n_el = (int64_t) (n_copy / el);
|
||||
|
||||
auto * src_v = ggml_view_1d(ctx_scratch, src_t, n_el, src_off);
|
||||
ggml_backend_view_init(src_v);
|
||||
auto * dst_v = ggml_view_1d(ctx_scratch, dst_t, n_el, dst_off);
|
||||
ggml_backend_view_init(dst_v);
|
||||
|
||||
ggml_backend_tensor_copy(src_v, dst_v);
|
||||
|
||||
src_pos += n_copy;
|
||||
dst_pos += n_copy;
|
||||
|
||||
if (src_pos - src_base == src_size) {
|
||||
src_base = src_pos;
|
||||
++src_j;
|
||||
}
|
||||
if (dst_pos - dst_base == dst_size) {
|
||||
dst_base = dst_pos;
|
||||
++dst_i;
|
||||
}
|
||||
}
|
||||
|
||||
GGML_ASSERT(src_pos == total && dst_pos == total);
|
||||
// any tensors left unvisited hold no data
|
||||
for (size_t i = src_j; i < mbuf_cur.cpy.size(); ++i) {
|
||||
GGML_ASSERT(ggml_nbytes(mbuf_cur.cpy[i]) == 0);
|
||||
}
|
||||
for (size_t i = dst_i; i < mbuf.org.size(); ++i) {
|
||||
GGML_ASSERT(ggml_nbytes(mbuf.org[i]) == 0);
|
||||
}
|
||||
|
||||
ggml_free(ctx_scratch);
|
||||
}
|
||||
|
||||
GGML_ASSERT(buf_size == 0);
|
||||
|
||||
+1
-22
@@ -201,11 +201,7 @@ uint32_t llama_hparams::n_embd_r() const {
|
||||
// TODO: maybe support other convolution strides than 1
|
||||
// NOTE: since the first column of the conv_state is shifted out each time, it's not actually needed
|
||||
// Corresponds to Mamba's conv_states size
|
||||
const uint32_t n_conv = (ssm_d_conv > 0 ? ssm_d_conv - 1 : 0) * (ssm_d_inner + 2*ssm_n_group*ssm_d_state);
|
||||
|
||||
// PLE conv history needs its own row: Meta splits cache_r_l by head, so a history packed behind the first is unaddressable
|
||||
// it lives in cache_ple_r_l instead, mirrored like the rest of the PLE module
|
||||
return n_conv;
|
||||
return (ssm_d_conv > 0 ? ssm_d_conv - 1 : 0) * (ssm_d_inner + 2*ssm_n_group*ssm_d_state);
|
||||
}
|
||||
|
||||
uint32_t llama_hparams::n_embd_s() const {
|
||||
@@ -240,23 +236,6 @@ bool llama_hparams::is_recr(uint32_t il) const {
|
||||
GGML_ABORT("%s: il (%u) out of bounds (n_layer_all: %u)\n", __func__, il, n_layer_all);
|
||||
}
|
||||
|
||||
uint32_t llama_hparams::ple_conv_state() const {
|
||||
if (ple_n_heads == 0 || ple_conv_kernel == 0) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
// dilation equals the n-gram size, matching the reference module
|
||||
return (ple_conv_kernel - 1) * ple_ngram_size * dsv4_hc_mult * n_embd;
|
||||
}
|
||||
|
||||
bool llama_hparams::is_ple(uint32_t il) const {
|
||||
if (il < n_layer_all) {
|
||||
return is_ple_impl[il];
|
||||
}
|
||||
|
||||
GGML_ABORT("%s: il (%u) out of bounds (n_layer_all: %u)\n", __func__, il, n_layer_all);
|
||||
}
|
||||
|
||||
uint32_t llama_hparams::n_pos_per_embd() const {
|
||||
return rope_type == LLAMA_ROPE_TYPE_MROPE || rope_type == LLAMA_ROPE_TYPE_IMROPE ? 4 : 1;
|
||||
}
|
||||
|
||||
@@ -3,15 +3,12 @@
|
||||
#include "llama.h"
|
||||
|
||||
#include <array>
|
||||
#include <bitset>
|
||||
#include <cassert>
|
||||
#include <cmath>
|
||||
|
||||
// bump if necessary
|
||||
#define LLAMA_MAX_LAYERS 512
|
||||
#define LLAMA_MAX_EXPERTS 1024 // Kimi K3
|
||||
#define LLAMA_MAX_PLE_NGRAM 8 // qwen4exp
|
||||
#define LLAMA_MAX_PLE_HEADS 64 // qwen4exp
|
||||
|
||||
enum llama_expert_gating_func_type {
|
||||
LLAMA_EXPERT_GATING_FUNC_TYPE_NONE = 0,
|
||||
@@ -279,30 +276,6 @@ struct llama_hparams {
|
||||
float dsv4_hc_eps = 0.0f;
|
||||
std::array<uint32_t, LLAMA_MAX_LAYERS> dsv4_compress_ratios;
|
||||
|
||||
// 0 = full rank (DeepSeek-V4)
|
||||
uint32_t hc_low_rank = 0;
|
||||
|
||||
uint32_t ple_ngram_size = 0;
|
||||
uint32_t ple_heads_per_ngram = 0;
|
||||
uint32_t ple_conv_kernel = 0;
|
||||
uint32_t ple_n_heads = 0; // (ngram_size - 1) * heads_per_ngram
|
||||
uint32_t ple_head_dim = 0;
|
||||
uint32_t ple_eos_token_id = 0;
|
||||
// the id the PLE hash stands in at image positions; 0 makes the loader fall back to EOS
|
||||
uint32_t ple_image_token_id = 0;
|
||||
// the file lists PLE layer indices, so this is never a per-layer gguf array and can hold one bit per layer
|
||||
std::bitset<LLAMA_MAX_LAYERS> is_ple_impl;
|
||||
// the hash multipliers reach ~2e13 and have to stay 64-bit
|
||||
std::array<uint64_t, LLAMA_MAX_PLE_NGRAM> ple_layer_multipliers;
|
||||
// head offsets and vocab sizes are token-space indices; the gather truncates them to int32 anyway
|
||||
std::array<uint32_t, LLAMA_MAX_PLE_HEADS> ple_head_offsets;
|
||||
std::array<uint32_t, LLAMA_MAX_PLE_HEADS> ple_head_vocab_sizes;
|
||||
|
||||
bool is_ple(uint32_t il) const;
|
||||
|
||||
// PLE conv history rows: (kernel - 1) * ngram_size; 0 without a PLE module
|
||||
uint32_t ple_conv_state() const;
|
||||
|
||||
// qwen3vl deepstack
|
||||
// When parsed from GGUF, this implies the first N layers consume the first
|
||||
// N deepstack embeddings. Use deepstack_mapping_arr if you need a more
|
||||
|
||||
+19
-136
@@ -6,7 +6,6 @@
|
||||
#include "llama-context.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <array>
|
||||
#include <cassert>
|
||||
#include <cmath>
|
||||
#include <cstring>
|
||||
@@ -79,8 +78,7 @@ llama_kv_cache::llama_kv_cache(
|
||||
llama_memory_t mem_other,
|
||||
const layer_filter_cb & filter,
|
||||
const layer_reuse_cb & reuse,
|
||||
const layer_share_cb & share,
|
||||
const char * name_tag) :
|
||||
const layer_share_cb & share) :
|
||||
model(model), hparams(hparams), v_trans(v_trans),
|
||||
n_seq_max(n_seq_max), n_stream(unified ? 1 : n_seq_max), n_pad(n_pad), n_swa(n_swa), swa_type(swa_type),
|
||||
other(static_cast<llama_kv_cache *>(mem_other)),
|
||||
@@ -234,8 +232,8 @@ llama_kv_cache::llama_kv_cache(
|
||||
ggml_tensor * k = has_k ? ggml_new_tensor_3d(ctx, type_k, n_embd_k_gqa, kv_size, n_stream) : nullptr;
|
||||
ggml_tensor * v = has_v ? ggml_new_tensor_3d(ctx, type_v, n_embd_v_gqa, kv_size, n_stream) : nullptr;
|
||||
|
||||
has_k && ggml_format_name(k, "cache_%sk_l%d", name_tag, il);
|
||||
has_v && ggml_format_name(v, "cache_%sv_l%d", name_tag, il);
|
||||
has_k && ggml_format_name(k, "cache_k_l%d", il);
|
||||
has_v && ggml_format_name(v, "cache_v_l%d", il);
|
||||
|
||||
std::vector<ggml_tensor *> k_stream;
|
||||
std::vector<ggml_tensor *> v_stream;
|
||||
@@ -1131,7 +1129,7 @@ void llama_kv_cache::apply_ubatch(const slot_info & sinfo, const llama_ubatch &
|
||||
|
||||
cells.pos_set(idx, ubatch.pos[i]);
|
||||
|
||||
if (ubatch.is_pos_2d() || ubatch.token || hparams.ple_n_heads > 0) {
|
||||
if (ubatch.is_pos_2d() || ubatch.token) {
|
||||
llama_kv_cell_ext ext;
|
||||
|
||||
if (ubatch.is_pos_2d()) {
|
||||
@@ -1141,12 +1139,6 @@ void llama_kv_cache::apply_ubatch(const slot_info & sinfo, const llama_ubatch &
|
||||
|
||||
if (ubatch.token) {
|
||||
ext.tok = ubatch.token[i];
|
||||
} else if (hparams.ple_n_heads > 0) {
|
||||
// embd batch (multimodal input) has no token ids, need to pad it with the correct ID for PLE layers
|
||||
// TODO @ngxson : check if we can do the same as gemma 3n / gemma 4
|
||||
ext.tok = hparams.ple_image_token_id != 0
|
||||
? (llama_token) hparams.ple_image_token_id
|
||||
: (llama_token) hparams.ple_eos_token_id;
|
||||
}
|
||||
|
||||
cells.ext_set(idx, ext);
|
||||
@@ -1822,8 +1814,7 @@ void llama_kv_cache::set_input_v_rot(ggml_tensor * dst) const {
|
||||
}
|
||||
|
||||
bool llama_kv_cache::has_cell_ext() const {
|
||||
// M-RoPE needs the 2D position, the PLE n-gram hash needs the token id
|
||||
return hparams.n_pos_per_embd() > 1 || hparams.ple_n_heads > 0;
|
||||
return hparams.n_pos_per_embd() > 1;
|
||||
}
|
||||
|
||||
void llama_kv_cache::get_prev_tokens(const llama_ubatch & ubatch, uint32_t n, std::vector<llama_token> & res) const {
|
||||
@@ -1852,8 +1843,6 @@ void llama_kv_cache::get_prev_tokens(const llama_ubatch & ubatch, uint32_t n, st
|
||||
seqs.set(ubatch.seq_id_unq[s]);
|
||||
}
|
||||
|
||||
const llama_pos w0 = p_min - (llama_pos) n;
|
||||
|
||||
// (seq_id, pos) -> token, for every cell that could be a predecessor of a ubatch token
|
||||
std::unordered_map<uint64_t, llama_token> hist;
|
||||
|
||||
@@ -1861,71 +1850,28 @@ void llama_kv_cache::get_prev_tokens(const llama_ubatch & ubatch, uint32_t n, st
|
||||
return ((uint64_t) seq_id << 32) | (uint32_t) pos;
|
||||
};
|
||||
|
||||
// handle M-RoPE gaps: multiple tokens share the same temporal pos
|
||||
// TODO @ngxson : improve this in the future
|
||||
std::array<std::pair<llama_pos, llama_token>, LLAMA_MAX_SEQ> below;
|
||||
below.fill({ -1, LLAMA_TOKEN_NULL });
|
||||
|
||||
for (uint32_t s = 0; s < n_stream; ++s) {
|
||||
// p_max inclusive: an embd token looks up cells at its own (shared) position
|
||||
v_cells[s].for_each_token_in(seqs, 0, p_max + 1,
|
||||
v_cells[s].for_each_token_in(seqs, p_min - (llama_pos) n, p_max,
|
||||
[&](llama_seq_id seq_id, llama_pos pos, llama_token tok) {
|
||||
if (pos >= w0) {
|
||||
hist[key(seq_id, pos)] = tok;
|
||||
} else if (pos > below[seq_id].first) {
|
||||
below[seq_id] = { pos, tok };
|
||||
}
|
||||
hist[key(seq_id, pos)] = tok;
|
||||
});
|
||||
}
|
||||
|
||||
// the token at pos p, or the nearest earlier one when p falls in an M-RoPE gap
|
||||
const auto lookup = [&](llama_seq_id seq_id, llama_pos p) -> llama_token {
|
||||
for (llama_pos q = p; q >= w0; --q) {
|
||||
const auto it = hist.find(key(seq_id, q));
|
||||
if (it != hist.end()) {
|
||||
return it->second;
|
||||
}
|
||||
}
|
||||
return below[seq_id].second;
|
||||
};
|
||||
|
||||
// an embd (multimodal) ubatch can repeat one position for a whole image, so positions
|
||||
// do not encode the token order; resolve its predecessors by ubatch order instead
|
||||
std::vector<uint32_t> ord; // index among the ubatch tokens of the same seq
|
||||
std::unordered_map<llama_seq_id, std::vector<uint32_t>> seq_idx;
|
||||
|
||||
if (!ubatch.token) {
|
||||
ord.resize(n_tokens);
|
||||
for (uint32_t i = 0; i < n_tokens; ++i) {
|
||||
auto & v = seq_idx[ubatch.seq_id[i][0]];
|
||||
ord[i] = v.size();
|
||||
v.push_back(i);
|
||||
}
|
||||
}
|
||||
|
||||
for (uint32_t i = 0; i < n_tokens; ++i) {
|
||||
// TODO: a token that belongs to more than one sequence has an ambiguous history.
|
||||
// the n-gram architectures have to reject such batches
|
||||
const llama_seq_id seq_id = ubatch.seq_id[i][0];
|
||||
|
||||
for (uint32_t j = 0; j < n; ++j) {
|
||||
const llama_pos d = (llama_pos) (n - j);
|
||||
|
||||
llama_pos p;
|
||||
if (!ubatch.token) {
|
||||
const auto & v = seq_idx[seq_id];
|
||||
const int64_t k = (int64_t) ord[i] - d;
|
||||
// k >= 0: an earlier token of this very ubatch; k < 0: before the chunk
|
||||
p = k >= 0 ? ubatch.pos[v[k]] : ubatch.pos[v[0]] + (llama_pos) k;
|
||||
} else {
|
||||
p = ubatch.pos[i] - d;
|
||||
}
|
||||
|
||||
const llama_pos p = ubatch.pos[i] - (llama_pos) (n - j);
|
||||
if (p < 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
res[i*n + j] = lookup(seq_id, p);
|
||||
const auto it = hist.find(key(seq_id, p));
|
||||
if (it != hist.end()) {
|
||||
res[i*n + j] = it->second;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -2162,15 +2108,6 @@ void llama_kv_cache::state_write(llama_io_write_i & io, llama_seq_id seq_id, lla
|
||||
}
|
||||
|
||||
void llama_kv_cache::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
|
||||
state_read_sinfo(io, seq_id, flags, nullptr, nullptr);
|
||||
}
|
||||
|
||||
void llama_kv_cache::state_read_sinfo(
|
||||
llama_io_read_i & io,
|
||||
llama_seq_id seq_id,
|
||||
llama_state_seq_flags flags,
|
||||
slot_info_vec_t * sinfos_out,
|
||||
const slot_info_vec_t * sinfos_in) {
|
||||
// TODO: refactor [TAG_KV_CACHE_SHARE_CELLS]
|
||||
if (other) {
|
||||
return;
|
||||
@@ -2181,35 +2118,17 @@ const slot_info_vec_t * sinfos_in) {
|
||||
// TODO: fix incosistent handling of `seq_id < 0` and `seq_id == -1` in the codebase [TAG_LLAMA_SEQ_ID_NEG]
|
||||
GGML_ASSERT(seq_id == -1 || (seq_id >= 0 && (size_t) seq_id < seq_to_stream.size()));
|
||||
|
||||
if (sinfos_out) {
|
||||
sinfos_out->assign(n_stream, slot_info{});
|
||||
}
|
||||
|
||||
if (sinfos_in && sinfos_in->size() != n_stream) {
|
||||
throw std::runtime_error("failed to restore kv cache: mirrored slot layout has the wrong stream count");
|
||||
}
|
||||
|
||||
uint32_t n_stream_cur;
|
||||
io.read(&n_stream_cur, sizeof(n_stream_cur));
|
||||
if (n_stream_cur != n_stream) {
|
||||
throw std::runtime_error("n_stream mismatch");
|
||||
}
|
||||
|
||||
// a whole-context restore replaces every stream, so the cache is emptied once here
|
||||
// clear() resets all streams at once, so doing it per stream below would keep only the last one
|
||||
if (seq_id == -1) {
|
||||
clear(true);
|
||||
}
|
||||
|
||||
for (uint32_t s = 0; s < n_stream; ++s) {
|
||||
uint32_t cell_count;
|
||||
io.read(&cell_count, sizeof(cell_count));
|
||||
|
||||
if (cell_count == 0) {
|
||||
// a mirrored cache must be empty here as well, or the two no longer agree cell for cell
|
||||
if (sinfos_in && !(*sinfos_in)[s].empty()) {
|
||||
throw std::runtime_error("failed to restore kv cache: mirrored cache holds cells this one does not");
|
||||
}
|
||||
continue;
|
||||
}
|
||||
|
||||
@@ -2218,7 +2137,7 @@ const slot_info_vec_t * sinfos_in) {
|
||||
slot_info sinfo;
|
||||
|
||||
bool res = true;
|
||||
res = res && state_read_meta(io, strm, cell_count, sinfo, seq_id, sinfos_in ? &(*sinfos_in)[s] : nullptr);
|
||||
res = res && state_read_meta(io, strm, cell_count, sinfo, seq_id);
|
||||
|
||||
try {
|
||||
res = res && state_read_data(io, strm, cell_count, sinfo);
|
||||
@@ -2234,10 +2153,6 @@ const slot_info_vec_t * sinfos_in) {
|
||||
}
|
||||
throw std::runtime_error("failed to restore kv cache");
|
||||
}
|
||||
|
||||
if (sinfos_out) {
|
||||
(*sinfos_out)[s] = sinfo;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -2373,7 +2288,7 @@ void llama_kv_cache::state_write_data(llama_io_write_i & io, const cell_ranges_t
|
||||
}
|
||||
}
|
||||
|
||||
bool llama_kv_cache::state_read_meta(llama_io_read_i & io, uint32_t strm, uint32_t cell_count, slot_info & sinfo, llama_seq_id dest_seq_id, const slot_info * sinfo_in) {
|
||||
bool llama_kv_cache::state_read_meta(llama_io_read_i & io, uint32_t strm, uint32_t cell_count, slot_info & sinfo, llama_seq_id dest_seq_id) {
|
||||
auto & cells = v_cells[strm];
|
||||
auto & head = v_heads[strm];
|
||||
|
||||
@@ -2423,37 +2338,10 @@ bool llama_kv_cache::state_read_meta(llama_io_read_i & io, uint32_t strm, uint32
|
||||
ubatch.seq_id[i] = &dest_seq_id;
|
||||
}
|
||||
|
||||
if (sinfo_in) {
|
||||
// this cache mirrors another one, so it takes that cache's layout instead of searching for its own cells
|
||||
if (sinfo_in->empty() || sinfo_in->n_stream() != 1 || sinfo_in->idxs[0].size() != cell_count) {
|
||||
LLAMA_LOG_ERROR("%s: mirrored slot layout holds %d cells, this cache restores %d\n", __func__,
|
||||
sinfo_in->empty() ? 0 : (int) sinfo_in->idxs[0].size(), cell_count);
|
||||
return false;
|
||||
}
|
||||
|
||||
sinfo = *sinfo_in;
|
||||
|
||||
// the layout is cell indices, so it means the same in both caches only while their streams line up
|
||||
sinfo.s0 = strm;
|
||||
sinfo.s1 = strm;
|
||||
sinfo.strm[0] = strm;
|
||||
|
||||
// seq_rm above freed exactly the cells this sequence held
|
||||
// anything else in the way is a cache that had already drifted, which this restore must not hide
|
||||
for (uint32_t i = 0; i < cell_count; ++i) {
|
||||
const uint32_t idx = sinfo.idxs[0][i];
|
||||
|
||||
if (idx >= cells.size() || !cells.is_empty(idx)) {
|
||||
LLAMA_LOG_ERROR("%s: cell %u of the mirrored slot layout is not free\n", __func__, idx);
|
||||
return false;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
sinfo = find_slot(ubatch, false);
|
||||
if (sinfo.empty()) {
|
||||
LLAMA_LOG_ERROR("%s: failed to find %d available cells in kv cache\n", __func__, cell_count);
|
||||
return false;
|
||||
}
|
||||
sinfo = find_slot(ubatch, false);
|
||||
if (sinfo.empty()) {
|
||||
LLAMA_LOG_ERROR("%s: failed to find %d available cells in kv cache\n", __func__, cell_count);
|
||||
return false;
|
||||
}
|
||||
|
||||
// note: apply_ubatch() rebuilds llama_kv_cell_ext from the ubatch
|
||||
@@ -2479,12 +2367,7 @@ bool llama_kv_cache::state_read_meta(llama_io_read_i & io, uint32_t strm, uint32
|
||||
return false;
|
||||
}
|
||||
|
||||
// the cells go in from 0, so a mirrored cache lands on the same ones as long as it restores the same count. the layout itself carries no more information here
|
||||
if (sinfo_in && (sinfo_in->empty() || sinfo_in->n_stream() != 1 || sinfo_in->idxs[0].size() != cell_count)) {
|
||||
LLAMA_LOG_ERROR("%s: mirrored slot layout holds %d cells, this cache restores %d\n", __func__,
|
||||
sinfo_in->empty() ? 0 : (int) sinfo_in->idxs[0].size(), cell_count);
|
||||
return false;
|
||||
}
|
||||
clear(true);
|
||||
|
||||
for (uint32_t i = 0; i < cell_count; ++i) {
|
||||
llama_pos pos;
|
||||
|
||||
+3
-20
@@ -112,9 +112,7 @@ public:
|
||||
llama_memory_t mem_other,
|
||||
const layer_filter_cb & filter,
|
||||
const layer_reuse_cb & reuse,
|
||||
const layer_share_cb & share,
|
||||
// a model can hold more than one cache, so the tensor names have to stay unique
|
||||
const char * name_tag = "");
|
||||
const layer_share_cb & share);
|
||||
|
||||
~llama_kv_cache() = default;
|
||||
|
||||
@@ -168,17 +166,6 @@ public:
|
||||
|
||||
const llama_kv_cells & get_cells(llama_seq_id seq_id) const;
|
||||
|
||||
// state_read, plus the cells the restored tokens were placed in
|
||||
// a cache that mirrors another one (the qwen4exp indexer) must not search for its own cells: two searches agree only by luck
|
||||
// sinfos_out: if set, filled with the layout used; a stream with no cells leaves an empty entry
|
||||
// sinfos_in : if set, the layout to use instead of searching. one entry per stream, cell count must match the blob
|
||||
void state_read_sinfo(
|
||||
llama_io_read_i & io,
|
||||
llama_seq_id seq_id,
|
||||
llama_state_seq_flags flags,
|
||||
slot_info_vec_t * sinfos_out,
|
||||
const slot_info_vec_t * sinfos_in);
|
||||
|
||||
//
|
||||
// graph_build API
|
||||
//
|
||||
@@ -236,10 +223,7 @@ public:
|
||||
bool has_cell_ext() const;
|
||||
|
||||
// for every token of the ubatch, the ids of the n tokens that precede it in its sequence
|
||||
// example for M-RoPE image case: tokens A B X X X C, where X is a 3-token image at pos 2 spanning positions 2..4:
|
||||
// tok: A B X X X C
|
||||
// pos: 0 1 2 2 2 5
|
||||
// prev, n=2: A -> [NULL, NULL], B -> [NULL, A], 3rd X -> [X, X], C -> [X, X]
|
||||
// entries with no matching cell are set to LLAMA_TOKEN_NULL
|
||||
// note: used by n-gram input embeddings
|
||||
void get_prev_tokens(const llama_ubatch & ubatch, uint32_t n, std::vector<llama_token> & res) const;
|
||||
|
||||
@@ -342,8 +326,7 @@ private:
|
||||
void state_write_meta(llama_io_write_i & io, const cell_ranges_t & cr, llama_seq_id seq_id = -1) const;
|
||||
void state_write_data(llama_io_write_i & io, const cell_ranges_t & cr) const;
|
||||
|
||||
// sinfo_in, when set, replaces the find_slot call: the cells are given by the caller
|
||||
bool state_read_meta(llama_io_read_i & io, uint32_t strm, uint32_t cell_count, slot_info & sinfo, llama_seq_id dest_seq_id = -1, const slot_info * sinfo_in = nullptr);
|
||||
bool state_read_meta(llama_io_read_i & io, uint32_t strm, uint32_t cell_count, slot_info & sinfo, llama_seq_id dest_seq_id = -1);
|
||||
bool state_read_data(llama_io_read_i & io, uint32_t strm, uint32_t cell_count, const slot_info & sinfo);
|
||||
};
|
||||
|
||||
|
||||
@@ -1,465 +0,0 @@
|
||||
#include "llama-memory-hybrid-idx.h"
|
||||
|
||||
#include "llama-impl.h"
|
||||
#include "llama-batch.h"
|
||||
#include "llama-io.h"
|
||||
#include "llama-model.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <cassert>
|
||||
#include <cmath>
|
||||
#include <iterator>
|
||||
#include <stdexcept>
|
||||
|
||||
//
|
||||
// llama_memory_hybrid_idx
|
||||
//
|
||||
|
||||
llama_memory_hybrid_idx::llama_memory_hybrid_idx(
|
||||
const llama_model & model,
|
||||
/* attn */
|
||||
ggml_type type_k,
|
||||
ggml_type type_v,
|
||||
bool v_trans,
|
||||
uint32_t kv_size,
|
||||
uint32_t n_pad,
|
||||
uint32_t n_swa,
|
||||
llama_swa_type swa_type,
|
||||
/* recurrent */
|
||||
ggml_type type_r,
|
||||
ggml_type type_s,
|
||||
uint32_t rs_size,
|
||||
/* common */
|
||||
uint32_t n_seq_max,
|
||||
uint32_t n_rs_seq,
|
||||
bool offload,
|
||||
bool unified,
|
||||
/* layer filters */
|
||||
const layer_filter_cb & filter_attn,
|
||||
const layer_filter_cb & filter_recr,
|
||||
const layer_filter_cb & filter_idx) :
|
||||
llama_memory_hybrid(
|
||||
model,
|
||||
type_k, type_v, v_trans, kv_size, n_pad, n_swa, swa_type,
|
||||
type_r, type_s, rs_size,
|
||||
n_seq_max, n_rs_seq, offload, unified,
|
||||
filter_attn, filter_recr),
|
||||
hparams_idx(model.hparams),
|
||||
mem_idx(filter_idx == nullptr ? nullptr : [&] {
|
||||
// MQA with a single key head of indexer_head_size, as llama_kv_cache_dsa shapes its own
|
||||
std::fill(hparams_idx.n_head_kv_arr.begin(), hparams_idx.n_head_kv_arr.end(), 1);
|
||||
hparams_idx.n_embd_head_k_full = model.hparams.indexer_head_size;
|
||||
|
||||
LLAMA_LOG_INFO("%s: creating indexer KV cache, size = %u cells\n", __func__, kv_size);
|
||||
|
||||
return new llama_kv_cache(
|
||||
model, hparams_idx, type_k, type_v, v_trans, offload, unified,
|
||||
kv_size, n_seq_max, n_pad, n_swa, swa_type,
|
||||
nullptr, filter_idx, nullptr, nullptr, "idx_");
|
||||
}()) {}
|
||||
|
||||
llama_memory_context_ptr llama_memory_hybrid_idx::init_batch(llama_batch_allocr & balloc, uint32_t n_ubatch, bool embd_all) {
|
||||
// note: repeats llama_memory_hybrid::init_batch, as the indexer needs the attention slot infos that the base context hides
|
||||
do {
|
||||
balloc.split_reset();
|
||||
|
||||
// follow the recurrent pattern for creating the ubatch splits
|
||||
std::vector<llama_ubatch> ubatches;
|
||||
|
||||
while (true) {
|
||||
llama_ubatch ubatch;
|
||||
|
||||
if (embd_all) {
|
||||
// if all tokens are output, split by sequence
|
||||
ubatch = balloc.split_seq(n_ubatch);
|
||||
} else {
|
||||
// Use non-sequential split when KV cache is unified (needed for hellaswag/winogrande/multiple-choice)
|
||||
const bool unified = (get_mem_attn()->get_n_stream() == 1);
|
||||
|
||||
// [TAG_RECURRENT_ROLLBACK_SPLITS]
|
||||
// the trailing (1 + n_rs_seq) tokens of each seq must stay in the same ubatch
|
||||
// so that the rollback snapshots remain valid
|
||||
const uint32_t n_rs_seq = get_mem_recr()->n_rs_seq;
|
||||
|
||||
ubatch = balloc.split_equal(n_ubatch, !unified, n_rs_seq > 0 ? n_rs_seq + 1 : 0);
|
||||
}
|
||||
|
||||
if (ubatch.n_tokens == 0) {
|
||||
break;
|
||||
}
|
||||
|
||||
ubatches.push_back(std::move(ubatch)); // NOLINT
|
||||
}
|
||||
|
||||
if (balloc.get_n_used() < balloc.get_n_tokens()) {
|
||||
// failed to find a suitable split
|
||||
break;
|
||||
}
|
||||
|
||||
// prepare the recurrent batches first
|
||||
if (!get_mem_recr()->prepare(ubatches)) {
|
||||
// TODO: will the recurrent cache be in an undefined context at this point?
|
||||
LLAMA_LOG_ERROR("%s: failed to prepare recurrent ubatches\n", __func__);
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(LLAMA_MEMORY_STATUS_FAILED_PREPARE);
|
||||
}
|
||||
|
||||
// prepare the attention cache
|
||||
auto heads_attn = get_mem_attn()->prepare(ubatches);
|
||||
if (heads_attn.empty()) {
|
||||
LLAMA_LOG_ERROR("%s: failed to prepare attention ubatches\n", __func__);
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(LLAMA_MEMORY_STATUS_FAILED_PREPARE);
|
||||
}
|
||||
|
||||
// the indexer uses the attention cache's slot layout; a separate one can drift from it
|
||||
llama_kv_cache::slot_info_vec_t heads_idx;
|
||||
if (mem_idx) {
|
||||
heads_idx = heads_attn;
|
||||
}
|
||||
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(
|
||||
this, std::move(heads_attn), std::move(heads_idx), std::move(ubatches));
|
||||
} while(false);
|
||||
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(LLAMA_MEMORY_STATUS_FAILED_PREPARE);
|
||||
}
|
||||
|
||||
llama_memory_context_ptr llama_memory_hybrid_idx::init_full() {
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(this);
|
||||
}
|
||||
|
||||
llama_memory_context_ptr llama_memory_hybrid_idx::init_update(llama_context * lctx, bool optimize) {
|
||||
return std::make_unique<llama_memory_hybrid_idx_context>(this, lctx, optimize);
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::clear(bool data) {
|
||||
llama_memory_hybrid::clear(data);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->clear(data);
|
||||
}
|
||||
}
|
||||
|
||||
bool llama_memory_hybrid_idx::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) {
|
||||
// same order as llama_memory_hybrid::seq_rm: the recurrent cache can refuse, so try it first
|
||||
if (!get_mem_recr()->seq_rm(seq_id, p0, p1)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_rm(seq_id, p0, p1);
|
||||
}
|
||||
|
||||
return get_mem_attn()->seq_rm(seq_id, p0, p1);
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) {
|
||||
llama_memory_hybrid::seq_cp(seq_id_src, seq_id_dst, p0, p1);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_cp(seq_id_src, seq_id_dst, p0, p1);
|
||||
}
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::seq_keep(llama_seq_id seq_id) {
|
||||
llama_memory_hybrid::seq_keep(seq_id);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_keep(seq_id);
|
||||
}
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::seq_add(llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) {
|
||||
llama_memory_hybrid::seq_add(seq_id, p0, p1, shift);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_add(seq_id, p0, p1, shift);
|
||||
}
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::seq_div(llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) {
|
||||
llama_memory_hybrid::seq_div(seq_id, p0, p1, d);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_div(seq_id, p0, p1, d);
|
||||
}
|
||||
}
|
||||
|
||||
std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid_idx::memory_breakdown() const {
|
||||
std::map<ggml_backend_buffer_type_t, size_t> mb = llama_memory_hybrid::memory_breakdown();
|
||||
|
||||
if (mem_idx) {
|
||||
for (const auto & buft_size : mem_idx->memory_breakdown()) {
|
||||
mb[buft_size.first] += buft_size.second;
|
||||
}
|
||||
}
|
||||
|
||||
return mb;
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::state_write(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) const {
|
||||
llama_memory_hybrid::state_write(io, seq_id, flags);
|
||||
|
||||
// [TAG_HYBRID_IDX_STATE] the indexer section goes last, so it is a pure suffix: an old reader stops early instead of misparsing it
|
||||
// The indexer mirrors the attention cache, so it uses the same PARTIAL_ONLY gate.
|
||||
if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) {
|
||||
if (mem_idx) {
|
||||
mem_idx->state_write(io, seq_id, flags);
|
||||
}
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
|
||||
// note: repeats llama_memory_hybrid::state_read
|
||||
// the indexer needs the attention cache's cells, and a half-failed restore must leave all three caches alike
|
||||
|
||||
// [TAG_HYBRID_IDX_SINFO]
|
||||
// the indexer restore adopts the attention cache's layout instead of searching for cells of its own
|
||||
// two find_slot calls agree only while both caches see the same occupancy, which a restore cannot promise
|
||||
llama_kv_cache::slot_info_vec_t sinfos_attn;
|
||||
|
||||
try {
|
||||
if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) {
|
||||
get_mem_attn()->state_read_sinfo(io, seq_id, flags, mem_idx ? &sinfos_attn : nullptr, nullptr);
|
||||
}
|
||||
|
||||
get_mem_recr()->state_read(io, seq_id, flags);
|
||||
|
||||
// [TAG_HYBRID_IDX_STATE] must mirror the write order in state_write
|
||||
if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) {
|
||||
if (mem_idx) {
|
||||
mem_idx->state_read_sinfo(io, seq_id, flags, nullptr, &sinfos_attn);
|
||||
}
|
||||
}
|
||||
|
||||
} catch (...) {
|
||||
// a half-restored context is the one state the indexer cannot fix by itself: attention holds new cells, the indexer old ones
|
||||
// drop what was being restored from all of them, which is a state they do agree on.
|
||||
state_drop(seq_id);
|
||||
|
||||
throw;
|
||||
}
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx::state_drop(llama_seq_id seq_id) {
|
||||
// dropped directly, not via seq_rm: the recurrent cache may refuse it and then only the other two get cleared
|
||||
if (seq_id < 0) {
|
||||
clear(true);
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
get_mem_attn()->seq_rm(seq_id, -1, -1);
|
||||
get_mem_recr()->seq_rm(seq_id, -1, -1);
|
||||
|
||||
if (mem_idx) {
|
||||
mem_idx->seq_rm(seq_id, -1, -1);
|
||||
}
|
||||
}
|
||||
|
||||
llama_kv_cache * llama_memory_hybrid_idx::get_mem_idx() const {
|
||||
return mem_idx.get();
|
||||
}
|
||||
|
||||
//
|
||||
// llama_memory_hybrid_idx_context
|
||||
//
|
||||
|
||||
// streams in each ubatch's slot info, matching get_k/get_v's `ns`
|
||||
static std::vector<uint32_t> llama_memory_hybrid_idx_ns(const llama_kv_cache::slot_info_vec_t & sinfos) {
|
||||
std::vector<uint32_t> res;
|
||||
res.reserve(sinfos.size());
|
||||
|
||||
for (const auto & sinfo : sinfos) {
|
||||
res.push_back(sinfo.s1 - sinfo.s0 + 1);
|
||||
}
|
||||
|
||||
return res;
|
||||
}
|
||||
|
||||
llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(llama_memory_status status) :
|
||||
llama_memory_hybrid_context(status) {}
|
||||
|
||||
llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(llama_memory_hybrid_idx * mem) :
|
||||
llama_memory_hybrid_context(mem),
|
||||
mem(mem),
|
||||
// graph reservation walks a full context, and qwen4exp builds the sparse attention only when this is set
|
||||
// without it the reserved worst case is the dense graph, so ggml-alloc must grow the buffer on the first decode
|
||||
ns_ubatch(mem->get_mem_idx() == nullptr ?
|
||||
std::vector<uint32_t>() : std::vector<uint32_t>{ mem->get_mem_idx()->get_n_stream() }),
|
||||
ctx_idx(mem->get_mem_idx() == nullptr ? nullptr :
|
||||
new llama_kv_cache_context(mem->get_mem_idx())) {}
|
||||
|
||||
llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(
|
||||
llama_memory_hybrid_idx * mem,
|
||||
llama_context * lctx,
|
||||
bool optimize) :
|
||||
llama_memory_hybrid_context(mem, lctx, optimize),
|
||||
mem(mem) {}
|
||||
|
||||
llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(
|
||||
llama_memory_hybrid_idx * mem,
|
||||
slot_info_vec_t sinfos_attn,
|
||||
slot_info_vec_t sinfos_idx,
|
||||
std::vector<llama_ubatch> ubatches) :
|
||||
// note: the base copies the ubatches; ctx_idx gets a copy of its own
|
||||
llama_memory_hybrid_context(mem, std::move(sinfos_attn), ubatches),
|
||||
mem(mem),
|
||||
ns_ubatch(llama_memory_hybrid_idx_ns(sinfos_idx)),
|
||||
ctx_idx(mem->get_mem_idx() == nullptr ? nullptr :
|
||||
new llama_kv_cache_context(mem->get_mem_idx(), std::move(sinfos_idx), ubatches)) {}
|
||||
|
||||
bool llama_memory_hybrid_idx_context::next() {
|
||||
if (ctx_idx) {
|
||||
ctx_idx->next();
|
||||
}
|
||||
|
||||
++i_cur;
|
||||
|
||||
return llama_memory_hybrid_context::next();
|
||||
}
|
||||
|
||||
bool llama_memory_hybrid_idx_context::apply() {
|
||||
bool res = llama_memory_hybrid_context::apply();
|
||||
|
||||
if (ctx_idx) {
|
||||
res = res & ctx_idx->apply();
|
||||
}
|
||||
|
||||
return res;
|
||||
}
|
||||
|
||||
const llama_kv_cache_context * llama_memory_hybrid_idx_context::get_idx() const {
|
||||
return static_cast<const llama_kv_cache_context *>(ctx_idx.get());
|
||||
}
|
||||
|
||||
uint32_t llama_memory_hybrid_idx_context::get_n_stream() const {
|
||||
GGML_ASSERT(i_cur < ns_ubatch.size());
|
||||
|
||||
return ns_ubatch[i_cur];
|
||||
}
|
||||
|
||||
void llama_memory_hybrid_idx_context::set_input_qsa(
|
||||
ggml_tensor * cell_blk,
|
||||
ggml_tensor * blk_cells,
|
||||
ggml_tensor * blk_pos,
|
||||
ggml_tensor * bias,
|
||||
const llama_ubatch * ubatch,
|
||||
uint32_t ratio,
|
||||
bool blk_bias) const {
|
||||
GGML_ASSERT(ratio > 0);
|
||||
GGML_ASSERT(mem != nullptr && mem->get_mem_idx() != nullptr);
|
||||
|
||||
GGML_ASSERT(ggml_backend_buffer_is_host(cell_blk->buffer));
|
||||
|
||||
const int64_t n_kv = cell_blk->ne[0];
|
||||
const int64_t n_ns = cell_blk->ne[1]; // streams in this ubatch
|
||||
const int64_t n_blocks = blk_pos->ne[0]/(4*n_ns);
|
||||
const int64_t n_tokens = ubatch->n_tokens;
|
||||
const int64_t r = ratio;
|
||||
|
||||
GGML_ASSERT(n_tokens % n_ns == 0);
|
||||
const int64_t n_tps = n_tokens/n_ns; // tokens per stream
|
||||
|
||||
int32_t * dst_cell_blk = (int32_t *) cell_blk->data;
|
||||
int32_t * dst_blk_cells = (int32_t *) blk_cells->data;
|
||||
int32_t * dst_blk_pos = (int32_t *) blk_pos->data;
|
||||
float * dst_bias = (float *) bias->data;
|
||||
|
||||
// block b covers [b*ratio, (b+1)*ratio), so its first token is at b*ratio
|
||||
// all mrope sections carry it: exact for text, approximate for images
|
||||
for (int64_t sec = 0; sec < 4; ++sec) {
|
||||
for (int64_t s = 0; s < n_ns; ++s) {
|
||||
for (int64_t b = 0; b < n_blocks; ++b) {
|
||||
dst_blk_pos[sec*(n_blocks*n_ns) + s*n_blocks + b] = (int32_t) (b*r);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// one pass per stream: cell j is a different token in each, so no mapping is shared
|
||||
std::vector<int32_t> blk_of(n_kv);
|
||||
std::vector<int32_t> filled(n_blocks);
|
||||
|
||||
for (int64_t s = 0; s < n_ns; ++s) {
|
||||
// ubatch index s*n_tps belongs to this stream; ask which cells array it uses
|
||||
const llama_seq_id seq_of_stream = ubatch->seq_id[s*n_tps][0];
|
||||
const auto & cells = mem->get_mem_idx()->get_cells(seq_of_stream);
|
||||
|
||||
int32_t * cur_cell_blk = dst_cell_blk + s*n_kv;
|
||||
int32_t * cur_blk_cells = dst_blk_cells + s*(r*n_blocks);
|
||||
|
||||
// an incomplete block cannot be pooled; the bias below forces those tail cells in
|
||||
// -1 means no usable block, and block 0 only keeps the gather in range
|
||||
std::fill(blk_of.begin(), blk_of.end(), -1);
|
||||
std::fill(filled.begin(), filled.end(), 0);
|
||||
std::fill(cur_blk_cells, cur_blk_cells + r*n_blocks, 0);
|
||||
|
||||
// a cell no block covers needs its own -inf, which a per-block bias cannot carry
|
||||
// every cache path keeps the position below the cell window, so this stays false
|
||||
bool oor = false;
|
||||
|
||||
for (int64_t j = 0; j < n_kv; ++j) {
|
||||
if (cells.is_empty(j)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
const llama_pos p = cells.pos_get(j);
|
||||
const int64_t b = p/r;
|
||||
|
||||
if (b >= n_blocks) {
|
||||
oor = true;
|
||||
continue;
|
||||
}
|
||||
|
||||
blk_of[j] = (int32_t) b;
|
||||
cur_blk_cells[b*r + (p%r)] = (int32_t) j;
|
||||
filled[b]++;
|
||||
}
|
||||
|
||||
GGML_ASSERT((!blk_bias || !oor) && "qsa: cell position runs past the cell window");
|
||||
|
||||
// per-block mode keeps an unpooled cell's real block, so the block's own -inf reaches it
|
||||
// per-cell mode carries that -inf itself and only needs the gather in range
|
||||
for (int64_t j = 0; j < n_kv; ++j) {
|
||||
if (blk_of[j] >= 0 && filled[blk_of[j]] < r && !blk_bias) {
|
||||
blk_of[j] = -1;
|
||||
}
|
||||
cur_cell_blk[j] = blk_of[j] < 0 ? 0 : blk_of[j];
|
||||
}
|
||||
|
||||
for (int64_t ii = 0; ii < n_tps; ++ii) {
|
||||
const int64_t i = s*n_tps + ii;
|
||||
const llama_seq_id seq_id = ubatch->seq_id[i][0];
|
||||
const llama_pos q = ubatch->pos[i];
|
||||
|
||||
// the tail is an incomplete block and is always visible, as in the reference
|
||||
const llama_pos tail_start = (q + 1)/r*r;
|
||||
|
||||
if (blk_bias) {
|
||||
// a block sits wholly inside or outside the tail, so one value covers it
|
||||
// the caller adds the attention mask, which drops empty, foreign and future cells
|
||||
float * cur_blk_bias = dst_bias + i*n_blocks;
|
||||
|
||||
for (int64_t b = 0; b < n_blocks; ++b) {
|
||||
// finite, so it can never meet a -inf and produce a nan
|
||||
cur_blk_bias[b] = b*r >= tail_start ? 1e9f : (filled[b] < r ? -INFINITY : 0.0f);
|
||||
}
|
||||
|
||||
continue;
|
||||
}
|
||||
|
||||
float * cur_bias = dst_bias + i*n_kv;
|
||||
|
||||
for (int64_t j = 0; j < n_kv; ++j) {
|
||||
float v = -INFINITY;
|
||||
|
||||
if (!cells.is_empty(j) && cells.seq_has(j, seq_id) && cells.pos_get(j) <= q) {
|
||||
// finite, so it can never meet a -inf and produce a nan
|
||||
v = cells.pos_get(j) >= tail_start ? 1e9f : (blk_of[j] < 0 ? -INFINITY : 0.0f);
|
||||
}
|
||||
|
||||
cur_bias[j] = v;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -1,156 +0,0 @@
|
||||
#pragma once
|
||||
|
||||
#include "llama-memory-hybrid.h"
|
||||
|
||||
#include <memory>
|
||||
#include <vector>
|
||||
|
||||
//
|
||||
// llama_memory_hybrid_idx
|
||||
//
|
||||
|
||||
// llama_memory_hybrid plus a third cache with one indexer key per token, for block-sparse attention (qwen4exp QSA)
|
||||
// the indexer is a side buffer over the attention cells: same size, padding, streams and slots, so cell j is one token in both
|
||||
|
||||
class llama_memory_hybrid_idx : public llama_memory_hybrid {
|
||||
public:
|
||||
llama_memory_hybrid_idx(
|
||||
const llama_model & model,
|
||||
/* attn */
|
||||
ggml_type type_k,
|
||||
ggml_type type_v,
|
||||
bool v_trans,
|
||||
uint32_t kv_size,
|
||||
uint32_t n_pad,
|
||||
uint32_t n_swa,
|
||||
llama_swa_type swa_type,
|
||||
/* recurrent */
|
||||
ggml_type type_r,
|
||||
ggml_type type_s,
|
||||
uint32_t rs_size,
|
||||
/* common */
|
||||
uint32_t n_seq_max,
|
||||
uint32_t n_rs_seq,
|
||||
bool offload,
|
||||
bool unified,
|
||||
/* layer filters */
|
||||
const layer_filter_cb & filter_attn,
|
||||
const layer_filter_cb & filter_recr,
|
||||
/* the indexer cache exists only if this is given */
|
||||
const layer_filter_cb & filter_idx);
|
||||
|
||||
~llama_memory_hybrid_idx() = default;
|
||||
|
||||
//
|
||||
// llama_memory_i
|
||||
//
|
||||
|
||||
llama_memory_context_ptr init_batch(
|
||||
llama_batch_allocr & balloc,
|
||||
uint32_t n_ubatch,
|
||||
bool embd_all) override;
|
||||
|
||||
llama_memory_context_ptr init_full() override;
|
||||
|
||||
llama_memory_context_ptr init_update(llama_context * lctx, bool optimize) override;
|
||||
|
||||
void clear(bool data) override;
|
||||
|
||||
bool seq_rm (llama_seq_id seq_id, llama_pos p0, llama_pos p1) override;
|
||||
void seq_cp (llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) override;
|
||||
void seq_keep(llama_seq_id seq_id) override;
|
||||
void seq_add (llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) override;
|
||||
void seq_div (llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) override;
|
||||
|
||||
std::map<ggml_backend_buffer_type_t, size_t> memory_breakdown() const override;
|
||||
|
||||
// state write/load
|
||||
|
||||
void state_write(llama_io_write_i & io, llama_seq_id seq_id = -1, llama_state_seq_flags flags = 0) const override;
|
||||
void state_read (llama_io_read_i & io, llama_seq_id seq_id = -1, llama_state_seq_flags flags = 0) override;
|
||||
|
||||
//
|
||||
// llama_memory_hybrid_idx specific API
|
||||
//
|
||||
|
||||
llama_kv_cache * get_mem_idx() const; // nullptr when the model carries no indexer
|
||||
|
||||
private:
|
||||
// forget seq_id (all of it if seq_id < 0) in every cache at once, so a failed restore cannot leave the caches out of step
|
||||
// seq_id < 0 drops the whole context, as the caches themselves do on a failed restore
|
||||
void state_drop(llama_seq_id seq_id);
|
||||
|
||||
// the indexer cache holds one key head per layer, so it needs its own hparams:
|
||||
// llama_kv_cache keeps a reference to what it is given
|
||||
llama_hparams hparams_idx;
|
||||
|
||||
const std::unique_ptr<llama_kv_cache> mem_idx;
|
||||
};
|
||||
|
||||
class llama_memory_hybrid_idx_context : public llama_memory_hybrid_context {
|
||||
public:
|
||||
using slot_info_vec_t = llama_kv_cache::slot_info_vec_t;
|
||||
|
||||
// used for errors
|
||||
explicit llama_memory_hybrid_idx_context(llama_memory_status status);
|
||||
|
||||
// used to create a full-cache context
|
||||
explicit llama_memory_hybrid_idx_context(llama_memory_hybrid_idx * mem);
|
||||
|
||||
// used to create an update context
|
||||
llama_memory_hybrid_idx_context(
|
||||
llama_memory_hybrid_idx * mem,
|
||||
llama_context * lctx,
|
||||
bool optimize);
|
||||
|
||||
// used to create a batch processing context from a batch
|
||||
llama_memory_hybrid_idx_context(
|
||||
llama_memory_hybrid_idx * mem,
|
||||
slot_info_vec_t sinfos_attn,
|
||||
slot_info_vec_t sinfos_idx,
|
||||
std::vector<llama_ubatch> ubatches);
|
||||
|
||||
~llama_memory_hybrid_idx_context() = default;
|
||||
|
||||
//
|
||||
// llama_memory_context_i
|
||||
//
|
||||
|
||||
bool next() override;
|
||||
bool apply() override;
|
||||
|
||||
//
|
||||
// llama_memory_hybrid_idx_context specific API
|
||||
//
|
||||
|
||||
// nullptr with no indexer, and for the update context, which builds no sparse graph
|
||||
const llama_kv_cache_context * get_idx() const;
|
||||
|
||||
// streams in the current slot info, the `ns` of get_k/get_v; 1 if unified
|
||||
uint32_t get_n_stream() const;
|
||||
|
||||
// block-compressed sparse attention (qwen4exp QSA) over the cells of the indexer cache.
|
||||
// Blocks cut the position line, not the cell array, so no caller assumes a contiguous layout:
|
||||
// cell_blk I32 [n_kv, ns] block each cell belongs to
|
||||
// blk_cells I32 [ratio*n_blocks, ns] cells making up each block
|
||||
// blk_pos I32 [4*n_blocks*ns] mrope position rows of each block's first token
|
||||
// bias F32 [n_kv, n_tokens/ns, ns] -inf where invisible, large where always visible
|
||||
// blk_bias asks for the bias per block instead: [n_blocks, n_tokens/ns, ns]
|
||||
// the caller then adds the attention mask, the only part of the bias that varies within a block
|
||||
void set_input_qsa(ggml_tensor * cell_blk, ggml_tensor * blk_cells, ggml_tensor * blk_pos,
|
||||
ggml_tensor * bias, const llama_ubatch * ubatch, uint32_t ratio,
|
||||
bool blk_bias) const;
|
||||
|
||||
private:
|
||||
const llama_memory_hybrid_idx * mem = nullptr;
|
||||
|
||||
// streams per ubatch, read from the slot infos before ctx_idx takes them
|
||||
// declared first, so it is initialised while sinfos_idx is still intact
|
||||
const std::vector<uint32_t> ns_ubatch;
|
||||
|
||||
// null unless the model has an indexer and this is a batch or full context
|
||||
const llama_memory_context_ptr ctx_idx;
|
||||
|
||||
// mirrors the base class's ubatch cursor, which is private there
|
||||
size_t i_cur = 0;
|
||||
};
|
||||
@@ -51,8 +51,7 @@ llama_memory_recurrent::llama_memory_recurrent(
|
||||
auto it = ctx_map.find(buft);
|
||||
if (it == ctx_map.end()) {
|
||||
ggml_init_params params = {
|
||||
// r and s per layer, plus the separate PLE conv row where the model has one
|
||||
/*.mem_size =*/ size_t((hparams.ple_conv_state() > 0 ? 3u : 2u)*n_layer*ggml_tensor_overhead()),
|
||||
/*.mem_size =*/ size_t(2u*n_layer*ggml_tensor_overhead()),
|
||||
/*.mem_buffer =*/ NULL,
|
||||
/*.no_alloc =*/ true,
|
||||
};
|
||||
@@ -72,7 +71,6 @@ llama_memory_recurrent::llama_memory_recurrent(
|
||||
|
||||
r_l.resize(n_layer);
|
||||
s_l.resize(n_layer);
|
||||
p_l.resize(n_layer);
|
||||
|
||||
for (int i = 0; i < n_layer; i++) {
|
||||
if (filter && !filter(i)) {
|
||||
@@ -105,13 +103,6 @@ llama_memory_recurrent::llama_memory_recurrent(
|
||||
ggml_format_name(s, "cache_s_l%d", i);
|
||||
r_l[i] = r;
|
||||
s_l[i] = s;
|
||||
|
||||
// the PLE history needs its own row: Meta must mirror it while the delta-net conv state next door stays split
|
||||
if (hparams.ple_conv_state() > 0 && hparams.is_ple(i)) {
|
||||
ggml_tensor * p = ggml_new_tensor_2d(ctx, type_r, hparams.ple_conv_state(), n_rows);
|
||||
ggml_format_name(p, "cache_ple_r_l%d", i);
|
||||
p_l[i] = p;
|
||||
}
|
||||
}
|
||||
|
||||
// allocate tensors and initialize the buffers to avoid NaNs in the padding
|
||||
@@ -128,13 +119,11 @@ llama_memory_recurrent::llama_memory_recurrent(
|
||||
{
|
||||
const size_t memory_size_r = size_r_bytes();
|
||||
const size_t memory_size_s = size_s_bytes();
|
||||
const size_t memory_size_p = size_p_bytes();
|
||||
|
||||
LLAMA_LOG_INFO("%s: size = %7.2f MiB (%6u cells, %3d layers, %2u seqs %2u rs_seq), R (%s): %7.2f MiB, S (%s): %7.2f MiB, P (%s): %7.2f MiB\n", __func__,
|
||||
(float)(memory_size_r + memory_size_s + memory_size_p) / (1024.0f * 1024.0f), mem_size, n_layer, n_seq_max, n_rs_seq,
|
||||
LLAMA_LOG_INFO("%s: size = %7.2f MiB (%6u cells, %3d layers, %2u seqs %2u rs_seq), R (%s): %7.2f MiB, S (%s): %7.2f MiB\n", __func__,
|
||||
(float)(memory_size_r + memory_size_s) / (1024.0f * 1024.0f), mem_size, n_layer, n_seq_max, n_rs_seq,
|
||||
ggml_type_name(type_r), (float)memory_size_r / (1024.0f * 1024.0f),
|
||||
ggml_type_name(type_s), (float)memory_size_s / (1024.0f * 1024.0f),
|
||||
ggml_type_name(type_r), (float)memory_size_p / (1024.0f * 1024.0f));
|
||||
ggml_type_name(type_s), (float)memory_size_s / (1024.0f * 1024.0f));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -751,18 +740,6 @@ size_t llama_memory_recurrent::size_s_bytes() const {
|
||||
return size_s_bytes;
|
||||
}
|
||||
|
||||
size_t llama_memory_recurrent::size_p_bytes() const {
|
||||
size_t size_p_bytes = 0;
|
||||
|
||||
for (const auto & p : p_l) {
|
||||
if (p != nullptr) {
|
||||
size_p_bytes += ggml_nbytes(p);
|
||||
}
|
||||
}
|
||||
|
||||
return size_p_bytes;
|
||||
}
|
||||
|
||||
void llama_memory_recurrent::state_write(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) const {
|
||||
GGML_UNUSED(flags);
|
||||
|
||||
@@ -922,17 +899,6 @@ void llama_memory_recurrent::state_write_data(llama_io_write_i & io, const std::
|
||||
const size_t buf_size = range_size * r_size_row;
|
||||
io.write_tensor(r_l[il], range.first * r_size_row, buf_size);
|
||||
}
|
||||
|
||||
// the PLE conv history is a second recurrent row, so it has to travel with the first
|
||||
if (p_l[il] != nullptr) {
|
||||
const uint64_t p_size_row = ggml_row_size(p_l[il]->type, hparams.ple_conv_state());
|
||||
io.write(&p_size_row, sizeof(p_size_row));
|
||||
|
||||
for (const auto & range : cell_ranges) {
|
||||
const size_t range_size = range.second - range.first;
|
||||
io.write_tensor(p_l[il], range.first * p_size_row, range_size * p_size_row);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (!s_trans) {
|
||||
@@ -1131,20 +1097,6 @@ bool llama_memory_recurrent::state_read_data(llama_io_read_i & io, uint32_t cell
|
||||
// Read and set the keys for the whole cell range
|
||||
io.read_tensor(r_l[il], head * r_size_row, cell_count * r_size_row);
|
||||
}
|
||||
|
||||
if (p_l[il] != nullptr) {
|
||||
uint64_t p_size_row_ref;
|
||||
io.read(&p_size_row_ref, sizeof(p_size_row_ref));
|
||||
const size_t p_size_row = ggml_row_size(p_l[il]->type, hparams.ple_conv_state());
|
||||
if (p_size_row != p_size_row_ref) {
|
||||
LLAMA_LOG_ERROR("%s: mismatched ple row size (%zu != %zu, layer %d)\n", __func__, p_size_row, (size_t) p_size_row_ref, il);
|
||||
return false;
|
||||
}
|
||||
|
||||
if (cell_count) {
|
||||
io.read_tensor(p_l[il], head * p_size_row, cell_count * p_size_row);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (!s_trans) {
|
||||
@@ -1299,10 +1251,6 @@ ggml_tensor * llama_memory_recurrent_context::get_s_l(int32_t il) const {
|
||||
return mem->s_l[il];
|
||||
}
|
||||
|
||||
ggml_tensor * llama_memory_recurrent_context::get_p_l(int32_t il) const {
|
||||
return mem->p_l[il];
|
||||
}
|
||||
|
||||
int32_t llama_memory_recurrent_context::s_copy(int i) const {
|
||||
const uint32_t cell_idx = i + mem->head;
|
||||
const int32_t src0 = mem->cells[cell_idx].src0;
|
||||
|
||||
@@ -111,8 +111,6 @@ public:
|
||||
// per layer
|
||||
std::vector<ggml_tensor *> r_l;
|
||||
std::vector<ggml_tensor *> s_l;
|
||||
// a second conv history that must stay replicated across devices, so it cannot share the r row
|
||||
std::vector<ggml_tensor *> p_l;
|
||||
|
||||
private:
|
||||
//const llama_model & model;
|
||||
@@ -127,7 +125,6 @@ private:
|
||||
|
||||
size_t size_r_bytes() const;
|
||||
size_t size_s_bytes() const;
|
||||
size_t size_p_bytes() const;
|
||||
|
||||
void state_write_meta(llama_io_write_i & io, const std::vector<std::pair<uint32_t, uint32_t>> & cell_ranges, llama_seq_id seq_id = -1) const;
|
||||
void state_write_data(llama_io_write_i & io, const std::vector<std::pair<uint32_t, uint32_t>> & cell_ranges) const;
|
||||
@@ -173,7 +170,6 @@ public:
|
||||
|
||||
ggml_tensor * get_r_l(int32_t il) const;
|
||||
ggml_tensor * get_s_l(int32_t il) const;
|
||||
ggml_tensor * get_p_l(int32_t il) const;
|
||||
|
||||
int32_t s_copy(int i) const;
|
||||
|
||||
|
||||
@@ -321,11 +321,10 @@ namespace GGUFMeta {
|
||||
case GGUF_TYPE_UINT32:
|
||||
case GGUF_TYPE_INT32: type_ok = (std::is_same<T, int32_t>::value) ||
|
||||
(std::is_same<T, uint32_t>::value); break;
|
||||
case GGUF_TYPE_UINT64: type_ok = (std::is_same<T, uint64_t>::value); break;
|
||||
case GGUF_TYPE_FLOAT32: type_ok = (std::is_same<T, float>::value); break;
|
||||
case GGUF_TYPE_STRING: type_ok = (std::is_same<T, std::string>::value); break;
|
||||
default:
|
||||
throw std::runtime_error(format("%s is not a string/float32/uint32/int32/uint64 array", key.c_str()));
|
||||
throw std::runtime_error(format("%s is not a string/float32/uint32/int32 array", key.c_str()));
|
||||
}
|
||||
if (!type_ok) {
|
||||
throw std::runtime_error(format("%s has wrong array element type %s", key.c_str(), gguf_type_name(arr_info.gt)));
|
||||
@@ -368,11 +367,10 @@ namespace GGUFMeta {
|
||||
case GGUF_TYPE_UINT32:
|
||||
case GGUF_TYPE_INT32: type_ok = (std::is_same<T, int32_t>::value) ||
|
||||
(std::is_same<T, uint32_t>::value); break;
|
||||
case GGUF_TYPE_UINT64: type_ok = (std::is_same<T, uint64_t>::value); break;
|
||||
case GGUF_TYPE_FLOAT32: type_ok = (std::is_same<T, float>::value); break;
|
||||
case GGUF_TYPE_STRING: type_ok = (std::is_same<T, std::string>::value); break;
|
||||
default:
|
||||
throw std::runtime_error(format("%s is not a string/float32/uint32/int32/uint64 array", key.c_str()));
|
||||
throw std::runtime_error(format("%s is not a string/float32/uint32/int32 array", key.c_str()));
|
||||
}
|
||||
if (!type_ok) {
|
||||
throw std::runtime_error(format("%s has wrong array element type %s", key.c_str(), gguf_type_name(arr_info.gt)));
|
||||
@@ -412,9 +410,6 @@ namespace GGUFMeta {
|
||||
template bool llama_model_loader::get_arr<std::array<int32_t, 512>>(enum llm_kv kid, std::array<int32_t, 512> & result, bool required);
|
||||
template bool llama_model_loader::get_arr<std::vector<int32_t>>(enum llm_kv kid, std::vector<int32_t> & result, bool required);
|
||||
template bool llama_model_loader::get_arr<std::array<uint32_t, LLAMA_MAX_LAYERS>>(enum llm_kv kid, std::array<uint32_t, LLAMA_MAX_LAYERS> & result, bool required);
|
||||
template bool llama_model_loader::get_arr<std::vector<uint32_t>>(enum llm_kv kid, std::vector<uint32_t> & result, bool required);
|
||||
template bool llama_model_loader::get_arr<std::array<uint64_t, LLAMA_MAX_PLE_NGRAM>>(enum llm_kv kid, std::array<uint64_t, LLAMA_MAX_PLE_NGRAM> & result, bool required);
|
||||
template bool llama_model_loader::get_arr<std::array<uint64_t, LLAMA_MAX_PLE_HEADS>>(enum llm_kv kid, std::array<uint64_t, LLAMA_MAX_PLE_HEADS> & result, bool required);
|
||||
|
||||
template<typename T>
|
||||
bool llama_model_loader::get_key(const std::string & key, T & result, bool required) {
|
||||
|
||||
@@ -60,10 +60,6 @@ void llama_model_saver::add_kv(const enum llm_kv key, const int32_t value) {
|
||||
gguf_set_val_i32(gguf_ctx, llm_kv(key).c_str(), value);
|
||||
}
|
||||
|
||||
void llama_model_saver::add_kv(const enum llm_kv key, const uint64_t value) {
|
||||
gguf_set_val_u64(gguf_ctx, llm_kv(key).c_str(), value);
|
||||
}
|
||||
|
||||
void llama_model_saver::add_kv(const enum llm_kv key, const float value) {
|
||||
gguf_set_val_f32(gguf_ctx, llm_kv(key).c_str(), value);
|
||||
}
|
||||
@@ -117,8 +113,6 @@ void llama_model_saver::add_kv(const enum llm_kv key, const Container & value, c
|
||||
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_BOOL, value.data(), n_values);
|
||||
} else if (std::is_same<typename Container::value_type, int32_t>::value) {
|
||||
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_INT32, value.data(), n_values);
|
||||
} else if (std::is_same<typename Container::value_type, uint64_t>::value) {
|
||||
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_UINT64, value.data(), n_values);
|
||||
} else if (std::is_same<typename Container::value_type, float>::value) {
|
||||
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_FLOAT32, value.data(), n_values);
|
||||
} else if (std::is_same<Container, std::string>::value) {
|
||||
@@ -130,7 +124,6 @@ void llama_model_saver::add_kv(const enum llm_kv key, const Container & value, c
|
||||
// instantiate for external usage:
|
||||
template void llama_model_saver::add_kv<std::vector<uint32_t>>(const enum llm_kv, const std::vector<uint32_t> &, const bool);
|
||||
template void llama_model_saver::add_kv<std::vector<float>>(const enum llm_kv, const std::vector<float> &, const bool);
|
||||
template void llama_model_saver::add_kv<std::vector<uint64_t>>(const enum llm_kv, const std::vector<uint64_t> &, const bool);
|
||||
|
||||
void llama_model_saver::add_kv(const enum llm_kv key, const std::vector<std::string> & value) {
|
||||
std::vector<const char *> tmp(value.size());
|
||||
@@ -315,32 +308,6 @@ void llama_model_saver::add_kv_from_model() {
|
||||
add_kv(LLM_KV_HYPER_CONNECTION_SINKHORN_ITERATIONS, hparams.dsv4_hc_sinkhorn_iters);
|
||||
add_kv(LLM_KV_HYPER_CONNECTION_EPSILON, hparams.dsv4_hc_eps);
|
||||
add_kv(LLM_KV_HASH_LAYER_COUNT, hparams.dsv4_hash_layer_count);
|
||||
add_kv(LLM_KV_HYPER_CONNECTION_LOW_RANK, hparams.hc_low_rank);
|
||||
|
||||
// the PLE group only means anything whole: write all of it or none
|
||||
if (hparams.ple_n_heads > 0) {
|
||||
std::vector<uint32_t> ple_layers;
|
||||
for (uint32_t il = 0; il < hparams.n_layer_all; ++il) {
|
||||
if (hparams.is_ple_impl[il]) {
|
||||
ple_layers.push_back(il);
|
||||
}
|
||||
}
|
||||
add_kv(LLM_KV_PLE_LAYERS, ple_layers);
|
||||
add_kv(LLM_KV_PLE_NGRAM_SIZE, hparams.ple_ngram_size);
|
||||
add_kv(LLM_KV_PLE_HEADS_PER_NGRAM, hparams.ple_heads_per_ngram);
|
||||
add_kv(LLM_KV_PLE_CONV_KERNEL, hparams.ple_conv_kernel);
|
||||
add_kv(LLM_KV_PLE_EOS_TOKEN_ID, hparams.ple_eos_token_id);
|
||||
add_kv(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, hparams.ple_head_dim);
|
||||
add_kv(LLM_KV_PLE_LAYER_MULTIPLIERS, std::vector<uint64_t>(
|
||||
hparams.ple_layer_multipliers.begin(),
|
||||
hparams.ple_layer_multipliers.begin() + hparams.ple_ngram_size));
|
||||
add_kv(LLM_KV_PLE_HEAD_OFFSETS, std::vector<uint64_t>(
|
||||
hparams.ple_head_offsets.begin(),
|
||||
hparams.ple_head_offsets.begin() + hparams.ple_n_heads));
|
||||
add_kv(LLM_KV_PLE_HEAD_VOCAB_SIZES, std::vector<uint64_t>(
|
||||
hparams.ple_head_vocab_sizes.begin(),
|
||||
hparams.ple_head_vocab_sizes.begin() + hparams.ple_n_heads));
|
||||
}
|
||||
|
||||
const float rope_scaling_factor = hparams.rope_freq_scale_train == 1.0f ? 0.0f : 1.0f/hparams.rope_freq_scale_train;
|
||||
|
||||
@@ -475,10 +442,6 @@ void llama_model_saver::add_tensors_from_model() {
|
||||
add_tensor(model->hc_head_fn);
|
||||
add_tensor(model->hc_head_base);
|
||||
add_tensor(model->hc_head_scale);
|
||||
add_tensor(model->per_layer_tok_embd);
|
||||
add_tensor(model->hc_head_norm);
|
||||
add_tensor(model->hc_head_down);
|
||||
add_tensor(model->hc_head_up);
|
||||
|
||||
for (const struct llama_layer & layer : model->layers) {
|
||||
for (size_t i = 0; i < sizeof(layer)/sizeof(struct ggml_tensor *); ++i) {
|
||||
|
||||
@@ -21,7 +21,6 @@ struct llama_model_saver {
|
||||
|
||||
void add_kv(enum llm_kv key, uint32_t value);
|
||||
void add_kv(enum llm_kv key, int32_t value);
|
||||
void add_kv(enum llm_kv key, uint64_t value);
|
||||
void add_kv(enum llm_kv key, float value);
|
||||
void add_kv(enum llm_kv key, bool value);
|
||||
void add_kv(enum llm_kv key, const char * value);
|
||||
|
||||
+3
-54
@@ -16,7 +16,6 @@
|
||||
#include "llama-kv-cache-dsv4.h"
|
||||
#include "llama-memory-hybrid.h"
|
||||
#include "llama-memory-hybrid-iswa.h"
|
||||
#include "llama-memory-hybrid-idx.h"
|
||||
#include "llama-memory-recurrent.h"
|
||||
|
||||
#include "llama.h"
|
||||
@@ -320,8 +319,6 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
|
||||
return new llama_model_qwen35(params);
|
||||
case LLM_ARCH_QWEN35MOE:
|
||||
return new llama_model_qwen35moe(params);
|
||||
case LLM_ARCH_QWEN4EXP:
|
||||
return new llama_model_qwen4exp(params);
|
||||
case LLM_ARCH_MISTRAL3:
|
||||
return new llama_model_mistral3(params);
|
||||
case LLM_ARCH_EAGLE3:
|
||||
@@ -379,7 +376,6 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
|
||||
static const std::regex pattern_qkv_bias ("blk\\.\\d*\\.attn_qkv.bias");
|
||||
static const std::regex pattern_qk_norm ("blk\\.\\d*\\.attn_(q|k)_norm\\.weight");
|
||||
static const std::regex pattern_kv_cache ("cache_(k|v)_l\\d*");
|
||||
static const std::regex pattern_idx_cache ("cache_idx_(k|v)_l\\d*");
|
||||
static const std::regex pattern_dsv4_state ("dsv4_(csa|hca|lid)_state_(kv|score)_l\\d*");
|
||||
static const std::regex pattern_attn_sinks ("blk\\.\\d*\\.attn_sinks.weight");
|
||||
static const std::regex pattern_attn_out_weight ("blk\\.\\d*\\.attn_output.weight");
|
||||
@@ -395,7 +391,6 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
|
||||
static const std::regex pattern_ssm_beta ("blk\\.\\d*\\.ssm_beta.weight");
|
||||
static const std::regex pattern_ssm_beta_alpha ("blk\\.\\d*\\.ssm_ba.weight");
|
||||
static const std::regex pattern_r_cache ("cache_r_l\\d*");
|
||||
static const std::regex pattern_ple_r_cache ("cache_ple_r_l\\d*");
|
||||
static const std::regex pattern_s_cache ("cache_s_l\\d*");
|
||||
static const std::regex pattern_ssm_conv1d ("blk\\.\\d*\\.ssm_conv1d.weight");
|
||||
static const std::regex pattern_ssm_out_weight ("blk\\.\\d*\\.ssm_out.weight");
|
||||
@@ -493,16 +488,6 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
|
||||
}
|
||||
}
|
||||
|
||||
// the qsa indexer has one key head and its projections are mirrored, so its cache cannot be split
|
||||
if (std::regex_match(tensor_name, pattern_idx_cache)) {
|
||||
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
|
||||
}
|
||||
|
||||
// the PLE table is model-level and its conv is mirrored, so every device runs the whole conv and needs the whole history
|
||||
if (std::regex_match(tensor_name, pattern_ple_r_cache)) {
|
||||
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_MIRRORED);
|
||||
}
|
||||
|
||||
// standard attention
|
||||
if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_kv_weight)) {
|
||||
return get_tensor_config_impl(GGML_BACKEND_SPLIT_AXIS_1, "attn_output.weight", "ssm_out.weight");
|
||||
@@ -591,8 +576,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
|
||||
};
|
||||
|
||||
auto get_split_segments = [&](int axis, uint32_t il) -> std::vector<std::pair<int64_t, uint32_t>> {
|
||||
if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE ||
|
||||
ud->model->arch == LLM_ARCH_QWEN4EXP) {
|
||||
if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE) {
|
||||
const int64_t head_k_dim = hparams.ssm_d_state;
|
||||
const int64_t head_v_dim = hparams.ssm_d_state;
|
||||
const int64_t n_k_heads = hparams.ssm_n_group;
|
||||
@@ -730,8 +714,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
|
||||
if (std::regex_match(tensor_name, pattern_q_weight) || std::regex_match(tensor_name, pattern_q_bias)) {
|
||||
GGML_ASSERT(segments.size() == 1);
|
||||
// some models have Q gate tensors, for those cases the granularity needs to be doubled:
|
||||
if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE ||
|
||||
ud->model->arch == LLM_ARCH_QWEN4EXP) {
|
||||
if (ud->model->arch == LLM_ARCH_QWEN3NEXT || ud->model->arch == LLM_ARCH_QWEN35 || ud->model->arch == LLM_ARCH_QWEN35MOE) {
|
||||
return {std::lcm(2*n_embd_q, blck_size_perf)};
|
||||
}
|
||||
return {granularity_q};
|
||||
@@ -944,7 +927,6 @@ const char * llm_type_name(llm_type type) {
|
||||
case LLM_TYPE_35B_A3B: return "35B.A3B";
|
||||
case LLM_TYPE_48B_A3B: return "48B.A3B";
|
||||
case LLM_TYPE_80B_A3B: return "80B.A3B";
|
||||
case LLM_TYPE_A3B: return "A3B";
|
||||
case LLM_TYPE_100B_A6B: return "100B.A6B";
|
||||
case LLM_TYPE_102B_A12B: return "102B.A12B";
|
||||
case LLM_TYPE_106B_A12B: return "106B.A12B";
|
||||
@@ -2449,10 +2431,6 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params,
|
||||
// layer filters, so pick the right one here
|
||||
llama_memory_hybrid::layer_filter_cb filter_attn = nullptr;
|
||||
llama_memory_hybrid::layer_filter_cb filter_recr = nullptr;
|
||||
// only the sparse-attention architectures use llama_memory_hybrid_idx
|
||||
// a null filter_idx means the GGUF has no indexer tensors
|
||||
llama_memory_hybrid::layer_filter_cb filter_idx = nullptr;
|
||||
const bool needs_mem_idx = (arch == LLM_ARCH_QWEN4EXP);
|
||||
if (arch == LLM_ARCH_FALCON_H1) {
|
||||
filter_attn = [&](uint32_t) { return true; };
|
||||
filter_recr = [&](uint32_t) { return true; };
|
||||
@@ -2463,20 +2441,13 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params,
|
||||
filter_recr = [&](uint32_t il) {
|
||||
return hparams.is_recr(il) && hparams.n_ff(il) == 0;
|
||||
};
|
||||
} else if (arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE || arch == LLM_ARCH_QWEN4EXP || arch == LLM_ARCH_MINIMAX_01) {
|
||||
} else if (arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE || arch == LLM_ARCH_MINIMAX_01) {
|
||||
filter_attn = [&](uint32_t il) {
|
||||
return il < hparams.n_layer() && !hparams.is_recr(il);
|
||||
};
|
||||
filter_recr = [&](uint32_t il) {
|
||||
return il < hparams.n_layer() && hparams.is_recr(il);
|
||||
};
|
||||
|
||||
if (arch == LLM_ARCH_QWEN4EXP && hparams.indexer_head_size > 0) {
|
||||
// QSA runs on the dense-attention layers only
|
||||
filter_idx = [&](uint32_t il) {
|
||||
return il < hparams.n_layer() && !hparams.is_recr(il);
|
||||
};
|
||||
}
|
||||
}
|
||||
|
||||
if (hparams.swa_type != LLAMA_SWA_TYPE_NONE) {
|
||||
@@ -2499,27 +2470,6 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params,
|
||||
/* unified */ cparams.kv_unified,
|
||||
/* filter_attn */ std::move(filter_attn),
|
||||
/* filter_recr */ std::move(filter_recr));
|
||||
} else if (needs_mem_idx) {
|
||||
// sparse attention over a per-token indexer cache, in its own memory type
|
||||
res = new llama_memory_hybrid_idx(
|
||||
/* model */ *this,
|
||||
/* attn_type_k */ params.type_k,
|
||||
/* attn_type_v */ params.type_v,
|
||||
/* attn_v_trans */ !cparams.flash_attn,
|
||||
/* attn_kv_size */ cparams.n_ctx_seq,
|
||||
/* attn_n_pad */ 1,
|
||||
/* attn_n_swa */ hparams.n_swa,
|
||||
/* attn_swa_type */ hparams.swa_type,
|
||||
/* recurrent_type_k */ GGML_TYPE_F32,
|
||||
/* recurrent_type_v */ GGML_TYPE_F32,
|
||||
/* recurrent_kv_size */ std::max((uint32_t) 1, cparams.n_seq_max),
|
||||
/* n_seq_max */ cparams.n_seq_max,
|
||||
/* n_rs_seq */ cparams.n_rs_seq,
|
||||
/* offload */ cparams.offload_kqv,
|
||||
/* unified */ cparams.kv_unified,
|
||||
/* filter_attn */ std::move(filter_attn),
|
||||
/* filter_recr */ std::move(filter_recr),
|
||||
/* filter_idx */ std::move(filter_idx));
|
||||
} else {
|
||||
res = new llama_memory_hybrid(
|
||||
/* model */ *this,
|
||||
@@ -2950,7 +2900,6 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
|
||||
case LLM_ARCH_QWEN3VLMOE:
|
||||
case LLM_ARCH_QWEN35:
|
||||
case LLM_ARCH_QWEN35MOE:
|
||||
case LLM_ARCH_QWEN4EXP:
|
||||
case LLM_ARCH_QWEN3TTS:
|
||||
return LLAMA_ROPE_TYPE_IMROPE;
|
||||
|
||||
|
||||
@@ -129,7 +129,6 @@ enum llm_type {
|
||||
LLM_TYPE_35B_A3B, // Qwen3.5
|
||||
LLM_TYPE_48B_A3B, // Kimi Linear
|
||||
LLM_TYPE_80B_A3B, // Qwen3 Next
|
||||
LLM_TYPE_A3B, // Qwen3.8 Flash Next
|
||||
LLM_TYPE_100B_A6B,
|
||||
LLM_TYPE_102B_A12B, // Solar-Open
|
||||
LLM_TYPE_106B_A12B, // GLM-4.5-Air
|
||||
@@ -561,22 +560,6 @@ struct llama_layer {
|
||||
struct ggml_tensor * index_q_norm = nullptr;
|
||||
struct ggml_tensor * index_k_norm = nullptr;
|
||||
|
||||
struct ggml_tensor * hc_attn_norm = nullptr;
|
||||
struct ggml_tensor * hc_attn_down = nullptr;
|
||||
struct ggml_tensor * hc_attn_up = nullptr;
|
||||
struct ggml_tensor * hc_attn_inject = nullptr;
|
||||
struct ggml_tensor * hc_ffn_norm = nullptr;
|
||||
struct ggml_tensor * hc_ffn_down = nullptr;
|
||||
struct ggml_tensor * hc_ffn_up = nullptr;
|
||||
struct ggml_tensor * hc_ffn_inject = nullptr;
|
||||
|
||||
struct ggml_tensor * ple_key = nullptr;
|
||||
struct ggml_tensor * ple_value = nullptr;
|
||||
struct ggml_tensor * ple_norm_key = nullptr;
|
||||
struct ggml_tensor * ple_norm_query = nullptr;
|
||||
struct ggml_tensor * ple_norm_conv = nullptr;
|
||||
struct ggml_tensor * ple_conv1d = nullptr;
|
||||
|
||||
// gemma4 layer output scale, reused for talkie embedding skip scale
|
||||
struct ggml_tensor * out_scale = nullptr;
|
||||
|
||||
@@ -657,10 +640,6 @@ struct llama_model {
|
||||
struct ggml_tensor * altup_proj = nullptr;
|
||||
struct ggml_tensor * altup_unembd_proj = nullptr;
|
||||
struct ggml_tensor * per_layer_tok_embd = nullptr;
|
||||
|
||||
struct ggml_tensor * hc_head_norm = nullptr;
|
||||
struct ggml_tensor * hc_head_down = nullptr;
|
||||
struct ggml_tensor * hc_head_up = nullptr;
|
||||
struct ggml_tensor * per_layer_model_proj = nullptr;
|
||||
struct ggml_tensor * per_layer_proj_norm = nullptr;
|
||||
|
||||
@@ -672,7 +651,6 @@ struct llama_model {
|
||||
// dspark
|
||||
struct ggml_tensor * dspark_markov_w1 = nullptr;
|
||||
struct ggml_tensor * dspark_markov_w2 = nullptr;
|
||||
struct ggml_tensor * dspark_markov_w2_s = nullptr;
|
||||
struct ggml_tensor * dspark_conf_proj = nullptr;
|
||||
struct ggml_tensor * dspark_conf_proj_b = nullptr;
|
||||
|
||||
|
||||
+1
-21
@@ -399,12 +399,6 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso
|
||||
case GGML_TYPE_Q5_K: return_type = GGML_TYPE_Q5_1; break;
|
||||
case GGML_TYPE_Q6_K: return_type = GGML_TYPE_Q8_0; break;
|
||||
default:
|
||||
if (qk_k <= 32) {
|
||||
// the target is already a 32-block type, so there is no smaller block to demote to
|
||||
// the check below turns it into F16, as a 256-block type does when its fallback does not fit
|
||||
return_type = target_type;
|
||||
break;
|
||||
}
|
||||
throw std::runtime_error(format("no tensor type fallback is defined for type %s",
|
||||
ggml_type_name(target_type)));
|
||||
}
|
||||
@@ -685,21 +679,7 @@ static ggml_type llama_tensor_get_type(quantize_state_impl & qs, const llama_mod
|
||||
return tensor->type;
|
||||
}
|
||||
if (params->token_embedding_type < GGML_TYPE_COUNT && tm.category == tensor_category::TOKEN_EMBD) {
|
||||
// per_layer_token_embd follows --token-embedding-type by default, but it is a large
|
||||
// separate table, so let an explicit --tensor-type name it
|
||||
bool named = false;
|
||||
if (std::strcmp(tensor->name, "per_layer_token_embd.weight") == 0) {
|
||||
const std::string tensor_name(tensor->name);
|
||||
for (const auto & [pattern, qtype] : qs.tensor_type_patterns) {
|
||||
if (std::regex_search(tensor_name, pattern)) {
|
||||
named = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (!named) {
|
||||
return params->token_embedding_type;
|
||||
}
|
||||
return params->token_embedding_type;
|
||||
}
|
||||
if (params->output_tensor_type < GGML_TYPE_COUNT && tm.category == tensor_category::OUTPUT) {
|
||||
return params->output_tensor_type;
|
||||
|
||||
+24
-34
@@ -115,11 +115,10 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
|
||||
if (markov_meta) {
|
||||
const int64_t dspark_markov_rank = markov_meta->ne[0];
|
||||
|
||||
dspark_markov_w1 = create_tensor(tn(LLM_TENSOR_DSPARK_MARKOV_W1, "weight"), { dspark_markov_rank, n_vocab }, 0);
|
||||
dspark_markov_w2 = create_tensor(tn(LLM_TENSOR_DSPARK_MARKOV_W2, "weight"), { dspark_markov_rank, n_vocab_draft }, 0);
|
||||
dspark_markov_w2_s = create_tensor(tn(LLM_TENSOR_DSPARK_MARKOV_W2, "scale"), { 1 }, TENSOR_NOT_REQUIRED);
|
||||
dspark_markov_w1 = create_tensor(tn(LLM_TENSOR_DSPARK_MARKOV_W1, "weight"), { dspark_markov_rank, n_vocab }, 0);
|
||||
dspark_markov_w2 = create_tensor(tn(LLM_TENSOR_DSPARK_MARKOV_W2, "weight"), { dspark_markov_rank, n_vocab_draft }, 0);
|
||||
|
||||
dspark_conf_proj = create_tensor(tn(LLM_TENSOR_DSPARK_CONF_PROJ, "weight"), { n_embd + dspark_markov_rank, 1 }, TENSOR_NOT_REQUIRED);
|
||||
dspark_conf_proj = create_tensor(tn(LLM_TENSOR_DSPARK_CONF_PROJ, "weight"), { n_embd + dspark_markov_rank, 1 }, 0);
|
||||
dspark_conf_proj_b = create_tensor(tn(LLM_TENSOR_DSPARK_CONF_PROJ, "bias"), { 1 }, TENSOR_NOT_REQUIRED);
|
||||
|
||||
LLAMA_LOG_INFO("%s: DFlash with DSpark markov head (rank = %lld)\n", __func__, (long long) dspark_markov_rank);
|
||||
@@ -220,9 +219,6 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
|
||||
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), { n_embd_head_k }, 0);
|
||||
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), { n_embd_head_k }, 0);
|
||||
|
||||
// optional per-head attention sinks (e.g. Nemotron DSpark)
|
||||
layer.attn_sinks = create_tensor(tn(LLM_TENSOR_ATTN_SINKS, "weight", i), { n_head }, TENSOR_NOT_REQUIRED);
|
||||
|
||||
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), { n_embd }, 0);
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), { n_embd, n_ff }, 0);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd }, 0);
|
||||
@@ -294,10 +290,7 @@ static void build_dspark_markov_head(llm_graph_context & g, const llama_model &
|
||||
|
||||
ggml_tensor * w1 = model.dspark_markov_w1;
|
||||
ggml_tensor * w2 = model.dspark_markov_w2;
|
||||
GGML_ASSERT(w1 && w2 && "DSpark markov weights not loaded");
|
||||
|
||||
// confidence head is optional
|
||||
const bool has_conf = model.dspark_conf_proj != nullptr;
|
||||
GGML_ASSERT(w1 && w2 && model.dspark_conf_proj && "DSpark markov/confidence weights not loaded");
|
||||
|
||||
ggml_tensor * base = res->t_logits; // [n_vocab, n_tokens]
|
||||
const int64_t n_vocab = base->ne[0];
|
||||
@@ -328,22 +321,23 @@ static void build_dspark_markov_head(llm_graph_context & g, const llama_model &
|
||||
ggml_tensor * prev = ggml_view_2d(ctx0, tokens, 1, n_blocks, token_stride, 0);
|
||||
prev = ggml_cont_1d(ctx0, prev, n_blocks);
|
||||
|
||||
// confidence head input: predicts per-position acceptance
|
||||
ggml_tensor * conf_inp = res->t_embd; // [n_embd, n_tok]
|
||||
|
||||
ggml_tensor * cat = nullptr;
|
||||
ggml_tensor * cat_conf = nullptr;
|
||||
|
||||
if (!sample_from_anchor) {
|
||||
// bonus anchor slot: pass the logits through unbiased, pad the (unread) confidence column
|
||||
cat = ggml_cont(ctx0, ggml_view_2d(ctx0, base, n_vocab, n_blocks, base_stride, 0));
|
||||
if (has_conf) {
|
||||
cat_conf = ggml_sigmoid(ctx0, ggml_cont(ctx0, ggml_view_2d(ctx0, base, 1, n_blocks, base_stride, 0)));
|
||||
}
|
||||
cat = ggml_cont(ctx0, ggml_view_2d(ctx0, base, n_vocab, n_blocks, base_stride, 0));
|
||||
cat_conf = ggml_sigmoid(ctx0, ggml_cont(ctx0, ggml_view_2d(ctx0, base, 1, n_blocks, base_stride, 0)));
|
||||
}
|
||||
|
||||
// TODO: the in-graph chain is greedy (argmax); sampling params affect only the final
|
||||
// token pick, not the Markov conditioning path
|
||||
for (int64_t i = i_draft_beg; i < block_drafts; ++i) {
|
||||
ggml_tensor * w1_prev = ggml_get_rows(ctx0, w1, prev); // [R, n_blocks]
|
||||
ggml_tensor * bias = g.build_lora_mm(w2, w1_prev, model.dspark_markov_w2_s); // [n_vocab_draft, n_blocks]
|
||||
ggml_tensor * w1_prev = ggml_get_rows(ctx0, w1, prev); // [R, n_blocks]
|
||||
ggml_tensor * bias = ggml_mul_mat(ctx0, w2, w1_prev); // [n_vocab_draft, n_blocks]
|
||||
if (model.d2t) {
|
||||
// reduced draft vocab: scatter the bias to the target rows (base is -inf on the others)
|
||||
const int64_t n_draft_vocab = bias->ne[0];
|
||||
@@ -360,21 +354,17 @@ static void build_dspark_markov_head(llm_graph_context & g, const llama_model &
|
||||
|
||||
cat = cat ? ggml_concat(ctx0, cat, col, 1) : col;
|
||||
|
||||
if (has_conf) {
|
||||
// confidence head input: predicts per-position acceptance
|
||||
ggml_tensor * conf_inp = res->t_embd; // [n_embd, n_tok]
|
||||
// conf(i) = sigmoid(conf_proj . [conf_inp(i); markov_w1[prev(i)]] + b) -- [1, n_blocks]
|
||||
ggml_tensor * conf_inp_i = ggml_view_2d(ctx0, conf_inp, conf_inp->ne[0], n_blocks,
|
||||
(size_t) block_drafts * conf_inp->nb[1], i*conf_inp->nb[1]);
|
||||
ggml_tensor * feat = ggml_concat(ctx0, ggml_cont(ctx0, conf_inp_i), w1_prev, 0);
|
||||
ggml_tensor * conf = ggml_mul_mat(ctx0, model.dspark_conf_proj, feat);
|
||||
if (model.dspark_conf_proj_b) {
|
||||
conf = ggml_add(ctx0, conf, model.dspark_conf_proj_b);
|
||||
}
|
||||
conf = ggml_sigmoid(ctx0, conf);
|
||||
|
||||
cat_conf = cat_conf ? ggml_concat(ctx0, cat_conf, conf, 1) : conf;
|
||||
// conf(i) = sigmoid(conf_proj . [conf_inp(i); markov_w1[prev(i)]] + b) -- [1, n_blocks]
|
||||
ggml_tensor * conf_inp_i = ggml_view_2d(ctx0, conf_inp, conf_inp->ne[0], n_blocks,
|
||||
(size_t) block_drafts * conf_inp->nb[1], i*conf_inp->nb[1]);
|
||||
ggml_tensor * feat = ggml_concat(ctx0, ggml_cont(ctx0, conf_inp_i), w1_prev, 0);
|
||||
ggml_tensor * conf = ggml_mul_mat(ctx0, model.dspark_conf_proj, feat);
|
||||
if (model.dspark_conf_proj_b) {
|
||||
conf = ggml_add(ctx0, conf, model.dspark_conf_proj_b);
|
||||
}
|
||||
conf = ggml_sigmoid(ctx0, conf);
|
||||
|
||||
cat_conf = cat_conf ? ggml_concat(ctx0, cat_conf, conf, 1) : conf;
|
||||
|
||||
if (i + 1 < block_drafts) {
|
||||
prev = ggml_argmax(ctx0, col);
|
||||
@@ -386,7 +376,7 @@ static void build_dspark_markov_head(llm_graph_context & g, const llama_model &
|
||||
out = ggml_cont(ctx0, ggml_permute(ctx0, out, 0, 2, 1, 3)); // [n_vocab, block_drafts, n_blocks]
|
||||
out = ggml_reshape_2d(ctx0, out, n_vocab, n_tok);
|
||||
|
||||
if (has_conf) {
|
||||
{
|
||||
ggml_tensor * conf = ggml_reshape_3d(ctx0, cat_conf, 1, n_blocks, block_drafts);
|
||||
conf = ggml_cont(ctx0, ggml_permute(ctx0, conf, 0, 2, 1, 3));
|
||||
conf = ggml_reshape_2d(ctx0, conf, 1, n_tok);
|
||||
@@ -717,8 +707,8 @@ llama_model_dflash::graph<false>::graph(const llama_model & model, const llm_gra
|
||||
|
||||
// cache-aware, non-causal attention
|
||||
ggml_tensor * cur = use_iswa
|
||||
? build_attn(inp_attn_iswa, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, layer.attn_sinks, nullptr, kq_scale, il)
|
||||
: build_attn(inp_attn, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, layer.attn_sinks, nullptr, kq_scale, il);
|
||||
? build_attn(inp_attn_iswa, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il)
|
||||
: build_attn(inp_attn, layer.wo, NULL, NULL, Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
|
||||
|
||||
if (attn_dynamic) {
|
||||
cur = build_dflash2_conv(*this, cur, attn_dynamic, layer.dflash_attn_conv_base, 1);
|
||||
|
||||
@@ -181,7 +181,7 @@ public:
|
||||
return res;
|
||||
}
|
||||
|
||||
const llama_hparams hparams;
|
||||
const llama_hparams & hparams;
|
||||
|
||||
ggml_tensor * inp_slopes = nullptr; // F32 [n_head]
|
||||
ggml_tensor * inp_q_decay = nullptr; // F32 [1, n_head, n_batch]
|
||||
|
||||
@@ -6,9 +6,6 @@
|
||||
|
||||
// note: almost all graphs require at least sqrtf, so include cmath globally
|
||||
#include <cmath>
|
||||
#include <map>
|
||||
|
||||
class llama_memory_hybrid_idx_context;
|
||||
|
||||
//
|
||||
// base classes
|
||||
@@ -2275,108 +2272,6 @@ struct llama_model_qwen35 : public llama_model_base {
|
||||
};
|
||||
|
||||
|
||||
struct llama_model_qwen4exp : public llama_model_base {
|
||||
llama_model_qwen4exp(const struct llama_model_params & params) : llama_model_base(params) {}
|
||||
|
||||
class llm_graph_input_qsa;
|
||||
|
||||
void load_arch_hparams(llama_model_loader & ml) override;
|
||||
void load_arch_tensors(llama_model_loader & ml) override;
|
||||
|
||||
struct graph : public llm_build_delta_net_base {
|
||||
graph(const llama_model & model, const llm_graph_params & params);
|
||||
private:
|
||||
// HC replaces every layer norm: residual is [n_embd, hc, n_tokens]
|
||||
ggml_tensor * build_hc_mix(
|
||||
ggml_tensor * x,
|
||||
ggml_tensor * w_norm,
|
||||
ggml_tensor * w_down,
|
||||
ggml_tensor * w_up,
|
||||
ggml_tensor * w_inject,
|
||||
ggml_tensor ** inject,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_hc_combine(
|
||||
ggml_tensor * residual,
|
||||
ggml_tensor * block_out,
|
||||
ggml_tensor * inject,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_layer_attn(
|
||||
llm_graph_input_attn_kv * inp_attn,
|
||||
const llama_memory_hybrid_idx_context * mctx_hyb,
|
||||
ggml_tensor * cur,
|
||||
ggml_tensor * inp_pos,
|
||||
int * sections,
|
||||
int il);
|
||||
|
||||
// dense self-attention restricted to the cells that top_k names
|
||||
ggml_tensor * build_attn_qsa(
|
||||
llm_graph_input_attn_kv * inp,
|
||||
ggml_tensor * q_cur,
|
||||
ggml_tensor * k_cur,
|
||||
ggml_tensor * v_cur,
|
||||
ggml_tensor * top_k,
|
||||
float kq_scale,
|
||||
int il);
|
||||
|
||||
// the QSA cache layout inputs do not depend on the layer, only on its compress ratio,
|
||||
// so the layers sharing a ratio share one input set
|
||||
std::map<uint32_t, llm_graph_input_qsa *> qsa_inps;
|
||||
|
||||
// QSA: token indices this layer's queries may attend to, or nullptr for dense
|
||||
ggml_tensor * build_qsa_top_k(
|
||||
const llama_memory_hybrid_idx_context * mctx_hyb,
|
||||
ggml_tensor * cur,
|
||||
ggml_tensor * inp_pos,
|
||||
ggml_tensor * kq_mask,
|
||||
int * sections,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_layer_attn_linear(
|
||||
llm_graph_input_rs * inp,
|
||||
ggml_tensor * cur,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_layer_ffn(
|
||||
ggml_tensor * cur,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_norm_gated(
|
||||
ggml_tensor * input,
|
||||
ggml_tensor * weights,
|
||||
ggml_tensor * gate,
|
||||
int layer);
|
||||
|
||||
// build_rs writes the state tensor in place, so one gather per cache tensor is reused
|
||||
std::map<ggml_tensor *, ggml_tensor *> rs_rows;
|
||||
|
||||
// one conv history per cache tensor: delta-net and PLE each have their own
|
||||
ggml_tensor * build_conv_state_at(
|
||||
llm_graph_input_rs * inp,
|
||||
ggml_tensor * conv_states_all,
|
||||
ggml_tensor * x,
|
||||
int64_t state_cols,
|
||||
int64_t channels,
|
||||
int il);
|
||||
|
||||
ggml_tensor * build_ple(
|
||||
llm_graph_input_rs * inp,
|
||||
const llama_memory_hybrid_idx_context * mctx_hyb,
|
||||
ggml_tensor * hidden,
|
||||
int il);
|
||||
|
||||
// returns pair of qkv, z
|
||||
std::pair<ggml_tensor *, ggml_tensor *> build_qkvz(
|
||||
ggml_tensor * input,
|
||||
int il);
|
||||
|
||||
const llama_model & model;
|
||||
};
|
||||
|
||||
std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
|
||||
};
|
||||
|
||||
struct llama_model_qwen35moe : public llama_model_base {
|
||||
llama_model_qwen35moe(const struct llama_model_params & params) : llama_model_base(params) {}
|
||||
void load_arch_hparams(llama_model_loader & ml) override;
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -149,7 +149,6 @@ if (LLAMA_LLGUIDANCE)
|
||||
endif ()
|
||||
|
||||
llama_build(test-recurrent-state-rollback.cpp)
|
||||
llama_build(test-save-load-state.cpp)
|
||||
|
||||
if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
|
||||
# these tests are disabled on Windows because they use internal functions not exported with LLAMA_API (when building with shared libraries)
|
||||
@@ -238,14 +237,6 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
|
||||
set_tests_properties(test-recurrent-state-rollback-dsv4 PROPERTIES
|
||||
FIXTURES_REQUIRED generate-models
|
||||
)
|
||||
|
||||
# Test state save/load functionality across all architectures, using the generated dummy models
|
||||
llama_test(
|
||||
test-save-load-state
|
||||
LABEL main
|
||||
ARGS --models "${MODEL_DIR}"
|
||||
)
|
||||
set_tests_properties(test-save-load-state PROPERTIES FIXTURES_REQUIRED generate-models)
|
||||
endif()
|
||||
|
||||
llama_build_and_test(test-chat-peg-parser.cpp peg-parser/simple-tokenize.cpp)
|
||||
@@ -308,6 +299,10 @@ llama_build_and_test(test-backend-sampler.cpp LABEL "model")
|
||||
llama_build_and_test(test-state-restore-fragmented.cpp LABEL "model" ARGS -m "${MODEL_DEST}")
|
||||
set_tests_properties(test-state-restore-fragmented PROPERTIES FIXTURES_REQUIRED test-download-model)
|
||||
|
||||
# Test state save/load functionality
|
||||
llama_build_and_test(test-save-load-state.cpp LABEL "model" ARGS -m "${MODEL_DEST}")
|
||||
set_tests_properties(test-save-load-state PROPERTIES FIXTURES_REQUIRED test-download-model)
|
||||
|
||||
if (APPLE)
|
||||
llama_build(test-rset-release.cpp)
|
||||
endif()
|
||||
|
||||
@@ -65,7 +65,7 @@ static void set_tensor_data(struct ggml_tensor * tensor, void * userdata) {
|
||||
}
|
||||
|
||||
static void usage(char ** argv) {
|
||||
printf("Usage: %s [-a/--arch arch] [-s/--seed seed] [-o/--out dir] [-v/--verbose] [-h/--help]\n", argv[0]);
|
||||
printf("Usage: %s [-a/--arch arch] [-s/--seed seed] [-v/--verbose]\n", argv[0]);
|
||||
}
|
||||
|
||||
static std::vector<llama_token> get_tokens(const uint32_t n_tokens, const uint32_t n_vocab, const size_t seed){
|
||||
@@ -82,7 +82,7 @@ static std::vector<llama_token> get_tokens(const uint32_t n_tokens, const uint32
|
||||
static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
gguf_context_ptr ret(gguf_init_empty());
|
||||
llama_model_saver ms(arch, ret.get());
|
||||
const uint32_t n_ctx = 256;
|
||||
const uint32_t n_ctx = 128;
|
||||
|
||||
uint32_t n_vocab = 128;
|
||||
uint32_t n_embd = 256;
|
||||
@@ -249,19 +249,8 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
|
||||
// MSA requires one indexer head per GQA (KV) head, unlike the DSA archs where the
|
||||
// indexer head count is independent of the main attention head count.
|
||||
if (arch == LLM_ARCH_QWEN4EXP) {
|
||||
ms.add_kv(LLM_KV_HYPER_CONNECTION_COUNT, uint32_t(4));
|
||||
ms.add_kv(LLM_KV_HYPER_CONNECTION_LOW_RANK, uint32_t(8));
|
||||
// without this the QSA layers fall back to dense and go uncovered
|
||||
ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, std::vector<uint32_t>(n_layer, 4));
|
||||
}
|
||||
|
||||
// minimax-m3 keeps one indexer head per GQA head; the rest use a fixed 64 to match the fused
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, arch == LLM_ARCH_MINIMAX_M3 ? n_head : uint32_t(64));
|
||||
// qwen4exp ropes indexer keys with the main rotary width, so its head can't be < n_rot
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH,
|
||||
arch == LLM_ARCH_QWEN4EXP ? n_embd_head : uint32_t(128));
|
||||
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, arch == LLM_ARCH_MINIMAX_M3 || arch == LLM_ARCH_DEEPSEEK4 ? n_head : uint32_t(1));
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, uint32_t(64));
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_TOP_K, uint32_t(8));
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_BLOCK_SIZE, uint32_t(4));
|
||||
ms.add_kv(LLM_KV_ATTENTION_INDEXER_LOCAL_BLOCKS, uint32_t(1));
|
||||
@@ -305,7 +294,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
ms.add_kv(LLM_KV_XIELU_ALPHA_P, 1.0f);
|
||||
ms.add_kv(LLM_KV_XIELU_BETA, 1.0f);
|
||||
ms.add_kv(LLM_KV_XIELU_EPS, 1.0e-7f);
|
||||
ms.add_kv(LLM_KV_SSM_INNER_SIZE, arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE || arch == LLM_ARCH_QWEN4EXP ? 256 : 2*n_embd);
|
||||
ms.add_kv(LLM_KV_SSM_INNER_SIZE, arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE ? 256 : 2*n_embd);
|
||||
ms.add_kv(LLM_KV_SSM_CONV_KERNEL, uint32_t(4));
|
||||
ms.add_kv(LLM_KV_SSM_STATE_SIZE, uint32_t(128));
|
||||
ms.add_kv(LLM_KV_SSM_TIME_STEP_RANK, n_head);
|
||||
@@ -422,7 +411,6 @@ static bool moe_mandatory(const llm_arch arch) {
|
||||
case LLM_ARCH_QWEN3NEXT:
|
||||
case LLM_ARCH_QWEN3VLMOE:
|
||||
case LLM_ARCH_QWEN35MOE:
|
||||
case LLM_ARCH_QWEN4EXP:
|
||||
case LLM_ARCH_PHIMOE:
|
||||
case LLM_ARCH_DBRX:
|
||||
case LLM_ARCH_OLMOE:
|
||||
@@ -519,7 +507,7 @@ static bool arch_supported(const llm_arch arch) {
|
||||
}
|
||||
// FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
|
||||
#ifdef GGML_USE_WEBGPU
|
||||
if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_QWEN4EXP) {
|
||||
if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE) {
|
||||
return false;
|
||||
}
|
||||
#endif // GGML_USE_WEBGPU
|
||||
@@ -764,10 +752,6 @@ int main(int argc, char ** argv) {
|
||||
std::string out;
|
||||
|
||||
for (int i = 1; i < argc; i++) {
|
||||
if (strcmp(argv[i], "-h") == 0 || strcmp(argv[i], "--help") == 0) {
|
||||
usage(argv);
|
||||
return 0;
|
||||
}
|
||||
if (strcmp(argv[i], "-a") == 0 || strcmp(argv[i], "--arch") == 0) {
|
||||
if (i + 1 < argc) {
|
||||
const std::string arch_name = argv[++i];
|
||||
|
||||
+35
-120
@@ -3,12 +3,8 @@
|
||||
#include "log.h"
|
||||
#include "llama-cpp.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <clocale>
|
||||
#include <cstring>
|
||||
#include <filesystem>
|
||||
#include <random>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
struct llama_batch_ptr {
|
||||
@@ -57,9 +53,7 @@ static llama_tokens generate_tokens(llama_context * ctx, llama_sampler * smpl, i
|
||||
// - decode the last token
|
||||
// - generate n_predict tokens
|
||||
static llama_tokens test_baseline(struct llama_model * model, const struct common_params & params, const llama_tokens & tokens) {
|
||||
auto params_ctx = common_context_params_to_llama(params);
|
||||
params_ctx.n_seq_max = 2;
|
||||
auto ctx = llama_context_ptr{llama_init_from_model(model, params_ctx)};
|
||||
auto ctx = llama_context_ptr{llama_init_from_model(model, common_context_params_to_llama(params))};
|
||||
|
||||
auto sparams = llama_sampler_chain_default_params();
|
||||
auto smpl = llama_sampler_ptr{llama_sampler_chain_init(sparams)};
|
||||
@@ -167,9 +161,7 @@ static bool test_seq_rm_isolated(
|
||||
// - replay the last prompt token
|
||||
// - generate n_predict tokens and compare against expected result
|
||||
static bool test_state_load(struct llama_model * model, const struct common_params & params, const llama_tokens & tokens, const llama_tokens & expected_result) {
|
||||
auto params_ctx = common_context_params_to_llama(params);
|
||||
params_ctx.n_seq_max = 2;
|
||||
auto ctx = llama_context_ptr{llama_init_from_model(model, params_ctx)};
|
||||
auto ctx = llama_context_ptr{llama_init_from_model(model, common_context_params_to_llama(params))};
|
||||
|
||||
auto sparams = llama_sampler_chain_default_params();
|
||||
auto smpl = llama_sampler_ptr{llama_sampler_chain_init(sparams)};
|
||||
@@ -355,18 +347,38 @@ static bool test_seq_cp_device(struct llama_model * model, const struct common_p
|
||||
}
|
||||
|
||||
|
||||
// Run the full save/load test suite (tests 1-5) for a single model.
|
||||
// Returns true if all tests pass, false otherwise.
|
||||
static bool run_save_load_tests_for_model(const std::string & model_path, const struct common_params & base_params) {
|
||||
struct common_params params = base_params;
|
||||
params.model.path = model_path;
|
||||
int main(int argc, char ** argv) {
|
||||
std::setlocale(LC_NUMERIC, "C");
|
||||
|
||||
common_params params;
|
||||
params.prompt = "";
|
||||
params.n_batch = 100;
|
||||
params.out_file = "dump_state.bin";
|
||||
params.sampling.seed = 1234;
|
||||
|
||||
common_init();
|
||||
|
||||
if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_COMMON)) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
if (params.n_parallel == 1) {
|
||||
LOG_TRC("%s: n_parallel == 1, enabling unified kv cache\n", __func__);
|
||||
params.kv_unified = true;
|
||||
}
|
||||
|
||||
if (params.n_predict < 0) {
|
||||
params.n_predict = 16;
|
||||
}
|
||||
|
||||
ggml_backend_load_all();
|
||||
|
||||
auto llama_init = common_init_from_params(params, true);
|
||||
auto * model = llama_init->model();
|
||||
|
||||
if (model == nullptr) {
|
||||
LOG_ERR("%s: failed to init model '%s'\n", __func__, model_path.c_str());
|
||||
return false;
|
||||
LOG_ERR("%s: failed to init\n", __func__);
|
||||
return 1;
|
||||
}
|
||||
|
||||
GGML_ASSERT(llama_init->context() == nullptr);
|
||||
@@ -399,127 +411,30 @@ static bool run_save_load_tests_for_model(const std::string & model_path, const
|
||||
// Test 1: baseline (saves state to disk)
|
||||
auto result_baseline = test_baseline(model, params, tokens);
|
||||
if (result_baseline.empty()) {
|
||||
return false;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Test 2: sequence removal isolation
|
||||
if (!test_seq_rm_isolated(model, params, tokens)) {
|
||||
return false;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Test 3: state load
|
||||
if (!test_state_load(model, params, tokens, result_baseline)) {
|
||||
return false;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Test 4: seq copy (host)
|
||||
if (!test_seq_cp_host(model, params, tokens, result_baseline)) {
|
||||
return false;
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Test 5: seq copy (device)
|
||||
if (!test_seq_cp_device(model, params, tokens, result_baseline)) {
|
||||
return false;
|
||||
return 1;
|
||||
}
|
||||
|
||||
LOG("\nAll tests passed.\n");
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
|
||||
int main(int argc, char ** argv) {
|
||||
std::setlocale(LC_NUMERIC, "C");
|
||||
|
||||
common_params params;
|
||||
params.prompt = "";
|
||||
params.n_batch = 100;
|
||||
params.out_file = "dump_state.bin";
|
||||
params.sampling.seed = 1234;
|
||||
|
||||
common_init();
|
||||
|
||||
// extract our own --models DIR option before handing the rest to the common arg parser
|
||||
std::string models_dir;
|
||||
std::vector<char *> filtered_argv;
|
||||
filtered_argv.push_back(argv[0]);
|
||||
for (int i = 1; i < argc; i++) {
|
||||
if (strcmp(argv[i], "--models") == 0) {
|
||||
if (i + 1 >= argc) {
|
||||
LOG_ERR("%s: --models requires a directory argument\n", __func__);
|
||||
return 1;
|
||||
}
|
||||
models_dir = argv[i + 1];
|
||||
i++;
|
||||
} else {
|
||||
filtered_argv.push_back(argv[i]);
|
||||
}
|
||||
}
|
||||
filtered_argv.push_back(nullptr);
|
||||
const int fargc = (int)filtered_argv.size() - 1;
|
||||
|
||||
// in --models mode there is no single model; set a placeholder so the common parser's
|
||||
// "--model is required" check passes (each model is set individually inside the loop)
|
||||
if (!models_dir.empty()) {
|
||||
params.model.path = models_dir;
|
||||
}
|
||||
|
||||
if (!common_params_parse(fargc, filtered_argv.data(), params, LLAMA_EXAMPLE_COMMON)) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
if (params.n_parallel == 1) {
|
||||
LOG_TRC("%s: n_parallel == 1, enabling unified kv cache\n", __func__);
|
||||
params.kv_unified = true;
|
||||
}
|
||||
|
||||
if (params.n_predict < 0) {
|
||||
params.n_predict = 16;
|
||||
}
|
||||
|
||||
ggml_backend_load_all();
|
||||
|
||||
if (!models_dir.empty()) {
|
||||
// run the suite over every dummy model in the directory
|
||||
if (!std::filesystem::exists(models_dir) || !std::filesystem::is_directory(models_dir)) {
|
||||
LOG_ERR("%s: models directory '%s' does not exist\n", __func__, models_dir.c_str());
|
||||
return 1;
|
||||
}
|
||||
|
||||
std::vector<std::string> models;
|
||||
for (const auto & entry : std::filesystem::directory_iterator(models_dir)) {
|
||||
if (entry.is_regular_file() && entry.path().extension() == ".gguf") {
|
||||
models.push_back(entry.path().string());
|
||||
}
|
||||
}
|
||||
std::sort(models.begin(), models.end());
|
||||
|
||||
if (models.empty()) {
|
||||
LOG_ERR("%s: no .gguf models found in '%s'\n", __func__, models_dir.c_str());
|
||||
return 1;
|
||||
}
|
||||
|
||||
LOG_INF("%s: running save/load tests over %zu models in '%s'\n", __func__, models.size(), models_dir.c_str());
|
||||
|
||||
size_t n_pass = 0;
|
||||
size_t n_fail = 0;
|
||||
for (const auto & model_path : models) {
|
||||
LOG("\n================================================================\n");
|
||||
LOG_INF("%s: model %s\n", __func__, model_path.c_str());
|
||||
|
||||
if (run_save_load_tests_for_model(model_path, params)) {
|
||||
n_pass++;
|
||||
} else {
|
||||
n_fail++;
|
||||
}
|
||||
}
|
||||
|
||||
LOG("\n================================================================\n");
|
||||
LOG_INF("%s: summary: %zu passed, %zu failed (of %zu)\n", __func__, n_pass, n_fail, models.size());
|
||||
|
||||
return n_fail == 0 ? 0 : 1;
|
||||
}
|
||||
|
||||
// single-model mode
|
||||
return run_save_load_tests_for_model(params.model.path, params) ? 0 : 1;
|
||||
return 0;
|
||||
}
|
||||
|
||||
@@ -163,7 +163,6 @@ For the full list of features, please refer to [server's changelog](https://gith
|
||||
| -------- | ----------- |
|
||||
| `-lcs, --lookup-cache-static FNAME` | path to static lookup cache to use for lookup decoding (not updated by generation) |
|
||||
| `-lcd, --lookup-cache-dynamic FNAME` | path to dynamic lookup cache to use for lookup decoding (updated by generation) |
|
||||
| `--kv-unified-per-slot N` | context limit per parallel slot (default: unset, behavior unchanged).<br/>when set without -c/--ctx-size, the shared KV pool is sized to n_parallel*N<br/>(env: LLAMA_ARG_KV_UNIFIED_PER_SLOT) |
|
||||
| `-ctxcp, --ctx-checkpoints, --swa-checkpoints N` | max number of context checkpoints to create per slot (default: 32)[(more info)](https://github.com/ggml-org/llama.cpp/pull/15293)<br/>(env: LLAMA_ARG_CTX_CHECKPOINTS) |
|
||||
| `-cms, --checkpoint-min-step N` | minimum spacing between context checkpoints in tokens (default: 8192, 0 = no minimum)<br/>(env: LLAMA_ARG_CHECKPOINT_MIN_SPACING_NT) |
|
||||
| `-cram, --cache-ram N` | set the maximum cache size in MiB (default: 8192, -1 - no limit, 0 - disable)[(more info)](https://github.com/ggml-org/llama.cpp/pull/16391)<br/>(env: LLAMA_ARG_CACHE_RAM) |
|
||||
|
||||
@@ -1208,31 +1208,10 @@ private:
|
||||
|
||||
const int n_ctx_train = llama_model_n_ctx_train(model_tgt);
|
||||
|
||||
{
|
||||
// note: the capping itself is done in n_ctx_slot(), here we only report it
|
||||
const int n_ctx_seq = llama_n_ctx_seq(ctx_tgt);
|
||||
|
||||
if (params_base.kv_unified_per_slot > 0) {
|
||||
if (n_ctx_seq > params_base.kv_unified_per_slot) {
|
||||
SRV_INF("capping per-slot context (%d) to --kv-unified-per-slot (%d)\n",
|
||||
n_ctx_seq, params_base.kv_unified_per_slot);
|
||||
} else if (params_base.kv_unified_per_slot > n_ctx_seq) {
|
||||
// cap is above the per-slot pool capacity, so it can never bind
|
||||
SRV_WRN(
|
||||
"--kv-unified-per-slot (%d) exceeds the per-slot pool capacity (%d) - cap has no effect, "
|
||||
"slots are limited to %d (raise the KV pool with -c, or unset -c to size it to "
|
||||
"n_parallel * kv_unified_per_slot)\n",
|
||||
params_base.kv_unified_per_slot, n_ctx_seq, n_ctx_seq);
|
||||
}
|
||||
}
|
||||
|
||||
const int n_ctx_capped = params_base.kv_unified_per_slot > 0 ?
|
||||
std::min(n_ctx_seq, params_base.kv_unified_per_slot) : n_ctx_seq;
|
||||
|
||||
if (n_ctx_capped > n_ctx_train) {
|
||||
SRV_WRN("the slot context (%d) exceeds the training context of the model (%d) - capping\n",
|
||||
n_ctx_capped, n_ctx_train);
|
||||
}
|
||||
int n_ctx_slot = llama_n_ctx_seq(ctx_tgt);
|
||||
if (n_ctx_slot > n_ctx_train) {
|
||||
SRV_WRN("the slot context (%d) exceeds the training context of the model (%d) - capping\n", n_ctx_slot, n_ctx_train);
|
||||
n_ctx_slot = n_ctx_train;
|
||||
}
|
||||
|
||||
slots.clear();
|
||||
@@ -1248,7 +1227,7 @@ private:
|
||||
|
||||
// setup slots
|
||||
SRV_INF("initializing, n_slots = %d, n_ctx_slot = %d, kv_unified = '%s'\n",
|
||||
params_base.n_parallel, n_ctx_slot(), params_base.kv_unified ? "true" : "false");
|
||||
params_base.n_parallel, n_ctx_slot, params_base.kv_unified ? "true" : "false");
|
||||
|
||||
// initialize slots
|
||||
for (int i = 0; i < params_base.n_parallel; i++) {
|
||||
@@ -1292,7 +1271,7 @@ private:
|
||||
slot.ctx_dft = ctx_dft;
|
||||
slot.mem.init(ctx_tgt, ctx_dft);
|
||||
slot.spec = spec.get();
|
||||
slot.n_ctx = n_ctx_slot();
|
||||
slot.n_ctx = n_ctx_slot;
|
||||
|
||||
slot.mctx = mctx;
|
||||
slot.prompt.tokens.has_mtmd = mctx != nullptr;
|
||||
@@ -3996,15 +3975,8 @@ private:
|
||||
});
|
||||
}
|
||||
|
||||
// context size of a single slot, capped by --kv-unified-per-slot and by the training context of the model
|
||||
int n_ctx_slot() const {
|
||||
int res = llama_n_ctx_seq(ctx_tgt);
|
||||
|
||||
if (params_base.kv_unified_per_slot > 0) {
|
||||
res = std::min(res, params_base.kv_unified_per_slot);
|
||||
}
|
||||
|
||||
return std::min(res, llama_model_n_ctx_train(model_tgt));
|
||||
int get_slot_n_ctx() {
|
||||
return slots.back().n_ctx;
|
||||
}
|
||||
|
||||
server_response_reader get_response_reader() {
|
||||
@@ -4170,7 +4142,7 @@ server_context_meta server_context::get_meta() const {
|
||||
/* has_inp_audio */ impl->chat_params.allow_audio,
|
||||
/* has_inp_video */ impl->chat_params.allow_video,
|
||||
/* json_ui_settings */ impl->json_ui_settings,
|
||||
/* slot_n_ctx */ impl->n_ctx_slot(),
|
||||
/* slot_n_ctx */ impl->get_slot_n_ctx(),
|
||||
/* pooling_type */ llama_pooling_type(impl->ctx_tgt),
|
||||
|
||||
/* chat_params */ impl->chat_params,
|
||||
|
||||
@@ -157,18 +157,6 @@ int llama_server(common_params & params, int argc, char ** argv) {
|
||||
}
|
||||
}
|
||||
|
||||
// size the KV pool from --kv-unified-per-slot, unless the user pinned it with -c
|
||||
// or with -c 0 for max context
|
||||
const bool ctx_pool_auto_sized = params.kv_unified_per_slot > 0 &&
|
||||
params.n_ctx == 0 &&
|
||||
(uint32_t) params.fit_params_min_ctx != UINT32_MAX;
|
||||
|
||||
if (ctx_pool_auto_sized) {
|
||||
params.n_ctx = params.n_parallel * params.kv_unified_per_slot;
|
||||
SRV_INF("--kv-unified-per-slot: sizing KV pool to n_parallel * kv_unified_per_slot = %d * %d = %d\n", params.n_parallel,
|
||||
params.kv_unified_per_slot, params.n_ctx);
|
||||
}
|
||||
|
||||
// for consistency between server router mode and single-model mode, we set the same model name as alias
|
||||
auto model_name = params.model.get_name();
|
||||
if (params.model_alias.empty() && !model_name.empty()) {
|
||||
|
||||
Reference in New Issue
Block a user