Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-09 14:29:15 +02:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit 'd646c9d15500a702425e8a90c19d2400deecbd0c' into concedo_experimental

Browse Source 
# Conflicts:
#	.github/actions/windows-setup-rocm/action.yml
#	.github/workflows/build-apple.yml
#	.github/workflows/release.yml
#	.github/workflows/server-self-hosted.yml
#	examples/training/README.md
#	ggml/src/ggml-hexagon/ggml-hexagon.cpp
#	ggml/src/ggml-hexagon/htp/hvx-arith.h
#	ggml/src/ggml-hexagon/htp/hvx-log.h
#	ggml/src/ggml-hexagon/htp/hvx-norm.h
#	ggml/src/ggml-hexagon/htp/hvx-scale.h
#	ggml/src/ggml-hexagon/htp/hvx-sqrt.h
#	ggml/src/ggml-hexagon/htp/unary-ops.c
#	ggml/src/ggml-hexagon/htp/unary-ops.h
#	ggml/src/ggml-sycl/mmvq.cpp
#	ggml/src/ggml-sycl/vecdotq.hpp
#	tests/test-backend-ops.cpp
#	tests/test-mtmd-c-api.c
#	tests/test-mtmd-impl.cpp
This commit is contained in:
Concedo
2026-09-04 16:05:40 +08:00
parent 99be3d0990 d646c9d155
commit a3e6006530
78 changed files with 693 additions and 221 deletions
Download Patch File Download Diff File Expand all files Collapse all files
common/arg.cpp
+11 -1
View File
@@ -961,6 +961,11 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
));
}
// if the preserve_reasoning kwarg was not specified explicitly, enable it by default
if (!params.default_template_kwargs.count("preserve_reasoning")) {
params.default_template_kwargs["preserve_reasoning"] = "true";
}
return true;
}
@@ -3554,6 +3559,10 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
LOG_WRN("Setting 'enable_thinking' via --chat-template-kwargs is deprecated. "
"Use --reasoning on / --reasoning off instead.\n");
}
if (item.key() == "preserve_reasoning") {
LOG_WRN("Setting 'preserve_reasoning' via --chat-template-kwargs is deprecated. "
"Use --reasoning-preserve / --no-reasoning-preserve instead.\n");
}
params.default_template_kwargs[item.key()] = item.value().dump();
}
}
@@ -3744,7 +3753,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
add_opt(common_arg(
{"--reasoning-preserve"},
{"--no-reasoning-preserve"},
"preserve reasoning trace in the full history, not just the last assistant message (default: template default)\n"
"preserve reasoning trace in the full history, not just the last assistant message (default: enabled)\n"
"compatible with certain templates having 'supports_preserve_reasoning' capability\n"
"example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking",
[](common_params & params, bool value) {
@@ -3753,6 +3762,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
} else {
params.default_template_kwargs["preserve_reasoning"] = "false";
}
params.preserve_reasoning_specified = true;
}
).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_REASONING_PRESERVE"));
add_opt(common_arg(
common/common.h
+2 -1
View File
@@ -271,7 +271,7 @@ struct common_params_sampling {
COMMON_SAMPLER_TYPE_TEMPERATURE,
};
common_grammar grammar; // optional grammar constraint (user / output-format / tool-calls)
common_grammar grammar; // optional grammar constraint (user / output-format / tool-calls)
bool grammar_lazy = false;
std::vector<common_grammar_trigger> grammar_triggers; // optional triggers (for lazy grammars)
std::set<llama_token> preserved_tokens;
@@ -658,6 +658,7 @@ struct common_params {
std::string ssl_file_cert = ""; // NOLINT
std::map<std::string, std::string> default_template_kwargs;
bool preserve_reasoning_specified = false;
// CLI params
std::string server_base; // if set, connect to this server instead of starting a new one
conversion/__init__.py
+1
View File
@@ -188,6 +188,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
"NanbeigeForCausalLM": "nanbeige",
"NemotronForCausalLM": "nemotron",
"NemotronHForCausalLM": "nemotron",
"NemotronHPuzzleForCausalLM": "nemotron",
"NeoBERT": "bert",
"NeoBERTForSequenceClassification": "bert",
"NeoBERTLMHead": "bert",
conversion/deepseek.py
+13 -2
View File
@@ -578,8 +578,7 @@ class DeepseekV4Model(TextModel):
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if (name.startswith(("aligner.", "image_"))
or name.endswith(".ffn.gate.bias_vl")):
if name.startswith(("aligner.", "image_")):
return None
if name.startswith("mtp."):
if not cls.mtp_only:
@@ -856,6 +855,7 @@ class DeepseekV4Model(TextModel):
"ffn_norm.weight": (gguf.MODEL_TENSOR.FFN_NORM, ".weight"),
"ffn.gate.weight": (gguf.MODEL_TENSOR.FFN_GATE_INP, ".weight"),
"ffn.gate.bias": (gguf.MODEL_TENSOR.FFN_EXP_PROBS_B, ".bias"),
"ffn.gate.bias_vl": (gguf.MODEL_TENSOR.FFN_EXP_PROBS_B_VL, ".bias"),
"ffn.gate.tid2eid": (gguf.MODEL_TENSOR.FFN_GATE_TID2EID, ".weight"),
"ffn.shared_experts.w1.weight": (gguf.MODEL_TENSOR.FFN_GATE_SHEXP, ".weight"),
"ffn.shared_experts.w2.weight": (gguf.MODEL_TENSOR.FFN_DOWN_SHEXP, ".weight"),
@@ -881,6 +881,10 @@ class DeepseekV4Model(TextModel):
if re.match(r"layers\.\d+\.ffn\.experts\.\d+\.w[123]\.(weight|scale)$", name):
return []
# hash layers route text tokens via tid2eid and image tokens via bias_vl; gate.bias is unused
if name.endswith(".ffn.gate.bias") and bid is not None and bid < self.hparams["num_hash_layers"]:
return []
tensor_key, suffix = self._map_dsv4_tensor_name(name, bid)
if tensor_key == gguf.MODEL_TENSOR.FFN_GATE_TID2EID:
return []
@@ -1003,6 +1007,13 @@ class DeepseekV4DSparkModel(DeepseekV4Model):
return self._DSPARK_ROOT_MAP[name]
return super()._map_dsv4_tensor_name(name, bid)
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
# the DFlash draft uses the plain exp-probs bias (ffn.gate.bias -> FFN_EXP_PROBS_B);
# the mtmd-only hash routing tensors (bias_vl, tid2eid) are not part of the DFLASH arch
if name.endswith(".ffn.gate.bias_vl"):
return
yield from super().modify_tensors(data_torch, name, bid)
def set_vocab(self):
if self.target_model_dir is None:
raise ValueError("DeepSeek-V4 DSpark requires --target-model-dir with the target tokenizer")
conversion/nemotron.py
+87
View File
@@ -5,6 +5,7 @@ from typing import Any, Callable, Iterable, TYPE_CHECKING
import torch
if TYPE_CHECKING:
from pathlib import Path
from torch import Tensor
from .base import MmprojModel, ModelBase, TextModel, gguf, logger
@@ -201,6 +202,7 @@ class NemotronHModel(GraniteHybridModel):
model_arch = gguf.MODEL_ARCH.NEMOTRON_H
is_moe: bool = False
supports_mtp_export = True
_experts: list[dict[str, Tensor]] | None = None
_SSM_LAYER_TYPES = {"mamba", "linear_attention"}
_ATTN_LAYER_TYPES = {"attention", "full_attention"}
@@ -513,3 +515,88 @@ class NemotronHModel(GraniteHybridModel):
experts = [k for d in self._experts for k in d.keys()]
if len(experts) > 0:
raise ValueError(f"Unprocessed experts: {experts}")
@ModelBase.register("NemotronHPuzzleForCausalLM")
@ModelBase.example("nvidia/NVIDIA-Nemotron-Labs-3-Puzzle-75B-A9B-BF16")
class NemotronHPuzzleModel(NemotronHModel):
"""NVIDIA Puzzle: NemotronH with a per-block MoE config (block_configs).
The checkpoint also ships an MTP draft head (mtp.safetensors). It is skipped
here: there is no Puzzle MTP inference path in tree, and the head is laid out
by mtp_block_configs rather than the mtp.layers.* form NemotronHModel maps."""
model_arch = gguf.MODEL_ARCH.NEMOTRON_H_MOE
is_moe: bool = True
supports_mtp_export = False
def __init__(self, dir_model: "Path", *args, **kwargs):
hparams = dict(kwargs.pop("hparams", None) or ModelBase.load_hparams(dir_model, self.is_mistral_format))
self.block_configs: list[dict] = hparams["block_configs"]
self.n_layer_trunk = len(self.block_configs)
# block_configs carries the per-block MoE shape, and is the authority on the
# block pattern too: the layers_block_type the HF config wrapper computes is
# not sized to it.
hparams["num_hidden_layers"] = self.n_layer_trunk
hparams["layers_block_type"] = [bc["block_type"] for bc in self.block_configs]
self.model_arch = gguf.MODEL_ARCH.NEMOTRON_H_MOE
# Bypass NemotronHModel.__init__: it assumes a flat num_experts_per_tok /
# moe_intermediate_size and a layers_block_type sized to block_count, neither
# of which hold for Puzzle's per-block config.
GraniteHybridModel.__init__(self, dir_model, *args, hparams=hparams, **kwargs)
self.head_dim = self.find_hparam(["head_dim", "attention_head_dim"])
self.d_inner = self.find_hparam(["num_heads"]) * self.d_model
# NemotronHModel.__init__ folds an MTP block into block_count when the
# config carries num_nextn_predict_layers; Puzzle's config does, but its
# head has a different layout and no inference path, so stay opted out.
self._mtp_bid = None
def set_gguf_parameters(self):
GraniteHybridModel.set_gguf_parameters(self)
head_dim = self.head_dim
if head_dim is None:
raise ValueError("Could not find the attention head dim in config")
self.gguf_writer.add_key_length(head_dim)
self.gguf_writer.add_value_length(head_dim)
ffn_lengths = [bc.get("moe_intermediate_size") or 0 for bc in self.block_configs]
experts_used = [bc.get("num_experts_per_tok") or 0 for bc in self.block_configs]
self.gguf_writer.add_feed_forward_length(ffn_lengths)
self.gguf_writer.add_expert_feed_forward_length(ffn_lengths)
self.gguf_writer.add_expert_used_count(experts_used)
self.gguf_writer.add_expert_shared_feed_forward_length(self.hparams["moe_shared_expert_intermediate_size"])
self.gguf_writer.add_expert_count(self.hparams["n_routed_experts"])
self.gguf_writer.add_expert_shared_count(self.hparams["n_shared_experts"])
self.gguf_writer.add_expert_weights_norm(self.hparams["norm_topk_prob"])
self.gguf_writer.add_expert_weights_scale(self.hparams["routed_scaling_factor"])
self.gguf_writer.add_expert_group_count(self.hparams["n_group"])
self.gguf_writer.add_moe_latent_size(self.hparams["moe_latent_size"])
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
# The official BF16 checkpoint (NVIDIA-Nemotron-Labs-3-Puzzle-75B-A9B-BF16)
# names the trunk "model.*" (model.layers.*, model.embeddings, model.norm_f)
# where the original release used the NemotronH-style "backbone.*", and spells
# the router bias "e_score_correction_bias" instead of "e_score_correction.bias";
# normalize so both convert identically.
if name.startswith("model."):
name = "backbone." + name[len("model."):]
if name.endswith("mixer.gate.e_score_correction_bias"):
name = name[: -len("e_score_correction_bias")] + "e_score_correction.bias"
yield from super().modify_tensors(data_torch, name, bid)
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
# Drop the MTP head unconditionally; see the class docstring.
if item[0].startswith("mtp."):
return None
return super().filter_tensors(item)
ggml/src/ggml-cuda/mmq-vec-dot.cuh
-11
View File
@@ -148,7 +148,6 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
typedef tile<16, 8, int, input_layout> tile_B;
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -204,7 +203,6 @@ static __device__ __forceinline__ void ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma(
typedef tile< 8, 8, int> tile_B;
typedef tile<16, 8, int> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -320,7 +318,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile<16, 8, int, input_layout> tile_B;
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -371,7 +368,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile< 8, 8, int> tile_B;
typedef tile<16, 8, int> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -486,7 +482,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile<16, 4, int, input_layout> tile_B;
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -537,7 +532,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile< 8, 4, int> tile_B;
typedef tile<16, 8, int> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -686,7 +680,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile<16, 4, int, input_layout> tile_B;
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -756,7 +749,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile< 8, 4, int> tile_B;
typedef tile<16, 8, int> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -1023,7 +1015,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile<16, 4, int, input_layout> tile_B;
typedef tile<16, 16, int, DATA_LAYOUT_J_MAJOR> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -1075,7 +1066,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile< 8, 4, int> tile_B;
typedef tile<16, 8, int> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -1190,7 +1180,6 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
typedef tile<8, 8, int> tile_B;
typedef tile<16, 8, float> tile_C;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp / tile_C::I;
ggml/src/ggml-cuda/mmq.cuh
-5
View File
@@ -482,9 +482,6 @@ static __device__ __forceinline__ void ggml_cuda_mmq_write_back_mma(
typedef tile<16, 8, int> tile_C;
#endif // defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
constexpr int warp_size = ggml_cuda_get_physical_warp_size();
constexpr int nwarps = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
constexpr int rows_per_warp = ggml_cuda_mmq_get_rows_per_warp(type, J, fallback);
constexpr int ntx = rows_per_warp/tile_C::I; // Number of x minitiles per warp.
@@ -541,8 +538,6 @@ struct ggml_cuda_mmq_util_funcs {
template <ggml_type type, int J, bool fallback>
static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_funcs() {
constexpr int I = ggml_cuda_mmq_get_I(type, J, fallback);
if (!ggml_cuda_mmq_get_config(type, J, fallback).use_mma_data_layout()) {
switch (type) {
case GGML_TYPE_Q1_0:
ggml/src/ggml-metal/ggml-metal-device.m
+4 -2
View File
@@ -1477,8 +1477,10 @@ void ggml_metal_device_event_synchronize(ggml_metal_device_t dev, ggml_metal_eve
void ggml_metal_device_get_memory(ggml_metal_device_t dev, size_t * free, size_t * total) {
if (@available(macOS 10.12, iOS 16.0, *)) {
*total = dev->mtl_device.recommendedMaxWorkingSetSize;
*free = *total - dev->mtl_device.currentAllocatedSize;
*total = dev->mtl_device.recommendedMaxWorkingSetSize;
size_t cur = dev->mtl_device.currentAllocatedSize;
// it's possible to allocate more than `recommendedMaxWorkingSetSize`
*free = *total > cur ? *total - cur : 0;
} else {
*free = 0;
*total = 0;
ggml/src/ggml-metal/ggml-metal-tuning.cpp
+101
View File
@@ -1468,6 +1468,107 @@ 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, GGML_TYPE_F16, 32, 32, 1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 1, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 2, 2 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 2, 3 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 3, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 32, 32, 3, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, 3, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, 1, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 64, 64, 3, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, 1, 4 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, 2, 4 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 96, 96, 3, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 128, 128, 1, 3 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 128, 128, 2, 3 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 128, 128, 3, 1 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 128, 128, 3, 3 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 192, 3, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 192, -1, 1 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 192, 1, 2 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 192, 1, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 192, 2, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 2, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 3, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 1, 2 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 1, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 2, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 3, 1 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 192, 128, 3, 2 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 3, 0 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, -1, 1 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 1, 1 }, { 1, 1 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 1, 2 }, { 1, 1 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 1, 4 }, { 1, 1 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 256, 256, 2, 2 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 320, 256, 2, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 320, 256, 3, 0 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 320, 256, -1, 1 }, { 2, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 320, 256, 1, 2 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 320, 256, 1, 4 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 512, 512, 2, 0 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 512, 512, 3, 0 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 512, 512, 3, 1 }, { 4, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 512, 512, 3, 3 }, { 4, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 2, 0 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 2, 1 }, { 4, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 2, 2 }, { 4, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 3, 1 }, { 4, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, 3, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, 2, 2 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, 2, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 64, 64, 3, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 96, 96, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 96, 96, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 96, 96, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 96, 96, 3, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, 3, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, 3, 3 }, { 4, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 128, 128, 3, 4 }, { 1, 1 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, 2, 1 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, 2, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 192, 3, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 128, 1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 128, 2, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 128, 1, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 192, 128, 3, 2 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 256, 256, 3, 0 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 320, 256, 1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 320, 256, 3, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 320, 256, -1, 1 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 320, 256, 1, 1 }, { 1, 2 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
{ { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 1, 1 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 1, 3 }, { 2, 4 } },
{ { GGML_METAL_DEVICE_M3_PRO, GGML_TYPE_F16, 32, 32, 2, 1 }, { 2, 4 } },
ggml/src/ggml.c
+1 -1
View File
@@ -7351,7 +7351,7 @@ void ggml_build_backward_expand(
}
// inplace operations are currently not supported
GGML_ASSERT(!node->view_src || node->op == GGML_OP_CPY || node->op == GGML_OP_VIEW ||
GGML_ASSERT(!node->view_src || node->op == GGML_OP_CPY || node->op == GGML_OP_SET_ROWS || node->op == GGML_OP_VIEW ||
node->op == GGML_OP_RESHAPE || node->op == GGML_OP_PERMUTE || node->op == GGML_OP_TRANSPOSE);
const size_t ihash = ggml_hash_find(&cgraph->visited_hash_set, node);
gguf-py/gguf/constants.py
+3
View File
@@ -697,6 +697,7 @@ class MODEL_TENSOR(IntEnum):
FFN_DOWN_CHEXP = auto()
FFN_UP_CHEXP = auto()
FFN_EXP_PROBS_B = auto()
FFN_EXP_PROBS_B_VL = auto() # deepseek4 vision (bias for image tokens)
FFN_GATE_TID2EID = auto()
MOE_LATENT_DOWN = auto() # nemotron 3 super
MOE_LATENT_UP = auto() # nemotron 3 super
@@ -1449,6 +1450,7 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
MODEL_TENSOR.FFN_UP_EXP: "blk.{bid}.ffn_up_exps",
MODEL_TENSOR.FFN_GATE_UP_EXP: "blk.{bid}.ffn_gate_up_exps",
MODEL_TENSOR.FFN_EXP_PROBS_B: "blk.{bid}.exp_probs_b",
MODEL_TENSOR.FFN_EXP_PROBS_B_VL: "blk.{bid}.exp_probs_b_vl",
MODEL_TENSOR.FFN_GATE_TID2EID: "blk.{bid}.ffn_gate_tid2eid",
MODEL_TENSOR.MOE_LATENT_DOWN: "blk.{bid}.ffn_latent_down", # nemotron 3 super
MODEL_TENSOR.MOE_LATENT_UP: "blk.{bid}.ffn_latent_up", # nemotron 3 super
@@ -3839,6 +3841,7 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_GATE_INP,
MODEL_TENSOR.FFN_GATE_TID2EID,
MODEL_TENSOR.FFN_EXP_PROBS_B,
MODEL_TENSOR.FFN_EXP_PROBS_B_VL,
MODEL_TENSOR.FFN_NORM,
MODEL_TENSOR.FFN_GATE_EXP,
MODEL_TENSOR.FFN_DOWN_EXP,
gguf-py/gguf/gguf_writer.py
+10 -4
View File
@@ -733,8 +733,11 @@ class GGUFWriter:
else:
self.add_array(Keys.LLM.FEED_FORWARD_LENGTH.format(arch=self.arch), length)
def add_expert_feed_forward_length(self, length: int) -> None:
self.add_uint32(Keys.LLM.EXPERT_FEED_FORWARD_LENGTH.format(arch=self.arch), length)
def add_expert_feed_forward_length(self, length: int | Sequence[int]) -> None:
if isinstance(length, int):
self.add_uint32(Keys.LLM.EXPERT_FEED_FORWARD_LENGTH.format(arch=self.arch), length)
else:
self.add_array(Keys.LLM.EXPERT_FEED_FORWARD_LENGTH.format(arch=self.arch), length)
def add_expert_shared_feed_forward_length(self, length: int) -> None:
self.add_uint32(Keys.LLM.EXPERT_SHARED_FEED_FORWARD_LENGTH.format(arch=self.arch), length)
@@ -860,8 +863,11 @@ class GGUFWriter:
def add_expert_count(self, count: int) -> None:
self.add_uint32(Keys.LLM.EXPERT_COUNT.format(arch=self.arch), count)
def add_expert_used_count(self, count: int) -> None:
self.add_uint32(Keys.LLM.EXPERT_USED_COUNT.format(arch=self.arch), count)
def add_expert_used_count(self, count: int | Sequence[int]) -> None:
if isinstance(count, int):
self.add_uint32(Keys.LLM.EXPERT_USED_COUNT.format(arch=self.arch), count)
else:
self.add_array(Keys.LLM.EXPERT_USED_COUNT.format(arch=self.arch), count)
def add_expert_shared_count(self, count: int) -> None:
self.add_uint32(Keys.LLM.EXPERT_SHARED_COUNT.format(arch=self.arch), count)
src/llama-arch.cpp
+2
View File
@@ -457,6 +457,7 @@ static const std::map<llm_tensor, const char *> LLM_TENSOR_NAMES = {
{ LLM_TENSOR_FFN_UP_SHEXP, "blk.%d.ffn_up_shexp" },
{ LLM_TENSOR_FFN_DOWN_SHEXP, "blk.%d.ffn_down_shexp" },
{ LLM_TENSOR_FFN_EXP_PROBS_B, "blk.%d.exp_probs_b" },
{ LLM_TENSOR_FFN_EXP_PROBS_B_VL, "blk.%d.exp_probs_b_vl" },
{ LLM_TENSOR_FFN_LATENT_DOWN, "blk.%d.ffn_latent_down" },
{ LLM_TENSOR_FFN_LATENT_UP, "blk.%d.ffn_latent_up" },
{ LLM_TENSOR_ATTN_NORM_2, "blk.%d.attn_norm_2" },
@@ -896,6 +897,7 @@ static const std::map<llm_tensor, llm_tensor_info> LLM_TENSOR_INFOS = {
{LLM_TENSOR_FFN_GATE_CHEXPS, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT_ID}},
{LLM_TENSOR_FFN_UP_CHEXPS, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT_ID}},
{LLM_TENSOR_FFN_EXP_PROBS_B, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_ADD}},
{LLM_TENSOR_FFN_EXP_PROBS_B_VL, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_ADD}},
// altup / laurel (gemma 3n)
{LLM_TENSOR_PER_LAYER_TOKEN_EMBD, {LLM_TENSOR_LAYER_INPUT, GGML_OP_GET_ROWS}},
{LLM_TENSOR_PER_LAYER_MODEL_PROJ, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}},
src/llama-arch.h
+1
View File
@@ -477,6 +477,7 @@ enum llm_tensor {
LLM_TENSOR_FFN_GATE_CHEXPS,
LLM_TENSOR_FFN_UP_CHEXPS,
LLM_TENSOR_FFN_EXP_PROBS_B,
LLM_TENSOR_FFN_EXP_PROBS_B_VL,
LLM_TENSOR_FFN_LATENT_DOWN,
LLM_TENSOR_FFN_LATENT_UP,
LLM_TENSOR_ATTN_Q_NORM,
src/llama-context.cpp
+11 -1
View File
@@ -490,7 +490,8 @@ llama_context::~llama_context() {
// wait for any pending asynchronous copies into the output buffers before they are freed
synchronize();
if (!model.hparams.no_alloc) {
// when training, ggml_opt allocates extra buffers through the scheduler, so the sizes no longer match the expectation
if (!model.hparams.no_alloc && !opt_ctx) {
for (size_t i = 0; i < backend_ptrs.size(); ++i) {
ggml_backend_t backend = backend_ptrs[i];
ggml_backend_buffer_type_t buft = backend_buft[i];
@@ -3418,6 +3419,15 @@ void llama_context::opt_init(struct llama_model * model, struct llama_opt_params
GGML_ASSERT(model->hparams.n_ctx_train % n_batch == 0);
GGML_ASSERT(n_batch % n_ubatch == 0);
if (cparams.flash_attn) {
LLAMA_LOG_INFO("%s: disabling flash attention, FLASH_ATTN_EXT has no backward pass\n", __func__);
cparams.flash_attn = false;
// the graph changes without flash attention, need to reserve again
sched_need_reserve = true;
sched_reserve();
}
ggml_opt_params opt_params = ggml_opt_default_params(sched.get(), GGML_OPT_LOSS_TYPE_CROSS_ENTROPY);
opt_params.opt_period = n_batch / n_ubatch;
opt_params.get_opt_pars = lopt_params.get_opt_pars;
src/llama-graph.cpp
+8 -7
View File
@@ -1467,7 +1467,7 @@ llm_graph_context::llm_graph_context(const llm_graph_params & params) :
n_embd_head_v (hparams.n_embd_head_v()),
n_embd_v_gqa (hparams.n_embd_v_gqa()),
n_expert (hparams.n_expert),
n_expert_used (cparams.warmup ? hparams.n_expert : hparams.n_expert_used),
n_expert_used (cparams.warmup ? hparams.n_expert : hparams.n_expert_used()),
freq_base (cparams.rope_freq_base),
freq_scale (cparams.rope_freq_scale),
ext_factor (cparams.yarn_ext_factor),
@@ -2271,25 +2271,26 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
assert(n_expert_used > 0);
// order the views before the adds
for (uint32_t i = 0; i < hparams.n_expert_used; ++i) {
// Use per-layer n_expert_used to bound the graph even during warmup (avoids
// the large-add-nodes issue for uniform arches; for Puzzle the per-layer
// value is correct). ref: https://github.com/ggml-org/llama.cpp/pull/14753
const uint32_t n_expert_used_il = hparams.n_expert_used(il);
for (uint32_t i = 0; i < n_expert_used_il; ++i) {
cur_experts[i] = ggml_view_2d(ctx0, experts, n_embd, n_tokens, experts->nb[2], i*experts->nb[1]);
ggml_build_forward_expand(gf, cur_experts[i]);
}
// aggregate experts
// note: here we explicitly use hparams.n_expert_used instead of n_expert_used
// to avoid potentially a large number of add nodes during warmup
// ref: https://github.com/ggml-org/llama.cpp/pull/14753
ggml_tensor * moe_out = cur_experts[0];
for (uint32_t i = 1; i < hparams.n_expert_used; ++i) {
for (uint32_t i = 1; i < n_expert_used_il; ++i) {
moe_out = ggml_add(ctx0, moe_out, cur_experts[i]);
ggml_build_forward_expand(gf, moe_out);
}
if (hparams.n_expert_used == 1) {
if (n_expert_used_il == 1) {
// avoid returning a non-contiguous tensor
moe_out = ggml_cont(ctx0, moe_out);
}
src/llama-hparams.cpp
+16
View File
@@ -71,6 +71,22 @@ uint32_t llama_hparams::n_ff(uint32_t il) const {
GGML_ABORT("fatal error");
}
uint32_t llama_hparams::n_ff_exp(uint32_t il) const {
if (il < n_layer_all) {
return n_ff_exp_arr[il];
}
GGML_ABORT("fatal error");
}
uint32_t llama_hparams::n_expert_used(uint32_t il) const {
if (il < n_layer_all) {
return n_expert_used_arr[il];
}
GGML_ABORT("fatal error");
}
uint32_t llama_hparams::n_gqa(uint32_t il) const {
const uint32_t n_head = this->n_head(il);
const uint32_t n_head_kv = this->n_head_kv(il);
src/llama-hparams.h
+13 -2
View File
@@ -62,7 +62,6 @@ struct llama_hparams {
// per-token adapter selection. -1 when the model has no such layer.
int32_t router_layer = -1;
uint32_t n_expert = 0;
uint32_t n_expert_used = 0;
uint32_t n_rel_attn_bkts = 0;
// TODO: this needs to be reworked
@@ -92,10 +91,14 @@ struct llama_hparams {
std::array<uint32_t, LLAMA_MAX_LAYERS> n_head_kv_arr;
std::array<uint32_t, LLAMA_MAX_LAYERS> n_ff_arr;
// per-layer expert feed-forward size
std::array<uint32_t, LLAMA_MAX_LAYERS> n_ff_exp_arr;
// per-layer top-k expert routing count
std::array<uint32_t, LLAMA_MAX_LAYERS> n_expert_used_arr;
uint32_t n_layer_dense_lead = 0;
uint32_t n_lora_q = 0;
uint32_t n_lora_kv = 0;
uint32_t n_ff_exp = 0;
uint32_t n_ff_shexp = 0;
uint32_t n_ff_chexp = 0;
uint32_t n_expert_shared = 0;
@@ -161,6 +164,10 @@ struct llama_hparams {
// the size of the sliding window (0 - no SWA)
uint32_t n_swa = 0;
// deepseek4 vision: when decoding non-causally (multimodal input), SWA is not applied between tokens of the current ubatch (the image span); older tokens are still window-clipped
// for other models (like gemma 3, gemma 4): SWA is always applied to match transformers implementation
bool swa_full_non_causal = false;
// if is_swa_impl[il] == 1, then layer il is SWA
// if is_swa_impl[il] == 0, then layer il is dense (i.e. non-SWA)
// by default, all layers are dense
@@ -381,6 +388,10 @@ struct llama_hparams {
uint32_t n_ff(uint32_t il = 0) const;
uint32_t n_ff_exp(uint32_t il = 0) const;
uint32_t n_expert_used(uint32_t il = 0) const;
uint32_t n_gqa(uint32_t il = 0) const;
uint32_t n_rot(uint32_t il = 0) const;
src/llama-kv-cache.cpp
+3 -1
View File
@@ -1686,7 +1686,9 @@ static void set_input_kq_mask_impl(const args_set_input_kq_mask & args, T * data
// apply SWA if any
if (swa) {
if (llama_hparams::is_masked_swa(n_swa, swa_type, p0, p1)) {
// see llama_hparams::swa_full_non_causal
const bool in_span = !causal && args.hparams.swa_full_non_causal && p0 >= seq_pos_min[seq_id];
if (!in_span && llama_hparams::is_masked_swa(n_swa, swa_type, p0, p1)) {
goto skip;
}
}
src/llama-model-loader.cpp
+22 -4
View File
@@ -952,7 +952,7 @@ static bool weight_buft_supported(const llama_hparams & hparams, ggml_tensor * w
case GGML_OP_MUL_MAT_ID:
{
// Used for either MoE expert routing or embedded adapter routing
const int n_ids_used = hparams.router_layer >= 0 ? 1 : hparams.n_expert_used;
const int n_ids_used = hparams.router_layer >= 0 ? 1 : hparams.n_expert_used();
GGML_ASSERT(n_ids_used > 0);
ggml_tensor * b = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, w->ne[0], n_ids_used, 512);
ggml_tensor * ids = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_ids_used, 512);
@@ -965,7 +965,7 @@ static bool weight_buft_supported(const llama_hparams & hparams, ggml_tensor * w
} break;
case GGML_OP_ADD_ID:
{
const int n_expert_used = hparams.n_expert_used;
const int n_expert_used = hparams.n_expert_used();
GGML_ASSERT(n_expert_used > 0);
ggml_tensor * a = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, w->ne[0], n_expert_used, 512);
ggml_tensor * c = ggml_new_tensor_2d(ctx, GGML_TYPE_I32, n_expert_used, 512);
@@ -1515,7 +1515,6 @@ bool llama_model_loader::load_all_data(
}
GGML_ASSERT(size_data != 0 && "call init_mappings() first");
std::vector<no_init<uint8_t>> read_buf;
std::vector<std::future<std::pair<ggml_tensor *, bool>>> validation_result;
// 4 staging buffers for async uploads, each sized 1MB seems to be a good default for single NVMe drives.
@@ -1616,7 +1615,25 @@ bool llama_model_loader::load_all_data(
ggml_backend_name(upload_backend));
}
std::vector<ggml_tensor *> tensors;
for (struct ggml_tensor * cur = ggml_get_first_tensor(ctx); cur != NULL; cur = ggml_get_next_tensor(ctx, cur)) {
tensors.push_back(cur);
}
// without mmap, tensors in non-host buffers are staged through a temporary buffer sized like the tensor
// load them biggest-first so the largest staging buffer is allocated while the fewest weights are resident
if (!use_mmap) {
std::stable_sort(tensors.begin(), tensors.end(), [](const ggml_tensor * a, const ggml_tensor * b) {
const bool staged_a = a->buffer && !ggml_backend_buffer_is_host(a->buffer);
const bool staged_b = b->buffer && !ggml_backend_buffer_is_host(b->buffer);
if (staged_a != staged_b) {
return staged_a;
}
return staged_a && ggml_nbytes(a) > ggml_nbytes(b);
});
}
for (struct ggml_tensor * cur : tensors) {
const auto * weight = get_weight(ggml_get_name(cur));
if (weight == nullptr) {
// this can happen with split experts models
@@ -1729,7 +1746,8 @@ bool llama_model_loader::load_all_data(
buffer_idx %= n_buffers;
}
} else {
read_buf.resize(n_size);
// scoped to one tensor so only one staging buffer is alive at a time
std::vector<no_init<uint8_t>> read_buf(n_size);
file->seek(weight->offs, SEEK_SET);
file->read_raw(read_buf.data(), n_size);
ggml_backend_tensor_set(cur, read_buf.data(), 0, n_size);
src/llama-model-saver.cpp
+2 -2
View File
@@ -222,7 +222,7 @@ void llama_model_saver::add_kv_from_model() {
add_kv(LLM_KV_BLOCK_COUNT, hparams.n_layer_all);
add_kv(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
add_kv(LLM_KV_FEED_FORWARD_LENGTH, hparams.n_ff_arr, true);
add_kv(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
add_kv(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp());
add_kv(LLM_KV_EXPERT_LATENT_LENGTH, hparams.n_expert_latent);
add_kv(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp);
add_kv(LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, hparams.n_ff_chexp);
@@ -233,7 +233,7 @@ void llama_model_saver::add_kv_from_model() {
add_kv(LLM_KV_USE_PARALLEL_RESIDUAL, hparams.use_par_res);
// add_kv(LLM_KV_TENSOR_DATA_LAYOUT, ???);
add_kv(LLM_KV_EXPERT_COUNT, hparams.n_expert);
add_kv(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used);
add_kv(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used());
add_kv(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
add_kv(LLM_KV_EXPERT_GROUP_COUNT, hparams.n_expert_groups);
add_kv(LLM_KV_EXPERT_GROUP_USED_COUNT, hparams.n_group_used);
src/llama-model.cpp
+38 -21
View File
@@ -786,7 +786,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
// the FFN is the same for Qwen 3 Next and Qwen 3.5:
if (std::regex_match(tensor_name, pattern_ffn_gate_up_weight)) {
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp(il);
GGML_ASSERT(tensor->ne[axis] == 2*n_ff_exp);
return {{n_ff_exp, 2}};
}
@@ -809,7 +809,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
return {{tensor->ne[axis], 1}};
}
if (std::regex_match(tensor_name, pattern_ffn_gate_up_weight)) {
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp(il);
GGML_ASSERT(tensor->ne[axis] == 2*n_ff_exp);
return {{n_ff_exp, 2}};
}
@@ -1095,6 +1095,7 @@ const char * llm_type_name(llm_type type) {
case LLM_TYPE_31B_A3_5B: return "31B.A3.5B";
case LLM_TYPE_35B_A3B: return "35B.A3B";
case LLM_TYPE_48B_A3B: return "48B.A3B";
case LLM_TYPE_75B_A9B: return "75B.A9B";
case LLM_TYPE_80B_A3B: return "80B.A3B";
case LLM_TYPE_A3B: return "A3B";
case LLM_TYPE_100B_A6B: return "100B.A6B";
@@ -1378,14 +1379,15 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_NEXTN_PREDICT_LAYERS, hparams.n_layer_nextn, false);
GGML_ASSERT(hparams.n_layer_nextn <= hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_COUNT, hparams.n_expert, false);
ml.get_key(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used, false);
std::fill(hparams.n_expert_used_arr.begin(), hparams.n_expert_used_arr.end(), 0);
ml.get_key_or_arr(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_GROUP_COUNT, hparams.n_expert_groups, false);
ml.get_key(LLM_KV_EXPERT_GROUP_USED_COUNT, hparams.n_group_used, false);
if (arch == LLM_ARCH_HUNYUAN_VL || arch == LLM_ARCH_HUNYUAN_DENSE) {
if (hparams.n_expert <= 1) {
hparams.n_expert = 0;
hparams.n_expert_used = 0;
hparams.n_expert = 0;
std::fill(hparams.n_expert_used_arr.begin(), hparams.n_expert_used_arr.end(), 0);
}
}
@@ -1403,10 +1405,16 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
GGML_ASSERT(hparams.convnext.n_layer <= hparams.n_layer_all);
}
// models may route a different number of experts per layer, so validate the maximum
uint32_t n_expert_used_max = 0;
for (uint32_t il = 0; il < hparams.n_layer_all; ++il) {
n_expert_used_max = std::max(n_expert_used_max, hparams.n_expert_used(il));
}
GGML_ASSERT(hparams.n_expert <= LLAMA_MAX_EXPERTS);
GGML_ASSERT(hparams.n_expert_used <= hparams.n_expert);
GGML_ASSERT(n_expert_used_max <= hparams.n_expert);
if (hparams.n_expert > 0) {
GGML_ASSERT(hparams.n_expert_used > 0);
GGML_ASSERT(n_expert_used_max > 0);
GGML_ASSERT(hparams.n_expert_groups < hparams.n_expert);
if (hparams.n_expert_groups > 1) {
GGML_ASSERT(hparams.n_expert % hparams.n_expert_groups == 0);
@@ -1414,13 +1422,14 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
GGML_ASSERT(hparams.n_group_used < hparams.n_expert_groups);
}
} else {
GGML_ASSERT(hparams.n_expert_used == 0);
GGML_ASSERT(n_expert_used_max == 0);
GGML_ASSERT(hparams.n_expert_groups == 0);
}
std::fill(hparams.n_head_arr.begin(), hparams.n_head_arr.end(), 0);
std::fill(hparams.n_head_kv_arr.begin(), hparams.n_head_kv_arr.end(), 0);
std::fill(hparams.n_ff_arr.begin(), hparams.n_ff_arr.end(), 0);
std::fill(hparams.n_head_arr.begin(), hparams.n_head_arr.end(), 0);
std::fill(hparams.n_head_kv_arr.begin(), hparams.n_head_kv_arr.end(), 0);
std::fill(hparams.n_ff_arr.begin(), hparams.n_ff_arr.end(), 0);
std::fill(hparams.n_ff_exp_arr.begin(), hparams.n_ff_exp_arr.end(), 0);
std::fill(hparams.rope_sections.begin(), hparams.rope_sections.end(), 0);
std::fill(hparams.rope_pattern.begin(), hparams.rope_pattern.end(), 1);
@@ -1653,7 +1662,7 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
const auto tn = LLM_TN(arch);
const int64_t n_expert = hparams.n_expert;
const int64_t n_expert_used = hparams.n_expert_used;
const int64_t n_expert_used = hparams.n_expert_used();
if (n_expert > 0 && n_expert_used == 0) {
throw std::runtime_error("model has expert layers but no expert layers are used");
@@ -1959,6 +1968,14 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
return true;
}
// without mmap, load non-host buffers first: their tensors go through a staging buffer, which is cheapest while the fewest weights are resident
if (!ml.use_mmap) {
std::stable_partition(ctx_buf_maps.begin(), ctx_buf_maps.end(), [](const auto & ctx_buf_map) {
const auto & buf_map = ctx_buf_map.second;
return !buf_map.empty() && !ggml_backend_buffer_is_host(buf_map.begin()->second);
});
}
// load tensor data
for (auto & [ctx, buf_map] : ctx_buf_maps) {
if (!ml.load_all_data(ctx, buf_map, use_mlock ? &pimpl->mlock_mmaps : NULL, params.progress_callback, params.progress_callback_user_data)) {
@@ -2109,7 +2126,7 @@ void llama_model::print_info() const {
LLAMA_LOG_INFO("%s: f_attn_value_scale = %.4f\n", __func__, hparams.f_attn_value_scale);
LLAMA_LOG_INFO("%s: n_ff = %s\n", __func__, print_f([&](uint32_t il) { return hparams.n_ff(il); }, hparams.n_layer_all).c_str());
LLAMA_LOG_INFO("%s: n_expert = %u\n", __func__, hparams.n_expert);
LLAMA_LOG_INFO("%s: n_expert_used = %u\n", __func__, hparams.n_expert_used);
LLAMA_LOG_INFO("%s: n_expert_used = %u\n", __func__, hparams.n_expert_used());
LLAMA_LOG_INFO("%s: n_expert_groups = %d\n", __func__, hparams.n_expert_groups);
LLAMA_LOG_INFO("%s: n_group_used = %d\n", __func__, hparams.n_group_used);
LLAMA_LOG_INFO("%s: causal attn = %d\n", __func__, hparams.causal_attn);
@@ -2184,7 +2201,7 @@ void llama_model::print_info() const {
if (arch == LLM_ARCH_DEEPSEEK) {
LLAMA_LOG_INFO("%s: n_layer_dense_lead = %d\n", __func__, hparams.n_layer_dense_lead);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_expert_shared = %d\n", __func__, hparams.n_expert_shared);
LLAMA_LOG_INFO("%s: expert_weights_scale = %.1f\n", __func__, hparams.expert_weights_scale);
}
@@ -2197,7 +2214,7 @@ void llama_model::print_info() const {
LLAMA_LOG_INFO("%s: n_lora_kv = %d\n", __func__, hparams.n_lora_kv);
LLAMA_LOG_INFO("%s: n_embd_head_k_mla = %d\n", __func__, hparams.n_embd_head_k_mla());
LLAMA_LOG_INFO("%s: n_embd_head_v_mla = %d\n", __func__, hparams.n_embd_head_v_mla());
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_expert_shared = %d\n", __func__, hparams.n_expert_shared);
LLAMA_LOG_INFO("%s: expert_weights_scale = %.1f\n", __func__, hparams.expert_weights_scale);
LLAMA_LOG_INFO("%s: expert_weights_norm = %d\n", __func__, hparams.expert_weights_norm);
@@ -2205,7 +2222,7 @@ void llama_model::print_info() const {
}
if (arch == LLM_ARCH_QWEN2MOE) {
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_ff_shexp = %d\n", __func__, hparams.n_ff_shexp);
}
@@ -2215,7 +2232,7 @@ void llama_model::print_info() const {
arch == LLM_ARCH_OPENAI_MOE ||
arch == LLM_ARCH_QWEN3VLMOE ||
arch == LLM_ARCH_RND1) {
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
}
if (arch == LLM_ARCH_MINICPM ||
@@ -2232,7 +2249,7 @@ void llama_model::print_info() const {
if (arch == LLM_ARCH_BAILINGMOE) {
LLAMA_LOG_INFO("%s: n_layer_dense_lead = %d\n", __func__, hparams.n_layer_dense_lead);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_expert_shared = %d\n", __func__, hparams.n_expert_shared);
LLAMA_LOG_INFO("%s: expert_weights_scale = %.1f\n", __func__, hparams.expert_weights_scale);
LLAMA_LOG_INFO("%s: expert_weights_norm = %d\n", __func__, hparams.expert_weights_norm);
@@ -2240,7 +2257,7 @@ void llama_model::print_info() const {
if (arch == LLM_ARCH_BAILINGMOE2 || arch == LLM_ARCH_BAILINGMOE3) {
LLAMA_LOG_INFO("%s: n_layer_dense_lead = %d\n", __func__, hparams.n_layer_dense_lead);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_ff_shexp = %d\n", __func__, hparams.n_ff_shexp);
LLAMA_LOG_INFO("%s: n_expert_shared = %d\n", __func__, hparams.n_expert_shared);
LLAMA_LOG_INFO("%s: expert_weights_scale = %.1f\n", __func__, hparams.expert_weights_scale);
@@ -2250,12 +2267,12 @@ void llama_model::print_info() const {
}
if (arch == LLM_ARCH_SMALLTHINKER || arch == LLM_ARCH_LFM2MOE) {
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: expert_gating_func = %s\n", __func__, llama_expert_gating_func_name((llama_expert_gating_func_type) hparams.expert_gating_func));
}
if (arch == LLM_ARCH_GROVEMOE) {
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp());
LLAMA_LOG_INFO("%s: n_ff_chexp = %d\n", __func__, hparams.n_ff_chexp);
LLAMA_LOG_INFO("%s: n_group_experts = %d\n", __func__, hparams.n_group_experts);
LLAMA_LOG_INFO("%s: expert_group_scale = %.2f\n", __func__, hparams.expert_group_scale);
src/llama-model.h
+3 -1
View File
@@ -128,6 +128,7 @@ enum llm_type {
LLM_TYPE_31B_A3_5B,
LLM_TYPE_35B_A3B, // Qwen3.5
LLM_TYPE_48B_A3B, // Kimi Linear
LLM_TYPE_75B_A9B, // Nemotron 3 Puzzle
LLM_TYPE_80B_A3B, // Qwen3 Next
LLM_TYPE_A3B, // Qwen3.8 Flash Next
LLM_TYPE_100B_A6B,
@@ -362,6 +363,7 @@ struct llama_layer {
struct ggml_tensor * ffn_up_b = nullptr; // b3
struct ggml_tensor * ffn_act = nullptr;
struct ggml_tensor * ffn_exp_probs_b = nullptr;
struct ggml_tensor * ffn_exp_probs_b_vl = nullptr; // deepseek4 vision (bias for image tokens)
struct ggml_tensor * ffn_gate_tid2eid = nullptr;
struct ggml_tensor * dflash_attn_conv_base = nullptr;
@@ -838,7 +840,7 @@ const char * llm_type_name(llm_type type);
const int64_t n_token_types = vocab.n_token_types(); GGML_UNUSED(n_token_types); \
const int64_t n_rot = hparams.n_rot(); GGML_UNUSED(n_rot); \
const int64_t n_expert = hparams.n_expert; GGML_UNUSED(n_expert); \
const int64_t n_expert_used = hparams.n_expert_used; GGML_UNUSED(n_expert_used); \
const int64_t n_expert_used = hparams.n_expert_used(); GGML_UNUSED(n_expert_used); \
const int64_t n_ctx_train = hparams.n_ctx_train; GGML_UNUSED(n_ctx_train);
// For internal test use
src/models/afmoe.cpp
+2 -2
View File
@@ -3,7 +3,7 @@
void llama_model_afmoe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -52,7 +52,7 @@ void llama_model_afmoe::load_arch_tensors(llama_model_loader &) {
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED);
}
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
src/models/bailingmoe.cpp
+2 -2
View File
@@ -3,7 +3,7 @@
void llama_model_bailingmoe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -19,7 +19,7 @@ void llama_model_bailingmoe::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/bailingmoe2.cpp
+2 -2
View File
@@ -3,7 +3,7 @@
void llama_model_bailingmoe2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -21,7 +21,7 @@ void llama_model_bailingmoe2::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/bailingmoe3.cpp
+8 -8
View File
@@ -15,7 +15,7 @@ void llama_model_bailingmoe3::load_arch_hparams(llama_model_loader & ml) {
hparams.kda_safe_gate = true;
}
ml.get_key(LLM_KV_KDA_GATE_LOWER_BOUND, hparams.kda_gate_lower_bound);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
@@ -26,7 +26,7 @@ void llama_model_bailingmoe3::load_arch_hparams(llama_model_loader & ml) {
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer_all, false);
if (hparams.n_ff_shexp == 0) {
hparams.n_ff_shexp = hparams.n_ff_exp * std::max(1u, hparams.n_expert_shared);
hparams.n_ff_shexp = hparams.n_ff_exp() * std::max(1u, hparams.n_expert_shared);
}
GGML_ASSERT(hparams.kda_safe_gate);
@@ -115,9 +115,9 @@ void llama_model_bailingmoe3::load_arch_tensors(llama_model_loader & ml) {
} else {
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", il), { n_embd, n_expert }, trunk_flags);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", il), { n_expert }, trunk_flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, trunk_flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, trunk_flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp, n_embd, n_expert }, trunk_flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp(), n_expert }, trunk_flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp(), n_expert }, trunk_flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp(), n_embd, n_expert }, trunk_flags);
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, trunk_flags);
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, trunk_flags);
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", il), { hparams.n_ff_shexp, n_embd }, trunk_flags);
@@ -145,9 +145,9 @@ void llama_model_bailingmoe3::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", il), { n_embd }, flags);
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", il), { n_embd, n_expert }, flags);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", il), { n_expert }, flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp, n_embd, n_expert }, flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp(), n_expert }, flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp(), n_expert }, flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp(), n_embd, n_expert }, flags);
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, flags);
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, flags);
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", il), { hparams.n_ff_shexp, n_embd }, flags);
src/models/bert.cpp
+1 -1
View File
@@ -182,7 +182,7 @@ llama_model_bert::graph::graph(const llama_model & model, const llm_graph_params
nullptr,
model.layers[il].ffn_down_exps,
nullptr,
hparams.n_expert, hparams.n_expert_used,
hparams.n_expert, hparams.n_expert_used(),
LLM_FFN_GELU, false,
hparams.expert_weights_scale,
LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
src/models/cohere2moe.cpp
+3 -3
View File
@@ -13,7 +13,7 @@ void llama_model_cohere2moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_LOGIT_SCALE, hparams.f_logit_scale);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -89,7 +89,7 @@ void llama_model_cohere2moe::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd }, flags);
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), { n_embd, n_ff }, flags);
} else {
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff;
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), { n_embd, n_expert }, flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff_exp, n_embd, n_expert }, flags);
@@ -113,7 +113,7 @@ void llama_model_cohere2moe::load_arch_tensors(llama_model_loader & ml) {
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head, n_embd_gqa, n_embd_gqa, flags);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), { n_embd_head_k * n_head, n_embd }, flags);
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff;
// Routed experts
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), { n_embd, n_expert }, flags);
src/models/deepseek.cpp
+3 -3
View File
@@ -3,11 +3,11 @@
void llama_model_deepseek::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
switch (hparams.n_ff_exp) {
switch (hparams.n_ff_exp()) {
case 1408: type = LLM_TYPE_16B; break;
case 1792: type = LLM_TYPE_20B; break;
default: type = LLM_TYPE_UNKNOWN;
@@ -19,7 +19,7 @@ void llama_model_deepseek::load_arch_tensors(llama_model_loader &) {
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/deepseek2.cpp
+2 -2
View File
@@ -15,7 +15,7 @@ void llama_model_deepseek2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK, hparams.n_lora_kv);
ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl, false);
ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, hparams.n_embd_head_v_mla_impl, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -79,7 +79,7 @@ void llama_model_deepseek2::load_arch_tensors(llama_model_loader & ml) {
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t kv_lora_rank = hparams.n_lora_kv;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/deepseek2ocr.cpp
+2 -2
View File
@@ -4,7 +4,7 @@ void llama_model_deepseek2ocr::load_arch_hparams(llama_model_loader & ml) {
// similar to deepseek2, but without MLA
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -25,7 +25,7 @@ void llama_model_deepseek2ocr::load_arch_tensors(llama_model_loader &) {
const int64_t n_expert_shared = hparams.n_expert_shared;
// similar to deepseek2, but without MLA
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/deepseek32.cpp
+3 -3
View File
@@ -4,7 +4,7 @@
#include "llama-kv-cache-dsa.h"
void llama_model_deepseek32::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
hparams.f_norm_eps = 1e-6; // eps for layer norm
ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false);
@@ -20,7 +20,7 @@ void llama_model_deepseek32::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK, hparams.n_lora_kv);
ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl, false);
ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, hparams.n_embd_head_v_mla_impl, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
// DSA parameters
@@ -71,7 +71,7 @@ void llama_model_deepseek32::load_arch_tensors(llama_model_loader & ml) {
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t kv_lora_rank = hparams.n_lora_kv;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_expert_shared = hparams.n_expert_shared;
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/deepseek4.cpp
+17 -5
View File
@@ -29,7 +29,7 @@ void llama_model_deepseek4::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_Q_LORA_RANK, hparams.n_lora_q);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm);
@@ -66,6 +66,9 @@ void llama_model_deepseek4::load_arch_hparams(llama_model_loader & ml) {
}
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
hparams.set_swa_pattern(0);
// tokens of an image span attend bidirectionally to the whole span, the window only applies to older tokens
// ref: get_window_topk_idxs_visible in the reference impl
hparams.swa_full_non_causal = true;
for (uint32_t il = hparams.n_layer(); il < hparams.n_layer_all; ++il) {
hparams.is_swa_impl[il] = true;
}
@@ -80,7 +83,7 @@ void llama_model_deepseek4::load_arch_tensors(llama_model_loader & ml) {
LLAMA_LOAD_LOCALS;
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_embd_head = hparams.n_embd_head_k();
@@ -156,6 +159,8 @@ void llama_model_deepseek4::load_arch_tensors(llama_model_loader & ml) {
} else {
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, flags);
}
// vision variant only: routing bias for image tokens
layer.ffn_exp_probs_b_vl = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B_VL, "bias", i), {n_expert}, flags | TENSOR_NOT_REQUIRED);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, flags);
@@ -1275,7 +1280,14 @@ llama_model_deepseek4::graph::graph(const llama_model & model, const llm_graph_p
const auto & layer = model.layers[il];
ggml_tensor * selected_experts = nullptr;
ggml_tensor * exp_probs_b = layer.ffn_exp_probs_b;
if ((uint32_t) il < hparams.dsv4_hash_layer_count) {
// may apply exp_probs_b_vl is input is from mtmd
const bool is_media = ubatch.embd != nullptr;
if (is_media) {
if (layer.ffn_exp_probs_b_vl) {
exp_probs_b = layer.ffn_exp_probs_b_vl;
}
} else if ((uint32_t) il < hparams.dsv4_hash_layer_count) {
selected_experts = ggml_get_rows(ctx0, layer.ffn_gate_tid2eid, res->t_inp_tokens);
exp_probs_b = nullptr;
}
@@ -1286,7 +1298,7 @@ llama_model_deepseek4::graph::graph(const llama_model & model, const llm_graph_p
layer.ffn_gate_exps,
layer.ffn_down_exps,
exp_probs_b,
n_expert, hparams.n_expert_used,
n_expert, hparams.n_expert_used(),
LLM_FFN_SILU, hparams.expert_weights_norm,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
@@ -1443,7 +1455,7 @@ llama_model_deepseek4::graph_mtp::graph_mtp(const llama_model & model, const llm
layer.ffn_gate_exps,
layer.ffn_down_exps,
layer.ffn_exp_probs_b,
n_expert, hparams.n_expert_used,
n_expert, hparams.n_expert_used(),
LLM_FFN_SILU, hparams.expert_weights_norm,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
src/models/dflash.cpp
+3 -3
View File
@@ -40,7 +40,7 @@ void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) {
if (hparams.dsv4_hc_mult > 0) {
ml.get_key(LLM_KV_ATTENTION_Q_LORA_RANK, hparams.n_lora_q);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm);
@@ -159,7 +159,7 @@ void llama_model_dflash::load_arch_tensors(llama_model_loader &) {
if (hparams.dsv4_hc_mult > 0) {
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_embd_head = hparams.n_embd_head_k();
const int64_t o_groups = hparams.dsv4_o_group_count;
@@ -948,7 +948,7 @@ llama_model_dflash::graph_dsv4::graph_dsv4(const llama_model & model, const llm_
layer.ffn_gate_exps,
layer.ffn_down_exps,
layer.ffn_exp_probs_b,
n_expert, hparams.n_expert_used,
n_expert, hparams.n_expert_used(),
LLM_FFN_SILU, hparams.expert_weights_norm,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
src/models/dots1.cpp
+2 -2
View File
@@ -3,7 +3,7 @@
void llama_model_dots1::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -19,7 +19,7 @@ void llama_model_dots1::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/dots3note.cpp
+2 -2
View File
@@ -11,7 +11,7 @@ void llama_model_dots3note::load_arch_hparams(llama_model_loader & ml) {
// MoE parameters
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -56,7 +56,7 @@ void llama_model_dots3note::load_arch_tensors(llama_model_loader & ml) {
const int64_t n_embd_head_qk_rope = hparams.n_rot();
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_expert_shared = hparams.n_expert_shared;
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/ernie4-5.cpp
+2 -2
View File
@@ -6,7 +6,7 @@ void llama_model_ernie4_5::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
if (arch == LLM_ARCH_ERNIE4_5_MOE) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_INTERLEAVE_MOE_LAYER_STEP, hparams.n_moe_layer_step);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
@@ -47,7 +47,7 @@ void llama_model_ernie4_5::load_arch_tensors(llama_model_loader &) {
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0);
if (arch == LLM_ARCH_ERNIE4_5_MOE && static_cast<uint32_t>(i) >= hparams.n_layer_dense_lead) { // MoE layers
int n_ff_exp = hparams.n_ff_exp;
int n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, TENSOR_NOT_REQUIRED);
src/models/exaone-moe.cpp
+2 -2
View File
@@ -13,7 +13,7 @@ void llama_model_exaone_moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -30,7 +30,7 @@ void llama_model_exaone_moe::load_arch_hparams(llama_model_loader & ml) {
void llama_model_exaone_moe::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_ff_shexp = hparams.n_ff_shexp > 0 ? hparams.n_ff_shexp : n_ff_exp;
const int64_t head_dim = hparams.n_embd_head_k();
const int64_t n_qo_dim = n_head * head_dim;
src/models/gemma4.cpp
+2 -2
View File
@@ -11,7 +11,7 @@ void llama_model_gemma4::load_arch_hparams(llama_model_loader & ml) {
hparams.f_attention_scale = 1.0f; // Gemma4 uses self.scaling = 1.0 (no pre-attn scaling)
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, hparams.n_embd_per_layer);
@@ -32,7 +32,7 @@ void llama_model_gemma4::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const uint32_t n_embd_per_layer = hparams.n_embd_per_layer;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
if (n_embd_head_k != n_embd_head_v) {
throw std::runtime_error("Gemma 4 requires n_embd_head_k == n_embd_head_v");
src/models/glm-dsa.cpp
+3 -3
View File
@@ -27,7 +27,7 @@ const std::array<uint32_t, LLAMA_MAX_LAYERS> GLM_5_2_DEFAULT_INDEXER_TYPES = {
};
void llama_model_glm_dsa::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false);
@@ -42,7 +42,7 @@ void llama_model_glm_dsa::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK, hparams.n_lora_kv);
ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl, false);
ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, hparams.n_embd_head_v_mla_impl, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
// DSA parameters
@@ -104,7 +104,7 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader & ml) {
const int64_t q_lora_rank = hparams.n_lora_q;
const int64_t kv_lora_rank = hparams.n_lora_kv;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/glm4-moe.cpp
+3 -3
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_glm4_moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false);
@@ -40,7 +40,7 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
}
GGML_ASSERT(hparams.n_expert > 0 && "n_expert must be > 0 for GLM4_MOE MoE layers");
GGML_ASSERT(hparams.n_expert_used > 0 && "n_expert_used must be > 0 for GLM4_MOE MoE layers");
GGML_ASSERT(hparams.n_expert_used() > 0 && "n_expert_used must be > 0 for GLM4_MOE MoE layers");
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);
@@ -82,7 +82,7 @@ void llama_model_glm4_moe::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), { n_expert }, flags);
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(
tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert }, flags);
src/models/granite-swa.cpp
+1 -1
View File
@@ -11,7 +11,7 @@ void llama_model_granite_swa::load_arch_hparams(llama_model_loader & ml) {
// MoE expert configuration
ml.get_key(LLM_KV_EXPERT_COUNT, hparams.n_expert, false);
ml.get_key(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used, false);
ml.get_key_or_arr(LLM_KV_EXPERT_USED_COUNT, hparams.n_expert_used_arr, hparams.n_layer_all, false);
// iSWA configuration
ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
src/models/grok.cpp
+2 -2
View File
@@ -12,7 +12,7 @@ void llama_model_grok::load_arch_hparams(llama_model_loader & ml) {
hparams.f_final_logit_softcapping = 0.0f;
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_LOGIT_SCALE, hparams.f_logit_scale, false);
ml.get_key(LLM_KV_EMBEDDING_SCALE, hparams.f_embedding_scale, false);
ml.get_key(LLM_KV_ATTENTION_OUTPUT_SCALE, hparams.f_attn_out_scale, false);
@@ -50,7 +50,7 @@ void llama_model_grok::load_arch_tensors(llama_model_loader &) {
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED);
}
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff/* / n_expert_used*/; // grok-1 n_ff_exp == n_ff
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff/* / n_expert_used*/; // grok-1 n_ff_exp == n_ff
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
src/models/grovemoe.cpp
+2 -2
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_grovemoe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, hparams.n_ff_chexp, false);
ml.get_key(LLM_KV_EXPERT_GROUP_SCALE, hparams.expert_group_scale);
ml.get_key(LLM_KV_EXPERTS_PER_GROUP, hparams.n_group_experts);
@@ -46,7 +46,7 @@ void llama_model_grovemoe::load_arch_tensors(llama_model_loader &) {
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
const int64_t n_ff_chexp = hparams.n_ff_chexp ? hparams.n_ff_chexp : n_embd_head_k;
const int64_t n_chunk_expert = n_expert / hparams.n_group_experts;
src/models/hunyuan-moe.cpp
+1 -1
View File
@@ -2,7 +2,7 @@
void llama_model_hunyuan_moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
switch (hparams.n_layer()) {
src/models/hy-v3.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
void llama_model_hy_v3::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -45,7 +45,7 @@ void llama_model_hy_v3::load_arch_tensors(llama_model_loader & ml) {
auto load_block = [&](int i, int flags) {
auto & layer = layers[i];
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / (n_expert_used > 0 ? n_expert_used : 1);
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / (n_expert_used > 0 ? n_expert_used : 1);
const int64_t n_ff_shexp = hparams.n_ff_shexp ? hparams.n_ff_shexp : n_ff_exp;
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, flags);
src/models/kimi-k3.cpp
+3 -3
View File
@@ -30,7 +30,7 @@ void llama_model_kimi_k3::load_arch_hparams(llama_model_loader & ml) {
hparams.is_recr_impl[i] = hparams.n_head_kv(i) == 0;
}
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -139,7 +139,7 @@ void llama_model_kimi_k3::load_arch_tensors(llama_model_loader &) {
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
} else {
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, 0);
@@ -584,7 +584,7 @@ ggml_tensor * llama_model_kimi_k3::graph::build_latent_moe(
layer.ffn_down_exps,
layer.ffn_exp_probs_b,
hparams.n_expert,
hparams.n_expert_used,
hparams.n_expert_used(),
LLM_FFN_SITU, hparams.expert_weights_norm,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
src/models/kimi-linear.cpp
+3 -3
View File
@@ -19,7 +19,7 @@ void llama_model_kimi_linear::load_arch_hparams(llama_model_loader & ml) {
}
// MoE parameters - Kimi uses moe_intermediate_size = 1024
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -137,7 +137,7 @@ void llama_model_kimi_linear::load_arch_tensors(llama_model_loader &) {
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0);
// MoE intermediate size (different from dense FFN)
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
// Kimi uses n_layer_dense_lead to determine which layers use dense FFN vs MoE
// first_k_dense_replace = 1 means layer 0 uses dense FFN, layers 1+ use MoE
@@ -504,7 +504,7 @@ llama_model_kimi_linear::graph::graph(const llama_model & model, const llm_graph
layer.ffn_down_exps,
layer.ffn_exp_probs_b,
hparams.n_expert,
hparams.n_expert_used,
hparams.n_expert_used(),
LLM_FFN_SILU, true,
hparams.expert_weights_scale,
(llama_expert_gating_func_type) hparams.expert_gating_func,
src/models/laguna.cpp
+3 -3
View File
@@ -9,7 +9,7 @@
void llama_model_laguna::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -24,7 +24,7 @@ void llama_model_laguna::load_arch_hparams(llama_model_loader & ml) {
// Weightless fixtures (test-llama-archs) omit this key; derive a nonzero
// size so the shared expert is still built. Real GGUFs always carry the
// exact value (routed and shared FF lengths may differ).
hparams.n_ff_shexp = hparams.n_ff_exp * hparams.n_expert_shared;
hparams.n_ff_shexp = hparams.n_ff_exp() * hparams.n_expert_shared;
}
// Sliding-window attention is OPTIONAL. XS.2 is hybrid (full / SWA / SWA /
@@ -76,7 +76,7 @@ void llama_model_laguna::load_arch_tensors(llama_model_loader & ml) {
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED);
}
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
const int64_t n_ff_shexp = hparams.n_ff_shexp;
for (int i = 0; i < n_layer; ++i) {
src/models/lfm2.cpp
+3 -3
View File
@@ -53,9 +53,9 @@ void llama_model_lfm2::load_arch_tensors(llama_model_loader &) {
if (is_moe_layer) {
GGML_ASSERT(n_expert && n_expert_used);
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, hparams.n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {hparams.n_ff_exp, n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, hparams.n_ff_exp, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, hparams.n_ff_exp(), n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {hparams.n_ff_exp(), n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, hparams.n_ff_exp(), n_expert}, 0);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, 0);
} else { // dense
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
src/models/lfm2moe.cpp
+4 -4
View File
@@ -6,7 +6,7 @@ void llama_model_lfm2moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SHORTCONV_L_CACHE, hparams.n_shortconv_l_cache);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);
for (uint32_t il = 0; il < hparams.n_layer(); ++il) {
@@ -42,9 +42,9 @@ void llama_model_lfm2moe::load_arch_tensors(llama_model_loader &) {
if (is_moe_layer) {
GGML_ASSERT(n_expert && n_expert_used);
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, hparams.n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {hparams.n_ff_exp, n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, hparams.n_ff_exp, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, hparams.n_ff_exp(), n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {hparams.n_ff_exp(), n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, hparams.n_ff_exp(), n_expert}, 0);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, 0);
} else { // dense
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
src/models/llada-moe.cpp
+2 -2
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_llada_moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
// diffusion language model uses non-causal attention
@@ -39,7 +39,7 @@ void llama_model_llada_moe::load_arch_tensors(llama_model_loader &) {
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/llama4.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
void llama_model_llama4::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_INTERLEAVE_MOE_LAYER_STEP, hparams.n_moe_layer_step);
const bool found_swa = ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false);
@@ -75,7 +75,7 @@ void llama_model_llama4::load_arch_tensors(llama_model_loader &) {
layer.rope_freqs = create_tensor(tn(LLM_TENSOR_ROPE_FREQS, "weight", i), {n_rot/2}, TENSOR_NOT_REQUIRED | (i != 0 ? TENSOR_DUPLICATED : 0));
if (is_moe_layer) {
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
src/models/mellum.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
void llama_model_mellum::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false);
if (hparams.n_swa > 0) {
@@ -61,7 +61,7 @@ void llama_model_mellum::load_arch_tensors(llama_model_loader &) {
throw std::runtime_error("n_expert_used must be > 0 for Mellum");
}
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/mimo2.cpp
+2 -2
View File
@@ -5,7 +5,7 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
@@ -62,7 +62,7 @@ void llama_model_mimo2::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED | flags);
// MoE branch
int64_t n_ff_exp = hparams.n_ff_exp;
int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED | flags);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED | flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED | flags);
src/models/minimax-m2.cpp
+1 -1
View File
@@ -2,7 +2,7 @@
void llama_model_minimax_m2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
switch (hparams.n_layer()) {
src/models/minimax-m3.cpp
+2 -2
View File
@@ -13,7 +13,7 @@
void llama_model_minimax_m3::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead, false);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -36,7 +36,7 @@ void llama_model_minimax_m3::load_arch_hparams(llama_model_loader & ml) {
void llama_model_minimax_m3::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_expert_shared = hparams.n_expert_shared;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/nemotron-h.cpp
+23 -7
View File
@@ -1,5 +1,7 @@
#include "models.h"
#include <algorithm> // std::max
void llama_model_nemotron_h::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SSM_CONV_KERNEL, hparams.ssm_d_conv);
ml.get_key(LLM_KV_SSM_INNER_SIZE, hparams.ssm_d_inner);
@@ -16,7 +18,8 @@ void llama_model_nemotron_h::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps); // MTP head final_layernorm
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
// Puzzle models set a different expert FFN size per layer
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
@@ -26,7 +29,17 @@ void llama_model_nemotron_h::load_arch_hparams(llama_model_loader & ml) {
switch (hparams.n_layer()) {
case 52: type = LLM_TYPE_31B_A3_5B; break; // Nemotron-H_MOE 31B
case 56: type = LLM_TYPE_9B; break;
case 88: type = LLM_TYPE_120B_A12B; break;
case 88:
{
// Nemotron 3 Super (uniform MoE) and Nemotron 3 Puzzle (per-layer
// heterogeneous MoE) both have 88 layers; the per-layer top-k array
// is the discriminator.
bool heterogeneous = false;
for (uint32_t i = 1; i < hparams.n_layer(); ++i) {
heterogeneous |= hparams.n_expert_used_arr[i] != hparams.n_expert_used_arr[0];
}
type = heterogeneous ? LLM_TYPE_75B_A9B : LLM_TYPE_120B_A12B;
} break;
default: type = LLM_TYPE_UNKNOWN;
}
}
@@ -94,7 +107,10 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
layer.wo_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
} else {
if (n_expert != 0) {
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
// Use per-layer n_ff_exp; fall back to n_ff/n_expert_used if absent (existing GGUFs).
const int64_t n_ff_exp_i = hparams.n_ff_exp(i)
? (int64_t)hparams.n_ff_exp(i)
: hparams.n_ff(i) / (int64_t)hparams.n_expert_used(i);
const int64_t n_ff_shexp = hparams.n_ff_shexp;
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), { n_embd, n_expert}, trunk_flags);
@@ -104,8 +120,8 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_latent_down = create_tensor(tn(LLM_TENSOR_FFN_LATENT_DOWN, "weight", i), {n_embd, moe_n_embd}, TENSOR_NOT_REQUIRED);
layer.ffn_latent_up = create_tensor(tn(LLM_TENSOR_FFN_LATENT_UP, "weight", i), {moe_n_embd, n_embd}, TENSOR_NOT_REQUIRED);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, moe_n_embd, n_expert}, trunk_flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {moe_n_embd, n_ff_exp, n_expert}, trunk_flags);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp_i, moe_n_embd, n_expert}, trunk_flags);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {moe_n_embd, n_ff_exp_i, n_expert}, trunk_flags);
// Shared expert branch
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", i), {n_ff_shexp, n_embd}, trunk_flags);
@@ -129,7 +145,7 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
const int64_t n_head_i = hparams.n_head(i);
const int64_t n_embd_k_gqa_i = hparams.n_embd_k_gqa(i);
const int64_t n_embd_v_gqa_i = hparams.n_embd_v_gqa(i);
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp(i) ? (int64_t)hparams.n_ff_exp(i) : n_ff / (int64_t)hparams.n_expert_used(i);
const int64_t n_ff_shexp = hparams.n_ff_shexp;
// NextN input-fusion tensors
@@ -280,7 +296,7 @@ ggml_tensor * llama_model_nemotron_h::graph::build_ffn_layer(ggml_tensor * cur,
nullptr, // no gate
model.layers[il].ffn_down_exps,
model.layers[il].ffn_exp_probs_b,
n_expert, n_expert_used,
n_expert, (int64_t)hparams.n_expert_used(il),
LLM_FFN_RELU_SQR, hparams.expert_weights_norm,
hparams.expert_weights_scale,
LLAMA_EXPERT_GATING_FUNC_TYPE_SIGMOID,
src/models/openai-moe.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
void llama_model_openai_moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
@@ -24,7 +24,7 @@ void llama_model_openai_moe::load_arch_hparams(llama_model_loader & ml) {
void llama_model_openai_moe::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
src/models/qwen2moe.cpp
+2 -2
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_qwen2moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -42,7 +42,7 @@ void llama_model_qwen2moe::load_arch_tensors(llama_model_loader &) {
}
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/qwen35moe.cpp
+3 -3
View File
@@ -2,7 +2,7 @@
#include "llama-memory-recurrent.h"
void llama_model_qwen35moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -54,7 +54,7 @@ void llama_model_qwen35moe::load_arch_tensors(llama_model_loader & ml) {
auto load_block_trunk = [&](int il, int flags) {
auto & layer = layers[il];
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
const int64_t n_ff_shexp = hparams.n_ff_shexp ? hparams.n_ff_shexp : n_ff;
// Calculate dimensions from hyperparameters
@@ -106,7 +106,7 @@ void llama_model_qwen35moe::load_arch_tensors(llama_model_loader & ml) {
auto load_block_mtp = [&](int il) {
auto & layer = layers[il];
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
const int64_t n_ff_shexp = hparams.n_ff_shexp ? hparams.n_ff_shexp : n_ff;
// MTP block looks like a full-attention Qwen3.5 decoder block with MoE FFN.
src/models/qwen3moe.cpp
+2 -2
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_qwen3moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
switch (hparams.n_layer()) {
@@ -47,7 +47,7 @@ void llama_model_qwen3moe::load_arch_tensors(llama_model_loader &) {
}
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/qwen3next.cpp
+2 -2
View File
@@ -2,7 +2,7 @@
#include "llama-memory-recurrent.h"
void llama_model_qwen3next::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -50,7 +50,7 @@ void llama_model_qwen3next::load_arch_tensors(llama_model_loader & ml) {
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, TENSOR_DUPLICATED);
}
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
// Calculate dimensions from hyperparameters
const int64_t head_k_dim = hparams.ssm_d_state;
src/models/qwen3vlmoe.cpp
+2 -2
View File
@@ -3,7 +3,7 @@
void llama_model_qwen3vlmoe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_NUM_DEEPSTACK_LAYERS, hparams.n_deepstack_layers, false);
ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, true);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
switch (hparams.n_layer()) {
@@ -49,7 +49,7 @@ void llama_model_qwen3vlmoe::load_arch_tensors(llama_model_loader &) {
}
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/qwen4exp.cpp
+2 -2
View File
@@ -24,7 +24,7 @@ static void qwen4exp_require_arr_len(llama_model_loader & ml, llm_kv kid, uint32
}
void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -191,7 +191,7 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
for (int il = 0; il < n_layer; ++il) {
auto & layer = layers[il];
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
const int64_t n_ff_shexp = hparams.n_ff_shexp ? hparams.n_ff_shexp : n_ff;
const int64_t head_k_dim = hparams.ssm_d_state;
src/models/rnd1.cpp
+2 -2
View File
@@ -1,7 +1,7 @@
#include "models.h"
void llama_model_rnd1::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
switch (hparams.n_layer()) {
@@ -49,7 +49,7 @@ void llama_model_rnd1::load_arch_tensors(llama_model_loader &) {
}
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
const int64_t n_ff_exp = hparams.n_ff_exp() ? hparams.n_ff_exp() : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
src/models/smallthinker.cpp
+2 -2
View File
@@ -18,7 +18,7 @@ void llama_model_smallthinker::load_arch_hparams(llama_model_loader & ml) {
hparams.n_no_rope_layer_step = hparams.n_layer();
}
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp, false);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all, false);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
@@ -57,7 +57,7 @@ void llama_model_smallthinker::load_arch_tensors(llama_model_loader &) {
GGML_ASSERT(n_expert_used > 0 && "n_expert_used must be > 0 for SMALLTHINKER");
// MoE branch
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), { n_embd, n_expert }, 0);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert }, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff_exp, n_embd, n_expert }, 0);
src/models/step35.cpp
+3 -3
View File
@@ -9,7 +9,7 @@ void llama_model_step35::load_arch_hparams(llama_model_loader & ml) {
hparams.n_rot_full = hparams.n_rot_full / 2;
// MoE + SWA parameters
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func, false);
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
@@ -99,7 +99,7 @@ void llama_model_step35::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED);
// MoE routed experts + selection bias (router_bias)
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED);
@@ -150,7 +150,7 @@ void llama_model_step35::load_arch_tensors(llama_model_loader & ml) {
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED);
// MoE routed experts + selection bias (router_bias)
const int64_t n_ff_exp = hparams.n_ff_exp;
const int64_t n_ff_exp = hparams.n_ff_exp();
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED);
tools/mtmd/README-dev.md
+1
View File
@@ -20,6 +20,7 @@ In short:
A typical pipeline of the core libmtmd is as follows:
- A bitmap (RGB image or PCM audio) is created
- Bitmap and the text prompt is provided to `mtmd_tokenize()` that breaks the input into chunks
- Alternatively, `mtmd_tokenize_from_parts()` takes a list of pre-split text/media parts instead of a marker-based prompt
- The tokenizer function first expands a "lazy" bitmap if it finds one. Typically, this is used by video, so that one media token corresponds to one input bitmap
- For models that support "fused" temporal frames like Qwen-VL, the tokenizer tries to merge pair of consecutive frames into one batch. Only bitmaps marked by `mtmd_bitmap_set_mergeable()` are merged
- The preprocessor will then be called, which produces a list of chunks
tools/mtmd/mtmd-cli.cpp
+43 -15
View File
@@ -109,16 +109,15 @@ struct mtmd_cli_context {
mtmd_cli_context(common_params & params) : llama_init(common_init_from_params(params)) {
model = llama_init->model();
lctx = llama_init->context();
if (!model || !lctx) {
exit(1);
}
vocab = llama_model_get_vocab(model);
smpl = common_sampler_init(model, params.sampling);
n_threads = params.cpuparams.n_threads;
batch = llama_batch_init(1, 0, 1); // batch for next token generation
n_batch = params.n_batch;
if (!model || !lctx) {
exit(1);
}
init_vision_context(params);
if (!mtmd_helper_model_can_chat(lctx, ctx_vision.get())) {
@@ -265,21 +264,50 @@ static int eval_message(mtmd_cli_context & ctx, common_chat_msg & msg) {
auto formatted_chat = chat_add_and_format(ctx, msg);
LOG_DBG("formatted_chat.prompt: %s\n", formatted_chat.c_str());
mtmd_input_text text;
text.text = formatted_chat.data();
text.text_len = formatted_chat.size();
text.add_special = add_bos;
text.parse_special = true;
if (g_is_interrupted) return 0;
mtmd::input_chunks chunks(mtmd_input_chunks_init());
// note: we replace the marker here instead of letting mtmd_tokenize() to do that
// because we want to demonstrate how to use mtmd_tokenize_from_parts()
// split the formatted chat on the media marker to get text segments
const std::string marker = mtmd_default_marker();
std::vector<std::string> segments;
size_t start = 0;
size_t pos;
while ((pos = formatted_chat.find(marker, start)) != std::string::npos) {
segments.push_back(formatted_chat.substr(start, pos - start));
start = pos + marker.size();
}
segments.push_back(formatted_chat.substr(start));
auto bitmaps_c_ptr = ctx.bitmaps.c_ptr();
int32_t res = mtmd_tokenize(ctx.ctx_vision.get(),
if (segments.size() - 1 != bitmaps_c_ptr.size()) {
LOG_ERR("Number of media markers (%zu) does not match number of loaded media (%zu)\n",
segments.size() - 1, bitmaps_c_ptr.size());
return 1;
}
// interleave text and media parts
std::vector<mtmd_input_text> texts(segments.size());
std::vector<mtmd_input_part> parts;
for (size_t i = 0; i < segments.size(); i++) {
texts[i] = {segments[i].data(), segments[i].size(), /* add_special */ false, /* parse_special */ true};
parts.push_back({&texts[i], nullptr});
if (i < bitmaps_c_ptr.size()) {
parts.push_back({nullptr, bitmaps_c_ptr[i]});
}
}
std::vector<const mtmd_input_part *> parts_ptr;
for (const auto & p : parts) {
parts_ptr.push_back(&p);
}
mtmd::input_chunks chunks(mtmd_input_chunks_init());
int32_t res = mtmd_tokenize_from_parts(ctx.ctx_vision.get(),
chunks.ptr.get(), // output
&text, // text
bitmaps_c_ptr.data(),
bitmaps_c_ptr.size());
parts_ptr.data(),
parts_ptr.size(),
add_bos);
if (res != 0) {
LOG_ERR("Unable to tokenize prompt, res = %d\n", res);
return 1;
tools/mtmd/mtmd-image.cpp
+50
View File
@@ -980,6 +980,56 @@ mtmd_image_preproc_out mtmd_image_preprocessor_idefics3::preprocess(const clip_i
//
// CITE: https://github.com/huggingface/transformers/blob/main/src/transformers/models/idefics3/image_processing_idefics3.py#L737
const clip_image_size original_size = img.get_size();
// old gguf files have no preprocessor longest size, custom token limits also need the generic size below
if (hparams.image_longest_edge > 0 && hparams.image_min_pixels <= 0 && hparams.image_max_pixels <= 0) {
const int tile_size = hparams.image_size;
const int longest_edge = hparams.image_longest_edge;
const double aspect_ratio = (double) original_size.width / original_size.height;
clip_image_size resized_size;
if (original_size.width >= original_size.height) {
resized_size.width = longest_edge;
resized_size.height = (int) (longest_edge / aspect_ratio);
resized_size.height += resized_size.height % 2;
} else {
resized_size.height = longest_edge;
resized_size.width = (int) (longest_edge * aspect_ratio);
resized_size.width += resized_size.width % 2;
}
const int grid_x = (resized_size.width + tile_size - 1) / tile_size;
const int grid_y = (resized_size.height + tile_size - 1) / tile_size;
const clip_image_size refined_size = clip_image_size{grid_x * tile_size, grid_y * tile_size};
clip_image_u8 resized_img;
img_tool::resize(img, resized_img, resized_size, hparams.image_resize_algo, PAD_NONE);
clip_image_u8 refined_img;
img_tool::resize(resized_img, refined_img, refined_size, hparams.image_resize_algo, PAD_NONE);
clip_image_u8 overview;
img_tool::resize(refined_img, overview, {tile_size, tile_size}, hparams.image_resize_algo, PAD_NONE);
std::vector<clip_image_u8> slices;
for (int y = 0; y < grid_y; y++) {
for (int x = 0; x < grid_x; x++) {
clip_image_u8 slice;
img_tool::crop(refined_img, slice, x * tile_size, y * tile_size, tile_size, tile_size);
slices.push_back(std::move(slice));
}
}
LOG_DBG("%s: grid size: %d x %d (%d tiles) + overview\n", __func__, grid_x, grid_y, grid_x * grid_y);
mtmd_image_preproc_out output;
output.append_overview(hparams, overview, true);
output.append(hparams, slices, true);
output.grid_x = grid_x;
output.grid_y = grid_y;
return output;
}
const clip_image_size refined_size = img_tool::calc_size_preserved_ratio(
original_size,
{ hparams.image_size, std::max(0, hparams.image_min_pixels), std::max(0, hparams.image_max_pixels), hparams.image_longest_edge });
tools/mtmd/mtmd-internal.h
+4 -2
View File
@@ -10,10 +10,12 @@
#define MTMD_INTERNAL_HEADER
// bitmap is null for text parts
struct mtmd_input_part {
struct mtmd_internal_part {
std::string text;
const mtmd_bitmap * bitmap;
// only used for text parts
bool parse_special = false;
};
// [QWEN_VIDEO] merged parts are erased from `parts`, so one group always maps to one part
std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_input_part> & parts, int n_merge);
std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_internal_part> & parts, int n_merge);
tools/mtmd/mtmd.cpp
+50 -5
View File
@@ -1097,7 +1097,7 @@ void mtmd_free(mtmd_context * ctx) {
delete ctx;
}
std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_input_part> & parts, int n_merge) {
std::vector<std::vector<const mtmd_bitmap *>> mtmd_group_mergeable_bitmaps(std::vector<mtmd_internal_part> & parts, int n_merge) {
std::vector<std::vector<const mtmd_bitmap *>> output;
for (size_t i = 0; i < parts.size(); i++) {
if (parts[i].bitmap == nullptr) {
@@ -1124,7 +1124,7 @@ struct mtmd_tokenizer {
bool parse_special;
const llama_vocab * vocab;
using part = mtmd_input_part;
using part = mtmd_internal_part;
std::vector<part> parts;
// these will be freed when mtmd_tokenizer finishes
std::vector<mtmd::bitmap> bm_from_lazy; // TODO @ngxson : refactor, free bm_from_lazy progressively
@@ -1160,7 +1160,7 @@ struct mtmd_tokenizer {
}
parts.push_back({"", bitmaps[i_bm++]});
} else {
parts.push_back({std::move(part), nullptr});
parts.push_back({std::move(part), nullptr, parse_special});
}
}
@@ -1177,6 +1177,26 @@ struct mtmd_tokenizer {
expand_lazy_bitmaps();
}
mtmd_tokenizer(mtmd_context * ctx,
const mtmd_input_part ** input_parts,
size_t n_parts,
bool add_special) : ctx(ctx) {
this->add_special = add_special;
parse_special = true; // only used for text returned by lazy bitmaps
vocab = ctx->vocab;
for (size_t i = 0; i < n_parts; i++) {
const mtmd_input_part * p = input_parts[i];
if (p->text != nullptr) {
parts.push_back({std::string(p->text->text, p->text->text_len), nullptr, p->text->parse_special});
} else {
parts.push_back({"", p->bitmap});
}
}
expand_lazy_bitmaps();
}
void expand_lazy_bitmaps() {
std::vector<part> expanded;
expanded.reserve(parts.size());
@@ -1201,7 +1221,7 @@ struct mtmd_tokenizer {
LOG_DBG("%s: lazy callback returned bitmap with dimensions %d x %d\n", __func__, out_bm->nx, out_bm->ny);
} else if (out_str) {
auto & ptr = text_from_lazy.emplace_back(out_str); // remember to free it later
expanded.push_back({ptr, nullptr});
expanded.push_back({ptr, nullptr, parse_special});
LOG_DBG("%s: lazy callback returned text: %s\n", __func__, out_str);
}
} else if (res == -1) {
@@ -1245,7 +1265,7 @@ struct mtmd_tokenizer {
return res;
}
} else {
add_text(p.text, parse_special);
add_text(p.text, p.parse_special);
}
}
@@ -1727,6 +1747,30 @@ int32_t mtmd_tokenize(mtmd_context * ctx,
}
}
int32_t mtmd_tokenize_from_parts(mtmd_context * ctx,
mtmd_input_chunks * output,
const mtmd_input_part ** parts,
size_t n_parts,
bool add_special) {
for (size_t i = 0; i < n_parts; i++) {
if ((parts[i]->text == nullptr) == (parts[i]->bitmap == nullptr)) {
LOG_ERR("%s: part %zu must have either text or bitmap set, not both\n", __func__, i);
return 1;
}
if (parts[i]->text != nullptr && parts[i]->text->text == nullptr) {
LOG_ERR("%s: part %zu has null text pointer\n", __func__, i);
return 1;
}
}
try {
mtmd_tokenizer tokenizer(ctx, parts, n_parts, add_special);
return tokenizer.tokenize(output);
} catch (const std::exception & e) {
LOG_ERR("%s: error: %s\n", __func__, e.what());
return 2;
}
}
static int32_t mtmd_encode_impl(mtmd_context * ctx, const mtmd_image_tokens * image_tokens, std::vector<float> & out_embd) {
clip_ctx * ctx_clip = ctx->ctx_v;
if (!ctx_clip) {
@@ -2132,6 +2176,7 @@ bool mtmd_decode_use_non_causal(const mtmd_context * ctx, const mtmd_input_chunk
case PROJECTOR_TYPE_GEMMA3:
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
case PROJECTOR_TYPE_DEEPSEEK4V:
return true;
default:
return false;
tools/mtmd/mtmd.h
+23 -4
View File
@@ -73,6 +73,12 @@ struct mtmd_input_text {
bool parse_special;
};
struct mtmd_input_part {
// only text or bitmap can be set, not both
const struct mtmd_input_text * text;
const struct mtmd_bitmap * bitmap;
};
//
// C API
//
@@ -83,6 +89,7 @@ typedef struct mtmd_image_tokens mtmd_image_tokens;
typedef struct mtmd_input_chunk mtmd_input_chunk;
typedef struct mtmd_input_chunks mtmd_input_chunks;
typedef struct mtmd_input_text mtmd_input_text;
typedef struct mtmd_input_part mtmd_input_part;
typedef struct mtmd_batch mtmd_batch;
typedef bool (*mtmd_progress_callback)(float progress, void * user_data);
@@ -276,10 +283,10 @@ struct mtmd_decoder_pos {
// return relative position (for example, embedding 0 will have position (0, 0, 0); remember to adjust it to the current absolute position)
MTMD_API struct mtmd_decoder_pos mtmd_image_tokens_get_decoder_pos(const mtmd_image_tokens * image_tokens, llama_pos pos_0, size_t i);
// tokenize an input text prompt and a list of bitmaps (images/audio)
// the prompt must have the input image marker (default: "<__media__>") in it
// tokenize an input text prompt and a list of bitmaps (image/audio)
// the prompt must have the input media marker (default: "<__media__>") in it
// the default marker is defined by mtmd_default_marker()
// the marker will be replaced with the image/audio chunk
// the marker will be replaced with the media chunk
// for example:
// "here is an image: <__media__>\ndescribe it in detail."
// this will gives 3 chunks:
@@ -291,13 +298,25 @@ MTMD_API struct mtmd_decoder_pos mtmd_image_tokens_get_decoder_pos(const mtmd_im
// return values:
// 0 on success
// 1 on number of bitmaps not matching the number of markers
// 2 on image preprocessing error
// 2 on media preprocessing error
MTMD_API int32_t mtmd_tokenize(mtmd_context * ctx,
mtmd_input_chunks * output,
const mtmd_input_text * text,
const mtmd_bitmap ** bitmaps,
size_t n_bitmaps);
// same as mtmd_tokenize(), but takes an array of mtmd_input_part
// use cases:
// - when you don't want to use media markers (they will be tokenized as normal text)
// - when you want to control parse_special for each text part
// note: per-part add_special will be ignored
// return 1 if a part has both text and bitmap set (or neither)
MTMD_API int32_t mtmd_tokenize_from_parts(mtmd_context * ctx,
mtmd_input_chunks * output,
const mtmd_input_part ** parts,
size_t n_parts,
bool add_special);
DEPRECATED(MTMD_API int32_t mtmd_encode(mtmd_context * ctx, const mtmd_image_tokens * image_tokens),
"use mtmd_encode_chunk() instead");
tools/server/server-common.cpp
+12 -11
View File
@@ -1062,8 +1062,7 @@ json oaicompat_completion_params_parse(const json & body) {
static void handle_media(
std::vector<raw_buffer> & out_files,
const std::string & url,
const std::string & media_path,
bool accept_base64_uri) {
const std::string & media_path) {
if (!media_path.empty()) {
// should already be enforced by arg.cpp, but checking just in case
GGML_ASSERT(media_path.back() == DIRECTORY_SEPARATOR);
@@ -1104,15 +1103,17 @@ static void handle_media(
data.assign((std::istreambuf_iterator<char>(file)), std::istreambuf_iterator<char>());
out_files.push_back(data);
} else if (accept_base64_uri && string_starts_with(url, "data:")) {
// try to decode base64 image
} else if (string_starts_with(url, "data:")) {
// try to decode base64 image, video, or audio
std::vector<std::string> parts = string_split<std::string>(url, /*separator*/ ',');
if (parts.size() != 2) {
throw std::runtime_error("Invalid uri-encoded base64 value");
} else if (!string_starts_with(parts[0], "data:image/")) {
throw std::runtime_error("Invalid uri format: " + parts[0]);
throw std::invalid_argument("Invalid uri-encoded base64 value");
} else if (!string_starts_with(parts[0], "data:image/")
&& !string_starts_with(parts[0], "data:video/")
&& !string_starts_with(parts[0], "data:audio/")) {
throw std::invalid_argument("Invalid uri format: " + parts[0]);
} else if (!string_ends_with(parts[0], "base64")) {
throw std::runtime_error("uri must be base64 encoded");
throw std::invalid_argument("uri must be base64 encoded");
} else {
auto base64_data = parts[1];
auto decoded_data = base64_decode(base64_data);
@@ -1219,7 +1220,7 @@ json oaicompat_chat_params_parse(
json image_url = json_value(p, "image_url", json::object());
std::string url = json_value(image_url, "url", std::string());
handle_media(out_files, url, opt.media_path, true);
handle_media(out_files, url, opt.media_path);
p["type"] = "media_marker";
p["text"] = get_media_marker();
@@ -1234,7 +1235,7 @@ json oaicompat_chat_params_parse(
json input_audio = json_value(p, "input_audio", json::object());
std::string url = json_value(input_audio, "data",
json_value(input_audio, "url", std::string()));
handle_media(out_files, url, opt.media_path, false);
handle_media(out_files, url, opt.media_path);
p["type"] = "media_marker";
p["text"] = get_media_marker();
@@ -1248,7 +1249,7 @@ json oaicompat_chat_params_parse(
json input_video = json_value(p, "input_video", json::object());
std::string url = json_value(input_video, "data",
json_value(input_video, "url", std::string()));
handle_media(out_files, url, opt.media_path, false);
handle_media(out_files, url, opt.media_path);
p["type"] = "media_marker";
p["text"] = get_media_marker();
tools/server/server-context.cpp
+13 -2
View File
@@ -1493,11 +1493,22 @@ private:
auto caps = common_chat_templates_get_caps(chat_params.tmpls.get());
auto it = params_base.default_template_kwargs.find("preserve_reasoning");
bool supported = caps.at("supports_preserve_reasoning");
bool enabled = it != params_base.default_template_kwargs.end();
bool specified = params_base.preserve_reasoning_specified;
// note: the kwarg is enabled by default if not specified explicitly, so check the value
bool enabled = it != params_base.default_template_kwargs.end() && it->second == "true";
if (supported) {
SRV_TRC("preserve_reasoning kwarg: %s\n",
it == params_base.default_template_kwargs.end() ? "unset (template default)" : it->second.c_str());
} else {
SRV_TRC("%s", "preserve_reasoning kwarg: not supported by template\n");
}
if (supported && !specified) {
SRV_WRN("%s", "chat template supports preserving reasoning, it is enabled by default (may use more tokens, disable via --no-reasoning-preserve)\n");
}
if (supported && !enabled) {
SRV_INF("%s", "chat template supports preserving reasoning, consider enabling it via --reasoning-preserve\n");
}
if (!supported && enabled) {
if (!supported && specified && enabled) {
SRV_WRN("%s", "chat template does NOT support preserving reasoning, --reasoning-preserve has no effect\n");
}
}
tools/server/tests/unit/test_vision_api.py
+1
View File
@@ -71,6 +71,7 @@ def test_v1_models_supports_multimodal_capability():
("What is this:\n", "malformed", False, None),
("What is this:\n", "https://google.com/404", False, None), # non-existent image
("What is this:\n", "https://ggml.ai", False, None), # non-image data
("What is this:\n", "data:text/html;base64,aGVsbG8=", False, None), # unsupported data uri mime
# TODO @ngxson : test with multiple images, no images and with audio
]
)