From 727cce06fd6af6a5e4c175ba590f55b9bfb28738 Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 00:20:26 +0800 Subject: [PATCH 1/8] [model] support ernie4-5-vl-moe model --- convert_hf_to_gguf.py | 165 ++++++++++++++++++++++ ggml/include/ggml.h | 1 + ggml/src/ggml-cpu/ops.cpp | 48 ++++++- ggml/src/ggml-cuda/rope.cu | 100 ++++++++++++- gguf-py/gguf/constants.py | 31 ++++ gguf-py/gguf/gguf_writer.py | 3 + include/llama.h | 1 + src/llama-arch.cpp | 1 + src/llama-arch.h | 1 + src/llama-hparams.cpp | 2 +- src/llama-hparams.h | 1 + src/llama-model.cpp | 95 +++++++++++++ src/models/ernie4-5-vl-moe.cpp | 185 ++++++++++++++++++++++++ tools/mtmd/clip.cpp | 92 +++++++++++- tools/mtmd/models/ernie45vlmoe.cpp | 219 +++++++++++++++++++++++++++++ tools/mtmd/mtmd.cpp | 5 + 16 files changed, 946 insertions(+), 4 deletions(-) create mode 100644 src/models/ernie4-5-vl-moe.cpp create mode 100644 tools/mtmd/models/ernie45vlmoe.cpp diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index e64756a74ab4..66eb93661734 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -3814,6 +3814,171 @@ def prepare_tensors(self): if len(experts) > 0: raise ValueError(f"Unprocessed experts: {experts}") +@ModelBase.register("Ernie4_5_VLMoeForConditionalGeneration") +class Ernie4_5VLMoeModel(Ernie4_5MoeModel): + model_arch = gguf.MODEL_ARCH.ERNIE4_5_VL_MOE + _experts: list[dict[str, Tensor]] | None = None + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + self._experts = [{} for _ in range(self.block_count)] + + def set_gguf_parameters(self): + super().set_gguf_parameters() + + # Handle list-based expert configurations by taking the first value + moe_num_experts = self.hparams["moe_num_experts"] + if isinstance(moe_num_experts, list): + moe_num_experts = moe_num_experts[0] + self.gguf_writer.add_expert_count(moe_num_experts) + + self.gguf_writer.add_expert_used_count(self.hparams["moe_k"]) + self.gguf_writer.add_interleave_moe_layer_step(self.hparams["moe_layer_interval"]) + + moe_layer_start_index = self.hparams["moe_layer_start_index"] + if isinstance(moe_layer_start_index, list): + moe_layer_start_index = moe_layer_start_index[0] + self.gguf_writer.add_leading_dense_block_count(moe_layer_start_index) + + if (moe_intermediate_size := self.hparams.get("moe_intermediate_size")) is not None: + if isinstance(moe_intermediate_size, list): + self.gguf_writer.add_expert_feed_forward_length(moe_intermediate_size[0]) + if len(moe_intermediate_size) > 1: + self.gguf_writer.add_vision_expert_feed_forward_length(moe_intermediate_size[1]) + else: + self.gguf_writer.add_expert_feed_forward_length(moe_intermediate_size) + + if (shared_expert_count := self.hparams.get('moe_num_shared_experts')) is not None: + self.gguf_writer.add_expert_shared_count(shared_expert_count) + if shared_expert_count > 0 and (shared_expert_intermediate_size := self.hparams.get('intermediate_size')) is not None and (num_key_value_heads := self.hparams.get('num_key_value_heads')) is not None: + self.gguf_writer.add_expert_shared_feed_forward_length(shared_expert_intermediate_size // num_key_value_heads) + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + # Skip vision and multimodal tensors - they are not part of the text model + if name.startswith("vision_model") or name.startswith("resampler_model") or \ + name.startswith("model.vision_model") or name.startswith("model.resampler_model"): + return [] + + # todo(megemini): gate_inp weight/weight_1 + # weight + if name.endswith(".mlp.gate.weight") or name.endswith(".mlp.gate.weight_1"): + if name.endswith(".mlp.gate.weight_1"): + name = name.replace(".mlp.gate.weight_1", ".mlp.gate.vision.weight") + + data_torch = data_torch.t() + # Extract bid from name if not provided + if bid is None: + match = re.search(r"model\.layers\.(\d+)", name) + if match: + bid = int(match.group(1)) + # todo(megemini): + logger.info("Processing gate.weight/weight_1: %s -> shape %s", name, data_torch.shape) + # Map the tensor name and ensure it has .weight suffix + mapped_name = self.map_tensor_name(name) + + return [(mapped_name, data_torch)] + + # todo(megemini): e_score_correction.bias/bias_1 for weight/weight_1 + if name.endswith(".mlp.moe_statics.e_score_correction_bias"): + name_text = name.replace("e_score_correction_bias", "e_score_correction.bias") + data_torch_text = data_torch[0, :] + + name_vision = name.replace("e_score_correction_bias", "e_score_correction.vision.bias") + data_torch_vision = data_torch[1, :] + + return [(self.map_tensor_name(name_text), data_torch_text), + (self.map_tensor_name(name_vision), data_torch_vision)] + + # process the experts separately + if name.find("mlp.experts") != -1: + n_experts = self.hparams["moe_num_experts"] + + # Handle n_experts being a list (for models with multiple expert groups) + if isinstance(n_experts, list): + total_experts = sum(n_experts) + else: + total_experts = n_experts + + assert bid is not None + if self._experts is None: + self._experts = [{} for _ in range(self.block_count)] + + self._experts[bid][name] = data_torch + + # Only merge routed experts (not shared experts) + # Total tensors = total_experts * 3 (gate, up, down) + if len(self._experts[bid]) >= total_experts * 3: + tensors: list[tuple[str, Tensor]] = [] + + # For models with multiple expert groups of different sizes, + for w_name in ["gate_proj", "up_proj", "down_proj"]: + # Collect all experts for this weight type + expert_data: dict[int, Tensor] = {} + for xid in range(total_experts): + ename = f"model.layers.{bid}.mlp.experts.{xid}.{w_name}.weight" + if ename in self._experts[bid]: + expert_data[xid] = self._experts[bid][ename] + del self._experts[bid][ename] + + if not expert_data: + continue + + # Group experts by shape (to handle different intermediate sizes) + shape_groups: dict[tuple[int, ...], list[tuple[int, Tensor]]] = {} + for xid, tensor in expert_data.items(): + shape_key = tuple(tensor.shape) + if shape_key not in shape_groups: + shape_groups[shape_key] = [] + shape_groups[shape_key].append((xid, tensor)) + + # For each shape group, stack the experts + # For ERNIE-4.5-VL with multiple expert groups of different sizes, + # we need to save them separately as llama.cpp doesn't support mixed sizes yet + if len(shape_groups) > 1: + # Sort shape groups by number of experts (descending) + sorted_groups = sorted(shape_groups.items(), key=lambda x: len(x[1]), reverse=True) + + for group_idx, (shape_key, expert_list) in enumerate(sorted_groups): + # Sort by expert ID to maintain order + expert_list.sort(key=lambda x: x[0]) + datas = [tensor for _, tensor in expert_list] + + data_torch = torch.stack(datas, dim=0) + + # Use group suffix for additional groups + if group_idx == 0: + merged_name = f"model.layers.{bid}.mlp.experts.{w_name}.weight" + else: + merged_name = f"model.vision.layers.{bid}.mlp.experts.{w_name}.weight" + + new_name = self.map_tensor_name(merged_name) + tensors.append((new_name, data_torch)) + else: + # Single shape - stack all experts + expert_list = list(shape_groups.values())[0] + expert_list.sort(key=lambda x: x[0]) + datas = [tensor for _, tensor in expert_list] + + data_torch = torch.stack(datas, dim=0) + + merged_name = f"model.layers.{bid}.mlp.experts.{w_name}.weight" + new_name = self.map_tensor_name(merged_name) + tensors.append((new_name, data_torch)) + + return tensors + else: + return [] + return [(self.map_tensor_name(name), data_torch)] + + def prepare_tensors(self): + super().prepare_tensors() + + if self._experts is not None: + # flatten `list[dict[str, Tensor]]` into `list[str]` + experts = [k for d in self._experts for k in d.keys()] + if len(experts) > 0: + raise ValueError(f"Unprocessed experts: {experts}") + @ModelBase.register( "Qwen2VLModel", diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index f759e2d5883e..2baf875fc7b5 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -252,6 +252,7 @@ #define GGML_ROPE_TYPE_MROPE 8 #define GGML_ROPE_TYPE_VISION 24 #define GGML_ROPE_TYPE_IMROPE 40 // binary: 101000 +#define GGML_ROPE_TYPE_ERNIE3D 72 // binary: 1001000, ERNIE-VL 3D RoPE (NORMAL rotation + interleaved h/w freq) #define GGML_MROPE_SECTIONS 4 diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index ce15b18ce0ee..3b608d682fa9 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -5651,6 +5651,43 @@ static void rotate_pairs(const int64_t n, const int64_t n_offset, const float * } } +static void ggml_ernie3d_rope_cache_init( + float theta_base_t, float theta_base_h, float theta_base_w, + int sections[4], + float freq_scale, const float * freq_factors, float corr_dims[2], int64_t ne0, float ext_factor, float mscale, + float * cache, float sin_sign, float theta_scale) { + // n_hw = sections[0] + sections[1] = total number of interleaved h/w frequencies + int n_hw = sections[0] + sections[1]; + + float theta_accum = 1.0f; // accumulated theta_scale^freq_idx + + for (int64_t i0 = 0; i0 < ne0; i0 += 2) { + int freq_idx = (int)(i0 / 2); + const float ff = freq_factors ? freq_factors[freq_idx] : 1.0f; + + float theta; + if (freq_idx < n_hw) { + if (freq_idx % 2 == 0) { + // even freq index -> height position + theta = theta_base_h * theta_accum; + } else { + // odd freq index -> width position + theta = theta_base_w * theta_accum; + } + } else { + // temporal position + theta = theta_base_t * theta_accum; + } + + rope_yarn( + theta/ff, freq_scale, corr_dims, i0, ext_factor, mscale, &cache[i0 + 0], &cache[i0 + 1] + ); + cache[i0 + 1] *= sin_sign; + + theta_accum *= theta_scale; + } +} + template //float or ggml_fp16_t static void ggml_compute_forward_rope_flt( const ggml_compute_params * params, @@ -5723,7 +5760,7 @@ static void ggml_compute_forward_rope_flt( if (is_vision) { GGML_ASSERT(n_dims == ne0/2); } - + const bool is_ernie3d = mode == GGML_ROPE_TYPE_ERNIE3D; const float * freq_factors = NULL; if (src2 != NULL) { GGML_ASSERT(src2->type == GGML_TYPE_F32); @@ -5745,6 +5782,14 @@ static void ggml_compute_forward_rope_flt( if (!mrope_used) { const int64_t p = pos[i2]; ggml_rope_cache_init(p, freq_scale, freq_factors, corr_dims, ne0, ext_factor, attn_factor, cache, sin_sign, theta_scale); + } else if (is_ernie3d) { + // ERNIE-VL 3D RoPE: interleaved h/w freq with NORMAL rotation + const int64_t p_t = pos[i2]; + const int64_t p_h = pos[i2 + ne2]; + const int64_t p_w = pos[i2 + ne2 * 2]; + ggml_ernie3d_rope_cache_init( + p_t, p_h, p_w, sections, + freq_scale, freq_factors, corr_dims, ne0, ext_factor, attn_factor, cache, sin_sign, theta_scale); } else { const int64_t p_t = pos[i2]; @@ -5765,6 +5810,7 @@ static void ggml_compute_forward_rope_flt( switch (mode) { case GGML_ROPE_TYPE_NORMAL: + case GGML_ROPE_TYPE_ERNIE3D: rotate_pairs(n_dims, 1, cache, src, dst_data, 1); break; case GGML_ROPE_TYPE_NEOX: diff --git a/ggml/src/ggml-cuda/rope.cu b/ggml/src/ggml-cuda/rope.cu index 45a49a5dc2a3..875e989704be 100644 --- a/ggml/src/ggml-cuda/rope.cu +++ b/ggml/src/ggml-cuda/rope.cu @@ -264,6 +264,68 @@ static __global__ void rope_multi(const T * x, dst[idst + n_dims/2] = x0*sin_theta + x1*cos_theta; } +template +static __global__ void rope_ernie3d( + const T * x, T * dst, const int ne0, const int ne1, const int ne2, const int s1, const int s2, + const int n_dims, const int32_t * pos, const float freq_scale, const float ext_factor, const float attn_factor, + const rope_corr_dims corr_dims, const float theta_scale, const float * freq_factors, const mrope_sections sections) { + const int i0 = 2*(blockDim.y*blockIdx.y + threadIdx.y); + + if (i0 >= ne0) { + return; + } + + const int row_dst = blockDim.x*blockIdx.x + threadIdx.x; + + const int row_x = row_dst % ne1; + const int channel_x = row_dst / ne1; + + // NORMAL rotation: pair (x[i0], x[i0+1]), stored at adjacent positions + const int idst = row_dst*ne0 + i0; + const int ix = channel_x*s2 + row_x*s1 + i0; + + if (i0 >= n_dims) { + dst[idst + 0] = x[ix + 0]; + dst[idst + 1] = x[ix + 1]; + return; + } + + // freq_idx = i0/2 (which frequency pair this is) + const int freq_idx = i0 / 2; + // n_hw = sections[0] + sections[1] = total number of h+w interleaved frequencies + const int n_hw = sections.v[0] + sections.v[1]; + + // Determine which position slot to use based on interleaved pattern + // Position slots: slot 0 = t_position, slot 1 = h_position, slot 2 = w_position + float theta_base = 0.0f; + if (freq_idx < n_hw) { + if (freq_idx % 2 == 0) { + // even freq index -> height position (slot 1) + theta_base = pos[channel_x + ne2 * 1] * powf(theta_scale, (float)freq_idx); + } else { + // odd freq index -> width position (slot 2) + theta_base = pos[channel_x + ne2 * 2] * powf(theta_scale, (float)freq_idx); + } + } else { + // temporal position (slot 0) + theta_base = pos[channel_x] * powf(theta_scale, (float)freq_idx); + } + + const float freq_factor = has_ff ? freq_factors[freq_idx] : 1.0f; + + float cos_theta; + float sin_theta; + + rope_yarn(theta_base/freq_factor, freq_scale, corr_dims, i0, ext_factor, attn_factor, cos_theta, sin_theta); + + // NORMAL (GPT-J) rotation: adjacent pair (x[i0], x[i0+1]) + const float x0 = x[ix + 0]; + const float x1 = x[ix + 1]; + + dst[idst + 0] = x0*cos_theta - x1*sin_theta; + dst[idst + 1] = x0*sin_theta + x1*cos_theta; +} + template static __global__ void rope_vision(const T * x, T * dst, @@ -453,6 +515,29 @@ static void rope_multi_cuda(const T * x, } } +template +static void rope_ernie3d_cuda( + const T * x, T * dst, const int ne0, const int ne1, const int ne2, const int s1, const int s2, const int n_dims, const int nr, + const int32_t * pos, const float freq_scale, const float freq_base, const float ext_factor, const float attn_factor, + const rope_corr_dims corr_dims, const float * freq_factors, const mrope_sections sections, cudaStream_t stream) { + GGML_ASSERT(ne0 % 2 == 0); + const dim3 block_dims(1, CUDA_ROPE_BLOCK_SIZE, 1); + const int n_blocks_x = (ne0 + 2*CUDA_ROPE_BLOCK_SIZE - 1) / (2*CUDA_ROPE_BLOCK_SIZE); + const dim3 block_nums(nr, n_blocks_x, 1); + + const float theta_scale = powf(freq_base, -2.0f/n_dims); + + if (freq_factors == nullptr) { + rope_ernie3d<<>>( + x, dst, ne0, ne1, ne2, s1, s2, n_dims, pos, freq_scale, ext_factor, + attn_factor, corr_dims, theta_scale, freq_factors, sections); + } else { + rope_ernie3d<<>>( + x, dst, ne0, ne1, ne2, s1, s2, n_dims, pos, freq_scale, ext_factor, + attn_factor, corr_dims, theta_scale, freq_factors, sections); + } +} + template static void rope_vision_cuda(const T * x, T * dst, @@ -603,7 +688,20 @@ void ggml_cuda_op_rope_impl(ggml_backend_cuda_context & ctx, s03, s1, s2, s3, n_dims, nr, pos, freq_scale, freq_base, ext_factor, attn_factor, corr_dims, freq_factors, row_indices, set_rows_stride, stream); - } else { + } else if (is_ernie3d) { + if (src0->type == GGML_TYPE_F32) { + rope_ernie3d_cuda( + (const float *) src0_d, (float *) dst_d, ne00, ne01, ne02, s01, s02, n_dims, nr, pos, freq_scale, + freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, stream); + } else if (src0->type == GGML_TYPE_F16) { + rope_ernie3d_cuda( + (const half *) src0_d, (half *) dst_d, ne00, ne01, ne02, s01, s02, n_dims, nr, pos, freq_scale, + freq_base, ext_factor, attn_factor, corr_dims, freq_factors, sections, stream); + } else { + GGML_ABORT("fatal error"); + } + } + else { GGML_ABORT("fatal error"); } } else if (is_mrope && !is_vision) { diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 8a3fab1e1c30..10298dffc0a7 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -110,6 +110,7 @@ class LLM: LEADING_DENSE_BLOCK_COUNT = "{arch}.leading_dense_block_count" FEED_FORWARD_LENGTH = "{arch}.feed_forward_length" EXPERT_FEED_FORWARD_LENGTH = "{arch}.expert_feed_forward_length" + VISION_EXPERT_FEED_FORWARD_LENGTH = "{arch}.vision_expert_feed_forward_length" EXPERT_SHARED_FEED_FORWARD_LENGTH = "{arch}.expert_shared_feed_forward_length" EXPERT_CHUNK_FEED_FORWARD_LENGTH = "{arch}.expert_chunk_feed_forward_length" USE_PARALLEL_RESIDUAL = "{arch}.use_parallel_residual" @@ -447,6 +448,7 @@ class MODEL_ARCH(IntEnum): AFMOE = auto() ERNIE4_5 = auto() ERNIE4_5_MOE = auto() + ERNIE4_5_VL_MOE = auto() HUNYUAN_MOE = auto() HUNYUAN_DENSE = auto() SMOLLM3 = auto() @@ -879,6 +881,7 @@ class MODEL_TENSOR(IntEnum): MODEL_ARCH.AFMOE: "afmoe", MODEL_ARCH.ERNIE4_5: "ernie4_5", MODEL_ARCH.ERNIE4_5_MOE: "ernie4_5-moe", + MODEL_ARCH.ERNIE4_5_VL_MOE: "ernie4_5-vl-moe", MODEL_ARCH.FALCON_H1: "falcon-h1", MODEL_ARCH.HUNYUAN_MOE: "hunyuan-moe", MODEL_ARCH.HUNYUAN_DENSE: "hunyuan-dense", @@ -2597,6 +2600,33 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.FFN_UP_SHEXP, MODEL_TENSOR.FFN_EXP_PROBS_B, ], + MODEL_ARCH.ERNIE4_5_VL_MOE: [ + MODEL_TENSOR.TOKEN_EMBD, + MODEL_TENSOR.OUTPUT_NORM, + MODEL_TENSOR.OUTPUT, + MODEL_TENSOR.ATTN_NORM, + MODEL_TENSOR.ATTN_Q, + MODEL_TENSOR.ATTN_K, + MODEL_TENSOR.ATTN_V, + MODEL_TENSOR.ATTN_OUT, + MODEL_TENSOR.FFN_NORM, + MODEL_TENSOR.FFN_GATE, + MODEL_TENSOR.FFN_DOWN, + MODEL_TENSOR.FFN_UP, + MODEL_TENSOR.FFN_GATE_INP, + MODEL_TENSOR.FFN_GATE_EXP, + MODEL_TENSOR.FFN_DOWN_EXP, + MODEL_TENSOR.FFN_UP_EXP, + MODEL_TENSOR.FFN_EXP_PROBS_B, + MODEL_TENSOR.V_FFN_GATE_INP, + MODEL_TENSOR.V_FFN_GATE_EXPS, + MODEL_TENSOR.V_FFN_DOWN_EXPS, + MODEL_TENSOR.V_FFN_UP_EXPS, + MODEL_TENSOR.V_FFN_EXP_PROBS_B, + MODEL_TENSOR.FFN_GATE_SHEXP, + MODEL_TENSOR.FFN_DOWN_SHEXP, + MODEL_TENSOR.FFN_UP_SHEXP, + ], MODEL_ARCH.PLM: [ MODEL_TENSOR.TOKEN_EMBD, MODEL_TENSOR.OUTPUT, @@ -3770,6 +3800,7 @@ class VisionProjectorType: MUSIC_FLAMINGO = "musicflamingo" # audio GLM4V = "glm4v" YOUTUVL = "youtuvl" + ERNIE45VLMOE = "ernie4.5vl_moe" # Items here are (block size, type size) diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index 62172b24c386..a8633afc2579 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -717,6 +717,9 @@ def add_feed_forward_length(self, length: int | Sequence[int]) -> None: 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_vision_expert_feed_forward_length(self, length: int) -> None: + self.add_uint32(Keys.LLM.VISION_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) diff --git a/include/llama.h b/include/llama.h index bf4e28a8be16..7ed25e7ab229 100644 --- a/include/llama.h +++ b/include/llama.h @@ -85,6 +85,7 @@ extern "C" { LLAMA_ROPE_TYPE_MROPE = GGML_ROPE_TYPE_MROPE, LLAMA_ROPE_TYPE_IMROPE = GGML_ROPE_TYPE_IMROPE, LLAMA_ROPE_TYPE_VISION = GGML_ROPE_TYPE_VISION, + LLAMA_ROPE_TYPE_ERNIE3D = GGML_ROPE_TYPE_ERNIE3D, }; enum llama_token_type { //TODO: remove, required until per token attributes are available from GGUF file diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index fce46772d7ed..8affd81694f6 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -163,6 +163,7 @@ static const std::map LLM_KV_NAMES = { { LLM_KV_LEADING_DENSE_BLOCK_COUNT, "%s.leading_dense_block_count" }, { LLM_KV_FEED_FORWARD_LENGTH, "%s.feed_forward_length" }, { LLM_KV_EXPERT_FEED_FORWARD_LENGTH, "%s.expert_feed_forward_length" }, + { LLM_KV_VISION_EXPERT_FEED_FORWARD_LENGTH, "%s.vision_expert_feed_forward_length" }, { LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, "%s.expert_shared_feed_forward_length" }, { LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, "%s.expert_chunk_feed_forward_length" }, { LLM_KV_SWIGLU_CLAMP_EXP, "%s.swiglu_clamp_exp" }, diff --git a/src/llama-arch.h b/src/llama-arch.h index a392ecce2b4c..dcd03b2b2129 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -167,6 +167,7 @@ enum llm_kv { LLM_KV_LEADING_DENSE_BLOCK_COUNT, LLM_KV_FEED_FORWARD_LENGTH, LLM_KV_EXPERT_FEED_FORWARD_LENGTH, + LLM_KV_VISION_EXPERT_FEED_FORWARD_LENGTH, LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, LLM_KV_SWIGLU_CLAMP_EXP, diff --git a/src/llama-hparams.cpp b/src/llama-hparams.cpp index 756dda1a7ab6..a584d0339493 100644 --- a/src/llama-hparams.cpp +++ b/src/llama-hparams.cpp @@ -178,7 +178,7 @@ bool llama_hparams::is_recurrent(uint32_t il) const { } uint32_t llama_hparams::n_pos_per_embd() const { - return rope_type == LLAMA_ROPE_TYPE_MROPE || rope_type == LLAMA_ROPE_TYPE_IMROPE ? 4 : 1; + return rope_type == LLAMA_ROPE_TYPE_MROPE || rope_type == LLAMA_ROPE_TYPE_IMROPE || rope_type == LLAMA_ROPE_TYPE_ERNIE3D ? 4 : 1; } bool llama_hparams::is_swa(uint32_t il) const { diff --git a/src/llama-hparams.h b/src/llama-hparams.h index 6c695bdbf662..3a2f6ca5f489 100644 --- a/src/llama-hparams.h +++ b/src/llama-hparams.h @@ -70,6 +70,7 @@ struct llama_hparams { uint32_t n_lora_q = 0; uint32_t n_lora_kv = 0; uint32_t n_ff_exp = 0; + uint32_t n_ff_v_exp = 0; uint32_t n_ff_shexp = 0; uint32_t n_ff_chexp = 0; uint32_t n_expert_shared = 0; diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 8fc61aee3727..e2be33e841e6 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -2195,6 +2195,34 @@ void llama_model::load_hparams(llama_model_loader & ml) { default: type = LLM_TYPE_UNKNOWN; } } break; + case LLM_ARCH_ERNIE4_5_VL_MOE: + { + 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(LLM_KV_VISION_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_v_exp, false); + 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); + if (!ml.get_key_or_arr(LLM_KV_ROPE_DIMENSION_SECTIONS, hparams.rope_sections, 4, false)) { + hparams.rope_sections[0] = 22; + hparams.rope_sections[1] = 22; + hparams.rope_sections[2] = 20; + hparams.rope_sections[3] = 0; + } + + LLAMA_LOG_INFO("%s: ERNIE-VL rope_sections=[%d,%d,%d,%d]\n", __func__, + hparams.rope_sections[0], hparams.rope_sections[1], + hparams.rope_sections[2], hparams.rope_sections[3]); + + if (hparams.n_ff_v_exp == 0) { + hparams.n_ff_v_exp = 512; // ERNIE-VL default + } + + switch (hparams.n_layer) { + case 28: type = LLM_TYPE_28B_A3B; break; + default: type = LLM_TYPE_UNKNOWN; + } + } break; case LLM_ARCH_FALCON_H1: { // Common parameters @@ -6397,6 +6425,67 @@ bool llama_model::load_tensors(llama_model_loader & ml) { layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff_exp, n_embd, n_expert}, 0); layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0); + // Shared expert (if present) + if (hparams.n_ff_shexp > 0) { + layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", i), { n_embd, hparams.n_ff_shexp}, 0); + layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", i), {hparams.n_ff_shexp, n_embd }, 0); + layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", i), { n_embd, hparams.n_ff_shexp}, 0); + } + } else { // Dense layers + layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0); + layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, 0); + layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0); + } + } + } break; + case LLM_ARCH_ERNIE4_5_VL_MOE: + { + tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0); + + // output + output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); + output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); + // if output is NULL, init from the input tok embed + if (output == NULL) { + output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED); + } + + for (int i = 0; i < n_layer; ++i) { + auto & layer = layers[i]; + + layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0); + + layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "weight", i), {n_embd, n_embd_head_k * n_head}, 0); + layer.wk = create_tensor(tn(LLM_TENSOR_ATTN_K, "weight", i), {n_embd, n_embd_gqa}, 0); + layer.wv = create_tensor(tn(LLM_TENSOR_ATTN_V, "weight", i), {n_embd, n_embd_gqa}, 0); + layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_k * n_head, n_embd}, 0); + + // optional bias tensors + layer.bq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED); + layer.bk = create_tensor(tn(LLM_TENSOR_ATTN_K, "bias", i), {n_embd_gqa}, TENSOR_NOT_REQUIRED); + layer.bv = create_tensor(tn(LLM_TENSOR_ATTN_V, "bias", i), {n_embd_gqa}, TENSOR_NOT_REQUIRED); + layer.bo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED); + + layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0); + + if (static_cast(i) >= hparams.n_layer_dense_lead) { // MoE layers + int n_ff_exp = hparams.n_ff_exp; + int n_ff_v_exp = hparams.n_ff_v_exp; // Vision expert intermediate size + + 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); + 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}, 0); + layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0); + + // Vision expert MoE tensors + layer.v_ffn_gate_inp = create_tensor(tn(LLM_TENSOR_V_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0); + layer.v_ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_V_FFN_EXP_PROBS_B, "bias", i), {n_expert}, TENSOR_NOT_REQUIRED); + layer.v_ffn_gate_exps = create_tensor(tn(LLM_TENSOR_V_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_v_exp, n_expert}, TENSOR_NOT_REQUIRED); + layer.v_ffn_down_exps = create_tensor(tn(LLM_TENSOR_V_FFN_DOWN_EXPS, "weight", i), { n_ff_v_exp, n_embd, n_expert}, 0); + layer.v_ffn_up_exps = create_tensor(tn(LLM_TENSOR_V_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_v_exp, n_expert}, 0); + + // Shared expert (if present) if (hparams.n_ff_shexp > 0) { layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", i), { n_embd, hparams.n_ff_shexp}, 0); @@ -8430,6 +8519,10 @@ ggml_cgraph * llama_model::build_graph(const llm_graph_params & params) const { { llm = std::make_unique(*this, params); } break; + case LLM_ARCH_ERNIE4_5_VL_MOE: + { + llm = std::make_unique(*this, params); + } break; case LLM_ARCH_HUNYUAN_MOE: { llm = std::make_unique(*this, params); @@ -8771,6 +8864,8 @@ llama_rope_type llama_model_rope_type(const llama_model * model) { return model->hparams.use_mrope() ? LLAMA_ROPE_TYPE_MROPE : LLAMA_ROPE_TYPE_NORM; case LLM_ARCH_GLM4_MOE: return model->hparams.use_mrope() ? LLAMA_ROPE_TYPE_MROPE : LLAMA_ROPE_TYPE_NEOX; + case LLM_ARCH_ERNIE4_5_VL_MOE: + return LLAMA_ROPE_TYPE_ERNIE3D; // all model arches should be listed explicitly here case LLM_ARCH_UNKNOWN: diff --git a/src/models/ernie4-5-vl-moe.cpp b/src/models/ernie4-5-vl-moe.cpp new file mode 100644 index 000000000000..f246085d2c01 --- /dev/null +++ b/src/models/ernie4-5-vl-moe.cpp @@ -0,0 +1,185 @@ +#include "models.h" + +llm_build_ernie4_5_vl_moe::llm_build_ernie4_5_vl_moe(const llama_model & model, const llm_graph_params & params) : + llm_graph_context(params) { + const int64_t n_embd_head = hparams.n_embd_head_v; + + GGML_ASSERT(n_embd_head == hparams.n_embd_head_k); + GGML_ASSERT(n_embd_head == hparams.n_rot); + + ggml_tensor * cur; + ggml_tensor * inpL; + + inpL = build_inp_embd(model.tok_embd); + + // inp_pos - contains the positions + ggml_tensor * inp_pos = build_inp_pos(); + + auto * inp_attn = build_attn_inp_kv(); + + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + GGML_ASSERT(hparams.n_moe_layer_step > 0 && "Ernie 4.5 MoE requires n_moe_layer_step > 0"); + + // Get MROPE sections from hparams + // ERNIE-VL uses [22, 22, 20, 0] for (n_h, n_w, n_t, extra) + // With ERNIE3D rope type: interleaved [h,w,h,w,...,t,t,...] frequency layout + int sections[4]; + std::copy(std::begin(hparams.rope_sections), std::end(hparams.rope_sections), sections); + + for (int il = 0; il < n_layer; ++il) { + ggml_tensor * inpSA = inpL; + + // norm + cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "attn_norm", il); + + // self-attention + { + // compute Q and K and RoPE them + ggml_tensor * Qcur = build_lora_mm(model.layers[il].wq, cur); + cb(Qcur, "Qcur", il); + if (model.layers[il].bq) { + Qcur = ggml_add(ctx0, Qcur, model.layers[il].bq); + cb(Qcur, "Qcur", il); + } + ggml_tensor * Kcur = build_lora_mm(model.layers[il].wk, cur); + cb(Kcur, "Kcur", il); + if (model.layers[il].bk) { + Kcur = ggml_add(ctx0, Kcur, model.layers[il].bk); + cb(Kcur, "Kcur", il); + } + ggml_tensor * Vcur = build_lora_mm(model.layers[il].wv, cur); + cb(Vcur, "Vcur", il); + if (model.layers[il].bv) { + Vcur = ggml_add(ctx0, Vcur, model.layers[il].bv); + cb(Vcur, "Vcur", il); + } + Qcur = ggml_reshape_3d(ctx0, Qcur, n_embd_head, n_head, n_tokens); + Kcur = ggml_reshape_3d(ctx0, Kcur, n_embd_head, n_head_kv, n_tokens); + Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head, n_head_kv, n_tokens); + + // ERNIE-VL uses ERNIE3D RoPE (NORMAL rotation + interleaved 3D frequency) + // sections [22, 22, 20, 0]: n_h=22, n_w=22, n_t=20 + // Frequency layout: [h0,w1,h2,w3,...,h42,w43,t44,...,t63] + // Position slots: [t, h, w, 0] + Qcur = ggml_rope_multi(ctx0, Qcur, inp_pos, nullptr, + n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + + Kcur = ggml_rope_multi(ctx0, Kcur, inp_pos, nullptr, + n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + + cb(Qcur, "Qcur", il); + cb(Kcur, "Kcur", il); + cb(Vcur, "Vcur", il); + + cur = build_attn(inp_attn, + model.layers[il].wo, NULL, + Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f / sqrtf(float(n_embd_head)), il); + cb(cur, "attn_out", il); + } + + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids); + } + + ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA); + cb(ffn_inp, "ffn_inp", il); + + // feed-forward network + bool is_moe_layer = + static_cast(il) >= hparams.n_layer_dense_lead && (il + 1) % hparams.n_moe_layer_step == 0; + + if (!is_moe_layer) { + // Dense layer + cur = build_norm(ffn_inp, model.layers[il].ffn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "ffn_norm", il); + + cur = build_ffn(cur, + model.layers[il].ffn_up, NULL, NULL, + model.layers[il].ffn_gate, NULL, NULL, + model.layers[il].ffn_down, NULL, NULL, + NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(cur, "ffn_out", il); + } else { + // MoE branch + cur = build_norm(ffn_inp, model.layers[il].ffn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "ffn_norm", il); + + ggml_tensor * moe_out = nullptr; + + // Use vision experts for vision tokens, text experts for text tokens + if (ubatch.embd) { + // Vision tokens: use vision MoE experts + moe_out = build_moe_ffn(cur, + model.layers[il].v_ffn_gate_inp, + model.layers[il].v_ffn_up_exps, + model.layers[il].v_ffn_gate_exps, + model.layers[il].v_ffn_down_exps, + model.layers[il].v_ffn_exp_probs_b, + n_expert, n_expert_used, + LLM_FFN_SILU, true, + false, 0.0, + LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX, + il); + cb(moe_out, "ffn_moe_out", il); + } else { + // Text tokens: use text MoE experts + moe_out = build_moe_ffn(cur, + model.layers[il].ffn_gate_inp, + model.layers[il].ffn_up_exps, + model.layers[il].ffn_gate_exps, + model.layers[il].ffn_down_exps, + model.layers[il].ffn_exp_probs_b, + n_expert, n_expert_used, + LLM_FFN_SILU, true, + false, 0.0, + LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX, + il); + cb(moe_out, "ffn_moe_out", il); + } + + // Shared expert (if present) + if (hparams.n_ff_shexp > 0) { + ggml_tensor * ffn_shexp = + build_ffn(cur, + model.layers[il].ffn_up_shexp, NULL, NULL, + model.layers[il].ffn_gate_shexp, NULL, NULL, + model.layers[il].ffn_down_shexp, NULL, NULL, + NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(ffn_shexp, "ffn_shexp", il); + + cur = ggml_add(ctx0, moe_out, ffn_shexp); + } else { + cur = moe_out; + } + cb(cur, "ffn_out", il); + } + + cur = ggml_add(ctx0, cur, ffn_inp); + cb(cur, "ffn_residual", il); + + cur = build_cvec(cur, il); + cb(cur, "l_out", il); + + // input for next layer + inpL = cur; + } + + cur = inpL; + + cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); + cb(cur, "result_norm", -1); + res->t_embd = cur; + + // lm_head + cur = build_lora_mm(model.output, cur); + + cb(cur, "result_output", -1); + res->t_logits = cur; + + ggml_build_forward_expand(gf, cur); +} diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index 9fa5afc390e3..3b2f0462cad2 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -829,6 +829,10 @@ static ggml_cgraph * clip_image_build_graph(clip_ctx * ctx, const clip_image_f32 { builder = std::make_unique(ctx, img); } break; + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + builder = std::make_unique(ctx, img); + } break; case PROJECTOR_TYPE_MLP: case PROJECTOR_TYPE_MLP_NORM: case PROJECTOR_TYPE_LDP: @@ -1139,6 +1143,15 @@ struct clip_model_loader { hparams.set_limit_image_tokens(8, 1024); hparams.set_warmup_n_tokens(256); // avoid OOM on warmup } break; + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + hparams.n_merge = 2; + hparams.spatial_conv_size = 2; + hparams.temporal_conv_size = 2; + hparams.use_temporal_conv = model.mm_temp_0_w != nullptr; + hparams.set_limit_image_tokens(8, 1024); + hparams.set_warmup_n_tokens(256); // avoid OOM on warmup + } break; case PROJECTOR_TYPE_GEMMA3: { // default value (used by all model sizes in gemma 3 family) @@ -1831,6 +1844,29 @@ struct clip_model_loader { layer.conv_pw2_b = get_tensor(string_format(TN_CONV_PW2, prefix, il, "bias")); } } break; + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + // spatial path + model.mm_spatial_0_w = get_tensor("mm.0.weight"); + model.mm_spatial_0_b = get_tensor("mm.0.bias"); + model.mm_spatial_2_w = get_tensor("mm.2.weight"); + model.mm_spatial_2_b = get_tensor("mm.2.bias"); + model.mm_spatial_norm_w = get_tensor("mm.3.weight"); + model.mm_spatial_norm_b = get_tensor("mm.3.bias", false); + + // temporal path (optional, not used for single images) + model.mm_temp_0_w = get_tensor("mm_temp.0.weight", false); + model.mm_temp_0_b = get_tensor("mm_temp.0.bias", false); + model.mm_temp_2_w = get_tensor("mm_temp.2.weight", false); + model.mm_temp_2_b = get_tensor("mm_temp.2.bias", false); + model.mm_temp_norm_w = get_tensor("mm_temp.3.weight", false); + model.mm_temp_norm_b = get_tensor("mm_temp.3.bias", false); + + // output + model.mm_mlp_w = get_tensor("mm.mlp.weight"); + model.mm_mlp_b = get_tensor("mm.mlp.bias"); + model.mm_after_norm_w = get_tensor("mm.norm.weight"); + } break; default: GGML_ASSERT(false && "unknown projector type"); } @@ -3003,7 +3039,21 @@ bool clip_image_preprocess(struct clip_ctx * ctx, const clip_image_u8 * img, str normalize_image_u8_to_f32(resized_image, *img_f32, params.image_mean, params.image_std); res_imgs->entries.push_back(std::move(img_f32)); } break; - + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + GGML_ASSERT(params.image_min_pixels > 0 && params.image_max_pixels > 0); + clip_image_u8 resized_image; + const int cur_merge = params.n_merge == 0 ? 1 : params.n_merge; + const clip_image_size target_size = img_tool::calc_size_preserved_ratio( + original_size, + params.patch_size * cur_merge, + params.image_min_pixels, + params.image_max_pixels); + img_tool::resize(*img, resized_image, target_size, img_tool::RESIZE_ALGO_BILINEAR); + clip_image_f32_ptr img_f32(clip_image_f32_init()); + normalize_image_u8_to_f32(resized_image, *img_f32, params.image_mean, params.image_std); + res_imgs->entries.push_back(std::move(img_f32)); + } break; case PROJECTOR_TYPE_LLAMA4: { GGML_ASSERT(!params.image_res_candidates.empty()); @@ -3145,6 +3195,8 @@ int clip_n_output_tokens_x(const struct clip_ctx * ctx, struct clip_image_f32 * case PROJECTOR_TYPE_GLM4V: case PROJECTOR_TYPE_YOUTUVL: return (img->nx / params.patch_size) / 2; + case PROJECTOR_TYPE_ERNIE45VLMOE: + return (img->nx / params.patch_size) / 2; default: break; } @@ -3161,6 +3213,8 @@ int clip_n_output_tokens_y(const struct clip_ctx * ctx, struct clip_image_f32 * case PROJECTOR_TYPE_GLM4V: case PROJECTOR_TYPE_YOUTUVL: return (img->ny / params.patch_size) / 2; + case PROJECTOR_TYPE_ERNIE45VLMOE: + return (img->nx / params.patch_size) / 2; default: break; } @@ -3230,6 +3284,13 @@ int clip_n_output_tokens(const struct clip_ctx * ctx, struct clip_image_f32 * im int y_patch = img->ny / (params.patch_size * 2); n_patches = x_patch * y_patch; } break; + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + // dynamic size (2 conv, so double patch size) + int x_patch = img->nx / (params.patch_size * 2); + int y_patch = img->ny / (params.patch_size * 2); + n_patches = x_patch * y_patch; + } break; case PROJECTOR_TYPE_GEMMA3: case PROJECTOR_TYPE_IDEFICS3: case PROJECTOR_TYPE_INTERNVL: @@ -3584,6 +3645,25 @@ bool clip_image_batch_encode(clip_ctx * ctx, const int n_threads, const clip_ima } } + set_input_i32("positions", positions); + } break; + case PROJECTOR_TYPE_ERNIE45VLMOE: + { + const int pw = image_size_width / patch_size; + const int ph = image_size_height / patch_size; + std::vector positions(n_pos * 4); + int ptr = 0; + + for (int y = 0; y < ph; y++) { + for (int x = 0; x < pw; x++) { + positions[ ptr] = y; + positions[ num_patches + ptr] = x; + positions[2 * num_patches + ptr] = 0; + positions[3 * num_patches + ptr] = 0; + ptr++; + } + } + set_input_i32("positions", positions); } break; case PROJECTOR_TYPE_PIXTRAL: @@ -3777,6 +3857,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { return ctx->model.position_embeddings->ne[0]; case PROJECTOR_TYPE_GLM4V: return ctx->model.mm_ffn_down_w->ne[1]; + case PROJECTOR_TYPE_ERNIE45VLMOE: + return ctx->model.mm_mlp_w->ne[0]; default: GGML_ABORT("Unknown projector type"); } @@ -3795,6 +3877,14 @@ bool clip_is_glm(const struct clip_ctx * ctx) { return ctx->proj_type() == PROJECTOR_TYPE_GLM_EDGE; } +bool clip_is_mrope(const struct clip_ctx * ctx) { + return ctx->proj_type() == PROJECTOR_TYPE_QWEN2VL + || ctx->proj_type() == PROJECTOR_TYPE_QWEN25VL + || ctx->proj_type() == PROJECTOR_TYPE_QWEN3VL + || ctx->proj_type() == PROJECTOR_TYPE_ERNIE45VLMOE + || ctx->proj_type() == PROJECTOR_TYPE_GLM4V; +} + bool clip_is_llava(const struct clip_ctx * ctx) { return ctx->model.hparams.has_llava_projector; } diff --git a/tools/mtmd/models/ernie45vlmoe.cpp b/tools/mtmd/models/ernie45vlmoe.cpp new file mode 100644 index 000000000000..cf2a3a9fbaa1 --- /dev/null +++ b/tools/mtmd/models/ernie45vlmoe.cpp @@ -0,0 +1,219 @@ +#include "models.h" + +ggml_cgraph * clip_graph_ernie45vlmoe::build() { + // ERNIE-4.5-VL-MoE Vision + Resampler: + // 1. ViT encoder with 2D position embeddings and M-RoPE support + // 2. Resampler with spatial conv (2x2 grouping) + optional temporal + MLP + RMS norm + + const int batch_size = 1; + const int n_pos = n_patches; + const int spatial_conv_size = hparams.spatial_conv_size; // 2 + const int temporal_conv_size = hparams.temporal_conv_size; // 2 + const bool use_temporal = hparams.use_temporal_conv; + + // GGML_ASSERT(n_patches_x == n_patches_y && "only square images supported"); + GGML_ASSERT(spatial_conv_size == 2 && "ERNIE-4.5-VL-MoE requires spatial_conv_size=2"); + + // ERNIE-VL Vision uses 2D position lookup RoPE: + // - Front half of frequencies use h_position + // - Back half of frequencies use w_position + // For d_head=80, n_dims=40, we need sections[0]=20 (for h) and sections[1]=20 (for w) + // GGML_ROPE_TYPE_VISION uses only 2 sections: sect_0 for first pos slot, sect_1 for second + int mrope_sections[4] = {d_head/4, d_head/4, 0, 0}; // [20, 20, 0, 0] for d_head=80 + + const int num_position_ids = n_pos * 4; // m-rope requires 4 dim per position + ggml_tensor * positions = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, num_position_ids); + ggml_set_name(positions, "positions"); + ggml_set_input(positions); + + + + // Build vision encoder with patch embedding + // Note: patch_embeddings_0 is reshaped to 4D during export for conv2d compatibility + ggml_tensor * inp = build_inp(); + + + + // ERNIE-4.5-VL uses RoPE (Rotary Position Embedding), not learned position embeddings + // Position encoding is applied within attention layers via RoPE + // So we don't need to add position embeddings here + + ggml_tensor * inpL = inp; + + // Pre-layernorm + if (model.pre_ln_w) { + inpL = build_norm(inpL, model.pre_ln_w, model.pre_ln_b, NORM_TYPE_NORMAL, eps, -1); + cb(inpL, "pre_ln", -1); + } + + // Loop over encoder layers + for (int il = 0; il < n_layer; il++) { + const auto & layer = model.layers[il]; + ggml_tensor * cur = inpL; + + // Layernorm 1 + cur = build_norm(cur, layer.ln_1_w, layer.ln_1_b, NORM_TYPE_NORMAL, eps, il); + cb(cur, "ln1", il); + + // Self-attention + { + ggml_tensor * Qcur = ggml_mul_mat(ctx0, layer.q_w, cur); + if (layer.q_b) { + Qcur = ggml_add(ctx0, Qcur, layer.q_b); + } + + ggml_tensor * Kcur = ggml_mul_mat(ctx0, layer.k_w, cur); + if (layer.k_b) { + Kcur = ggml_add(ctx0, Kcur, layer.k_b); + } + + ggml_tensor * Vcur = ggml_mul_mat(ctx0, layer.v_w, cur); + if (layer.v_b) { + Vcur = ggml_add(ctx0, Vcur, layer.v_b); + } + + Qcur = ggml_reshape_3d(ctx0, Qcur, d_head, n_head, n_pos); + Kcur = ggml_reshape_3d(ctx0, Kcur, d_head, n_head, n_pos); + Vcur = ggml_reshape_3d(ctx0, Vcur, d_head, n_head, n_pos); + + cb(Qcur, "Qcur", il); + cb(Kcur, "Kcur", il); + cb(Vcur, "Vcur", il); + + + // apply M-RoPE + Qcur = ggml_rope_multi( + ctx0, Qcur, positions, nullptr, + d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, 32768, 10000, 1, 0, 1, 32, 1); + Kcur = ggml_rope_multi( + ctx0, Kcur, positions, nullptr, + d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, 32768, 10000, 1, 0, 1, 32, 1); + + cb(Qcur, "Qcur_rope", il); + cb(Kcur, "Kcur_rope", il); + + + + cur = build_attn(layer.o_w, layer.o_b, Qcur, Kcur, Vcur, nullptr, kq_scale, il); + cb(cur, "attn_out", il); + } + + // Residual + cur = ggml_add(ctx0, cur, inpL); + inpL = cur; + cb(cur, "ffn_inp", il); + + // Layernorm 2 + cur = build_norm(cur, layer.ln_2_w, layer.ln_2_b, NORM_TYPE_NORMAL, eps, il); + cb(cur, "ffn_inp_normed", il); + + // FFN + cur = build_ffn(cur, layer.ff_up_w, layer.ff_up_b, layer.ff_gate_w, layer.ff_gate_b, layer.ff_down_w, + layer.ff_down_b, hparams.ffn_op, il); + + cb(cur, "ffn_out", il); + + // Residual 2 + cur = ggml_add(ctx0, inpL, cur); + cb(cur, "layer_out", il); + + inpL = cur; + } + + // Post-layernorm + if (model.post_ln_w) { + inpL = build_norm(inpL, model.post_ln_w, model.post_ln_b, NORM_TYPE_NORMAL, eps, -1); + } + + ggml_tensor * embeddings = inpL; + cb(embeddings, "vision_output", -1); + + // ------------------------------------------- + // Resampler projection + // ------------------------------------------- + // Input shape: [n_embd, n_patches] = [1280, 1600] + // We need to group 2x2 patches: 40x40 patches -> 20x20 groups + // Output shape: [n_embd*4, n_groups] = [5120, 400] + + const int n_groups_x = n_patches_x / spatial_conv_size; // 40/2 = 20 + const int n_groups_y = n_patches_y / spatial_conv_size; // 40/2 = 20 + const int n_groups = n_groups_x * n_groups_y; // 400 + + // Use patch_merge_permute to group 2x2 patches + // Note: build_patch_merge_permute expects 2D input [n_embd, n_patches] + embeddings = build_patch_merge_permute(embeddings, spatial_conv_size); + + // embeddings is now [n_embd*4, n_groups] = [5120, 400] + cb(embeddings, "spatial_reshape", -1); + + // Spatial linear path: Linear -> GELU -> Linear -> LayerNorm + // Note: weights were transposed (.t()) during GGUF conversion, so we must + // undo that with ggml_transpose before ggml_mul_mat + ggml_tensor * spatial_out = embeddings; + + // First linear + ggml_tensor * mm_spatial_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_spatial_0_w)); + spatial_out = ggml_mul_mat(ctx0, mm_spatial_0_w, spatial_out); + spatial_out = ggml_add(ctx0, spatial_out, model.mm_spatial_0_b); + cb(spatial_out, "spatial_linear_0", -1); + + // GELU + spatial_out = ggml_gelu(ctx0, spatial_out); + cb(spatial_out, "spatial_gelu", -1); + + // Second linear + ggml_tensor * mm_spatial_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_spatial_2_w)); + spatial_out = ggml_mul_mat(ctx0, mm_spatial_2_w, spatial_out); + spatial_out = ggml_add(ctx0, spatial_out, model.mm_spatial_2_b); + cb(spatial_out, "spatial_linear_2", -1); + + // LayerNorm + spatial_out = build_norm(spatial_out, model.mm_spatial_norm_w, model.mm_spatial_norm_b, NORM_TYPE_NORMAL, eps, -1); + cb(spatial_out, "spatial_norm", -1); + + ggml_tensor * resampler_out = spatial_out; + + // Temporal processing for single images (t=1): + // Following ERNIE-VL original: when t=1, slice_offsets and slice_offsets2 both point to the same frame + // So we concat(x, x, dim=-1) which in GGML's [hidden, seq] layout is dim=0 + // This doubles the hidden dimension: [5120, 400] -> [10240, 400] + resampler_out = ggml_concat(ctx0, resampler_out, resampler_out, 0); + + // Temporal linear path: Linear -> GELU -> Linear -> LayerNorm + // Weights were transposed (.t()) during GGUF conversion, undo with ggml_transpose + + // First temporal linear + ggml_tensor * mm_temp_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_temp_0_w)); + resampler_out = ggml_mul_mat(ctx0, mm_temp_0_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_temp_0_b); + cb(resampler_out, "temporal_linear_0", -1); + + // GELU + resampler_out = ggml_gelu(ctx0, resampler_out); + cb(resampler_out, "temporal_gelu", -1); + + // Second temporal linear + ggml_tensor * mm_temp_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_temp_2_w)); + resampler_out = ggml_mul_mat(ctx0, mm_temp_2_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_temp_2_b); + cb(resampler_out, "temporal_linear_2", -1); + + // LayerNorm + resampler_out = build_norm(resampler_out, model.mm_temp_norm_w, model.mm_temp_norm_b, NORM_TYPE_NORMAL, eps, -1); + cb(resampler_out, "temporal_norm", -1); + + // Final MLP + ggml_tensor * mm_mlp_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_mlp_w)); + resampler_out = ggml_mul_mat(ctx0, mm_mlp_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_mlp_b); + cb(resampler_out, "mlp", -1); + + // RMS norm (final output normalization) + resampler_out = build_norm(resampler_out, model.mm_after_norm_w, nullptr, NORM_TYPE_RMS, eps, -1); + cb(resampler_out, "after_norm", -1); + + // Build the graph + ggml_build_forward_expand(gf, resampler_out); + + return gf; +} diff --git a/tools/mtmd/mtmd.cpp b/tools/mtmd/mtmd.cpp index d037e834f3bc..48018e85ac5c 100644 --- a/tools/mtmd/mtmd.cpp +++ b/tools/mtmd/mtmd.cpp @@ -315,6 +315,11 @@ struct mtmd_context { img_end = "<|end_of_image|>"; } + else if (proj == PROJECTOR_TYPE_ERNIE45VLMOE) { + img_beg = "<|IMAGE_START|>"; + img_end = "<|IMAGE_END|>"; + + } } void init_audio() { From 59774d232f3c76c0e90896e15ba5ec5d9f5c6141 Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 01:12:13 +0800 Subject: [PATCH 2/8] [model] support ernie4-5-vl-moe model --- src/llama-arch.h | 1118 +++++++++++++++++++++---------------------- src/llama-model.cpp | 1 + src/models/models.h | 4 + 3 files changed, 542 insertions(+), 581 deletions(-) diff --git a/src/llama-arch.h b/src/llama-arch.h index dcd03b2b2129..8485f44cacc7 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -1,605 +1,561 @@ #pragma once -#include "ggml.h" // ggml_op - +#include "llama.h" +#include "llama-arch.h" +#include "llama-graph.h" +#include "llama-hparams.h" +#include "llama-memory.h" +#include "llama-vocab.h" + +#include +#include #include -#include - -// -// gguf constants (sync with gguf.py) -// - -enum llm_arch { - LLM_ARCH_CLIP, - LLM_ARCH_LLAMA, - LLM_ARCH_LLAMA4, - LLM_ARCH_DECI, - LLM_ARCH_FALCON, - LLM_ARCH_BAICHUAN, - LLM_ARCH_GROK, - LLM_ARCH_GPT2, - LLM_ARCH_GPTJ, - LLM_ARCH_GPTNEOX, - LLM_ARCH_MPT, - LLM_ARCH_STARCODER, - LLM_ARCH_REFACT, - LLM_ARCH_BERT, - LLM_ARCH_MODERN_BERT, - LLM_ARCH_NOMIC_BERT, - LLM_ARCH_NOMIC_BERT_MOE, - LLM_ARCH_NEO_BERT, - LLM_ARCH_JINA_BERT_V2, - LLM_ARCH_JINA_BERT_V3, - LLM_ARCH_BLOOM, - LLM_ARCH_STABLELM, - LLM_ARCH_QWEN, - LLM_ARCH_QWEN2, - LLM_ARCH_QWEN2MOE, - LLM_ARCH_QWEN2VL, - LLM_ARCH_QWEN3, - LLM_ARCH_QWEN3MOE, - LLM_ARCH_QWEN3NEXT, - LLM_ARCH_QWEN3_5, - LLM_ARCH_QWEN3_5_MOE, - LLM_ARCH_QWEN3VL, - LLM_ARCH_QWEN3VLMOE, - LLM_ARCH_PHI2, - LLM_ARCH_PHI3, - LLM_ARCH_PHIMOE, - LLM_ARCH_PLAMO, - LLM_ARCH_PLAMO2, - LLM_ARCH_PLAMO3, - LLM_ARCH_CODESHELL, - LLM_ARCH_ORION, - LLM_ARCH_INTERNLM2, - LLM_ARCH_MINICPM, - LLM_ARCH_MINICPM3, - LLM_ARCH_GEMMA, - LLM_ARCH_GEMMA2, - LLM_ARCH_GEMMA3, - LLM_ARCH_GEMMA3N, - LLM_ARCH_GEMMA_EMBEDDING, - LLM_ARCH_STARCODER2, - LLM_ARCH_MAMBA, - LLM_ARCH_MAMBA2, - LLM_ARCH_JAMBA, - LLM_ARCH_FALCON_H1, - LLM_ARCH_XVERSE, - LLM_ARCH_COMMAND_R, - LLM_ARCH_COHERE2, - LLM_ARCH_DBRX, - LLM_ARCH_OLMO, - LLM_ARCH_OLMO2, - LLM_ARCH_OLMOE, - LLM_ARCH_OPENELM, - LLM_ARCH_ARCTIC, - LLM_ARCH_DEEPSEEK, - LLM_ARCH_DEEPSEEK2, - LLM_ARCH_CHATGLM, - LLM_ARCH_GLM4, - LLM_ARCH_GLM4_MOE, - LLM_ARCH_BITNET, - LLM_ARCH_T5, - LLM_ARCH_T5ENCODER, - LLM_ARCH_JAIS, - LLM_ARCH_NEMOTRON, - LLM_ARCH_NEMOTRON_H, - LLM_ARCH_NEMOTRON_H_MOE, - LLM_ARCH_EXAONE, - LLM_ARCH_EXAONE4, - LLM_ARCH_EXAONE_MOE, - LLM_ARCH_RWKV6, - LLM_ARCH_RWKV6QWEN2, - LLM_ARCH_RWKV7, - LLM_ARCH_ARWKV7, - LLM_ARCH_GRANITE, - LLM_ARCH_GRANITE_MOE, - LLM_ARCH_GRANITE_HYBRID, - LLM_ARCH_CHAMELEON, - LLM_ARCH_WAVTOKENIZER_DEC, - LLM_ARCH_PLM, - LLM_ARCH_BAILINGMOE, - LLM_ARCH_BAILINGMOE2, - LLM_ARCH_DOTS1, - LLM_ARCH_ARCEE, - LLM_ARCH_AFMOE, - LLM_ARCH_ERNIE4_5, - LLM_ARCH_ERNIE4_5_MOE, - LLM_ARCH_HUNYUAN_MOE, - LLM_ARCH_HUNYUAN_DENSE, - LLM_ARCH_SMOLLM3, - LLM_ARCH_OPENAI_MOE, - LLM_ARCH_LFM2, - LLM_ARCH_LFM2MOE, - LLM_ARCH_DREAM, - LLM_ARCH_SMALLTHINKER, - LLM_ARCH_LLADA, - LLM_ARCH_LLADA_MOE, - LLM_ARCH_SEED_OSS, - LLM_ARCH_GROVEMOE, - LLM_ARCH_APERTUS, - LLM_ARCH_MINIMAX_M2, - LLM_ARCH_COGVLM, - LLM_ARCH_RND1, - LLM_ARCH_PANGU_EMBED, - LLM_ARCH_MISTRAL3, - LLM_ARCH_MIMO2, - LLM_ARCH_STEP35, - LLM_ARCH_LLAMA_EMBED, - LLM_ARCH_MAINCODER, - LLM_ARCH_KIMI_LINEAR, - LLM_ARCH_UNKNOWN, +#include +#include +#include + +struct llama_cparams; +struct llama_ubatch; +struct llama_model_loader; + +// available models +enum llm_type { + LLM_TYPE_UNKNOWN, + LLM_TYPE_14M, + LLM_TYPE_17M, + LLM_TYPE_22M, + LLM_TYPE_33M, + LLM_TYPE_47M, + LLM_TYPE_60M, + LLM_TYPE_70M, + LLM_TYPE_80M, + LLM_TYPE_109M, + LLM_TYPE_137M, + LLM_TYPE_140M, + LLM_TYPE_149M, + LLM_TYPE_160M, + LLM_TYPE_190M, + LLM_TYPE_220M, + LLM_TYPE_250M, + LLM_TYPE_256M, + LLM_TYPE_270M, + LLM_TYPE_335M, + LLM_TYPE_350M, + LLM_TYPE_360M, + LLM_TYPE_395M, + LLM_TYPE_410M, + LLM_TYPE_450M, + LLM_TYPE_475M, + LLM_TYPE_558M, + LLM_TYPE_700M, + LLM_TYPE_770M, + LLM_TYPE_780M, + LLM_TYPE_950M, + LLM_TYPE_0_3B, + LLM_TYPE_0_5B, + LLM_TYPE_0_6B, + LLM_TYPE_1B, + LLM_TYPE_1_2B, + LLM_TYPE_1_3B, + LLM_TYPE_1_4B, + LLM_TYPE_1_5B, + LLM_TYPE_1_6B, + LLM_TYPE_1_7B, + LLM_TYPE_1_8B, + LLM_TYPE_2B, + LLM_TYPE_2_6B, + LLM_TYPE_2_8B, + LLM_TYPE_2_9B, + LLM_TYPE_3B, + LLM_TYPE_4B, + LLM_TYPE_6B, + LLM_TYPE_6_9B, + LLM_TYPE_7B, + LLM_TYPE_8B, + LLM_TYPE_9B, + LLM_TYPE_11B, + LLM_TYPE_12B, + LLM_TYPE_13B, + LLM_TYPE_14B, + LLM_TYPE_15B, + LLM_TYPE_16B, + LLM_TYPE_20B, + LLM_TYPE_26B, + LLM_TYPE_27B, + LLM_TYPE_30B, + LLM_TYPE_32B, + LLM_TYPE_34B, + LLM_TYPE_35B, + LLM_TYPE_36B, + LLM_TYPE_40B, + LLM_TYPE_65B, + LLM_TYPE_70B, + LLM_TYPE_120B, + LLM_TYPE_142B, + LLM_TYPE_236B, + LLM_TYPE_290B, + LLM_TYPE_314B, + LLM_TYPE_405B, + LLM_TYPE_671B, + LLM_TYPE_SMALL, + LLM_TYPE_MEDIUM, + LLM_TYPE_LARGE, + LLM_TYPE_XL, + LLM_TYPE_A1_7B, + LLM_TYPE_A2_7B, + LLM_TYPE_8x7B, + LLM_TYPE_8x22B, + LLM_TYPE_16x12B, + LLM_TYPE_16x3_8B, + LLM_TYPE_10B_128x3_66B, + LLM_TYPE_57B_A14B, + LLM_TYPE_17B_16E, // llama4 Scout + LLM_TYPE_17B_128E, // llama4 Maverick + LLM_TYPE_A13B, + LLM_TYPE_7B_A1B, + LLM_TYPE_8B_A1B, // lfm2moe + LLM_TYPE_16B_A1B, + LLM_TYPE_21B_A3B, // Ernie MoE small + LLM_TYPE_28B_A3B, // Ernie MoE vl small + LLM_TYPE_30B_A3B, + LLM_TYPE_31B_A3_5B, + LLM_TYPE_48B_A3B, // Kimi Linear + LLM_TYPE_80B_A3B, // Qwen3 Next + LLM_TYPE_100B_A6B, + LLM_TYPE_102B_A12B, // Solar-Open + LLM_TYPE_106B_A12B, // GLM-4.5-Air + LLM_TYPE_196B_A11B, // Step3.5-Flash + LLM_TYPE_230B_A10B, // Minimax M2 + LLM_TYPE_235B_A22B, + LLM_TYPE_300B_A47B, // Ernie MoE big + LLM_TYPE_310B_A15B, // /MiMo-V2-Flash + LLM_TYPE_355B_A32B, // GLM-4.5 + LLM_TYPE_E2B, + LLM_TYPE_E4B, }; -enum llm_kv { - LLM_KV_GENERAL_TYPE, - LLM_KV_GENERAL_ARCHITECTURE, - LLM_KV_GENERAL_QUANTIZATION_VERSION, - LLM_KV_GENERAL_ALIGNMENT, - LLM_KV_GENERAL_FILE_TYPE, - LLM_KV_GENERAL_SAMPLING_SEQUENCE, - LLM_KV_GENERAL_SAMPLING_TOP_K, - LLM_KV_GENERAL_SAMPLING_TOP_P, - LLM_KV_GENERAL_SAMPLING_MIN_P, - LLM_KV_GENERAL_SAMPLING_XTC_PROBABILITY, - LLM_KV_GENERAL_SAMPLING_XTC_THRESHOLD, - LLM_KV_GENERAL_SAMPLING_TEMP, - LLM_KV_GENERAL_SAMPLING_PENALTY_LAST_N, - LLM_KV_GENERAL_SAMPLING_PENALTY_REPEAT, - LLM_KV_GENERAL_SAMPLING_MIROSTAT, - LLM_KV_GENERAL_SAMPLING_MIROSTAT_TAU, - LLM_KV_GENERAL_SAMPLING_MIROSTAT_ETA, - LLM_KV_GENERAL_NAME, - LLM_KV_GENERAL_AUTHOR, - LLM_KV_GENERAL_VERSION, - LLM_KV_GENERAL_URL, - LLM_KV_GENERAL_DESCRIPTION, - LLM_KV_GENERAL_LICENSE, - LLM_KV_GENERAL_SOURCE_URL, - LLM_KV_GENERAL_SOURCE_HF_REPO, - - LLM_KV_VOCAB_SIZE, - LLM_KV_CONTEXT_LENGTH, - LLM_KV_EMBEDDING_LENGTH, - LLM_KV_EMBEDDING_LENGTH_OUT, - LLM_KV_FEATURES_LENGTH, - LLM_KV_BLOCK_COUNT, - LLM_KV_LEADING_DENSE_BLOCK_COUNT, - LLM_KV_FEED_FORWARD_LENGTH, - LLM_KV_EXPERT_FEED_FORWARD_LENGTH, - LLM_KV_VISION_EXPERT_FEED_FORWARD_LENGTH, - LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, - LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, - LLM_KV_SWIGLU_CLAMP_EXP, - LLM_KV_SWIGLU_CLAMP_SHEXP, - LLM_KV_USE_PARALLEL_RESIDUAL, - LLM_KV_TENSOR_DATA_LAYOUT, - LLM_KV_EXPERT_COUNT, - LLM_KV_EXPERT_USED_COUNT, - LLM_KV_EXPERT_SHARED_COUNT, - LLM_KV_EXPERT_GROUP_COUNT, - LLM_KV_EXPERT_GROUP_USED_COUNT, - LLM_KV_EXPERT_WEIGHTS_SCALE, - LLM_KV_EXPERT_WEIGHTS_NORM, - LLM_KV_EXPERT_GATING_FUNC, - LLM_KV_EXPERT_GROUP_SCALE, - LLM_KV_EXPERTS_PER_GROUP, - LLM_KV_MOE_EVERY_N_LAYERS, - LLM_KV_NEXTN_PREDICT_LAYERS, - LLM_KV_NUM_DEEPSTACK_LAYERS, - LLM_KV_POOLING_TYPE, - LLM_KV_LOGIT_SCALE, - LLM_KV_DECODER_START_TOKEN_ID, - LLM_KV_DECODER_BLOCK_COUNT, - LLM_KV_ATTN_LOGIT_SOFTCAPPING, - LLM_KV_ROUTER_LOGIT_SOFTCAPPING, - LLM_KV_FINAL_LOGIT_SOFTCAPPING, - LLM_KV_SWIN_NORM, - LLM_KV_RESCALE_EVERY_N_LAYERS, - LLM_KV_TIME_MIX_EXTRA_DIM, - LLM_KV_TIME_DECAY_EXTRA_DIM, - LLM_KV_RESIDUAL_SCALE, - LLM_KV_EMBEDDING_SCALE, - LLM_KV_TOKEN_SHIFT_COUNT, - LLM_KV_INTERLEAVE_MOE_LAYER_STEP, - - LLM_KV_ATTENTION_HEAD_COUNT, - LLM_KV_ATTENTION_HEAD_COUNT_KV, - LLM_KV_ATTENTION_MAX_ALIBI_BIAS, - LLM_KV_ATTENTION_CLAMP_KQV, - LLM_KV_ATTENTION_KEY_LENGTH, - LLM_KV_ATTENTION_VALUE_LENGTH, - LLM_KV_ATTENTION_LAYERNORM_EPS, - LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, - LLM_KV_ATTENTION_GROUPNORM_EPS, - LLM_KV_ATTENTION_GROUPNORM_GROUPS, - LLM_KV_ATTENTION_CAUSAL, - LLM_KV_ATTENTION_Q_LORA_RANK, - LLM_KV_ATTENTION_KV_LORA_RANK, - LLM_KV_ATTENTION_DECAY_LORA_RANK, - LLM_KV_ATTENTION_ICLR_LORA_RANK, - LLM_KV_ATTENTION_VALUE_RESIDUAL_MIX_LORA_RANK, - LLM_KV_ATTENTION_GATE_LORA_RANK, - LLM_KV_ATTENTION_RELATIVE_BUCKETS_COUNT, - LLM_KV_ATTENTION_SLIDING_WINDOW, - LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, - LLM_KV_ATTENTION_SCALE, - LLM_KV_ATTENTION_OUTPUT_SCALE, - LLM_KV_ATTENTION_TEMPERATURE_LENGTH, - LLM_KV_ATTENTION_TEMPERATURE_SCALE, - LLM_KV_ATTENTION_KEY_LENGTH_MLA, - LLM_KV_ATTENTION_VALUE_LENGTH_MLA, - - LLM_KV_ROPE_DIMENSION_COUNT, - LLM_KV_ROPE_DIMENSION_SECTIONS, - LLM_KV_ROPE_FREQ_BASE, - LLM_KV_ROPE_FREQ_BASE_SWA, - LLM_KV_ROPE_SCALE_LINEAR, - LLM_KV_ROPE_SCALING_TYPE, - LLM_KV_ROPE_SCALING_FACTOR, - LLM_KV_ROPE_SCALING_ATTN_FACTOR, - LLM_KV_ROPE_SCALING_ORIG_CTX_LEN, - LLM_KV_ROPE_SCALING_FINETUNED, - LLM_KV_ROPE_SCALING_YARN_LOG_MUL, - LLM_KV_ROPE_SCALING_YARN_EXT_FACTOR, - LLM_KV_ROPE_SCALING_YARN_ATTN_FACTOR, - LLM_KV_ROPE_SCALING_YARN_BETA_FAST, - LLM_KV_ROPE_SCALING_YARN_BETA_SLOW, - - LLM_KV_SPLIT_NO, - LLM_KV_SPLIT_COUNT, - LLM_KV_SPLIT_TENSORS_COUNT, - - LLM_KV_SSM_INNER_SIZE, - LLM_KV_SSM_CONV_KERNEL, - LLM_KV_SSM_STATE_SIZE, - LLM_KV_SSM_TIME_STEP_RANK, - LLM_KV_SSM_GROUP_COUNT, - LLM_KV_SSM_DT_B_C_RMS, - - LLM_KV_KDA_HEAD_DIM, - - LLM_KV_WKV_HEAD_SIZE, - - LLM_KV_TOKENIZER_MODEL, - LLM_KV_TOKENIZER_PRE, - LLM_KV_TOKENIZER_LIST, - LLM_KV_TOKENIZER_TOKEN_TYPE, - LLM_KV_TOKENIZER_TOKEN_TYPE_COUNT, - LLM_KV_TOKENIZER_SCORES, - LLM_KV_TOKENIZER_MERGES, - LLM_KV_TOKENIZER_BOS_ID, - LLM_KV_TOKENIZER_EOS_ID, - LLM_KV_TOKENIZER_EOT_ID, - LLM_KV_TOKENIZER_EOM_ID, - LLM_KV_TOKENIZER_UNK_ID, - LLM_KV_TOKENIZER_SEP_ID, - LLM_KV_TOKENIZER_PAD_ID, - LLM_KV_TOKENIZER_CLS_ID, - LLM_KV_TOKENIZER_MASK_ID, - LLM_KV_TOKENIZER_ADD_BOS, - LLM_KV_TOKENIZER_ADD_EOS, - LLM_KV_TOKENIZER_ADD_SEP, - LLM_KV_TOKENIZER_ADD_PREFIX, - LLM_KV_TOKENIZER_REMOVE_EXTRA_WS, - LLM_KV_TOKENIZER_PRECOMPILED_CHARSMAP, - LLM_KV_TOKENIZER_HF_JSON, - LLM_KV_TOKENIZER_RWKV, - LLM_KV_TOKENIZER_CHAT_TEMPLATE, - LLM_KV_TOKENIZER_FIM_PRE_ID, - LLM_KV_TOKENIZER_FIM_SUF_ID, - LLM_KV_TOKENIZER_FIM_MID_ID, - LLM_KV_TOKENIZER_FIM_PAD_ID, - LLM_KV_TOKENIZER_FIM_REP_ID, - LLM_KV_TOKENIZER_FIM_SEP_ID, - - LLM_KV_ADAPTER_TYPE, - LLM_KV_ADAPTER_LORA_ALPHA, - LLM_KV_ADAPTER_LORA_TASK_NAME, - LLM_KV_ADAPTER_LORA_PROMPT_PREFIX, - LLM_KV_ADAPTER_ALORA_INVOCATION_TOKENS, - - LLM_KV_POSNET_EMBEDDING_LENGTH, - LLM_KV_POSNET_BLOCK_COUNT, - - LLM_KV_CONVNEXT_EMBEDDING_LENGTH, - LLM_KV_CONVNEXT_BLOCK_COUNT, - - LLM_KV_CLASSIFIER_OUTPUT_LABELS, - - LLM_KV_SHORTCONV_L_CACHE, - - LLM_KV_XIELU_ALPHA_N, - LLM_KV_XIELU_ALPHA_P, - LLM_KV_XIELU_BETA, - LLM_KV_XIELU_EPS, - - // deprecated: - LLM_KV_TOKENIZER_PREFIX_ID, - LLM_KV_TOKENIZER_SUFFIX_ID, - LLM_KV_TOKENIZER_MIDDLE_ID, - - // sentence-transformers dense layers in and out features - LLM_KV_DENSE_2_FEAT_IN, - LLM_KV_DENSE_2_FEAT_OUT, - LLM_KV_DENSE_3_FEAT_IN, - LLM_KV_DENSE_3_FEAT_OUT, -}; +std::string llama_rope_scaling_type_name(llama_rope_scaling_type rope_scaling_type); -enum llm_tensor { - LLM_TENSOR_TOKEN_EMBD, - LLM_TENSOR_TOKEN_EMBD_NORM, - LLM_TENSOR_TOKEN_TYPES, - LLM_TENSOR_POS_EMBD, - LLM_TENSOR_DENSE_2_OUT, - LLM_TENSOR_DENSE_3_OUT, - LLM_TENSOR_OUTPUT, - LLM_TENSOR_OUTPUT_NORM, - LLM_TENSOR_OUTPUT_NORM_LFM2, // fix for wrong tensor name - LLM_TENSOR_ROPE_FREQS, - LLM_TENSOR_ROPE_FACTORS_LONG, - LLM_TENSOR_ROPE_FACTORS_SHORT, - LLM_TENSOR_ATTN_Q, - LLM_TENSOR_ATTN_K, - LLM_TENSOR_ATTN_V, - LLM_TENSOR_ATTN_QKV, - LLM_TENSOR_ATTN_OUT, - LLM_TENSOR_ATTN_NORM, - LLM_TENSOR_ATTN_NORM_2, - LLM_TENSOR_ATTN_OUT_NORM, - LLM_TENSOR_ATTN_POST_NORM, - LLM_TENSOR_ATTN_ROT_EMBD, - LLM_TENSOR_ATTN_SINKS, - LLM_TENSOR_ATTN_GATE, - LLM_TENSOR_FFN_GATE_INP, - LLM_TENSOR_FFN_GATE_INP_SHEXP, - LLM_TENSOR_FFN_NORM, - LLM_TENSOR_FFN_POST_NORM, - LLM_TENSOR_FFN_GATE, - LLM_TENSOR_FFN_DOWN, - LLM_TENSOR_FFN_UP, - LLM_TENSOR_FFN_ACT, - LLM_TENSOR_FFN_DOWN_EXP, // split experts for backward compatibility - LLM_TENSOR_FFN_GATE_EXP, - LLM_TENSOR_FFN_UP_EXP, - LLM_TENSOR_FFN_NORM_EXPS, - LLM_TENSOR_FFN_DOWN_EXPS, // merged experts - LLM_TENSOR_FFN_GATE_EXPS, - LLM_TENSOR_FFN_UP_EXPS, - LLM_TENSOR_FFN_DOWN_SHEXP, - LLM_TENSOR_FFN_GATE_SHEXP, - LLM_TENSOR_FFN_UP_SHEXP, - LLM_TENSOR_FFN_DOWN_CHEXPS, - LLM_TENSOR_FFN_GATE_CHEXPS, - LLM_TENSOR_FFN_UP_CHEXPS, - LLM_TENSOR_FFN_EXP_PROBS_B, - LLM_TENSOR_ATTN_Q_NORM, - LLM_TENSOR_ATTN_K_NORM, - LLM_TENSOR_LAYER_OUT_NORM, - LLM_TENSOR_POST_ATTN_NORM, - LLM_TENSOR_POST_MLP_NORM, - LLM_TENSOR_PER_LAYER_TOKEN_EMBD, // gemma3n - LLM_TENSOR_PER_LAYER_MODEL_PROJ, // gemma3n - LLM_TENSOR_PER_LAYER_INP_GATE, // gemma3n - LLM_TENSOR_PER_LAYER_PROJ, // gemma3n - LLM_TENSOR_PER_LAYER_PROJ_NORM, // gemma3n - LLM_TENSOR_PER_LAYER_POST_NORM, // gemma3n - LLM_TENSOR_ALTUP_PROJ, // gemma3n - LLM_TENSOR_ALTUP_UNEMBD_PROJ, // gemma3n - LLM_TENSOR_ALTUP_CORRECT_COEF, // gemma3n - LLM_TENSOR_ALTUP_CORRECT_SCALE, // gemma3n - LLM_TENSOR_ALTUP_PREDICT_COEF, // gemma3n - LLM_TENSOR_ALTUP_ROUTER, // gemma3n - LLM_TENSOR_ALTUP_ROUTER_NORM, // gemma3n - LLM_TENSOR_LAUREL_L, // gemma3n - LLM_TENSOR_LAUREL_R, // gemma3n - LLM_TENSOR_LAUREL_POST_NORM, // gemma3n - LLM_TENSOR_SSM_IN, - LLM_TENSOR_SSM_CONV1D, - LLM_TENSOR_SSM_X, - LLM_TENSOR_SSM_DT, - LLM_TENSOR_SSM_DT_NORM, - LLM_TENSOR_SSM_A, - LLM_TENSOR_SSM_A_NOSCAN, // qwen3next special case with MUL instead of SSM_SCAN - LLM_TENSOR_SSM_B_NORM, - LLM_TENSOR_SSM_C_NORM, - LLM_TENSOR_SSM_D, - LLM_TENSOR_SSM_NORM, - LLM_TENSOR_SSM_OUT, - LLM_TENSOR_SSM_BETA_ALPHA, // qwen3next - // Kimi Linear KDA (using SSM_ prefix for consistency) - LLM_TENSOR_SSM_CONV1D_Q, // kimi: Q conv1d weight - LLM_TENSOR_SSM_CONV1D_K, // kimi: K conv1d weight - LLM_TENSOR_SSM_CONV1D_V, // kimi: V conv1d weight - LLM_TENSOR_SSM_F_A, // kimi: forget gate projection A - LLM_TENSOR_SSM_F_B, // kimi: forget gate projection B - LLM_TENSOR_SSM_BETA, // kimi: beta mixing coefficient - LLM_TENSOR_SSM_G_A, // kimi: output gate projection A - LLM_TENSOR_SSM_G_B, // kimi: output gate projection B - LLM_TENSOR_TIME_MIX_W0, - LLM_TENSOR_TIME_MIX_W1, - LLM_TENSOR_TIME_MIX_W2, - LLM_TENSOR_TIME_MIX_A0, - LLM_TENSOR_TIME_MIX_A1, - LLM_TENSOR_TIME_MIX_A2, - LLM_TENSOR_TIME_MIX_V0, - LLM_TENSOR_TIME_MIX_V1, - LLM_TENSOR_TIME_MIX_V2, - LLM_TENSOR_TIME_MIX_G1, - LLM_TENSOR_TIME_MIX_G2, - LLM_TENSOR_TIME_MIX_K_K, - LLM_TENSOR_TIME_MIX_K_A, - LLM_TENSOR_TIME_MIX_R_K, - LLM_TENSOR_TIME_MIX_LERP_X, - LLM_TENSOR_TIME_MIX_LERP_W, - LLM_TENSOR_TIME_MIX_LERP_K, - LLM_TENSOR_TIME_MIX_LERP_V, - LLM_TENSOR_TIME_MIX_LERP_R, - LLM_TENSOR_TIME_MIX_LERP_G, - LLM_TENSOR_TIME_MIX_LERP_FUSED, - LLM_TENSOR_TIME_MIX_FIRST, - LLM_TENSOR_TIME_MIX_DECAY, - LLM_TENSOR_TIME_MIX_DECAY_W1, - LLM_TENSOR_TIME_MIX_DECAY_W2, - LLM_TENSOR_TIME_MIX_KEY, - LLM_TENSOR_TIME_MIX_VALUE, - LLM_TENSOR_TIME_MIX_RECEPTANCE, - LLM_TENSOR_TIME_MIX_GATE, - LLM_TENSOR_TIME_MIX_LN, - LLM_TENSOR_TIME_MIX_OUTPUT, - LLM_TENSOR_CHANNEL_MIX_LERP_K, - LLM_TENSOR_CHANNEL_MIX_LERP_R, - LLM_TENSOR_CHANNEL_MIX_KEY, - LLM_TENSOR_CHANNEL_MIX_RECEPTANCE, - LLM_TENSOR_CHANNEL_MIX_VALUE, - LLM_TENSOR_ATTN_Q_A, - LLM_TENSOR_ATTN_Q_B, - LLM_TENSOR_ATTN_KV_A_MQA, - LLM_TENSOR_ATTN_KV_B, - LLM_TENSOR_ATTN_K_B, - LLM_TENSOR_ATTN_V_B, - LLM_TENSOR_ATTN_Q_A_NORM, - LLM_TENSOR_ATTN_KV_A_NORM, - LLM_TENSOR_ATTN_SUB_NORM, - LLM_TENSOR_FFN_SUB_NORM, - LLM_TENSOR_DEC_ATTN_NORM, - LLM_TENSOR_DEC_ATTN_Q, - LLM_TENSOR_DEC_ATTN_K, - LLM_TENSOR_DEC_ATTN_V, - LLM_TENSOR_DEC_ATTN_OUT, - LLM_TENSOR_DEC_ATTN_REL_B, - LLM_TENSOR_DEC_CROSS_ATTN_NORM, - LLM_TENSOR_DEC_CROSS_ATTN_Q, - LLM_TENSOR_DEC_CROSS_ATTN_K, - LLM_TENSOR_DEC_CROSS_ATTN_V, - LLM_TENSOR_DEC_CROSS_ATTN_OUT, - LLM_TENSOR_DEC_CROSS_ATTN_REL_B, - LLM_TENSOR_DEC_FFN_NORM, - LLM_TENSOR_DEC_FFN_GATE, - LLM_TENSOR_DEC_FFN_DOWN, - LLM_TENSOR_DEC_FFN_UP, - LLM_TENSOR_DEC_OUTPUT_NORM, - LLM_TENSOR_ENC_ATTN_NORM, - LLM_TENSOR_ENC_ATTN_Q, - LLM_TENSOR_ENC_ATTN_K, - LLM_TENSOR_ENC_ATTN_V, - LLM_TENSOR_ENC_ATTN_OUT, - LLM_TENSOR_ENC_ATTN_REL_B, - LLM_TENSOR_ENC_FFN_NORM, - LLM_TENSOR_ENC_FFN_GATE, - LLM_TENSOR_ENC_FFN_DOWN, - LLM_TENSOR_ENC_FFN_UP, - LLM_TENSOR_ENC_OUTPUT_NORM, - LLM_TENSOR_CLS, - LLM_TENSOR_CLS_OUT, - LLM_TENSOR_CONV1D, - LLM_TENSOR_CONVNEXT_DW, - LLM_TENSOR_CONVNEXT_NORM, - LLM_TENSOR_CONVNEXT_PW1, - LLM_TENSOR_CONVNEXT_PW2, - LLM_TENSOR_CONVNEXT_GAMMA, - LLM_TENSOR_POS_NET_CONV1, - LLM_TENSOR_POS_NET_CONV2, - LLM_TENSOR_POS_NET_NORM, - LLM_TENSOR_POS_NET_NORM1, - LLM_TENSOR_POS_NET_NORM2, - LLM_TENSOR_POS_NET_ATTN_NORM, - LLM_TENSOR_POS_NET_ATTN_Q, - LLM_TENSOR_POS_NET_ATTN_K, - LLM_TENSOR_POS_NET_ATTN_V, - LLM_TENSOR_POS_NET_ATTN_OUT, - LLM_TENSOR_SHORTCONV_CONV, - LLM_TENSOR_SHORTCONV_INPROJ, - LLM_TENSOR_SHORTCONV_OUTPROJ, - LLM_TENSOR_VISEXP_ATTN_QKV, - LLM_TENSOR_VISEXP_ATTN_OUT, - LLM_TENSOR_VISEXP_FFN_GATE, - LLM_TENSOR_VISEXP_FFN_DOWN, - LLM_TENSOR_VISEXP_FFN_UP, - LLM_TENSOR_NEXTN_EH_PROJ, - LLM_TENSOR_NEXTN_EMBED_TOKENS, - LLM_TENSOR_NEXTN_ENORM, - LLM_TENSOR_NEXTN_HNORM, - LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, - LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, -}; +struct llama_layer_posnet { + // resnet + struct ggml_tensor * norm1 = nullptr; + struct ggml_tensor * norm1_b = nullptr; -enum llm_tensor_layer { - LLM_TENSOR_LAYER_INPUT, - LLM_TENSOR_LAYER_REPEATING, - LLM_TENSOR_LAYER_OUTPUT, -}; + struct ggml_tensor * conv1 = nullptr; + struct ggml_tensor * conv1_b = nullptr; -struct LLM_KV { - LLM_KV(llm_arch arch, const char * suffix = nullptr); + struct ggml_tensor * norm2 = nullptr; + struct ggml_tensor * norm2_b = nullptr; - llm_arch arch; - const char * suffix; + struct ggml_tensor * conv2 = nullptr; + struct ggml_tensor * conv2_b = nullptr; - std::string operator()(llm_kv kv) const; -}; + // attention + struct ggml_tensor * attn_norm = nullptr; + struct ggml_tensor * attn_norm_b = nullptr; + + struct ggml_tensor * attn_q = nullptr; + struct ggml_tensor * attn_q_b = nullptr; -// helper to handle gguf constants -// usage: -// -// const auto tn = LLM_TN(LLM_ARCH_LLAMA); -// -// std::string name = tn(LLM_TENSOR_OUTPUT); -> "output" -// std::string name = tn(LLM_TENSOR_TOKEN_EMBD, "bias"); -> "token_embd.bias" -// std::string name = tn(LLM_TENSOR_ATTN_NORM, "weight", 3); -> "blk.3.attn_norm.weight" -// -struct LLM_TN_IMPL { - const llm_arch arch; - const llm_tensor tensor; - const char * const suffix; - const int bid; - const int xid; - - const std::set model_tensors; - - LLM_TN_IMPL(llm_arch arch, llm_tensor tensor, const char * suffix, int bid, int xid); - - std::string str() const; - - operator std::string() const { - return str(); - } - - friend bool operator==(const std::string & str, const LLM_TN_IMPL & tn) { - return str == tn.str(); - } - - friend bool operator!=(const std::string & str, const LLM_TN_IMPL & tn) { - return str != tn.str(); - } + struct ggml_tensor * attn_k = nullptr; + struct ggml_tensor * attn_k_b = nullptr; + + struct ggml_tensor * attn_v = nullptr; + struct ggml_tensor * attn_v_b = nullptr; + + struct ggml_tensor * attn_o = nullptr; + struct ggml_tensor * attn_o_b = nullptr; + + // normalize + struct ggml_tensor * norm = nullptr; + struct ggml_tensor * norm_b = nullptr; }; -struct LLM_TN { - LLM_TN(llm_arch arch) : arch(arch) {} +struct llama_layer_convnext { + struct ggml_tensor * dw = nullptr; + struct ggml_tensor * dw_b = nullptr; - llm_arch arch; + struct ggml_tensor * norm = nullptr; + struct ggml_tensor * norm_b = nullptr; - LLM_TN_IMPL operator()(llm_tensor tensor, const char * suffix, int bid = -1, int xid = -1) const { - return LLM_TN_IMPL(arch, tensor, suffix, bid, xid); - } + struct ggml_tensor * pw1 = nullptr; + struct ggml_tensor * pw1_b = nullptr; + + struct ggml_tensor * pw2 = nullptr; + struct ggml_tensor * pw2_b = nullptr; + + struct ggml_tensor * gamma = nullptr; +}; - LLM_TN_IMPL operator()(llm_tensor tensor, int bid = -1, int xid = -1) const { - return LLM_TN_IMPL(arch, tensor, nullptr, bid, xid); - } +struct llama_layer_shortconv { + struct ggml_tensor * in_proj = nullptr; + struct ggml_tensor * conv = nullptr; + struct ggml_tensor * out_proj = nullptr; }; +struct llama_layer_nextn { + struct ggml_tensor * eh_proj = nullptr; + struct ggml_tensor * embed_tokens = nullptr; + struct ggml_tensor * enorm = nullptr; + struct ggml_tensor * hnorm = nullptr; + struct ggml_tensor * shared_head_head = nullptr; + struct ggml_tensor * shared_head_norm = nullptr; +}; -struct llm_tensor_info { - llm_tensor_layer layer; - ggml_op op; +struct llama_layer { + // normalization + struct ggml_tensor * attn_norm = nullptr; + struct ggml_tensor * attn_norm_b = nullptr; + struct ggml_tensor * attn_norm_2 = nullptr; + struct ggml_tensor * attn_norm_2_b = nullptr; + struct ggml_tensor * attn_q_norm = nullptr; + struct ggml_tensor * attn_q_norm_b = nullptr; + struct ggml_tensor * attn_k_norm = nullptr; + struct ggml_tensor * attn_k_norm_b = nullptr; + struct ggml_tensor * attn_out_norm = nullptr; + struct ggml_tensor * attn_out_norm_b = nullptr; + struct ggml_tensor * attn_q_a_norm = nullptr; + struct ggml_tensor * attn_kv_a_norm = nullptr; + struct ggml_tensor * attn_sub_norm = nullptr; + struct ggml_tensor * attn_post_norm = nullptr; + struct ggml_tensor * ffn_sub_norm = nullptr; + struct ggml_tensor * attn_norm_cross = nullptr; + struct ggml_tensor * attn_norm_enc = nullptr; + struct ggml_tensor * ssm_norm = nullptr; + struct ggml_tensor * ssm_dt_norm = nullptr; + struct ggml_tensor * ssm_b_norm = nullptr; + struct ggml_tensor * ssm_c_norm = nullptr; + + // attention + struct ggml_tensor * wq = nullptr; + struct ggml_tensor * wk = nullptr; + struct ggml_tensor * wv = nullptr; + struct ggml_tensor * wo = nullptr; + struct ggml_tensor * wqkv = nullptr; + struct ggml_tensor * wq_a = nullptr; + struct ggml_tensor * wq_b = nullptr; + struct ggml_tensor * wkv_a_mqa = nullptr; + struct ggml_tensor * wkv_b = nullptr; + struct ggml_tensor * wk_b = nullptr; + struct ggml_tensor * wv_b = nullptr; + struct ggml_tensor * wq_cross = nullptr; + struct ggml_tensor * wk_cross = nullptr; + struct ggml_tensor * wv_cross = nullptr; + struct ggml_tensor * wo_cross = nullptr; + struct ggml_tensor * wq_enc = nullptr; + struct ggml_tensor * wk_enc = nullptr; + struct ggml_tensor * wv_enc = nullptr; + struct ggml_tensor * wo_enc = nullptr; + struct ggml_tensor * wqkv_gate = nullptr; + + // attention bias + struct ggml_tensor * bq = nullptr; + struct ggml_tensor * bk = nullptr; + struct ggml_tensor * bv = nullptr; + struct ggml_tensor * bo = nullptr; + struct ggml_tensor * bqkv = nullptr; + + // relative position bias + struct ggml_tensor * attn_rel_b = nullptr; + struct ggml_tensor * attn_rel_b_enc = nullptr; + struct ggml_tensor * attn_rel_b_cross = nullptr; + + // normalization + struct ggml_tensor * ffn_norm = nullptr; + struct ggml_tensor * ffn_norm_b = nullptr; + struct ggml_tensor * ffn_post_norm = nullptr; + struct ggml_tensor * layer_out_norm = nullptr; + struct ggml_tensor * layer_out_norm_b = nullptr; + struct ggml_tensor * ffn_norm_exps = nullptr; + struct ggml_tensor * ffn_norm_enc = nullptr; + + // ff + struct ggml_tensor * ffn_gate = nullptr; // w1 + struct ggml_tensor * ffn_down = nullptr; // w2 + struct ggml_tensor * ffn_up = nullptr; // w3 + struct ggml_tensor * ffn_gate_enc = nullptr; + struct ggml_tensor * ffn_down_enc = nullptr; + struct ggml_tensor * ffn_up_enc = nullptr; + + // ff MoE + struct ggml_tensor * ffn_gate_inp = nullptr; + struct ggml_tensor * ffn_gate_exps = nullptr; + struct ggml_tensor * ffn_down_exps = nullptr; + struct ggml_tensor * ffn_up_exps = nullptr; + struct ggml_tensor * ffn_gate_inp_b = nullptr; + struct ggml_tensor * ffn_gate_exps_b = nullptr; + struct ggml_tensor * ffn_down_exps_b = nullptr; + struct ggml_tensor * ffn_up_exps_b = nullptr; + + // ff shared expert (shexp) + struct ggml_tensor * ffn_gate_inp_shexp = nullptr; + struct ggml_tensor * ffn_gate_shexp = nullptr; + struct ggml_tensor * ffn_down_shexp = nullptr; + struct ggml_tensor * ffn_up_shexp = nullptr; + + // ff adjugate experts (chexps) + struct ggml_tensor * ffn_gate_chexps = nullptr; + struct ggml_tensor * ffn_down_chexps = nullptr; + struct ggml_tensor * ffn_up_chexps = nullptr; + + // ff bias + struct ggml_tensor * ffn_gate_b = nullptr; + struct ggml_tensor * ffn_down_b = nullptr; // b2 + 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 * v_ffn_exp_probs_b = nullptr; + + // mamba proj + struct ggml_tensor * ssm_in = nullptr; + struct ggml_tensor * ssm_x = nullptr; + struct ggml_tensor * ssm_dt = nullptr; + struct ggml_tensor * ssm_out = nullptr; + + // mamba + struct ggml_tensor * ssm_conv1d = nullptr; + struct ggml_tensor * ssm_a = nullptr; + struct ggml_tensor * ssm_d = nullptr; + + // mamba bias + struct ggml_tensor * ssm_conv1d_b = nullptr; + struct ggml_tensor * ssm_dt_b = nullptr; + + // qwen3next + struct ggml_tensor * ssm_beta_alpha = nullptr; + + // rwkv + struct ggml_tensor * time_mix_w1 = nullptr; + struct ggml_tensor * time_mix_w2 = nullptr; + struct ggml_tensor * time_mix_lerp_x = nullptr; + struct ggml_tensor * time_mix_lerp_w = nullptr; + struct ggml_tensor * time_mix_lerp_k = nullptr; + struct ggml_tensor * time_mix_lerp_v = nullptr; + struct ggml_tensor * time_mix_lerp_r = nullptr; + struct ggml_tensor * time_mix_lerp_g = nullptr; + struct ggml_tensor * time_mix_lerp_fused = nullptr; + + struct ggml_tensor * time_mix_first = nullptr; + struct ggml_tensor * time_mix_decay = nullptr; + struct ggml_tensor * time_mix_decay_w1 = nullptr; + struct ggml_tensor * time_mix_decay_w2 = nullptr; + struct ggml_tensor * time_mix_key = nullptr; + struct ggml_tensor * time_mix_key_b = nullptr; + struct ggml_tensor * time_mix_value = nullptr; + struct ggml_tensor * time_mix_value_b = nullptr; + struct ggml_tensor * time_mix_receptance = nullptr; + struct ggml_tensor * time_mix_receptance_b = nullptr; + struct ggml_tensor * time_mix_gate = nullptr; + + // rwkv7 + struct ggml_tensor * time_mix_w0 = nullptr; + struct ggml_tensor * time_mix_a0 = nullptr; + struct ggml_tensor * time_mix_a1 = nullptr; + struct ggml_tensor * time_mix_a2 = nullptr; + struct ggml_tensor * time_mix_v0 = nullptr; + struct ggml_tensor * time_mix_v1 = nullptr; + struct ggml_tensor * time_mix_v2 = nullptr; + struct ggml_tensor * time_mix_g1 = nullptr; + struct ggml_tensor * time_mix_g2 = nullptr; + struct ggml_tensor * time_mix_k_k = nullptr; + struct ggml_tensor * time_mix_k_a = nullptr; + struct ggml_tensor * time_mix_r_k = nullptr; + + struct ggml_tensor * time_mix_ln = nullptr; + struct ggml_tensor * time_mix_ln_b = nullptr; + struct ggml_tensor * time_mix_output = nullptr; + + struct ggml_tensor * channel_mix_lerp_k = nullptr; + struct ggml_tensor * channel_mix_lerp_r = nullptr; + + struct ggml_tensor * channel_mix_key = nullptr; + struct ggml_tensor * channel_mix_receptance = nullptr; + struct ggml_tensor * channel_mix_value = nullptr; + + // long rope factors + struct ggml_tensor * rope_long = nullptr; + struct ggml_tensor * rope_short = nullptr; + struct ggml_tensor * rope_freqs = nullptr; + + // bitnet scale + struct ggml_tensor * wq_scale = nullptr; + struct ggml_tensor * wk_scale = nullptr; + struct ggml_tensor * wv_scale = nullptr; + struct ggml_tensor * wo_scale = nullptr; + struct ggml_tensor * ffn_gate_scale = nullptr; + struct ggml_tensor * ffn_up_scale = nullptr; + struct ggml_tensor * ffn_down_scale = nullptr; + + // altup & laurel + struct ggml_tensor * per_layer_inp_gate = nullptr; + struct ggml_tensor * per_layer_proj = nullptr; + struct ggml_tensor * per_layer_post_norm = nullptr; + struct ggml_tensor * altup_correct_coef = nullptr; + struct ggml_tensor * altup_correct_scale = nullptr; + struct ggml_tensor * altup_predict_coef = nullptr; + struct ggml_tensor * altup_router = nullptr; + struct ggml_tensor * altup_router_norm = nullptr; + struct ggml_tensor * laurel_l = nullptr; + struct ggml_tensor * laurel_r = nullptr; + struct ggml_tensor * laurel_post_norm = nullptr; + + // openai-moe + struct ggml_tensor * attn_sinks = nullptr; + + // cogvlm + struct ggml_tensor * visexp_attn_wqkv = nullptr; + struct ggml_tensor * visexp_attn_wo = nullptr; + struct ggml_tensor * visexp_ffn_gate = nullptr; + struct ggml_tensor * visexp_ffn_down = nullptr; + struct ggml_tensor * visexp_ffn_up = nullptr; + + // xIELU activation parameters for Apertus + struct ggml_tensor * ffn_act_alpha_n = nullptr; + struct ggml_tensor * ffn_act_alpha_p = nullptr; + struct ggml_tensor * ffn_act_beta = nullptr; + struct ggml_tensor * ffn_act_eps = nullptr; + + // Kimi Linear KDA (using ssm_ prefix for consistency) + // Note: ssm_dt_b already exists above (mamba bias), reused for Kimi dt_bias + struct ggml_tensor * ssm_q_conv = nullptr; + struct ggml_tensor * ssm_k_conv = nullptr; + struct ggml_tensor * ssm_v_conv = nullptr; + struct ggml_tensor * ssm_f_a = nullptr; + struct ggml_tensor * ssm_f_b = nullptr; + struct ggml_tensor * ssm_beta = nullptr; + struct ggml_tensor * ssm_g_a = nullptr; + struct ggml_tensor * ssm_g_b = nullptr; + struct ggml_tensor * ssm_o_norm = nullptr; + + struct llama_layer_posnet posnet; + + struct llama_layer_convnext convnext; + + struct llama_layer_shortconv shortconv; + + struct llama_layer_nextn nextn; }; -const char * llm_arch_name(llm_arch arch); +struct llama_model { + llm_type type = LLM_TYPE_UNKNOWN; + llm_arch arch = LLM_ARCH_UNKNOWN; + + std::string name = "n/a"; + + llama_hparams hparams = {}; + llama_vocab vocab; + + // for classifier models + std::vector classifier_labels; + + struct ggml_tensor * tok_embd = nullptr; + struct ggml_tensor * type_embd = nullptr; + struct ggml_tensor * pos_embd = nullptr; + struct ggml_tensor * tok_norm = nullptr; + struct ggml_tensor * tok_norm_b = nullptr; + + struct ggml_tensor * output_norm = nullptr; + struct ggml_tensor * output_norm_b = nullptr; + struct ggml_tensor * output = nullptr; + struct ggml_tensor * output_b = nullptr; + struct ggml_tensor * output_norm_enc = nullptr; + + // classifier + struct ggml_tensor * cls = nullptr; + struct ggml_tensor * cls_b = nullptr; + struct ggml_tensor * cls_out = nullptr; + struct ggml_tensor * cls_out_b = nullptr; + + struct ggml_tensor * conv1d = nullptr; + struct ggml_tensor * conv1d_b = nullptr; + + // gemma3n altup + struct ggml_tensor * tok_embd_per_layer = nullptr; + struct ggml_tensor * altup_proj = nullptr; + struct ggml_tensor * altup_unembd_proj = nullptr; + struct ggml_tensor * per_layer_model_proj = nullptr; + struct ggml_tensor * per_layer_proj_norm = nullptr; + + std::vector layers; + + //Dense linear projections for SentenceTransformers models like embeddinggemma + // For Sentence Transformers models structure see + // https://sbert.net/docs/sentence_transformer/usage/custom_models.html#structure-of-sentence-transformer-models + struct ggml_tensor * dense_2_out_layers = nullptr; + struct ggml_tensor * dense_3_out_layers = nullptr; + + // gguf metadata + std::unordered_map gguf_kv; -llm_arch llm_arch_from_string(const std::string & name); + // list of devices used in this model + std::vector devices; + + // for quantize-stats only + std::vector> tensors_by_name; + + // for keeping track of associated LoRA adapters + std::unordered_set loras; + + int64_t t_load_us = 0; + int64_t t_start_us = 0; + + explicit llama_model(const struct llama_model_params & params); + ~llama_model(); + + void load_stats (llama_model_loader & ml); + void load_arch (llama_model_loader & ml); + void load_hparams(llama_model_loader & ml); + void load_vocab (llama_model_loader & ml); + bool load_tensors(llama_model_loader & ml); // returns false if cancelled by progress_callback + + std::string arch_name() const; + std::string type_name() const; + + std::string desc() const; + + size_t size() const; // file size + size_t n_tensors() const; + size_t n_devices() const; + + uint32_t n_gpu_layers() const; + llama_split_mode split_mode() const; + + std::map memory_breakdown() const; + + // total number of parameters in the model + uint64_t n_elements() const; + + void print_info() const; + + ggml_backend_dev_t dev_layer(int il) const; + ggml_backend_dev_t dev_output() const; + + ggml_backend_buffer_type_t select_buft(int il) const; + + bool has_tensor_overrides() const; + + const struct ggml_tensor * get_tensor(const char * name) const; + + float get_rope_freq_base (const llama_cparams & cparams, int il) const; + float get_rope_freq_scale(const llama_cparams & cparams, int il) const; + + ggml_tensor * get_rope_factors(const llama_cparams & cparams, int il) const; + + // TODO: move this to new llm_arch_model_i interface + llama_memory_i * create_memory(const llama_memory_params & params, const llama_cparams & cparams) const; + + // TODO: move this to new llm_arch_model_i interface + ggml_cgraph * build_graph(const llm_graph_params & params) const; + +private: + llama_model_params params; + + struct impl; + std::unique_ptr pimpl; +}; -const llm_tensor_info & llm_tensor_info_for(llm_tensor tensor); +const char * llm_type_name(llm_type type); -bool llm_arch_is_recurrent(const llm_arch & arch); -bool llm_arch_is_hybrid (const llm_arch & arch); -bool llm_arch_is_diffusion(const llm_arch & arch); +// For internal test use +// TODO: remove +const std::vector> & llama_internal_get_tensor_map(const llama_model * model); diff --git a/src/llama-model.cpp b/src/llama-model.cpp index e2be33e841e6..d409c68b5286 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -123,6 +123,7 @@ const char * llm_type_name(llm_type type) { case LLM_TYPE_8B_A1B: return "8B.A1B"; case LLM_TYPE_16B_A1B: return "16B.A1B"; case LLM_TYPE_21B_A3B: return "21B.A3B"; + case LLM_TYPE_28B_A3B: return "28B.A3B"; case LLM_TYPE_30B_A3B: return "30B.A3B"; case LLM_TYPE_31B_A3_5B: return "31B.A3.5B"; case LLM_TYPE_48B_A3B: return "48B.A3B"; diff --git a/src/models/models.h b/src/models/models.h index 2a750c168ea9..dbb75dc9571a 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -205,6 +205,10 @@ struct llm_build_ernie4_5_moe : public llm_graph_context { llm_build_ernie4_5_moe(const llama_model & model, const llm_graph_params & params); }; +struct llm_build_ernie4_5_vl_moe : public llm_graph_context { + llm_build_ernie4_5_vl_moe(const llama_model & model, const llm_graph_params & params); +}; + template struct llm_build_exaone4 : public llm_graph_context { llm_build_exaone4(const llama_model & model, const llm_graph_params & params); From b35f2dbd219cc692e92d91389f55f59151b0a977 Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 11:22:28 +0800 Subject: [PATCH 3/8] [model] support ernie4-5-vl-moe model --- src/CMakeLists.txt | 1 + src/llama-model.h | 8 ++++++++ tools/mtmd/CMakeLists.txt | 1 + tools/mtmd/clip-impl.h | 2 ++ tools/mtmd/clip-model.h | 22 ++++++++++++++++++++++ tools/mtmd/models/models.h | 5 +++++ 6 files changed, 39 insertions(+) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 0c164617a12e..2d3f13209b01 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -61,6 +61,7 @@ add_library(llama models/dots1.cpp models/dream.cpp models/ernie4-5-moe.cpp + models/ernie4-5-vl-moe.cpp models/ernie4-5.cpp models/exaone.cpp models/exaone4.cpp diff --git a/src/llama-model.h b/src/llama-model.h index 7b580043b337..876cd30280fd 100644 --- a/src/llama-model.h +++ b/src/llama-model.h @@ -116,6 +116,7 @@ enum llm_type { LLM_TYPE_8B_A1B, // lfm2moe LLM_TYPE_16B_A1B, LLM_TYPE_21B_A3B, // Ernie MoE small + LLM_TYPE_28B_A3B, // Ernie MoE vl small LLM_TYPE_30B_A3B, LLM_TYPE_31B_A3_5B, LLM_TYPE_48B_A3B, // Kimi Linear @@ -286,6 +287,12 @@ struct llama_layer { struct ggml_tensor * ffn_down_exps_b = nullptr; struct ggml_tensor * ffn_up_exps_b = nullptr; + // ff Vision expert MoE + struct ggml_tensor * v_ffn_gate_inp = nullptr; + struct ggml_tensor * v_ffn_gate_exps = nullptr; + struct ggml_tensor * v_ffn_down_exps = nullptr; + struct ggml_tensor * v_ffn_up_exps = nullptr; + // ff shared expert (shexp) struct ggml_tensor * ffn_gate_inp_shexp = nullptr; struct ggml_tensor * ffn_gate_shexp = nullptr; @@ -303,6 +310,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 * v_ffn_exp_probs_b = nullptr; // mamba proj struct ggml_tensor * ssm_in = nullptr; diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt index 751440af323c..ac840a23ed0b 100644 --- a/tools/mtmd/CMakeLists.txt +++ b/tools/mtmd/CMakeLists.txt @@ -29,6 +29,7 @@ add_library(mtmd models/whisper-enc.cpp models/mobilenetv5.cpp models/youtuvl.cpp + models/ernie45vlmoe.cpp ) set_target_properties(mtmd PROPERTIES diff --git a/tools/mtmd/clip-impl.h b/tools/mtmd/clip-impl.h index ad232178bf4b..af54b8be793f 100644 --- a/tools/mtmd/clip-impl.h +++ b/tools/mtmd/clip-impl.h @@ -235,6 +235,7 @@ enum projector_type { PROJECTOR_TYPE_LFM2A, PROJECTOR_TYPE_GLM4V, PROJECTOR_TYPE_YOUTUVL, + PROJECTOR_TYPE_ERNIE45VLMOE, PROJECTOR_TYPE_UNKNOWN, }; @@ -268,6 +269,7 @@ static std::map PROJECTOR_TYPE_NAMES = { { PROJECTOR_TYPE_LFM2A, "lfm2a"}, { PROJECTOR_TYPE_GLM4V, "glm4v"}, { PROJECTOR_TYPE_YOUTUVL, "youtuvl"}, + { PROJECTOR_TYPE_ERNIE45VLMOE, "ernie45vlmoe"}, }; static projector_type clip_projector_type_from_string(const std::string & str) { diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index d4ff9151bb08..c4e54323c77c 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -79,6 +79,11 @@ struct clip_hparams { int minicpmv_version = 0; int32_t minicpmv_query_num = 0; // MiniCPM-V query number + // ernie4.5-vl-moe + int32_t spatial_conv_size = 0; + int32_t temporal_conv_size = 0; + bool use_temporal_conv = false; + // custom value provided by user, can be undefined if not set int32_t custom_image_min_tokens = -1; int32_t custom_image_max_tokens = -1; @@ -359,6 +364,23 @@ struct clip_model { ggml_tensor * mm_norm_pre_b = nullptr; ggml_tensor * mm_norm_mid_w = nullptr; + // ernie4.5-vl-moe + ggml_tensor * mm_spatial_0_w = nullptr; + ggml_tensor * mm_spatial_0_b = nullptr; + ggml_tensor * mm_spatial_2_w = nullptr; + ggml_tensor * mm_spatial_2_b = nullptr; + ggml_tensor * mm_spatial_norm_w = nullptr; + ggml_tensor * mm_spatial_norm_b = nullptr; + ggml_tensor * mm_temp_0_w = nullptr; + ggml_tensor * mm_temp_0_b = nullptr; + ggml_tensor * mm_temp_2_w = nullptr; + ggml_tensor * mm_temp_2_b = nullptr; + ggml_tensor * mm_temp_norm_w = nullptr; + ggml_tensor * mm_temp_norm_b = nullptr; + ggml_tensor * mm_mlp_w = nullptr; + ggml_tensor * mm_mlp_b = nullptr; + ggml_tensor * mm_after_norm_w = nullptr; + // cogvlm ggml_tensor * mm_post_fc_norm_w = nullptr; ggml_tensor * mm_post_fc_norm_b = nullptr; diff --git a/tools/mtmd/models/models.h b/tools/mtmd/models/models.h index 9970980c7bc7..2e7966c40f49 100644 --- a/tools/mtmd/models/models.h +++ b/tools/mtmd/models/models.h @@ -109,3 +109,8 @@ struct clip_graph_mobilenetv5 : clip_graph { ggml_tensor * inp, const mobilenetv5_block & block); }; + +struct clip_graph_ernie45vlmoe : clip_graph { + clip_graph_ernie45vlmoe(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} + ggml_cgraph * build() override; +}; From 790be8a22d7cdd8a4ffc25fd74e489e60042246a Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 20:48:31 +0800 Subject: [PATCH 4/8] [model] support ernie4-5-vl-moe model --- src/llama-arch.cpp | 7 + src/llama-arch.h | 1124 +++++++++++++++++++++++--------------------- 2 files changed, 594 insertions(+), 537 deletions(-) diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 8affd81694f6..6dd427274736 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -101,6 +101,7 @@ static const std::map LLM_ARCH_NAMES = { { LLM_ARCH_AFMOE, "afmoe" }, { LLM_ARCH_ERNIE4_5, "ernie4_5" }, { LLM_ARCH_ERNIE4_5_MOE, "ernie4_5-moe" }, + { LLM_ARCH_ERNIE4_5_VL_MOE, "ernie4_5-vl-moe" }, { LLM_ARCH_HUNYUAN_MOE, "hunyuan-moe" }, { LLM_ARCH_HUNYUAN_DENSE, "hunyuan-dense" }, { LLM_ARCH_SMOLLM3, "smollm3" }, @@ -2107,6 +2108,7 @@ static std::set llm_get_tensor_names(llm_arch arch) { LLM_TENSOR_FFN_EXP_PROBS_B, }; case LLM_ARCH_ERNIE4_5_MOE: + case LLM_ARCH_ERNIE4_5_VL_MOE: return { LLM_TENSOR_TOKEN_EMBD, LLM_TENSOR_OUTPUT_NORM, @@ -2128,6 +2130,11 @@ static std::set llm_get_tensor_names(llm_arch arch) { LLM_TENSOR_FFN_DOWN_EXPS, LLM_TENSOR_FFN_UP_EXPS, LLM_TENSOR_FFN_EXP_PROBS_B, + LLM_TENSOR_V_FFN_GATE_INP, + LLM_TENSOR_V_FFN_GATE_EXPS, + LLM_TENSOR_V_FFN_DOWN_EXPS, + LLM_TENSOR_V_FFN_UP_EXPS, + LLM_TENSOR_V_FFN_EXP_PROBS_B, }; case LLM_ARCH_HUNYUAN_MOE: return { diff --git a/src/llama-arch.h b/src/llama-arch.h index 8485f44cacc7..bed5f5317488 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -1,561 +1,611 @@ #pragma once -#include "llama.h" -#include "llama-arch.h" -#include "llama-graph.h" -#include "llama-hparams.h" -#include "llama-memory.h" -#include "llama-vocab.h" - -#include -#include -#include -#include -#include -#include - -struct llama_cparams; -struct llama_ubatch; -struct llama_model_loader; - -// available models -enum llm_type { - LLM_TYPE_UNKNOWN, - LLM_TYPE_14M, - LLM_TYPE_17M, - LLM_TYPE_22M, - LLM_TYPE_33M, - LLM_TYPE_47M, - LLM_TYPE_60M, - LLM_TYPE_70M, - LLM_TYPE_80M, - LLM_TYPE_109M, - LLM_TYPE_137M, - LLM_TYPE_140M, - LLM_TYPE_149M, - LLM_TYPE_160M, - LLM_TYPE_190M, - LLM_TYPE_220M, - LLM_TYPE_250M, - LLM_TYPE_256M, - LLM_TYPE_270M, - LLM_TYPE_335M, - LLM_TYPE_350M, - LLM_TYPE_360M, - LLM_TYPE_395M, - LLM_TYPE_410M, - LLM_TYPE_450M, - LLM_TYPE_475M, - LLM_TYPE_558M, - LLM_TYPE_700M, - LLM_TYPE_770M, - LLM_TYPE_780M, - LLM_TYPE_950M, - LLM_TYPE_0_3B, - LLM_TYPE_0_5B, - LLM_TYPE_0_6B, - LLM_TYPE_1B, - LLM_TYPE_1_2B, - LLM_TYPE_1_3B, - LLM_TYPE_1_4B, - LLM_TYPE_1_5B, - LLM_TYPE_1_6B, - LLM_TYPE_1_7B, - LLM_TYPE_1_8B, - LLM_TYPE_2B, - LLM_TYPE_2_6B, - LLM_TYPE_2_8B, - LLM_TYPE_2_9B, - LLM_TYPE_3B, - LLM_TYPE_4B, - LLM_TYPE_6B, - LLM_TYPE_6_9B, - LLM_TYPE_7B, - LLM_TYPE_8B, - LLM_TYPE_9B, - LLM_TYPE_11B, - LLM_TYPE_12B, - LLM_TYPE_13B, - LLM_TYPE_14B, - LLM_TYPE_15B, - LLM_TYPE_16B, - LLM_TYPE_20B, - LLM_TYPE_26B, - LLM_TYPE_27B, - LLM_TYPE_30B, - LLM_TYPE_32B, - LLM_TYPE_34B, - LLM_TYPE_35B, - LLM_TYPE_36B, - LLM_TYPE_40B, - LLM_TYPE_65B, - LLM_TYPE_70B, - LLM_TYPE_120B, - LLM_TYPE_142B, - LLM_TYPE_236B, - LLM_TYPE_290B, - LLM_TYPE_314B, - LLM_TYPE_405B, - LLM_TYPE_671B, - LLM_TYPE_SMALL, - LLM_TYPE_MEDIUM, - LLM_TYPE_LARGE, - LLM_TYPE_XL, - LLM_TYPE_A1_7B, - LLM_TYPE_A2_7B, - LLM_TYPE_8x7B, - LLM_TYPE_8x22B, - LLM_TYPE_16x12B, - LLM_TYPE_16x3_8B, - LLM_TYPE_10B_128x3_66B, - LLM_TYPE_57B_A14B, - LLM_TYPE_17B_16E, // llama4 Scout - LLM_TYPE_17B_128E, // llama4 Maverick - LLM_TYPE_A13B, - LLM_TYPE_7B_A1B, - LLM_TYPE_8B_A1B, // lfm2moe - LLM_TYPE_16B_A1B, - LLM_TYPE_21B_A3B, // Ernie MoE small - LLM_TYPE_28B_A3B, // Ernie MoE vl small - LLM_TYPE_30B_A3B, - LLM_TYPE_31B_A3_5B, - LLM_TYPE_48B_A3B, // Kimi Linear - LLM_TYPE_80B_A3B, // Qwen3 Next - LLM_TYPE_100B_A6B, - LLM_TYPE_102B_A12B, // Solar-Open - LLM_TYPE_106B_A12B, // GLM-4.5-Air - LLM_TYPE_196B_A11B, // Step3.5-Flash - LLM_TYPE_230B_A10B, // Minimax M2 - LLM_TYPE_235B_A22B, - LLM_TYPE_300B_A47B, // Ernie MoE big - LLM_TYPE_310B_A15B, // /MiMo-V2-Flash - LLM_TYPE_355B_A32B, // GLM-4.5 - LLM_TYPE_E2B, - LLM_TYPE_E4B, -}; - -std::string llama_rope_scaling_type_name(llama_rope_scaling_type rope_scaling_type); - -struct llama_layer_posnet { - // resnet - struct ggml_tensor * norm1 = nullptr; - struct ggml_tensor * norm1_b = nullptr; - - struct ggml_tensor * conv1 = nullptr; - struct ggml_tensor * conv1_b = nullptr; - - struct ggml_tensor * norm2 = nullptr; - struct ggml_tensor * norm2_b = nullptr; - - struct ggml_tensor * conv2 = nullptr; - struct ggml_tensor * conv2_b = nullptr; - - // attention - struct ggml_tensor * attn_norm = nullptr; - struct ggml_tensor * attn_norm_b = nullptr; - - struct ggml_tensor * attn_q = nullptr; - struct ggml_tensor * attn_q_b = nullptr; - - struct ggml_tensor * attn_k = nullptr; - struct ggml_tensor * attn_k_b = nullptr; - - struct ggml_tensor * attn_v = nullptr; - struct ggml_tensor * attn_v_b = nullptr; - - struct ggml_tensor * attn_o = nullptr; - struct ggml_tensor * attn_o_b = nullptr; +#include "ggml.h" // ggml_op - // normalize - struct ggml_tensor * norm = nullptr; - struct ggml_tensor * norm_b = nullptr; -}; - -struct llama_layer_convnext { - struct ggml_tensor * dw = nullptr; - struct ggml_tensor * dw_b = nullptr; - - struct ggml_tensor * norm = nullptr; - struct ggml_tensor * norm_b = nullptr; - - struct ggml_tensor * pw1 = nullptr; - struct ggml_tensor * pw1_b = nullptr; - - struct ggml_tensor * pw2 = nullptr; - struct ggml_tensor * pw2_b = nullptr; - - struct ggml_tensor * gamma = nullptr; +#include +#include + +// +// gguf constants (sync with gguf.py) +// + +enum llm_arch { + LLM_ARCH_CLIP, + LLM_ARCH_LLAMA, + LLM_ARCH_LLAMA4, + LLM_ARCH_DECI, + LLM_ARCH_FALCON, + LLM_ARCH_BAICHUAN, + LLM_ARCH_GROK, + LLM_ARCH_GPT2, + LLM_ARCH_GPTJ, + LLM_ARCH_GPTNEOX, + LLM_ARCH_MPT, + LLM_ARCH_STARCODER, + LLM_ARCH_REFACT, + LLM_ARCH_BERT, + LLM_ARCH_MODERN_BERT, + LLM_ARCH_NOMIC_BERT, + LLM_ARCH_NOMIC_BERT_MOE, + LLM_ARCH_NEO_BERT, + LLM_ARCH_JINA_BERT_V2, + LLM_ARCH_JINA_BERT_V3, + LLM_ARCH_BLOOM, + LLM_ARCH_STABLELM, + LLM_ARCH_QWEN, + LLM_ARCH_QWEN2, + LLM_ARCH_QWEN2MOE, + LLM_ARCH_QWEN2VL, + LLM_ARCH_QWEN3, + LLM_ARCH_QWEN3MOE, + LLM_ARCH_QWEN3NEXT, + LLM_ARCH_QWEN3_5, + LLM_ARCH_QWEN3_5_MOE, + LLM_ARCH_QWEN3VL, + LLM_ARCH_QWEN3VLMOE, + LLM_ARCH_PHI2, + LLM_ARCH_PHI3, + LLM_ARCH_PHIMOE, + LLM_ARCH_PLAMO, + LLM_ARCH_PLAMO2, + LLM_ARCH_PLAMO3, + LLM_ARCH_CODESHELL, + LLM_ARCH_ORION, + LLM_ARCH_INTERNLM2, + LLM_ARCH_MINICPM, + LLM_ARCH_MINICPM3, + LLM_ARCH_GEMMA, + LLM_ARCH_GEMMA2, + LLM_ARCH_GEMMA3, + LLM_ARCH_GEMMA3N, + LLM_ARCH_GEMMA_EMBEDDING, + LLM_ARCH_STARCODER2, + LLM_ARCH_MAMBA, + LLM_ARCH_MAMBA2, + LLM_ARCH_JAMBA, + LLM_ARCH_FALCON_H1, + LLM_ARCH_XVERSE, + LLM_ARCH_COMMAND_R, + LLM_ARCH_COHERE2, + LLM_ARCH_DBRX, + LLM_ARCH_OLMO, + LLM_ARCH_OLMO2, + LLM_ARCH_OLMOE, + LLM_ARCH_OPENELM, + LLM_ARCH_ARCTIC, + LLM_ARCH_DEEPSEEK, + LLM_ARCH_DEEPSEEK2, + LLM_ARCH_CHATGLM, + LLM_ARCH_GLM4, + LLM_ARCH_GLM4_MOE, + LLM_ARCH_BITNET, + LLM_ARCH_T5, + LLM_ARCH_T5ENCODER, + LLM_ARCH_JAIS, + LLM_ARCH_NEMOTRON, + LLM_ARCH_NEMOTRON_H, + LLM_ARCH_NEMOTRON_H_MOE, + LLM_ARCH_EXAONE, + LLM_ARCH_EXAONE4, + LLM_ARCH_EXAONE_MOE, + LLM_ARCH_RWKV6, + LLM_ARCH_RWKV6QWEN2, + LLM_ARCH_RWKV7, + LLM_ARCH_ARWKV7, + LLM_ARCH_GRANITE, + LLM_ARCH_GRANITE_MOE, + LLM_ARCH_GRANITE_HYBRID, + LLM_ARCH_CHAMELEON, + LLM_ARCH_WAVTOKENIZER_DEC, + LLM_ARCH_PLM, + LLM_ARCH_BAILINGMOE, + LLM_ARCH_BAILINGMOE2, + LLM_ARCH_DOTS1, + LLM_ARCH_ARCEE, + LLM_ARCH_AFMOE, + LLM_ARCH_ERNIE4_5, + LLM_ARCH_ERNIE4_5_MOE, + LLM_ARCH_ERNIE4_5_VL_MOE, + LLM_ARCH_HUNYUAN_MOE, + LLM_ARCH_HUNYUAN_DENSE, + LLM_ARCH_SMOLLM3, + LLM_ARCH_OPENAI_MOE, + LLM_ARCH_LFM2, + LLM_ARCH_LFM2MOE, + LLM_ARCH_DREAM, + LLM_ARCH_SMALLTHINKER, + LLM_ARCH_LLADA, + LLM_ARCH_LLADA_MOE, + LLM_ARCH_SEED_OSS, + LLM_ARCH_GROVEMOE, + LLM_ARCH_APERTUS, + LLM_ARCH_MINIMAX_M2, + LLM_ARCH_COGVLM, + LLM_ARCH_RND1, + LLM_ARCH_PANGU_EMBED, + LLM_ARCH_MISTRAL3, + LLM_ARCH_MIMO2, + LLM_ARCH_STEP35, + LLM_ARCH_LLAMA_EMBED, + LLM_ARCH_MAINCODER, + LLM_ARCH_KIMI_LINEAR, + LLM_ARCH_UNKNOWN, }; -struct llama_layer_shortconv { - struct ggml_tensor * in_proj = nullptr; - struct ggml_tensor * conv = nullptr; - struct ggml_tensor * out_proj = nullptr; +enum llm_kv { + LLM_KV_GENERAL_TYPE, + LLM_KV_GENERAL_ARCHITECTURE, + LLM_KV_GENERAL_QUANTIZATION_VERSION, + LLM_KV_GENERAL_ALIGNMENT, + LLM_KV_GENERAL_FILE_TYPE, + LLM_KV_GENERAL_SAMPLING_SEQUENCE, + LLM_KV_GENERAL_SAMPLING_TOP_K, + LLM_KV_GENERAL_SAMPLING_TOP_P, + LLM_KV_GENERAL_SAMPLING_MIN_P, + LLM_KV_GENERAL_SAMPLING_XTC_PROBABILITY, + LLM_KV_GENERAL_SAMPLING_XTC_THRESHOLD, + LLM_KV_GENERAL_SAMPLING_TEMP, + LLM_KV_GENERAL_SAMPLING_PENALTY_LAST_N, + LLM_KV_GENERAL_SAMPLING_PENALTY_REPEAT, + LLM_KV_GENERAL_SAMPLING_MIROSTAT, + LLM_KV_GENERAL_SAMPLING_MIROSTAT_TAU, + LLM_KV_GENERAL_SAMPLING_MIROSTAT_ETA, + LLM_KV_GENERAL_NAME, + LLM_KV_GENERAL_AUTHOR, + LLM_KV_GENERAL_VERSION, + LLM_KV_GENERAL_URL, + LLM_KV_GENERAL_DESCRIPTION, + LLM_KV_GENERAL_LICENSE, + LLM_KV_GENERAL_SOURCE_URL, + LLM_KV_GENERAL_SOURCE_HF_REPO, + + LLM_KV_VOCAB_SIZE, + LLM_KV_CONTEXT_LENGTH, + LLM_KV_EMBEDDING_LENGTH, + LLM_KV_EMBEDDING_LENGTH_OUT, + LLM_KV_FEATURES_LENGTH, + LLM_KV_BLOCK_COUNT, + LLM_KV_LEADING_DENSE_BLOCK_COUNT, + LLM_KV_FEED_FORWARD_LENGTH, + LLM_KV_EXPERT_FEED_FORWARD_LENGTH, + LLM_KV_VISION_EXPERT_FEED_FORWARD_LENGTH, + LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, + LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, + LLM_KV_SWIGLU_CLAMP_EXP, + LLM_KV_SWIGLU_CLAMP_SHEXP, + LLM_KV_USE_PARALLEL_RESIDUAL, + LLM_KV_TENSOR_DATA_LAYOUT, + LLM_KV_EXPERT_COUNT, + LLM_KV_EXPERT_USED_COUNT, + LLM_KV_EXPERT_SHARED_COUNT, + LLM_KV_EXPERT_GROUP_COUNT, + LLM_KV_EXPERT_GROUP_USED_COUNT, + LLM_KV_EXPERT_WEIGHTS_SCALE, + LLM_KV_EXPERT_WEIGHTS_NORM, + LLM_KV_EXPERT_GATING_FUNC, + LLM_KV_EXPERT_GROUP_SCALE, + LLM_KV_EXPERTS_PER_GROUP, + LLM_KV_MOE_EVERY_N_LAYERS, + LLM_KV_NEXTN_PREDICT_LAYERS, + LLM_KV_NUM_DEEPSTACK_LAYERS, + LLM_KV_POOLING_TYPE, + LLM_KV_LOGIT_SCALE, + LLM_KV_DECODER_START_TOKEN_ID, + LLM_KV_DECODER_BLOCK_COUNT, + LLM_KV_ATTN_LOGIT_SOFTCAPPING, + LLM_KV_ROUTER_LOGIT_SOFTCAPPING, + LLM_KV_FINAL_LOGIT_SOFTCAPPING, + LLM_KV_SWIN_NORM, + LLM_KV_RESCALE_EVERY_N_LAYERS, + LLM_KV_TIME_MIX_EXTRA_DIM, + LLM_KV_TIME_DECAY_EXTRA_DIM, + LLM_KV_RESIDUAL_SCALE, + LLM_KV_EMBEDDING_SCALE, + LLM_KV_TOKEN_SHIFT_COUNT, + LLM_KV_INTERLEAVE_MOE_LAYER_STEP, + + LLM_KV_ATTENTION_HEAD_COUNT, + LLM_KV_ATTENTION_HEAD_COUNT_KV, + LLM_KV_ATTENTION_MAX_ALIBI_BIAS, + LLM_KV_ATTENTION_CLAMP_KQV, + LLM_KV_ATTENTION_KEY_LENGTH, + LLM_KV_ATTENTION_VALUE_LENGTH, + LLM_KV_ATTENTION_LAYERNORM_EPS, + LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, + LLM_KV_ATTENTION_GROUPNORM_EPS, + LLM_KV_ATTENTION_GROUPNORM_GROUPS, + LLM_KV_ATTENTION_CAUSAL, + LLM_KV_ATTENTION_Q_LORA_RANK, + LLM_KV_ATTENTION_KV_LORA_RANK, + LLM_KV_ATTENTION_DECAY_LORA_RANK, + LLM_KV_ATTENTION_ICLR_LORA_RANK, + LLM_KV_ATTENTION_VALUE_RESIDUAL_MIX_LORA_RANK, + LLM_KV_ATTENTION_GATE_LORA_RANK, + LLM_KV_ATTENTION_RELATIVE_BUCKETS_COUNT, + LLM_KV_ATTENTION_SLIDING_WINDOW, + LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, + LLM_KV_ATTENTION_SCALE, + LLM_KV_ATTENTION_OUTPUT_SCALE, + LLM_KV_ATTENTION_TEMPERATURE_LENGTH, + LLM_KV_ATTENTION_TEMPERATURE_SCALE, + LLM_KV_ATTENTION_KEY_LENGTH_MLA, + LLM_KV_ATTENTION_VALUE_LENGTH_MLA, + + LLM_KV_ROPE_DIMENSION_COUNT, + LLM_KV_ROPE_DIMENSION_SECTIONS, + LLM_KV_ROPE_FREQ_BASE, + LLM_KV_ROPE_FREQ_BASE_SWA, + LLM_KV_ROPE_SCALE_LINEAR, + LLM_KV_ROPE_SCALING_TYPE, + LLM_KV_ROPE_SCALING_FACTOR, + LLM_KV_ROPE_SCALING_ATTN_FACTOR, + LLM_KV_ROPE_SCALING_ORIG_CTX_LEN, + LLM_KV_ROPE_SCALING_FINETUNED, + LLM_KV_ROPE_SCALING_YARN_LOG_MUL, + LLM_KV_ROPE_SCALING_YARN_EXT_FACTOR, + LLM_KV_ROPE_SCALING_YARN_ATTN_FACTOR, + LLM_KV_ROPE_SCALING_YARN_BETA_FAST, + LLM_KV_ROPE_SCALING_YARN_BETA_SLOW, + + LLM_KV_SPLIT_NO, + LLM_KV_SPLIT_COUNT, + LLM_KV_SPLIT_TENSORS_COUNT, + + LLM_KV_SSM_INNER_SIZE, + LLM_KV_SSM_CONV_KERNEL, + LLM_KV_SSM_STATE_SIZE, + LLM_KV_SSM_TIME_STEP_RANK, + LLM_KV_SSM_GROUP_COUNT, + LLM_KV_SSM_DT_B_C_RMS, + + LLM_KV_KDA_HEAD_DIM, + + LLM_KV_WKV_HEAD_SIZE, + + LLM_KV_TOKENIZER_MODEL, + LLM_KV_TOKENIZER_PRE, + LLM_KV_TOKENIZER_LIST, + LLM_KV_TOKENIZER_TOKEN_TYPE, + LLM_KV_TOKENIZER_TOKEN_TYPE_COUNT, + LLM_KV_TOKENIZER_SCORES, + LLM_KV_TOKENIZER_MERGES, + LLM_KV_TOKENIZER_BOS_ID, + LLM_KV_TOKENIZER_EOS_ID, + LLM_KV_TOKENIZER_EOT_ID, + LLM_KV_TOKENIZER_EOM_ID, + LLM_KV_TOKENIZER_UNK_ID, + LLM_KV_TOKENIZER_SEP_ID, + LLM_KV_TOKENIZER_PAD_ID, + LLM_KV_TOKENIZER_CLS_ID, + LLM_KV_TOKENIZER_MASK_ID, + LLM_KV_TOKENIZER_ADD_BOS, + LLM_KV_TOKENIZER_ADD_EOS, + LLM_KV_TOKENIZER_ADD_SEP, + LLM_KV_TOKENIZER_ADD_PREFIX, + LLM_KV_TOKENIZER_REMOVE_EXTRA_WS, + LLM_KV_TOKENIZER_PRECOMPILED_CHARSMAP, + LLM_KV_TOKENIZER_HF_JSON, + LLM_KV_TOKENIZER_RWKV, + LLM_KV_TOKENIZER_CHAT_TEMPLATE, + LLM_KV_TOKENIZER_FIM_PRE_ID, + LLM_KV_TOKENIZER_FIM_SUF_ID, + LLM_KV_TOKENIZER_FIM_MID_ID, + LLM_KV_TOKENIZER_FIM_PAD_ID, + LLM_KV_TOKENIZER_FIM_REP_ID, + LLM_KV_TOKENIZER_FIM_SEP_ID, + + LLM_KV_ADAPTER_TYPE, + LLM_KV_ADAPTER_LORA_ALPHA, + LLM_KV_ADAPTER_LORA_TASK_NAME, + LLM_KV_ADAPTER_LORA_PROMPT_PREFIX, + LLM_KV_ADAPTER_ALORA_INVOCATION_TOKENS, + + LLM_KV_POSNET_EMBEDDING_LENGTH, + LLM_KV_POSNET_BLOCK_COUNT, + + LLM_KV_CONVNEXT_EMBEDDING_LENGTH, + LLM_KV_CONVNEXT_BLOCK_COUNT, + + LLM_KV_CLASSIFIER_OUTPUT_LABELS, + + LLM_KV_SHORTCONV_L_CACHE, + + LLM_KV_XIELU_ALPHA_N, + LLM_KV_XIELU_ALPHA_P, + LLM_KV_XIELU_BETA, + LLM_KV_XIELU_EPS, + + // deprecated: + LLM_KV_TOKENIZER_PREFIX_ID, + LLM_KV_TOKENIZER_SUFFIX_ID, + LLM_KV_TOKENIZER_MIDDLE_ID, + + // sentence-transformers dense layers in and out features + LLM_KV_DENSE_2_FEAT_IN, + LLM_KV_DENSE_2_FEAT_OUT, + LLM_KV_DENSE_3_FEAT_IN, + LLM_KV_DENSE_3_FEAT_OUT, }; -struct llama_layer_nextn { - struct ggml_tensor * eh_proj = nullptr; - struct ggml_tensor * embed_tokens = nullptr; - struct ggml_tensor * enorm = nullptr; - struct ggml_tensor * hnorm = nullptr; - struct ggml_tensor * shared_head_head = nullptr; - struct ggml_tensor * shared_head_norm = nullptr; +enum llm_tensor { + LLM_TENSOR_TOKEN_EMBD, + LLM_TENSOR_TOKEN_EMBD_NORM, + LLM_TENSOR_TOKEN_TYPES, + LLM_TENSOR_POS_EMBD, + LLM_TENSOR_DENSE_2_OUT, + LLM_TENSOR_DENSE_3_OUT, + LLM_TENSOR_OUTPUT, + LLM_TENSOR_OUTPUT_NORM, + LLM_TENSOR_OUTPUT_NORM_LFM2, // fix for wrong tensor name + LLM_TENSOR_ROPE_FREQS, + LLM_TENSOR_ROPE_FACTORS_LONG, + LLM_TENSOR_ROPE_FACTORS_SHORT, + LLM_TENSOR_ATTN_Q, + LLM_TENSOR_ATTN_K, + LLM_TENSOR_ATTN_V, + LLM_TENSOR_ATTN_QKV, + LLM_TENSOR_ATTN_OUT, + LLM_TENSOR_ATTN_NORM, + LLM_TENSOR_ATTN_NORM_2, + LLM_TENSOR_ATTN_OUT_NORM, + LLM_TENSOR_ATTN_POST_NORM, + LLM_TENSOR_ATTN_ROT_EMBD, + LLM_TENSOR_ATTN_SINKS, + LLM_TENSOR_ATTN_GATE, + LLM_TENSOR_FFN_GATE_INP, + LLM_TENSOR_V_FFN_GATE_INP, + LLM_TENSOR_FFN_GATE_INP_SHEXP, + LLM_TENSOR_FFN_NORM, + LLM_TENSOR_FFN_POST_NORM, + LLM_TENSOR_FFN_GATE, + LLM_TENSOR_FFN_DOWN, + LLM_TENSOR_FFN_UP, + LLM_TENSOR_FFN_ACT, + LLM_TENSOR_FFN_DOWN_EXP, // split experts for backward compatibility + LLM_TENSOR_FFN_GATE_EXP, + LLM_TENSOR_FFN_UP_EXP, + LLM_TENSOR_FFN_NORM_EXPS, + LLM_TENSOR_FFN_DOWN_EXPS, // merged experts + LLM_TENSOR_V_FFN_DOWN_EXPS, // merged experts + LLM_TENSOR_FFN_GATE_EXPS, + LLM_TENSOR_V_FFN_GATE_EXPS, + LLM_TENSOR_FFN_UP_EXPS, + LLM_TENSOR_V_FFN_UP_EXPS, + LLM_TENSOR_FFN_DOWN_SHEXP, + LLM_TENSOR_FFN_GATE_SHEXP, + LLM_TENSOR_FFN_UP_SHEXP, + LLM_TENSOR_FFN_DOWN_CHEXPS, + LLM_TENSOR_FFN_GATE_CHEXPS, + LLM_TENSOR_FFN_UP_CHEXPS, + LLM_TENSOR_FFN_EXP_PROBS_B, + LLM_TENSOR_V_FFN_EXP_PROBS_B, + LLM_TENSOR_ATTN_Q_NORM, + LLM_TENSOR_ATTN_K_NORM, + LLM_TENSOR_LAYER_OUT_NORM, + LLM_TENSOR_POST_ATTN_NORM, + LLM_TENSOR_POST_MLP_NORM, + LLM_TENSOR_PER_LAYER_TOKEN_EMBD, // gemma3n + LLM_TENSOR_PER_LAYER_MODEL_PROJ, // gemma3n + LLM_TENSOR_PER_LAYER_INP_GATE, // gemma3n + LLM_TENSOR_PER_LAYER_PROJ, // gemma3n + LLM_TENSOR_PER_LAYER_PROJ_NORM, // gemma3n + LLM_TENSOR_PER_LAYER_POST_NORM, // gemma3n + LLM_TENSOR_ALTUP_PROJ, // gemma3n + LLM_TENSOR_ALTUP_UNEMBD_PROJ, // gemma3n + LLM_TENSOR_ALTUP_CORRECT_COEF, // gemma3n + LLM_TENSOR_ALTUP_CORRECT_SCALE, // gemma3n + LLM_TENSOR_ALTUP_PREDICT_COEF, // gemma3n + LLM_TENSOR_ALTUP_ROUTER, // gemma3n + LLM_TENSOR_ALTUP_ROUTER_NORM, // gemma3n + LLM_TENSOR_LAUREL_L, // gemma3n + LLM_TENSOR_LAUREL_R, // gemma3n + LLM_TENSOR_LAUREL_POST_NORM, // gemma3n + LLM_TENSOR_SSM_IN, + LLM_TENSOR_SSM_CONV1D, + LLM_TENSOR_SSM_X, + LLM_TENSOR_SSM_DT, + LLM_TENSOR_SSM_DT_NORM, + LLM_TENSOR_SSM_A, + LLM_TENSOR_SSM_A_NOSCAN, // qwen3next special case with MUL instead of SSM_SCAN + LLM_TENSOR_SSM_B_NORM, + LLM_TENSOR_SSM_C_NORM, + LLM_TENSOR_SSM_D, + LLM_TENSOR_SSM_NORM, + LLM_TENSOR_SSM_OUT, + LLM_TENSOR_SSM_BETA_ALPHA, // qwen3next + // Kimi Linear KDA (using SSM_ prefix for consistency) + LLM_TENSOR_SSM_CONV1D_Q, // kimi: Q conv1d weight + LLM_TENSOR_SSM_CONV1D_K, // kimi: K conv1d weight + LLM_TENSOR_SSM_CONV1D_V, // kimi: V conv1d weight + LLM_TENSOR_SSM_F_A, // kimi: forget gate projection A + LLM_TENSOR_SSM_F_B, // kimi: forget gate projection B + LLM_TENSOR_SSM_BETA, // kimi: beta mixing coefficient + LLM_TENSOR_SSM_G_A, // kimi: output gate projection A + LLM_TENSOR_SSM_G_B, // kimi: output gate projection B + LLM_TENSOR_TIME_MIX_W0, + LLM_TENSOR_TIME_MIX_W1, + LLM_TENSOR_TIME_MIX_W2, + LLM_TENSOR_TIME_MIX_A0, + LLM_TENSOR_TIME_MIX_A1, + LLM_TENSOR_TIME_MIX_A2, + LLM_TENSOR_TIME_MIX_V0, + LLM_TENSOR_TIME_MIX_V1, + LLM_TENSOR_TIME_MIX_V2, + LLM_TENSOR_TIME_MIX_G1, + LLM_TENSOR_TIME_MIX_G2, + LLM_TENSOR_TIME_MIX_K_K, + LLM_TENSOR_TIME_MIX_K_A, + LLM_TENSOR_TIME_MIX_R_K, + LLM_TENSOR_TIME_MIX_LERP_X, + LLM_TENSOR_TIME_MIX_LERP_W, + LLM_TENSOR_TIME_MIX_LERP_K, + LLM_TENSOR_TIME_MIX_LERP_V, + LLM_TENSOR_TIME_MIX_LERP_R, + LLM_TENSOR_TIME_MIX_LERP_G, + LLM_TENSOR_TIME_MIX_LERP_FUSED, + LLM_TENSOR_TIME_MIX_FIRST, + LLM_TENSOR_TIME_MIX_DECAY, + LLM_TENSOR_TIME_MIX_DECAY_W1, + LLM_TENSOR_TIME_MIX_DECAY_W2, + LLM_TENSOR_TIME_MIX_KEY, + LLM_TENSOR_TIME_MIX_VALUE, + LLM_TENSOR_TIME_MIX_RECEPTANCE, + LLM_TENSOR_TIME_MIX_GATE, + LLM_TENSOR_TIME_MIX_LN, + LLM_TENSOR_TIME_MIX_OUTPUT, + LLM_TENSOR_CHANNEL_MIX_LERP_K, + LLM_TENSOR_CHANNEL_MIX_LERP_R, + LLM_TENSOR_CHANNEL_MIX_KEY, + LLM_TENSOR_CHANNEL_MIX_RECEPTANCE, + LLM_TENSOR_CHANNEL_MIX_VALUE, + LLM_TENSOR_ATTN_Q_A, + LLM_TENSOR_ATTN_Q_B, + LLM_TENSOR_ATTN_KV_A_MQA, + LLM_TENSOR_ATTN_KV_B, + LLM_TENSOR_ATTN_K_B, + LLM_TENSOR_ATTN_V_B, + LLM_TENSOR_ATTN_Q_A_NORM, + LLM_TENSOR_ATTN_KV_A_NORM, + LLM_TENSOR_ATTN_SUB_NORM, + LLM_TENSOR_FFN_SUB_NORM, + LLM_TENSOR_DEC_ATTN_NORM, + LLM_TENSOR_DEC_ATTN_Q, + LLM_TENSOR_DEC_ATTN_K, + LLM_TENSOR_DEC_ATTN_V, + LLM_TENSOR_DEC_ATTN_OUT, + LLM_TENSOR_DEC_ATTN_REL_B, + LLM_TENSOR_DEC_CROSS_ATTN_NORM, + LLM_TENSOR_DEC_CROSS_ATTN_Q, + LLM_TENSOR_DEC_CROSS_ATTN_K, + LLM_TENSOR_DEC_CROSS_ATTN_V, + LLM_TENSOR_DEC_CROSS_ATTN_OUT, + LLM_TENSOR_DEC_CROSS_ATTN_REL_B, + LLM_TENSOR_DEC_FFN_NORM, + LLM_TENSOR_DEC_FFN_GATE, + LLM_TENSOR_DEC_FFN_DOWN, + LLM_TENSOR_DEC_FFN_UP, + LLM_TENSOR_DEC_OUTPUT_NORM, + LLM_TENSOR_ENC_ATTN_NORM, + LLM_TENSOR_ENC_ATTN_Q, + LLM_TENSOR_ENC_ATTN_K, + LLM_TENSOR_ENC_ATTN_V, + LLM_TENSOR_ENC_ATTN_OUT, + LLM_TENSOR_ENC_ATTN_REL_B, + LLM_TENSOR_ENC_FFN_NORM, + LLM_TENSOR_ENC_FFN_GATE, + LLM_TENSOR_ENC_FFN_DOWN, + LLM_TENSOR_ENC_FFN_UP, + LLM_TENSOR_ENC_OUTPUT_NORM, + LLM_TENSOR_CLS, + LLM_TENSOR_CLS_OUT, + LLM_TENSOR_CONV1D, + LLM_TENSOR_CONVNEXT_DW, + LLM_TENSOR_CONVNEXT_NORM, + LLM_TENSOR_CONVNEXT_PW1, + LLM_TENSOR_CONVNEXT_PW2, + LLM_TENSOR_CONVNEXT_GAMMA, + LLM_TENSOR_POS_NET_CONV1, + LLM_TENSOR_POS_NET_CONV2, + LLM_TENSOR_POS_NET_NORM, + LLM_TENSOR_POS_NET_NORM1, + LLM_TENSOR_POS_NET_NORM2, + LLM_TENSOR_POS_NET_ATTN_NORM, + LLM_TENSOR_POS_NET_ATTN_Q, + LLM_TENSOR_POS_NET_ATTN_K, + LLM_TENSOR_POS_NET_ATTN_V, + LLM_TENSOR_POS_NET_ATTN_OUT, + LLM_TENSOR_SHORTCONV_CONV, + LLM_TENSOR_SHORTCONV_INPROJ, + LLM_TENSOR_SHORTCONV_OUTPROJ, + LLM_TENSOR_VISEXP_ATTN_QKV, + LLM_TENSOR_VISEXP_ATTN_OUT, + LLM_TENSOR_VISEXP_FFN_GATE, + LLM_TENSOR_VISEXP_FFN_DOWN, + LLM_TENSOR_VISEXP_FFN_UP, + LLM_TENSOR_NEXTN_EH_PROJ, + LLM_TENSOR_NEXTN_EMBED_TOKENS, + LLM_TENSOR_NEXTN_ENORM, + LLM_TENSOR_NEXTN_HNORM, + LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, + LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, }; -struct llama_layer { - // normalization - struct ggml_tensor * attn_norm = nullptr; - struct ggml_tensor * attn_norm_b = nullptr; - struct ggml_tensor * attn_norm_2 = nullptr; - struct ggml_tensor * attn_norm_2_b = nullptr; - struct ggml_tensor * attn_q_norm = nullptr; - struct ggml_tensor * attn_q_norm_b = nullptr; - struct ggml_tensor * attn_k_norm = nullptr; - struct ggml_tensor * attn_k_norm_b = nullptr; - struct ggml_tensor * attn_out_norm = nullptr; - struct ggml_tensor * attn_out_norm_b = nullptr; - struct ggml_tensor * attn_q_a_norm = nullptr; - struct ggml_tensor * attn_kv_a_norm = nullptr; - struct ggml_tensor * attn_sub_norm = nullptr; - struct ggml_tensor * attn_post_norm = nullptr; - struct ggml_tensor * ffn_sub_norm = nullptr; - struct ggml_tensor * attn_norm_cross = nullptr; - struct ggml_tensor * attn_norm_enc = nullptr; - struct ggml_tensor * ssm_norm = nullptr; - struct ggml_tensor * ssm_dt_norm = nullptr; - struct ggml_tensor * ssm_b_norm = nullptr; - struct ggml_tensor * ssm_c_norm = nullptr; - - // attention - struct ggml_tensor * wq = nullptr; - struct ggml_tensor * wk = nullptr; - struct ggml_tensor * wv = nullptr; - struct ggml_tensor * wo = nullptr; - struct ggml_tensor * wqkv = nullptr; - struct ggml_tensor * wq_a = nullptr; - struct ggml_tensor * wq_b = nullptr; - struct ggml_tensor * wkv_a_mqa = nullptr; - struct ggml_tensor * wkv_b = nullptr; - struct ggml_tensor * wk_b = nullptr; - struct ggml_tensor * wv_b = nullptr; - struct ggml_tensor * wq_cross = nullptr; - struct ggml_tensor * wk_cross = nullptr; - struct ggml_tensor * wv_cross = nullptr; - struct ggml_tensor * wo_cross = nullptr; - struct ggml_tensor * wq_enc = nullptr; - struct ggml_tensor * wk_enc = nullptr; - struct ggml_tensor * wv_enc = nullptr; - struct ggml_tensor * wo_enc = nullptr; - struct ggml_tensor * wqkv_gate = nullptr; - - // attention bias - struct ggml_tensor * bq = nullptr; - struct ggml_tensor * bk = nullptr; - struct ggml_tensor * bv = nullptr; - struct ggml_tensor * bo = nullptr; - struct ggml_tensor * bqkv = nullptr; - - // relative position bias - struct ggml_tensor * attn_rel_b = nullptr; - struct ggml_tensor * attn_rel_b_enc = nullptr; - struct ggml_tensor * attn_rel_b_cross = nullptr; - - // normalization - struct ggml_tensor * ffn_norm = nullptr; - struct ggml_tensor * ffn_norm_b = nullptr; - struct ggml_tensor * ffn_post_norm = nullptr; - struct ggml_tensor * layer_out_norm = nullptr; - struct ggml_tensor * layer_out_norm_b = nullptr; - struct ggml_tensor * ffn_norm_exps = nullptr; - struct ggml_tensor * ffn_norm_enc = nullptr; - - // ff - struct ggml_tensor * ffn_gate = nullptr; // w1 - struct ggml_tensor * ffn_down = nullptr; // w2 - struct ggml_tensor * ffn_up = nullptr; // w3 - struct ggml_tensor * ffn_gate_enc = nullptr; - struct ggml_tensor * ffn_down_enc = nullptr; - struct ggml_tensor * ffn_up_enc = nullptr; - - // ff MoE - struct ggml_tensor * ffn_gate_inp = nullptr; - struct ggml_tensor * ffn_gate_exps = nullptr; - struct ggml_tensor * ffn_down_exps = nullptr; - struct ggml_tensor * ffn_up_exps = nullptr; - struct ggml_tensor * ffn_gate_inp_b = nullptr; - struct ggml_tensor * ffn_gate_exps_b = nullptr; - struct ggml_tensor * ffn_down_exps_b = nullptr; - struct ggml_tensor * ffn_up_exps_b = nullptr; - - // ff shared expert (shexp) - struct ggml_tensor * ffn_gate_inp_shexp = nullptr; - struct ggml_tensor * ffn_gate_shexp = nullptr; - struct ggml_tensor * ffn_down_shexp = nullptr; - struct ggml_tensor * ffn_up_shexp = nullptr; - - // ff adjugate experts (chexps) - struct ggml_tensor * ffn_gate_chexps = nullptr; - struct ggml_tensor * ffn_down_chexps = nullptr; - struct ggml_tensor * ffn_up_chexps = nullptr; - - // ff bias - struct ggml_tensor * ffn_gate_b = nullptr; - struct ggml_tensor * ffn_down_b = nullptr; // b2 - 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 * v_ffn_exp_probs_b = nullptr; - - // mamba proj - struct ggml_tensor * ssm_in = nullptr; - struct ggml_tensor * ssm_x = nullptr; - struct ggml_tensor * ssm_dt = nullptr; - struct ggml_tensor * ssm_out = nullptr; - - // mamba - struct ggml_tensor * ssm_conv1d = nullptr; - struct ggml_tensor * ssm_a = nullptr; - struct ggml_tensor * ssm_d = nullptr; - - // mamba bias - struct ggml_tensor * ssm_conv1d_b = nullptr; - struct ggml_tensor * ssm_dt_b = nullptr; - - // qwen3next - struct ggml_tensor * ssm_beta_alpha = nullptr; - - // rwkv - struct ggml_tensor * time_mix_w1 = nullptr; - struct ggml_tensor * time_mix_w2 = nullptr; - struct ggml_tensor * time_mix_lerp_x = nullptr; - struct ggml_tensor * time_mix_lerp_w = nullptr; - struct ggml_tensor * time_mix_lerp_k = nullptr; - struct ggml_tensor * time_mix_lerp_v = nullptr; - struct ggml_tensor * time_mix_lerp_r = nullptr; - struct ggml_tensor * time_mix_lerp_g = nullptr; - struct ggml_tensor * time_mix_lerp_fused = nullptr; - - struct ggml_tensor * time_mix_first = nullptr; - struct ggml_tensor * time_mix_decay = nullptr; - struct ggml_tensor * time_mix_decay_w1 = nullptr; - struct ggml_tensor * time_mix_decay_w2 = nullptr; - struct ggml_tensor * time_mix_key = nullptr; - struct ggml_tensor * time_mix_key_b = nullptr; - struct ggml_tensor * time_mix_value = nullptr; - struct ggml_tensor * time_mix_value_b = nullptr; - struct ggml_tensor * time_mix_receptance = nullptr; - struct ggml_tensor * time_mix_receptance_b = nullptr; - struct ggml_tensor * time_mix_gate = nullptr; - - // rwkv7 - struct ggml_tensor * time_mix_w0 = nullptr; - struct ggml_tensor * time_mix_a0 = nullptr; - struct ggml_tensor * time_mix_a1 = nullptr; - struct ggml_tensor * time_mix_a2 = nullptr; - struct ggml_tensor * time_mix_v0 = nullptr; - struct ggml_tensor * time_mix_v1 = nullptr; - struct ggml_tensor * time_mix_v2 = nullptr; - struct ggml_tensor * time_mix_g1 = nullptr; - struct ggml_tensor * time_mix_g2 = nullptr; - struct ggml_tensor * time_mix_k_k = nullptr; - struct ggml_tensor * time_mix_k_a = nullptr; - struct ggml_tensor * time_mix_r_k = nullptr; - - struct ggml_tensor * time_mix_ln = nullptr; - struct ggml_tensor * time_mix_ln_b = nullptr; - struct ggml_tensor * time_mix_output = nullptr; - - struct ggml_tensor * channel_mix_lerp_k = nullptr; - struct ggml_tensor * channel_mix_lerp_r = nullptr; - - struct ggml_tensor * channel_mix_key = nullptr; - struct ggml_tensor * channel_mix_receptance = nullptr; - struct ggml_tensor * channel_mix_value = nullptr; - - // long rope factors - struct ggml_tensor * rope_long = nullptr; - struct ggml_tensor * rope_short = nullptr; - struct ggml_tensor * rope_freqs = nullptr; - - // bitnet scale - struct ggml_tensor * wq_scale = nullptr; - struct ggml_tensor * wk_scale = nullptr; - struct ggml_tensor * wv_scale = nullptr; - struct ggml_tensor * wo_scale = nullptr; - struct ggml_tensor * ffn_gate_scale = nullptr; - struct ggml_tensor * ffn_up_scale = nullptr; - struct ggml_tensor * ffn_down_scale = nullptr; - - // altup & laurel - struct ggml_tensor * per_layer_inp_gate = nullptr; - struct ggml_tensor * per_layer_proj = nullptr; - struct ggml_tensor * per_layer_post_norm = nullptr; - struct ggml_tensor * altup_correct_coef = nullptr; - struct ggml_tensor * altup_correct_scale = nullptr; - struct ggml_tensor * altup_predict_coef = nullptr; - struct ggml_tensor * altup_router = nullptr; - struct ggml_tensor * altup_router_norm = nullptr; - struct ggml_tensor * laurel_l = nullptr; - struct ggml_tensor * laurel_r = nullptr; - struct ggml_tensor * laurel_post_norm = nullptr; - - // openai-moe - struct ggml_tensor * attn_sinks = nullptr; - - // cogvlm - struct ggml_tensor * visexp_attn_wqkv = nullptr; - struct ggml_tensor * visexp_attn_wo = nullptr; - struct ggml_tensor * visexp_ffn_gate = nullptr; - struct ggml_tensor * visexp_ffn_down = nullptr; - struct ggml_tensor * visexp_ffn_up = nullptr; - - // xIELU activation parameters for Apertus - struct ggml_tensor * ffn_act_alpha_n = nullptr; - struct ggml_tensor * ffn_act_alpha_p = nullptr; - struct ggml_tensor * ffn_act_beta = nullptr; - struct ggml_tensor * ffn_act_eps = nullptr; - - // Kimi Linear KDA (using ssm_ prefix for consistency) - // Note: ssm_dt_b already exists above (mamba bias), reused for Kimi dt_bias - struct ggml_tensor * ssm_q_conv = nullptr; - struct ggml_tensor * ssm_k_conv = nullptr; - struct ggml_tensor * ssm_v_conv = nullptr; - struct ggml_tensor * ssm_f_a = nullptr; - struct ggml_tensor * ssm_f_b = nullptr; - struct ggml_tensor * ssm_beta = nullptr; - struct ggml_tensor * ssm_g_a = nullptr; - struct ggml_tensor * ssm_g_b = nullptr; - struct ggml_tensor * ssm_o_norm = nullptr; - - struct llama_layer_posnet posnet; - - struct llama_layer_convnext convnext; - - struct llama_layer_shortconv shortconv; - - struct llama_layer_nextn nextn; +enum llm_tensor_layer { + LLM_TENSOR_LAYER_INPUT, + LLM_TENSOR_LAYER_REPEATING, + LLM_TENSOR_LAYER_OUTPUT, }; -struct llama_model { - llm_type type = LLM_TYPE_UNKNOWN; - llm_arch arch = LLM_ARCH_UNKNOWN; - - std::string name = "n/a"; - - llama_hparams hparams = {}; - llama_vocab vocab; - - // for classifier models - std::vector classifier_labels; - - struct ggml_tensor * tok_embd = nullptr; - struct ggml_tensor * type_embd = nullptr; - struct ggml_tensor * pos_embd = nullptr; - struct ggml_tensor * tok_norm = nullptr; - struct ggml_tensor * tok_norm_b = nullptr; - - struct ggml_tensor * output_norm = nullptr; - struct ggml_tensor * output_norm_b = nullptr; - struct ggml_tensor * output = nullptr; - struct ggml_tensor * output_b = nullptr; - struct ggml_tensor * output_norm_enc = nullptr; +struct LLM_KV { + LLM_KV(llm_arch arch, const char * suffix = nullptr); - // classifier - struct ggml_tensor * cls = nullptr; - struct ggml_tensor * cls_b = nullptr; - struct ggml_tensor * cls_out = nullptr; - struct ggml_tensor * cls_out_b = nullptr; + llm_arch arch; + const char * suffix; - struct ggml_tensor * conv1d = nullptr; - struct ggml_tensor * conv1d_b = nullptr; - - // gemma3n altup - struct ggml_tensor * tok_embd_per_layer = nullptr; - struct ggml_tensor * altup_proj = nullptr; - struct ggml_tensor * altup_unembd_proj = nullptr; - struct ggml_tensor * per_layer_model_proj = nullptr; - struct ggml_tensor * per_layer_proj_norm = nullptr; - - std::vector layers; - - //Dense linear projections for SentenceTransformers models like embeddinggemma - // For Sentence Transformers models structure see - // https://sbert.net/docs/sentence_transformer/usage/custom_models.html#structure-of-sentence-transformer-models - struct ggml_tensor * dense_2_out_layers = nullptr; - struct ggml_tensor * dense_3_out_layers = nullptr; - - // gguf metadata - std::unordered_map gguf_kv; - - // list of devices used in this model - std::vector devices; - - // for quantize-stats only - std::vector> tensors_by_name; - - // for keeping track of associated LoRA adapters - std::unordered_set loras; - - int64_t t_load_us = 0; - int64_t t_start_us = 0; - - explicit llama_model(const struct llama_model_params & params); - ~llama_model(); - - void load_stats (llama_model_loader & ml); - void load_arch (llama_model_loader & ml); - void load_hparams(llama_model_loader & ml); - void load_vocab (llama_model_loader & ml); - bool load_tensors(llama_model_loader & ml); // returns false if cancelled by progress_callback - - std::string arch_name() const; - std::string type_name() const; - - std::string desc() const; - - size_t size() const; // file size - size_t n_tensors() const; - size_t n_devices() const; - - uint32_t n_gpu_layers() const; - llama_split_mode split_mode() const; - - std::map memory_breakdown() const; - - // total number of parameters in the model - uint64_t n_elements() const; - - void print_info() const; - - ggml_backend_dev_t dev_layer(int il) const; - ggml_backend_dev_t dev_output() const; + std::string operator()(llm_kv kv) const; +}; - ggml_backend_buffer_type_t select_buft(int il) const; +// helper to handle gguf constants +// usage: +// +// const auto tn = LLM_TN(LLM_ARCH_LLAMA); +// +// std::string name = tn(LLM_TENSOR_OUTPUT); -> "output" +// std::string name = tn(LLM_TENSOR_TOKEN_EMBD, "bias"); -> "token_embd.bias" +// std::string name = tn(LLM_TENSOR_ATTN_NORM, "weight", 3); -> "blk.3.attn_norm.weight" +// +struct LLM_TN_IMPL { + const llm_arch arch; + const llm_tensor tensor; + const char * const suffix; + const int bid; + const int xid; + + const std::set model_tensors; + + LLM_TN_IMPL(llm_arch arch, llm_tensor tensor, const char * suffix, int bid, int xid); + + std::string str() const; + + operator std::string() const { + return str(); + } + + friend bool operator==(const std::string & str, const LLM_TN_IMPL & tn) { + return str == tn.str(); + } + + friend bool operator!=(const std::string & str, const LLM_TN_IMPL & tn) { + return str != tn.str(); + } +}; - bool has_tensor_overrides() const; +struct LLM_TN { + LLM_TN(llm_arch arch) : arch(arch) {} - const struct ggml_tensor * get_tensor(const char * name) const; + llm_arch arch; - float get_rope_freq_base (const llama_cparams & cparams, int il) const; - float get_rope_freq_scale(const llama_cparams & cparams, int il) const; + LLM_TN_IMPL operator()(llm_tensor tensor, const char * suffix, int bid = -1, int xid = -1) const { + return LLM_TN_IMPL(arch, tensor, suffix, bid, xid); + } - ggml_tensor * get_rope_factors(const llama_cparams & cparams, int il) const; + LLM_TN_IMPL operator()(llm_tensor tensor, int bid = -1, int xid = -1) const { + return LLM_TN_IMPL(arch, tensor, nullptr, bid, xid); + } +}; - // TODO: move this to new llm_arch_model_i interface - llama_memory_i * create_memory(const llama_memory_params & params, const llama_cparams & cparams) const; - // TODO: move this to new llm_arch_model_i interface - ggml_cgraph * build_graph(const llm_graph_params & params) const; +struct llm_tensor_info { + llm_tensor_layer layer; + ggml_op op; +}; -private: - llama_model_params params; +const char * llm_arch_name(llm_arch arch); - struct impl; - std::unique_ptr pimpl; -}; +llm_arch llm_arch_from_string(const std::string & name); -const char * llm_type_name(llm_type type); +const llm_tensor_info & llm_tensor_info_for(llm_tensor tensor); -// For internal test use -// TODO: remove -const std::vector> & llama_internal_get_tensor_map(const llama_model * model); +bool llm_arch_is_recurrent(const llm_arch & arch); +bool llm_arch_is_hybrid (const llm_arch & arch); +bool llm_arch_is_diffusion(const llm_arch & arch); From 75705b5bbd2bb1c0cac773868459537bd2635097 Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 21:24:16 +0800 Subject: [PATCH 5/8] [model] support ernie4-5-vl-moe model --- src/llama-arch.cpp | 10 ++++++++++ tools/mtmd/clip-impl.h | 2 +- 2 files changed, 11 insertions(+), 1 deletion(-) diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 6dd427274736..ab6f9082bb2d 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -335,6 +335,7 @@ static const std::map LLM_TENSOR_NAMES = { { LLM_TENSOR_ATTN_OUT, "blk.%d.attn_output" }, { LLM_TENSOR_ATTN_ROT_EMBD, "blk.%d.attn_rot_embd" }, { LLM_TENSOR_FFN_GATE_INP, "blk.%d.ffn_gate_inp" }, + { LLM_TENSOR_V_FFN_GATE_INP, "blk.%d.v_ffn_gate_inp" }, { LLM_TENSOR_FFN_NORM, "blk.%d.ffn_norm" }, { LLM_TENSOR_FFN_GATE, "blk.%d.ffn_gate" }, { LLM_TENSOR_FFN_DOWN, "blk.%d.ffn_down" }, @@ -343,8 +344,11 @@ static const std::map LLM_TENSOR_NAMES = { { LLM_TENSOR_FFN_DOWN_EXP, "blk.%d.ffn_down.%d" }, { LLM_TENSOR_FFN_UP_EXP, "blk.%d.ffn_up.%d" }, { LLM_TENSOR_FFN_GATE_EXPS, "blk.%d.ffn_gate_exps" }, + { LLM_TENSOR_V_FFN_GATE_EXPS, "blk.%d.v_ffn_gate_exps" }, { LLM_TENSOR_FFN_DOWN_EXPS, "blk.%d.ffn_down_exps" }, + { LLM_TENSOR_V_FFN_DOWN_EXPS, "blk.%d.v_ffn_down_exps" }, { LLM_TENSOR_FFN_UP_EXPS, "blk.%d.ffn_up_exps" }, + { LLM_TENSOR_V_FFN_UP_EXPS, "blk.%d.v_ffn_up_exps" }, { LLM_TENSOR_ATTN_POST_NORM, "blk.%d.post_attention_norm" }, { LLM_TENSOR_ATTN_Q_NORM, "blk.%d.attn_q_norm" }, { LLM_TENSOR_ATTN_K_NORM, "blk.%d.attn_k_norm" }, @@ -354,6 +358,7 @@ static const std::map 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_V_FFN_EXP_PROBS_B, "blk.%d.v_exp_probs_b" }, { LLM_TENSOR_ATTN_NORM_2, "blk.%d.attn_norm_2" }, { LLM_TENSOR_ATTN_QKV, "blk.%d.attn_qkv" }, { LLM_TENSOR_LAYER_OUT_NORM, "blk.%d.layer_output_norm" }, @@ -2519,6 +2524,10 @@ static const std::map LLM_TENSOR_INFOS = { {LLM_TENSOR_ENC_FFN_UP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, {LLM_TENSOR_FFN_GATE_INP_SHEXP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, {LLM_TENSOR_FFN_GATE_INP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_V_FFN_GATE_INP, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_V_FFN_GATE_EXPS, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_V_FFN_DOWN_EXPS, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, + {LLM_TENSOR_V_FFN_UP_EXPS, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, {LLM_TENSOR_SSM_IN, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, {LLM_TENSOR_SSM_X, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, {LLM_TENSOR_SSM_DT, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_MUL_MAT}}, @@ -2606,6 +2615,7 @@ static const std::map 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_V_FFN_EXP_PROBS_B, {LLM_TENSOR_LAYER_REPEATING, GGML_OP_ADD}}, // altup / laurel (gemma 3n) {LLM_TENSOR_PER_LAYER_TOKEN_EMBD, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_GET_ROWS}}, {LLM_TENSOR_PER_LAYER_MODEL_PROJ, {LLM_TENSOR_LAYER_OUTPUT, GGML_OP_MUL_MAT}}, diff --git a/tools/mtmd/clip-impl.h b/tools/mtmd/clip-impl.h index af54b8be793f..af7508364a4f 100644 --- a/tools/mtmd/clip-impl.h +++ b/tools/mtmd/clip-impl.h @@ -269,7 +269,7 @@ static std::map PROJECTOR_TYPE_NAMES = { { PROJECTOR_TYPE_LFM2A, "lfm2a"}, { PROJECTOR_TYPE_GLM4V, "glm4v"}, { PROJECTOR_TYPE_YOUTUVL, "youtuvl"}, - { PROJECTOR_TYPE_ERNIE45VLMOE, "ernie45vlmoe"}, + { PROJECTOR_TYPE_ERNIE45VLMOE, "ernie4.5vl_moe"}, }; static projector_type clip_projector_type_from_string(const std::string & str) { From 57133bae4171acd9cc1d07f586780744345325bb Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 10 Feb 2026 23:44:06 +0800 Subject: [PATCH 6/8] [model] support ernie4-5-vl-moe model --- tools/mtmd/clip-model.h | 17 ------------ tools/mtmd/clip.cpp | 44 +++++++++++++++--------------- tools/mtmd/models/ernie45vlmoe.cpp | 36 ++++++++++++------------ 3 files changed, 40 insertions(+), 57 deletions(-) diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index c4e54323c77c..3843bf334851 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -364,23 +364,6 @@ struct clip_model { ggml_tensor * mm_norm_pre_b = nullptr; ggml_tensor * mm_norm_mid_w = nullptr; - // ernie4.5-vl-moe - ggml_tensor * mm_spatial_0_w = nullptr; - ggml_tensor * mm_spatial_0_b = nullptr; - ggml_tensor * mm_spatial_2_w = nullptr; - ggml_tensor * mm_spatial_2_b = nullptr; - ggml_tensor * mm_spatial_norm_w = nullptr; - ggml_tensor * mm_spatial_norm_b = nullptr; - ggml_tensor * mm_temp_0_w = nullptr; - ggml_tensor * mm_temp_0_b = nullptr; - ggml_tensor * mm_temp_2_w = nullptr; - ggml_tensor * mm_temp_2_b = nullptr; - ggml_tensor * mm_temp_norm_w = nullptr; - ggml_tensor * mm_temp_norm_b = nullptr; - ggml_tensor * mm_mlp_w = nullptr; - ggml_tensor * mm_mlp_b = nullptr; - ggml_tensor * mm_after_norm_w = nullptr; - // cogvlm ggml_tensor * mm_post_fc_norm_w = nullptr; ggml_tensor * mm_post_fc_norm_b = nullptr; diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index 3b2f0462cad2..c7ee83b101fe 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -1148,7 +1148,7 @@ struct clip_model_loader { hparams.n_merge = 2; hparams.spatial_conv_size = 2; hparams.temporal_conv_size = 2; - hparams.use_temporal_conv = model.mm_temp_0_w != nullptr; + hparams.use_temporal_conv = model.mm_1_w != nullptr; hparams.set_limit_image_tokens(8, 1024); hparams.set_warmup_n_tokens(256); // avoid OOM on warmup } break; @@ -1846,26 +1846,26 @@ struct clip_model_loader { } break; case PROJECTOR_TYPE_ERNIE45VLMOE: { - // spatial path - model.mm_spatial_0_w = get_tensor("mm.0.weight"); - model.mm_spatial_0_b = get_tensor("mm.0.bias"); - model.mm_spatial_2_w = get_tensor("mm.2.weight"); - model.mm_spatial_2_b = get_tensor("mm.2.bias"); - model.mm_spatial_norm_w = get_tensor("mm.3.weight"); - model.mm_spatial_norm_b = get_tensor("mm.3.bias", false); - - // temporal path (optional, not used for single images) - model.mm_temp_0_w = get_tensor("mm_temp.0.weight", false); - model.mm_temp_0_b = get_tensor("mm_temp.0.bias", false); - model.mm_temp_2_w = get_tensor("mm_temp.2.weight", false); - model.mm_temp_2_b = get_tensor("mm_temp.2.bias", false); - model.mm_temp_norm_w = get_tensor("mm_temp.3.weight", false); - model.mm_temp_norm_b = get_tensor("mm_temp.3.bias", false); - - // output - model.mm_mlp_w = get_tensor("mm.mlp.weight"); - model.mm_mlp_b = get_tensor("mm.mlp.bias"); - model.mm_after_norm_w = get_tensor("mm.norm.weight"); + // spatial path: Linear -> GELU -> Linear -> LayerNorm + model.mm_0_w = get_tensor(string_format(TN_LLAVA_PROJ, 0, "weight")); + model.mm_0_b = get_tensor(string_format(TN_LLAVA_PROJ, 0, "bias")); + model.mm_2_w = get_tensor(string_format(TN_LLAVA_PROJ, 2, "weight")); + model.mm_2_b = get_tensor(string_format(TN_LLAVA_PROJ, 2, "bias")); + model.mm_post_norm_w = get_tensor(string_format(TN_LLAVA_PROJ, 3, "weight")); + model.mm_post_norm_b = get_tensor(string_format(TN_LLAVA_PROJ, 3, "bias"), false); + + // temporal path: Linear -> GELU -> Linear -> LayerNorm (optional, not used for single images) + model.mm_1_w = get_tensor("mm_temp.0.weight", false); + model.mm_1_b = get_tensor("mm_temp.0.bias", false); + model.mm_3_w = get_tensor("mm_temp.2.weight", false); + model.mm_3_b = get_tensor("mm_temp.2.bias", false); + model.mm_input_norm_w = get_tensor("mm_temp.3.weight", false); + model.mm_input_norm_b = get_tensor("mm_temp.3.bias", false); + + // output MLP + RMS norm + model.mm_fc_w = get_tensor("mm.mlp.weight"); + model.mm_fc_b = get_tensor("mm.mlp.bias"); + model.mm_norm_mid_w = get_tensor("mm.norm.weight"); } break; default: GGML_ASSERT(false && "unknown projector type"); @@ -3858,7 +3858,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { case PROJECTOR_TYPE_GLM4V: return ctx->model.mm_ffn_down_w->ne[1]; case PROJECTOR_TYPE_ERNIE45VLMOE: - return ctx->model.mm_mlp_w->ne[0]; + return ctx->model.mm_fc_w->ne[0]; default: GGML_ABORT("Unknown projector type"); } diff --git a/tools/mtmd/models/ernie45vlmoe.cpp b/tools/mtmd/models/ernie45vlmoe.cpp index cf2a3a9fbaa1..03182537d689 100644 --- a/tools/mtmd/models/ernie45vlmoe.cpp +++ b/tools/mtmd/models/ernie45vlmoe.cpp @@ -152,9 +152,9 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { ggml_tensor * spatial_out = embeddings; // First linear - ggml_tensor * mm_spatial_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_spatial_0_w)); - spatial_out = ggml_mul_mat(ctx0, mm_spatial_0_w, spatial_out); - spatial_out = ggml_add(ctx0, spatial_out, model.mm_spatial_0_b); + ggml_tensor * spatial_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_0_w)); + spatial_out = ggml_mul_mat(ctx0, spatial_0_w, spatial_out); + spatial_out = ggml_add(ctx0, spatial_out, model.mm_0_b); cb(spatial_out, "spatial_linear_0", -1); // GELU @@ -162,13 +162,13 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { cb(spatial_out, "spatial_gelu", -1); // Second linear - ggml_tensor * mm_spatial_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_spatial_2_w)); - spatial_out = ggml_mul_mat(ctx0, mm_spatial_2_w, spatial_out); - spatial_out = ggml_add(ctx0, spatial_out, model.mm_spatial_2_b); + ggml_tensor * spatial_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_2_w)); + spatial_out = ggml_mul_mat(ctx0, spatial_2_w, spatial_out); + spatial_out = ggml_add(ctx0, spatial_out, model.mm_2_b); cb(spatial_out, "spatial_linear_2", -1); // LayerNorm - spatial_out = build_norm(spatial_out, model.mm_spatial_norm_w, model.mm_spatial_norm_b, NORM_TYPE_NORMAL, eps, -1); + spatial_out = build_norm(spatial_out, model.mm_post_norm_w, model.mm_post_norm_b, NORM_TYPE_NORMAL, eps, -1); cb(spatial_out, "spatial_norm", -1); ggml_tensor * resampler_out = spatial_out; @@ -183,9 +183,9 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { // Weights were transposed (.t()) during GGUF conversion, undo with ggml_transpose // First temporal linear - ggml_tensor * mm_temp_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_temp_0_w)); - resampler_out = ggml_mul_mat(ctx0, mm_temp_0_w, resampler_out); - resampler_out = ggml_add(ctx0, resampler_out, model.mm_temp_0_b); + ggml_tensor * temp_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_1_w)); + resampler_out = ggml_mul_mat(ctx0, temp_0_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_1_b); cb(resampler_out, "temporal_linear_0", -1); // GELU @@ -193,23 +193,23 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { cb(resampler_out, "temporal_gelu", -1); // Second temporal linear - ggml_tensor * mm_temp_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_temp_2_w)); - resampler_out = ggml_mul_mat(ctx0, mm_temp_2_w, resampler_out); - resampler_out = ggml_add(ctx0, resampler_out, model.mm_temp_2_b); + ggml_tensor * temp_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_3_w)); + resampler_out = ggml_mul_mat(ctx0, temp_2_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_3_b); cb(resampler_out, "temporal_linear_2", -1); // LayerNorm - resampler_out = build_norm(resampler_out, model.mm_temp_norm_w, model.mm_temp_norm_b, NORM_TYPE_NORMAL, eps, -1); + resampler_out = build_norm(resampler_out, model.mm_input_norm_w, model.mm_input_norm_b, NORM_TYPE_NORMAL, eps, -1); cb(resampler_out, "temporal_norm", -1); // Final MLP - ggml_tensor * mm_mlp_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_mlp_w)); - resampler_out = ggml_mul_mat(ctx0, mm_mlp_w, resampler_out); - resampler_out = ggml_add(ctx0, resampler_out, model.mm_mlp_b); + ggml_tensor * mlp_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_fc_w)); + resampler_out = ggml_mul_mat(ctx0, mlp_w, resampler_out); + resampler_out = ggml_add(ctx0, resampler_out, model.mm_fc_b); cb(resampler_out, "mlp", -1); // RMS norm (final output normalization) - resampler_out = build_norm(resampler_out, model.mm_after_norm_w, nullptr, NORM_TYPE_RMS, eps, -1); + resampler_out = build_norm(resampler_out, model.mm_norm_mid_w, nullptr, NORM_TYPE_RMS, eps, -1); cb(resampler_out, "after_norm", -1); // Build the graph From 23a38d52f169cba8082ac05aca86837c9fc89b8c Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Wed, 11 Feb 2026 00:05:12 +0800 Subject: [PATCH 7/8] [model] support ernie4-5-vl-moe model --- tools/mtmd/models/ernie45vlmoe.cpp | 133 ++++------------------------- 1 file changed, 17 insertions(+), 116 deletions(-) diff --git a/tools/mtmd/models/ernie45vlmoe.cpp b/tools/mtmd/models/ernie45vlmoe.cpp index 03182537d689..ae0beefdeb62 100644 --- a/tools/mtmd/models/ernie45vlmoe.cpp +++ b/tools/mtmd/models/ernie45vlmoe.cpp @@ -5,13 +5,9 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { // 1. ViT encoder with 2D position embeddings and M-RoPE support // 2. Resampler with spatial conv (2x2 grouping) + optional temporal + MLP + RMS norm - const int batch_size = 1; - const int n_pos = n_patches; - const int spatial_conv_size = hparams.spatial_conv_size; // 2 - const int temporal_conv_size = hparams.temporal_conv_size; // 2 - const bool use_temporal = hparams.use_temporal_conv; + const int n_pos = n_patches; + const int spatial_conv_size = hparams.spatial_conv_size; // 2 - // GGML_ASSERT(n_patches_x == n_patches_y && "only square images supported"); GGML_ASSERT(spatial_conv_size == 2 && "ERNIE-4.5-VL-MoE requires spatial_conv_size=2"); // ERNIE-VL Vision uses 2D position lookup RoPE: @@ -26,124 +22,29 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { ggml_set_name(positions, "positions"); ggml_set_input(positions); - - - // Build vision encoder with patch embedding - // Note: patch_embeddings_0 is reshaped to 4D during export for conv2d compatibility ggml_tensor * inp = build_inp(); - - - // ERNIE-4.5-VL uses RoPE (Rotary Position Embedding), not learned position embeddings - // Position encoding is applied within attention layers via RoPE - // So we don't need to add position embeddings here - - ggml_tensor * inpL = inp; - - // Pre-layernorm - if (model.pre_ln_w) { - inpL = build_norm(inpL, model.pre_ln_w, model.pre_ln_b, NORM_TYPE_NORMAL, eps, -1); - cb(inpL, "pre_ln", -1); - } - - // Loop over encoder layers - for (int il = 0; il < n_layer; il++) { - const auto & layer = model.layers[il]; - ggml_tensor * cur = inpL; - - // Layernorm 1 - cur = build_norm(cur, layer.ln_1_w, layer.ln_1_b, NORM_TYPE_NORMAL, eps, il); - cb(cur, "ln1", il); - - // Self-attention - { - ggml_tensor * Qcur = ggml_mul_mat(ctx0, layer.q_w, cur); - if (layer.q_b) { - Qcur = ggml_add(ctx0, Qcur, layer.q_b); - } - - ggml_tensor * Kcur = ggml_mul_mat(ctx0, layer.k_w, cur); - if (layer.k_b) { - Kcur = ggml_add(ctx0, Kcur, layer.k_b); - } - - ggml_tensor * Vcur = ggml_mul_mat(ctx0, layer.v_w, cur); - if (layer.v_b) { - Vcur = ggml_add(ctx0, Vcur, layer.v_b); - } - - Qcur = ggml_reshape_3d(ctx0, Qcur, d_head, n_head, n_pos); - Kcur = ggml_reshape_3d(ctx0, Kcur, d_head, n_head, n_pos); - Vcur = ggml_reshape_3d(ctx0, Vcur, d_head, n_head, n_pos); - - cb(Qcur, "Qcur", il); - cb(Kcur, "Kcur", il); - cb(Vcur, "Vcur", il); - - - // apply M-RoPE - Qcur = ggml_rope_multi( - ctx0, Qcur, positions, nullptr, - d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, 32768, 10000, 1, 0, 1, 32, 1); - Kcur = ggml_rope_multi( - ctx0, Kcur, positions, nullptr, - d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, 32768, 10000, 1, 0, 1, 32, 1); - - cb(Qcur, "Qcur_rope", il); - cb(Kcur, "Kcur_rope", il); - - - - cur = build_attn(layer.o_w, layer.o_b, Qcur, Kcur, Vcur, nullptr, kq_scale, il); - cb(cur, "attn_out", il); - } - - // Residual - cur = ggml_add(ctx0, cur, inpL); - inpL = cur; - cb(cur, "ffn_inp", il); - - // Layernorm 2 - cur = build_norm(cur, layer.ln_2_w, layer.ln_2_b, NORM_TYPE_NORMAL, eps, il); - cb(cur, "ffn_inp_normed", il); - - // FFN - cur = build_ffn(cur, layer.ff_up_w, layer.ff_up_b, layer.ff_gate_w, layer.ff_gate_b, layer.ff_down_w, - layer.ff_down_b, hparams.ffn_op, il); - - cb(cur, "ffn_out", il); - - // Residual 2 - cur = ggml_add(ctx0, inpL, cur); - cb(cur, "layer_out", il); - - inpL = cur; - } - - // Post-layernorm - if (model.post_ln_w) { - inpL = build_norm(inpL, model.post_ln_w, model.post_ln_b, NORM_TYPE_NORMAL, eps, -1); - } - - ggml_tensor * embeddings = inpL; + // Build ViT encoder using the generic build_vit() with M-RoPE position encoding + auto add_pos = [&](ggml_tensor * cur, const clip_layer &) { + return ggml_rope_multi( + ctx0, cur, positions, nullptr, + d_head/2, mrope_sections, GGML_ROPE_TYPE_VISION, + 32768, 10000, 1, 0, 1, 32, 1); + }; + + ggml_tensor * embeddings = build_vit( + inp, n_pos, + NORM_TYPE_NORMAL, + hparams.ffn_op, + nullptr, // no learned position embeddings, using RoPE + add_pos); cb(embeddings, "vision_output", -1); // ------------------------------------------- // Resampler projection // ------------------------------------------- - // Input shape: [n_embd, n_patches] = [1280, 1600] - // We need to group 2x2 patches: 40x40 patches -> 20x20 groups - // Output shape: [n_embd*4, n_groups] = [5120, 400] - - const int n_groups_x = n_patches_x / spatial_conv_size; // 40/2 = 20 - const int n_groups_y = n_patches_y / spatial_conv_size; // 40/2 = 20 - const int n_groups = n_groups_x * n_groups_y; // 400 - - // Use patch_merge_permute to group 2x2 patches - // Note: build_patch_merge_permute expects 2D input [n_embd, n_patches] + // Group 2x2 patches: 40x40 -> 20x20, output shape [n_embd*4, n_groups] embeddings = build_patch_merge_permute(embeddings, spatial_conv_size); - - // embeddings is now [n_embd*4, n_groups] = [5120, 400] cb(embeddings, "spatial_reshape", -1); // Spatial linear path: Linear -> GELU -> Linear -> LayerNorm From c932d7ea4494de5ac550f8bdcf8bfbff90c0bbed Mon Sep 17 00:00:00 2001 From: isLinXu <2267379130@qq.com> Date: Tue, 17 Mar 2026 13:48:55 +0800 Subject: [PATCH 8/8] update&fix --- convert_hf_to_gguf.py | 162 +++++++++++++++++++++++++++-- gguf-py/gguf/constants.py | 16 +++ gguf-py/gguf/tensor_mapping.py | 42 ++++++++ src/llama-arch.h | 3 - tools/mtmd/clip-model.h | 5 - tools/mtmd/clip.cpp | 5 +- tools/mtmd/models/ernie45vlmoe.cpp | 66 ++++-------- 7 files changed, 235 insertions(+), 64 deletions(-) diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index 66eb93661734..7ad88314af90 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -552,7 +552,6 @@ def prepare_tensors(self): break for new_name, data_torch in (self.modify_tensors(data_torch, name, bid)): - # TODO: why do we squeeze here? # data = data_torch.squeeze().numpy() data = data_torch.numpy() @@ -646,6 +645,9 @@ def prepare_tensors(self): # n_dims is implicit in the shape logger.info(f"{f'%-{max_name_len}s' % f'{new_name},'} {old_dtype} --> {data_qtype.name}, shape = {shape_str}") + # Debug: print all tensors being added + print(f"DEBUG ADD: {new_name}, shape={data.shape}") + self.gguf_writer.add_tensor(new_name, data, raw_dtype=data_qtype) def set_type(self): @@ -3856,8 +3858,9 @@ def set_gguf_parameters(self): def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: # Skip vision and multimodal tensors - they are not part of the text model if name.startswith("vision_model") or name.startswith("resampler_model") or \ - name.startswith("model.vision_model") or name.startswith("model.resampler_model"): - return [] + name.startswith("model.vision_model") or name.startswith("model.resampler_model") or \ + name.endswith(".rotary_emb.original_inv_freq") or name.endswith(".rotary_emb.inv_freq"): + return # todo(megemini): gate_inp weight/weight_1 # weight @@ -3876,7 +3879,8 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter # Map the tensor name and ensure it has .weight suffix mapped_name = self.map_tensor_name(name) - return [(mapped_name, data_torch)] + yield (mapped_name, data_torch) + return # todo(megemini): e_score_correction.bias/bias_1 for weight/weight_1 if name.endswith(".mlp.moe_statics.e_score_correction_bias"): @@ -3886,8 +3890,9 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter name_vision = name.replace("e_score_correction_bias", "e_score_correction.vision.bias") data_torch_vision = data_torch[1, :] - return [(self.map_tensor_name(name_text), data_torch_text), - (self.map_tensor_name(name_vision), data_torch_vision)] + yield (self.map_tensor_name(name_text), data_torch_text) + yield (self.map_tensor_name(name_vision), data_torch_vision) + return # process the experts separately if name.find("mlp.experts") != -1: @@ -3965,10 +3970,13 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter new_name = self.map_tensor_name(merged_name) tensors.append((new_name, data_torch)) - return tensors + for tensor_tuple in tensors: + yield tensor_tuple + return else: - return [] - return [(self.map_tensor_name(name), data_torch)] + return + yield (self.map_tensor_name(name), data_torch) + return def prepare_tensors(self): super().prepare_tensors() @@ -3980,6 +3988,142 @@ def prepare_tensors(self): raise ValueError(f"Unprocessed experts: {experts}") +@ModelBase.register("Ernie4_5_VLMoeForConditionalGeneration") +class Ernie4_5VLMoeVisionModel(MmprojModel): + # Resampler tensor name mapping: HF name -> GGUF name + _resampler_mapping = { + "model.resampler_model.spatial_linear.0": "mm.0", + "model.resampler_model.spatial_linear.2": "mm.2", + "model.resampler_model.spatial_linear.3": "mm.3", + "model.resampler_model.temporal_linear.0": "mm_temp.0", + "model.resampler_model.temporal_linear.2": "mm_temp.2", + "model.resampler_model.temporal_linear.3": "mm_temp.3", + "model.resampler_model.mlp": "mm.mlp", + "model.resampler_model.after_norm": "mm.norm", + } + + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + assert self.hparams_vision is not None + # Set default vision parameters for ERNIE-4.5-VL + if "image_size" not in self.hparams_vision: + self.hparams_vision["image_size"] = 448 # default for ERNIE-4.5-VL + if "patch_size" not in self.hparams_vision: + self.hparams_vision["patch_size"] = 14 + if "hidden_size" not in self.hparams_vision: + self.hparams_vision["hidden_size"] = self.hparams_vision.get("embed_dim", 1280) + if "intermediate_size" not in self.hparams_vision: + self.hparams_vision["intermediate_size"] = self.hparams_vision.get("mlp_ratio", 4) * self.hparams_vision["hidden_size"] + if "num_attention_heads" not in self.hparams_vision: + self.hparams_vision["num_attention_heads"] = self.hparams_vision.get("num_heads", 16) + + def set_gguf_parameters(self): + super().set_gguf_parameters() + self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.ERNIE45VLMOE) + self.gguf_writer.add_vision_attention_layernorm_eps(self.hparams_vision.get("layer_norm_eps", 1e-6)) + # ERNIE-VL uses quick_gelu activation (C++ default when neither use_gelu nor use_silu is set) + ffn_op = self.hparams_vision.get("hidden_act", "quick_gelu") + if ffn_op == "gelu": + self.gguf_writer.add_vision_use_gelu(True) + elif ffn_op == "silu": + self.gguf_writer.add_vision_use_silu(True) + # quick_gelu: don't set either flag, C++ defaults to FFN_GELU_QUICK + + def tensor_force_quant(self, name, new_name, bid, n_dims): + # Handle resampler tensors: bias should be F32, weights F16 + # new_name is already mapped by modify_tensors (e.g., "mm.0.weight", "mm.0.bias") + if new_name.startswith("mm.") or new_name.startswith("mm_"): + if new_name.endswith(".bias"): + return gguf.GGMLQuantizationType.F32 + else: + return gguf.GGMLQuantizationType.F16 + # Let parent handle other tensors + return super().tensor_force_quant(name, new_name, bid, n_dims) + + def prepare_tensors(self): + # Call parent prepare_tensors - resampler tensors will be handled by modify_tensors + # and their types will be controlled by tensor_force_quant + super().prepare_tensors() + + def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + # Handle resampler tensors with manual mapping + for hf_prefix, gguf_prefix in self._resampler_mapping.items(): + if name.startswith(hf_prefix): + suffix = name[len(hf_prefix):] # e.g. ".weight" or ".bias" + new_name = gguf_prefix + suffix + print(f"DEBUG: Resampler mapping: {name} -> {new_name}, shape={data_torch.shape}") + # Yield the tensor - it will be handled by prepare_tensors + yield (new_name, data_torch) + return + + # Debug: print all model.* tensors that are being skipped + if name.startswith("model."): + print(f"DEBUG: Skipping model tensor: {name}") + + # Handle vision encoder tensors + if name.startswith("vision_model."): + # Split fused QKV into separate Q, K, V + if ".attn.qkv." in name: + if data_torch.ndim == 2: # weight + c3, _ = data_torch.shape + else: # bias + c3 = data_torch.shape[0] + assert c3 % 3 == 0 + c = c3 // 3 + wq = data_torch[:c] + wk = data_torch[c: c * 2] + wv = data_torch[c * 2:] + yield from super().modify_tensors(wq, name.replace("qkv", "q"), bid) + yield from super().modify_tensors(wk, name.replace("qkv", "k"), bid) + yield from super().modify_tensors(wv, name.replace("qkv", "v"), bid) + # Split Conv3D patch_embed into Conv2Ds (similar to QWEN2VL) + elif 'patch_embed.proj.weight' in name: + print(f"DEBUG: patch_embed.proj.weight shape = {data_torch.shape}") + if data_torch.ndim == 5: + # Conv3D: [out_channels, in_channels, 2, height, width] for spatial merge + c1, c2, kt, kh, kw = data_torch.shape + del c1, c2, kh, kw # unused + assert kt == 2, "Current implementation only supports spatial_merge_size of 2" + yield (gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.V_ENC_EMBD_PATCH] + ".weight", data_torch[:, :, 0, ...]) + yield (gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.V_ENC_EMBD_PATCH] + ".weight.1", data_torch[:, :, 1, ...]) + elif data_torch.ndim == 4: + # Conv2D: [out_channels, in_channels, height, width] - use as is + yield (gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.V_ENC_EMBD_PATCH] + ".weight", data_torch) + elif data_torch.ndim == 2: + # Linear projection: [out_features, in_features] = (1280, 588) + # Convert to Conv2D: (out_channels, in_channels, height, width) = (1280, 3, 14, 14) + # ERNIE-VL uses a linear layer, but we convert it to Conv2D for compatibility + out_ch, in_ch = data_torch.shape + patch_size = 14 + channels = 3 + assert in_ch == channels * patch_size * patch_size, \ + f"Expected in_features={channels * patch_size * patch_size}, got {in_ch}" + # Reshape: (out_ch, in_ch) -> (out_ch, channels, patch_size, patch_size) + # Note: data is stored as (out_ch, in_ch) = (1280, 588) + # We need to reshape to (out_ch, channels, patch_size, patch_size) = (1280, 3, 14, 14) + # The memory layout is contiguous, so we can view directly + data_conv = data_torch.view(out_ch, channels, patch_size, patch_size) + print(f"DEBUG: Converted linear to Conv2D: {data_conv.shape}") + yield (gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.V_ENC_EMBD_PATCH] + ".weight", data_conv) + else: + raise ValueError(f"Unexpected patch_embed.proj.weight shape: {data_torch.shape}") + # Handle patch_embed bias - it's used by the C++ code + # NOTE: The conv_2d output is f32 because inp_raw is created as f32 in build_inp_raw() + # So we must keep bias as f32 to match the output type of conv_2d + elif 'patch_embed.proj.bias' in name: + # Keep as f32 to match output type + if data_torch.dtype != torch.float32: + data_torch = data_torch.to(torch.float32) + yield (gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.V_ENC_EMBD_PATCH] + ".bias", data_torch) + else: + yield from super().modify_tensors(data_torch, name, bid) + # Skip text model tensors (model.* but not model.resampler_model.* which is handled above) + elif name.startswith("model.") or name.startswith("ernie."): + return + else: + yield from super().modify_tensors(data_torch, name, bid) + + @ModelBase.register( "Qwen2VLModel", "Qwen2VLForConditionalGeneration", diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 10298dffc0a7..454bd9b95082 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -725,6 +725,17 @@ class MODEL_TENSOR(IntEnum): V_DS_NORM = auto() # qwen3vl V_DS_FC1 = auto() # qwen3vl V_DS_FC2 = auto() # qwen3vl + V_FFN_GATE_INP = auto() # ernie45vlmoe + V_FFN_UP_EXPS = auto() # ernie45vlmoe + V_FFN_DOWN_EXPS = auto() # ernie45vlmoe + V_FFN_NORM_EXPS = auto() # ernie45vlmoe + V_FFN_GATE_EXPS = auto() # ernie45vlmoe + V_FFN_GATE_SHEXP = auto() # ernie45vlmoe + V_FFN_UP_SHEXP = auto() # ernie45vlmoe + V_FFN_DOWN_SHEXP = auto() # ernie45vlmoe + V_FFN_GATE_INP_SHEXP = auto() # ernie45vlmoe + V_FFN_NORM_SHEXP = auto() # ernie45vlmoe + V_FFN_EXP_PROBS_B = auto() # ernie45vlmoe V_MM_POST_FC_NORM = auto() # cogvlm V_MM_UP = auto() # cogvlm V_MM_DOWN = auto() # cogvlm @@ -1162,6 +1173,11 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.V_MM_GATE: "mm.gate", MODEL_TENSOR.V_TOK_BOI: "v.boi", MODEL_TENSOR.V_TOK_EOI: "v.eoi", + MODEL_TENSOR.V_FFN_GATE_INP: "blk.{bid}.v_ffn_gate_inp", + MODEL_TENSOR.V_FFN_GATE_EXPS: "blk.{bid}.v_ffn_gate_exps", + MODEL_TENSOR.V_FFN_DOWN_EXPS: "blk.{bid}.v_ffn_down_exps", + MODEL_TENSOR.V_FFN_UP_EXPS: "blk.{bid}.v_ffn_up_exps", + MODEL_TENSOR.V_FFN_EXP_PROBS_B: "blk.{bid}.v_exp_probs_b", # audio (mtmd) # note: all audio tensor names must use prefix "a." or "mm.a." MODEL_TENSOR.A_ENC_EMBD_POS: "a.position_embd", diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index 43f32c7b5223..26d6162c24ee 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -1324,6 +1324,8 @@ class TensorNameMap: "model.vision_tower.embeddings.cls_token", # Intern-S1 "vision_model.class_embedding", # llama 4 "model.vision.patch_embedding.cls_embedding", # cogvlm + "vision_model.embeddings.class_embedding", # ernie4.5-vl-moe + "vision_model.patch_embed.cls_embedding", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_EMBD_PATCH: ( @@ -1338,10 +1340,13 @@ class TensorNameMap: "vision_tower.patch_embed.proj", # kimi-vl "model.vision.patch_embedding.proj", # cogvlm "siglip2.vision_model.embeddings.patch_embedding", + "vision_model.embeddings.patch_embedding", # ernie4.5-vl-moe + "vision_model.patch_embed.proj", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_EMBD_NORM: ( "visual.post_conv_layernorm", # glm4v + "vision_model.ln", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_EMBD_POS: ( @@ -1354,11 +1359,14 @@ class TensorNameMap: "visual.pos_embed", # qwen3vl "model.vision.patch_embedding.position_embedding", # cogvlm "visual.embeddings.position_embedding", # glm4v + "vision_model.embeddings.position_embedding", # ernie4.5-vl-moe + "vision_model.patch_embed.pos_emb", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_ATTN_QKV: ( "visual.blocks.{bid}.attn.qkv", # qwen3vl "model.vision.transformer.layers.{bid}.attention.query_key_value", # cogvlm + "vision_model.blocks.{bid}.attn.qkv", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_ATTN_Q: ( @@ -1372,6 +1380,7 @@ class TensorNameMap: "visual.blocks.{bid}.attn.q", # qwen2vl, generated "vision_tower.encoder.blocks.{bid}.wq", # kimi-vl, generated "siglip2.vision_model.encoder.layers.{bid}.self_attn.q_proj", # youtuvl + "vision_model.blocks.{bid}.attn.q", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_ATTN_Q_NORM: ( @@ -1390,6 +1399,7 @@ class TensorNameMap: "visual.blocks.{bid}.attn.k", # qwen2vl, generated "vision_tower.encoder.blocks.{bid}.wk", # kimi-vl, generated "siglip2.vision_model.encoder.layers.{bid}.self_attn.k_proj", + "vision_model.blocks.{bid}.attn.k", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_ATTN_K_NORM: ( @@ -1408,6 +1418,7 @@ class TensorNameMap: "visual.blocks.{bid}.attn.v", # qwen2vl, generated "vision_tower.encoder.blocks.{bid}.wv", # kimi-vl, generated "siglip2.vision_model.encoder.layers.{bid}.self_attn.v_proj", + "vision_model.blocks.{bid}.attn.v", # ernie4.5-vl-moe ), MODEL_TENSOR.V_ENC_INPUT_NORM: ( @@ -1421,6 +1432,7 @@ class TensorNameMap: "vision_model.model.layers.{bid}.input_layernorm", # llama4 "visual.blocks.{bid}.norm1", # qwen2vl "vision_tower.encoder.blocks.{bid}.norm0", # kimi-vl (norm0/norm1) + "vision_model.blocks.{bid}.norm1", # ernie4.5-vl-moe "model.vision.transformer.layers.{bid}.input_layernorm", # cogvlm "siglip2.vision_model.encoder.layers.{bid}.layer_norm1", ), @@ -1428,6 +1440,7 @@ class TensorNameMap: MODEL_TENSOR.V_ENC_ATTN_O: ( "vision_tower.vision_model.encoder.layers.{bid}.self_attn.out_proj", "vision_tower.vision_model.encoder.layers.{bid}.attn.proj", # InternVL + "vision_model.blocks.{bid}.attn.proj", # ernie4.5-vl-moe "model.vision_tower.encoder.layer.{bid}.attention.projection_layer", # Intern-S1 "vpm.encoder.layers.{bid}.self_attn.out_proj", "model.vision_model.encoder.layers.{bid}.self_attn.out_proj", # SmolVLM @@ -1452,6 +1465,7 @@ class TensorNameMap: "vision_encoder.transformer.layers.{bid}.ffn_norm", # pixtral "visual.blocks.{bid}.norm2", # qwen2vl "vision_tower.encoder.blocks.{bid}.norm1", # kimi-vl (norm0/norm1) + "vision_model.blocks.{bid}.norm2", # ernie4.5-vl-moe "model.vision.transformer.layers.{bid}.post_attention_layernorm", # cogvlm "siglip2.vision_model.encoder.layers.{bid}.layer_norm2", ), @@ -1467,6 +1481,7 @@ class TensorNameMap: "visual.blocks.{bid}.mlp.fc1", # qwen2vl "visual.blocks.{bid}.mlp.up_proj", # qwen2.5vl "visual.blocks.{bid}.mlp.linear_fc1", # qwen3vl + "vision_model.blocks.{bid}.mlp.fc1", # ernie4.5-vl-moe "vision_tower.encoder.blocks.{bid}.mlp.fc0", # kimi-vl (fc0/fc1) "model.vision.transformer.layers.{bid}.mlp.fc1", # cogvlm "siglip2.vision_model.encoder.layers.{bid}.mlp.fc1", @@ -1489,6 +1504,7 @@ class TensorNameMap: "visual.blocks.{bid}.mlp.fc2", # qwen2vl "visual.blocks.{bid}.mlp.down_proj", # qwen2.5vl "visual.blocks.{bid}.mlp.linear_fc2", # qwen3vl + "vision_model.blocks.{bid}.mlp.fc2", # ernie4.5-vl-moe "vision_tower.encoder.blocks.{bid}.mlp.fc1", # kimi-vl (fc0/fc1) "model.vision.transformer.layers.{bid}.mlp.fc2", # cogvlm "siglip2.vision_model.encoder.layers.{bid}.mlp.fc2", @@ -1519,6 +1535,7 @@ class TensorNameMap: "vision_tower.encoder.final_layernorm", # kimi-vl "visual.post_layernorm", # glm4v "siglip2.vision_model.post_layernorm", + "vision_model.post_layernorm", # ernie4.5-vl-moe ), MODEL_TENSOR.V_MM_POST_NORM: ( @@ -1544,18 +1561,22 @@ class TensorNameMap: MODEL_TENSOR.V_RESMPL_POS_EMBD_K: ( "resampler.pos_embed_k", + "resampler_model.pos_embed_k", # ernie4.5-vl-moe ), MODEL_TENSOR.V_RESMPL_ATTN_Q: ( "resampler.attn.in_proj_q", # tensor generated from resampler.attn.in_proj + "resampler_model.attn.in_proj_q", # ernie4.5-vl-moe ), MODEL_TENSOR.V_RESMPL_ATTN_K: ( "resampler.attn.in_proj_k", # tensor generated from resampler.attn.in_proj + "resampler_model.attn.in_proj_k", # ernie4.5-vl-moe ), MODEL_TENSOR.V_RESMPL_ATTN_V: ( "resampler.attn.in_proj_v", # tensor generated from resampler.attn.in_proj + "resampler_model.attn.in_proj_v", # ernie4.5-vl-moe ), MODEL_TENSOR.V_RESMPL_ATTN_OUT: ( @@ -1584,6 +1605,7 @@ class TensorNameMap: MODEL_TENSOR.V_RESMPL_QUERY: ( "resampler.query", + "resampler_model.query", # ernie4.5-vl-moe ), MODEL_TENSOR.V_TOK_EMBD_IMG_BREAK: ( @@ -1635,6 +1657,26 @@ class TensorNameMap: "model.vision.eoi", # cogvlm ), + MODEL_TENSOR.V_FFN_GATE_INP: ( + "model.layers.{bid}.mlp.gate.vision", # ernie4.5-vl-moe + ), + + MODEL_TENSOR.V_FFN_GATE_EXPS: ( + "model.vision.layers.{bid}.mlp.experts.gate_proj", # ernie4.5-vl-moe + ), + + MODEL_TENSOR.V_FFN_DOWN_EXPS: ( + "model.vision.layers.{bid}.mlp.experts.down_proj", # ernie4.5-vl-moe + ), + + MODEL_TENSOR.V_FFN_UP_EXPS: ( + "model.vision.layers.{bid}.mlp.experts.up_proj", # ernie4.5-vl-moe + ), + + MODEL_TENSOR.V_FFN_EXP_PROBS_B: ( + "model.layers.{bid}.mlp.moe_statics.e_score_correction.vision", # ernie4.5-vl-moe + ), + # audio (mtmd) MODEL_TENSOR.A_ENC_EMBD_POS: ( diff --git a/src/llama-arch.h b/src/llama-arch.h index bed5f5317488..c6ff643ed01c 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -158,7 +158,6 @@ enum llm_kv { LLM_KV_GENERAL_LICENSE, LLM_KV_GENERAL_SOURCE_URL, LLM_KV_GENERAL_SOURCE_HF_REPO, - LLM_KV_VOCAB_SIZE, LLM_KV_CONTEXT_LENGTH, LLM_KV_EMBEDDING_LENGTH, @@ -203,7 +202,6 @@ enum llm_kv { LLM_KV_EMBEDDING_SCALE, LLM_KV_TOKEN_SHIFT_COUNT, LLM_KV_INTERLEAVE_MOE_LAYER_STEP, - LLM_KV_ATTENTION_HEAD_COUNT, LLM_KV_ATTENTION_HEAD_COUNT_KV, LLM_KV_ATTENTION_MAX_ALIBI_BIAS, @@ -230,7 +228,6 @@ enum llm_kv { LLM_KV_ATTENTION_TEMPERATURE_SCALE, LLM_KV_ATTENTION_KEY_LENGTH_MLA, LLM_KV_ATTENTION_VALUE_LENGTH_MLA, - LLM_KV_ROPE_DIMENSION_COUNT, LLM_KV_ROPE_DIMENSION_SECTIONS, LLM_KV_ROPE_FREQ_BASE, diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index 3843bf334851..d4ff9151bb08 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -79,11 +79,6 @@ struct clip_hparams { int minicpmv_version = 0; int32_t minicpmv_query_num = 0; // MiniCPM-V query number - // ernie4.5-vl-moe - int32_t spatial_conv_size = 0; - int32_t temporal_conv_size = 0; - bool use_temporal_conv = false; - // custom value provided by user, can be undefined if not set int32_t custom_image_min_tokens = -1; int32_t custom_image_max_tokens = -1; diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index c7ee83b101fe..43efae422bdc 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -1146,9 +1146,6 @@ struct clip_model_loader { case PROJECTOR_TYPE_ERNIE45VLMOE: { hparams.n_merge = 2; - hparams.spatial_conv_size = 2; - hparams.temporal_conv_size = 2; - hparams.use_temporal_conv = model.mm_1_w != nullptr; hparams.set_limit_image_tokens(8, 1024); hparams.set_warmup_n_tokens(256); // avoid OOM on warmup } break; @@ -3858,7 +3855,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { case PROJECTOR_TYPE_GLM4V: return ctx->model.mm_ffn_down_w->ne[1]; case PROJECTOR_TYPE_ERNIE45VLMOE: - return ctx->model.mm_fc_w->ne[0]; + return ctx->model.mm_fc_w->ne[1]; default: GGML_ABORT("Unknown projector type"); } diff --git a/tools/mtmd/models/ernie45vlmoe.cpp b/tools/mtmd/models/ernie45vlmoe.cpp index ae0beefdeb62..23679e0d0ef6 100644 --- a/tools/mtmd/models/ernie45vlmoe.cpp +++ b/tools/mtmd/models/ernie45vlmoe.cpp @@ -5,10 +5,11 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { // 1. ViT encoder with 2D position embeddings and M-RoPE support // 2. Resampler with spatial conv (2x2 grouping) + optional temporal + MLP + RMS norm - const int n_pos = n_patches; - const int spatial_conv_size = hparams.spatial_conv_size; // 2 + const int n_pos = n_patches; + // Use n_merge for patch merge size (same as spatial_conv_size = 2) + const int spatial_merge_size = hparams.n_merge > 0 ? hparams.n_merge : 2; - GGML_ASSERT(spatial_conv_size == 2 && "ERNIE-4.5-VL-MoE requires spatial_conv_size=2"); + GGML_ASSERT(spatial_merge_size == 2 && "ERNIE-4.5-VL-MoE requires n_merge=2"); // ERNIE-VL Vision uses 2D position lookup RoPE: // - Front half of frequencies use h_position @@ -22,6 +23,8 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { ggml_set_name(positions, "positions"); ggml_set_input(positions); + // Use the standard build_inp() which handles Conv2D patch embeddings + // The Python conversion now converts linear weights to Conv2D format ggml_tensor * inp = build_inp(); // Build ViT encoder using the generic build_vit() with M-RoPE position encoding @@ -44,29 +47,19 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { // Resampler projection // ------------------------------------------- // Group 2x2 patches: 40x40 -> 20x20, output shape [n_embd*4, n_groups] - embeddings = build_patch_merge_permute(embeddings, spatial_conv_size); + embeddings = build_patch_merge_permute(embeddings, spatial_merge_size); cb(embeddings, "spatial_reshape", -1); // Spatial linear path: Linear -> GELU -> Linear -> LayerNorm - // Note: weights were transposed (.t()) during GGUF conversion, so we must - // undo that with ggml_transpose before ggml_mul_mat + // Weights are expected to be already transposed in GGUF format ggml_tensor * spatial_out = embeddings; - // First linear - ggml_tensor * spatial_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_0_w)); - spatial_out = ggml_mul_mat(ctx0, spatial_0_w, spatial_out); - spatial_out = ggml_add(ctx0, spatial_out, model.mm_0_b); - cb(spatial_out, "spatial_linear_0", -1); - - // GELU - spatial_out = ggml_gelu(ctx0, spatial_out); - cb(spatial_out, "spatial_gelu", -1); - - // Second linear - ggml_tensor * spatial_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_2_w)); - spatial_out = ggml_mul_mat(ctx0, spatial_2_w, spatial_out); - spatial_out = ggml_add(ctx0, spatial_out, model.mm_2_b); - cb(spatial_out, "spatial_linear_2", -1); + spatial_out = build_ffn(spatial_out, + model.mm_0_w, model.mm_0_b, + nullptr, nullptr, + model.mm_2_w, model.mm_2_b, + FFN_GELU, + -1); // LayerNorm spatial_out = build_norm(spatial_out, model.mm_post_norm_w, model.mm_post_norm_b, NORM_TYPE_NORMAL, eps, -1); @@ -76,36 +69,23 @@ ggml_cgraph * clip_graph_ernie45vlmoe::build() { // Temporal processing for single images (t=1): // Following ERNIE-VL original: when t=1, slice_offsets and slice_offsets2 both point to the same frame - // So we concat(x, x, dim=-1) which in GGML's [hidden, seq] layout is dim=0 - // This doubles the hidden dimension: [5120, 400] -> [10240, 400] resampler_out = ggml_concat(ctx0, resampler_out, resampler_out, 0); // Temporal linear path: Linear -> GELU -> Linear -> LayerNorm - // Weights were transposed (.t()) during GGUF conversion, undo with ggml_transpose - - // First temporal linear - ggml_tensor * temp_0_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_1_w)); - resampler_out = ggml_mul_mat(ctx0, temp_0_w, resampler_out); - resampler_out = ggml_add(ctx0, resampler_out, model.mm_1_b); - cb(resampler_out, "temporal_linear_0", -1); - - // GELU - resampler_out = ggml_gelu(ctx0, resampler_out); - cb(resampler_out, "temporal_gelu", -1); - - // Second temporal linear - ggml_tensor * temp_2_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_3_w)); - resampler_out = ggml_mul_mat(ctx0, temp_2_w, resampler_out); - resampler_out = ggml_add(ctx0, resampler_out, model.mm_3_b); - cb(resampler_out, "temporal_linear_2", -1); + // Weights are expected to be already transposed in GGUF format + resampler_out = build_ffn(resampler_out, + model.mm_1_w, model.mm_1_b, + nullptr, nullptr, + model.mm_3_w, model.mm_3_b, + FFN_GELU, + -1); // LayerNorm resampler_out = build_norm(resampler_out, model.mm_input_norm_w, model.mm_input_norm_b, NORM_TYPE_NORMAL, eps, -1); cb(resampler_out, "temporal_norm", -1); - // Final MLP - ggml_tensor * mlp_w = ggml_cont(ctx0, ggml_transpose(ctx0, model.mm_fc_w)); - resampler_out = ggml_mul_mat(ctx0, mlp_w, resampler_out); + // Final MLP: Linear (weights are expected to be already transposed in GGUF format) + resampler_out = ggml_mul_mat(ctx0, model.mm_fc_w, resampler_out); resampler_out = ggml_add(ctx0, resampler_out, model.mm_fc_b); cb(resampler_out, "mlp", -1);