mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-08-17 10:12:34 +02:00
Compare commits
5 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Toby
|
3733366720 | ||
 Georgi Gerganov
|
4197155add | ||
|
Georgi Gerganov
|
cea66f4c5a | ||
 Fathi Boudra
|
4695f001fe | ||
 Titaniumtown
|
f275595dd1 |
@@ -193,6 +193,14 @@ static std::vector<std::function<void(const common_chat_template & tmpl, autopar
|
||||
LOG_DBG(ANSI_ORANGE "[Patch: Laguna]\n" ANSI_RESET);
|
||||
}
|
||||
},
|
||||
// Bailing V3
|
||||
[](const common_chat_template & tmpl, autoparser & analysis) -> void {
|
||||
if (tmpl.src.find("Bailing V3 chat template") != std::string::npos) {
|
||||
analysis.tools.arguments.value_suffix = trim_whitespace(analysis.tools.arguments.value_suffix);
|
||||
analysis.tools.arguments.tolerate_intertag_whitespace = true;
|
||||
LOG_DBG(ANSI_ORANGE "[Patch: Bailing V3]\n" ANSI_RESET);
|
||||
}
|
||||
},
|
||||
|
||||
});
|
||||
|
||||
|
||||
@@ -27,6 +27,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
|
||||
"BaichuanForCausalLM": "baichuan",
|
||||
"BailingMoeForCausalLM": "bailingmoe",
|
||||
"BailingMoeV2ForCausalLM": "bailingmoe",
|
||||
"BailingMoeV3ForCausalLM": "bailingmoe3",
|
||||
"BambaForCausalLM": "granite",
|
||||
"BertForMaskedLM": "bert",
|
||||
"BertForSequenceClassification": "bert",
|
||||
|
||||
@@ -0,0 +1,192 @@
|
||||
from __future__ import annotations
|
||||
|
||||
import re
|
||||
|
||||
from typing import Callable, Iterable, TYPE_CHECKING
|
||||
|
||||
import torch
|
||||
|
||||
if TYPE_CHECKING:
|
||||
from torch import Tensor
|
||||
|
||||
from .base import ModelBase, TextModel, gguf
|
||||
|
||||
|
||||
@ModelBase.register("BailingMoeV3ForCausalLM")
|
||||
class BailingMoeV3Model(TextModel):
|
||||
model_arch = gguf.MODEL_ARCH.BAILINGMOE3
|
||||
supports_mtp_export = True
|
||||
|
||||
_experts: list[dict[str, Tensor]] | None = None
|
||||
_main_layers: int | None = None
|
||||
|
||||
def __init__(self, *args, **kwargs):
|
||||
super().__init__(*args, **kwargs)
|
||||
nextn_layers = self.hparams.get("num_nextn_predict_layers", 0) or 0
|
||||
if self.no_mtp:
|
||||
nextn_layers = 0
|
||||
self.block_count = self.hparams["num_hidden_layers"] + nextn_layers
|
||||
self.tensor_map = gguf.get_tensor_name_map(self.model_arch, self.block_count)
|
||||
|
||||
def index_tensors(self, remote_hf_model_id: str | None = None):
|
||||
type(self)._main_layers = self.hparams["num_hidden_layers"]
|
||||
return super().index_tensors(remote_hf_model_id=remote_hf_model_id)
|
||||
|
||||
def set_vocab(self):
|
||||
self._set_vocab_gpt2()
|
||||
|
||||
def is_full_attention(self, bid: int) -> bool:
|
||||
n_layer = self.hparams["num_hidden_layers"]
|
||||
layer_group_size = self.hparams["layer_group_size"]
|
||||
return bid >= n_layer or (bid + 1) % layer_group_size == 0 or bid >= n_layer // layer_group_size * layer_group_size
|
||||
|
||||
def set_gguf_parameters(self):
|
||||
if not self.hparams.get("no_kda_lora", False):
|
||||
raise ValueError("BailingMoeV3 KDA LoRA projections are not supported")
|
||||
if not self.hparams.get("kda_safe_gate", False):
|
||||
raise ValueError("BailingMoeV3 non-safe KDA gates are not supported")
|
||||
if self.hparams.get("gated_attention_proj_granularity_type") != "head_wise":
|
||||
raise ValueError("BailingMoeV3 requires head-wise attention gates")
|
||||
|
||||
self.hparams["num_key_value_heads"] = 1
|
||||
super().set_gguf_parameters()
|
||||
|
||||
n_head_kv = [1 if self.is_full_attention(il) else 0 for il in range(self.block_count)]
|
||||
self.gguf_writer.add_head_count_kv(n_head_kv)
|
||||
|
||||
self.gguf_writer.add_vocab_size(self.hparams["vocab_size"])
|
||||
self.gguf_writer.add_ssm_conv_kernel(self.hparams["short_conv_kernel_size"])
|
||||
self.gguf_writer.add_kda_head_dim(self.hparams["head_dim"])
|
||||
self.gguf_writer.add_kda_safe_gate(self.hparams["kda_safe_gate"])
|
||||
self.gguf_writer.add_kda_gate_lower_bound(self.hparams["kda_lower_bound"])
|
||||
|
||||
kv_lora_rank = self.hparams["kv_lora_rank"]
|
||||
qk_nope_head_dim = self.hparams["qk_nope_head_dim"]
|
||||
qk_rope_head_dim = self.hparams["qk_rope_head_dim"]
|
||||
if (q_lora_rank := self.hparams.get("q_lora_rank")) is not None:
|
||||
self.gguf_writer.add_q_lora_rank(q_lora_rank)
|
||||
self.gguf_writer.add_kv_lora_rank(kv_lora_rank)
|
||||
self.gguf_writer.add_rope_dimension_count(qk_rope_head_dim)
|
||||
self.gguf_writer.add_key_length(kv_lora_rank + qk_rope_head_dim)
|
||||
self.gguf_writer.add_key_length_mla(qk_nope_head_dim + qk_rope_head_dim)
|
||||
self.gguf_writer.add_value_length_mla(self.hparams["v_head_dim"])
|
||||
|
||||
self.gguf_writer.add_expert_feed_forward_length(self.hparams["moe_intermediate_size"])
|
||||
self.gguf_writer.add_expert_shared_feed_forward_length(self.hparams["moe_shared_expert_intermediate_size"])
|
||||
self.gguf_writer.add_expert_shared_count(self.hparams["num_shared_experts"])
|
||||
self.gguf_writer.add_leading_dense_block_count(self.hparams["first_k_dense_replace"])
|
||||
self.gguf_writer.add_expert_weights_scale(self.hparams["routed_scaling_factor"])
|
||||
self.gguf_writer.add_expert_weights_norm(self.hparams["norm_topk_prob"])
|
||||
|
||||
def clamp_limits(key: str) -> list[float] | None:
|
||||
values = self.hparams.get(key)
|
||||
if values is None:
|
||||
return None
|
||||
values = [0.0 if value is None else float(value) for value in values[:self.block_count]]
|
||||
return values + [0.0] * (self.block_count - len(values))
|
||||
|
||||
if (values := clamp_limits("expert_swiglu_limit_list")) is not None:
|
||||
self.gguf_writer.add_swiglu_clamp_exp(values)
|
||||
if (values := clamp_limits("share_expert_swiglu_limit_list")) is not None:
|
||||
self.gguf_writer.add_swiglu_clamp_shexp(values)
|
||||
|
||||
if not self.no_mtp and (nextn_layers := self.hparams.get("num_nextn_predict_layers", 0)):
|
||||
self.gguf_writer.add_nextn_predict_layers(nextn_layers)
|
||||
|
||||
def prepare_metadata(self, vocab_only: bool):
|
||||
from_dir = self.fname_out.is_dir()
|
||||
super().prepare_metadata(vocab_only=vocab_only)
|
||||
|
||||
if not self.mtp_only or not from_dir:
|
||||
return
|
||||
|
||||
output_type: str = self.ftype.name.partition("_")[2]
|
||||
fname_default: str = gguf.naming_convention(
|
||||
self.metadata.name, self.metadata.basename, self.metadata.finetune,
|
||||
self.metadata.version, size_label=None, output_type=output_type, model_type=None)
|
||||
self.fname_out = self.fname_out.parent / f"mtp-{fname_default}.gguf"
|
||||
|
||||
@classmethod
|
||||
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
|
||||
name, gen = item
|
||||
if name.endswith(".expert_bias"):
|
||||
name += ".bias"
|
||||
|
||||
if cls._main_layers is None:
|
||||
return super().filter_tensors((name, gen))
|
||||
|
||||
m = re.match(r"model\.layers\.(\d+)\.", name)
|
||||
is_mtp = m is not None and int(m.group(1)) >= cls._main_layers
|
||||
|
||||
if is_mtp and cls.no_mtp:
|
||||
return None
|
||||
if cls.mtp_only and not is_mtp and name not in (
|
||||
"model.word_embeddings.weight", "model.norm.weight", "lm_head.weight",
|
||||
):
|
||||
return None
|
||||
|
||||
return super().filter_tensors((name, gen))
|
||||
|
||||
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
|
||||
if name.endswith((".q_conv1d.weight", ".k_conv1d.weight", ".v_conv1d.weight")) and data_torch.ndim in (2, 3):
|
||||
d_inner = data_torch.shape[0]
|
||||
d_conv = data_torch.shape[-1]
|
||||
data_torch = data_torch.reshape(1, d_inner, 1, d_conv)
|
||||
|
||||
if name.endswith(".A_log"):
|
||||
data_torch = torch.exp(data_torch).reshape(-1, 1)
|
||||
|
||||
if name.endswith(".dt_bias"):
|
||||
name = name.rpartition(".dt_bias")[0] + ".dt_proj.bias"
|
||||
|
||||
if name.endswith(".attention.f_proj.weight"):
|
||||
assert bid is not None
|
||||
if self.is_full_attention(bid):
|
||||
raise ValueError(f"unexpected f_proj on full-attention layer {bid}")
|
||||
name = self.format_tensor_name(gguf.MODEL_TENSOR.SSM_F_A, bid)
|
||||
|
||||
if name.endswith(".attention.g_proj.weight"):
|
||||
assert bid is not None
|
||||
tensor = gguf.MODEL_TENSOR.ATTN_GATE if self.is_full_attention(bid) else gguf.MODEL_TENSOR.SSM_G_A
|
||||
name = self.format_tensor_name(tensor, bid)
|
||||
|
||||
if ".mlp.experts." in name:
|
||||
n_experts = self.hparams["num_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
|
||||
if len(self._experts[bid]) >= n_experts * 3:
|
||||
for weight_name in ("down_proj", "gate_proj", "up_proj"):
|
||||
tensors = []
|
||||
for expert_id in range(n_experts):
|
||||
expert_name = f"model.layers.{bid}.mlp.experts.{expert_id}.{weight_name}.weight"
|
||||
tensors.append(self._experts[bid].pop(expert_name))
|
||||
merged_name = f"model.layers.{bid}.mlp.experts.{weight_name}.weight"
|
||||
yield from super().modify_tensors(torch.stack(tensors, dim=0), merged_name, bid)
|
||||
return
|
||||
|
||||
if name.endswith(".attention.kv_b_proj.weight"):
|
||||
assert bid is not None
|
||||
n_head = self.hparams["num_attention_heads"]
|
||||
v_head_dim = self.hparams["v_head_dim"]
|
||||
qk_nope_head_dim = self.hparams["qk_nope_head_dim"]
|
||||
assert data_torch.shape[0] == n_head * (v_head_dim + qk_nope_head_dim)
|
||||
kv_b = data_torch.view(n_head, v_head_dim + qk_nope_head_dim, data_torch.shape[-1])
|
||||
k_b, v_b = torch.split(kv_b, [qk_nope_head_dim, v_head_dim], dim=1)
|
||||
name_k = self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_K_B, bid)
|
||||
name_v = self.format_tensor_name(gguf.MODEL_TENSOR.ATTN_V_B, bid)
|
||||
yield from super().modify_tensors(k_b.transpose(1, 2), name_k, bid)
|
||||
yield from super().modify_tensors(v_b, name_v, bid)
|
||||
return
|
||||
|
||||
yield from super().modify_tensors(data_torch, name, bid)
|
||||
|
||||
def prepare_tensors(self):
|
||||
super().prepare_tensors()
|
||||
if self._experts is not None:
|
||||
experts = [name for layer in self._experts for name in layer]
|
||||
if experts:
|
||||
raise ValueError(f"Unprocessed experts: {experts}")
|
||||
+1
-1
@@ -5,7 +5,7 @@ project("ggml" C CXX ASM)
|
||||
### GGML Version
|
||||
set(GGML_VERSION_MAJOR 0)
|
||||
set(GGML_VERSION_MINOR 20)
|
||||
set(GGML_VERSION_PATCH 0)
|
||||
set(GGML_VERSION_PATCH 1)
|
||||
set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
|
||||
|
||||
+102
-56
@@ -349,8 +349,9 @@ static void ggml_cpy_f32_q8_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ne / QK8_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q8_0, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -361,8 +362,10 @@ static void ggml_cpy_q8_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q8_0_f32, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -373,9 +376,11 @@ static void ggml_cpy_q2_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK2_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK2_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_q_f32<cpy_blck_q2_0_f32, QK2_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11,
|
||||
ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -387,8 +392,9 @@ static void ggml_cpy_f32_q4_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ne / QK4_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -399,9 +405,11 @@ static void ggml_cpy_q4_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -414,8 +422,9 @@ static void ggml_cpy_f32_q4_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ne / QK4_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -426,9 +435,11 @@ static void ggml_cpy_q4_1_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -441,8 +452,9 @@ static void ggml_cpy_f32_q5_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ne / QK5_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_f32_q<cpy_blck_f32_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -453,9 +465,11 @@ static void ggml_cpy_q5_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -468,8 +482,9 @@ static void ggml_cpy_f32_q5_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ne / QK5_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q5_1, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -480,9 +495,11 @@ static void ggml_cpy_q5_1_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -494,9 +511,11 @@ static void ggml_cpy_mxfp4_f32_sycl(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK_MXFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_MXFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_mxfp4, QK_MXFP4>, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
|
||||
nb01, nb02, nb03, ne10, ne11, ne12,
|
||||
@@ -509,9 +528,10 @@ static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_NL == 0);
|
||||
const int num_blocks = ne / QK4_NL;
|
||||
const int num_blocks = ceil_div(ne / QK4_NL, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11,
|
||||
ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -556,8 +576,9 @@ static void ggml_cpy_f16_q4_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ne / QK4_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -570,8 +591,9 @@ static void ggml_cpy_f16_q4_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ne / QK4_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -584,8 +606,9 @@ static void ggml_cpy_f16_q5_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ne / QK5_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -849,7 +872,8 @@ static void ggml_cpy_q8_0_q8_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
@@ -863,7 +887,8 @@ static void ggml_cpy_q5_0_q5_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
@@ -877,7 +902,8 @@ static void ggml_cpy_q5_1_q5_1(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
@@ -892,7 +918,8 @@ static void ggml_cpy_q4_0_q4_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -906,8 +933,9 @@ static void ggml_cpy_q4_1_q4_1(const char * cx, char * cdst, const int ne, const
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_q<block_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -918,7 +946,8 @@ static void ggml_cpy_q1_0_q1_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK1_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK1_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
@@ -930,7 +959,8 @@ static void ggml_cpy_q2_0_q2_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK2_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK2_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -942,7 +972,8 @@ static void ggml_cpy_mxfp4_mxfp4(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_MXFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_MXFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
@@ -954,7 +985,8 @@ static void ggml_cpy_nvfp4_nvfp4(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_NVFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_NVFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -966,7 +998,8 @@ static void ggml_cpy_q2_K_q2_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -978,7 +1011,8 @@ static void ggml_cpy_q3_K_q3_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -990,7 +1024,8 @@ static void ggml_cpy_q4_K_q4_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1002,7 +1037,8 @@ static void ggml_cpy_q5_K_q5_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1014,7 +1050,8 @@ static void ggml_cpy_q6_K_q6_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1026,7 +1063,8 @@ static void ggml_cpy_iq2_xxs_iq2_xxs(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1038,7 +1076,8 @@ static void ggml_cpy_iq2_xs_iq2_xs(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1050,7 +1089,8 @@ static void ggml_cpy_iq2_s_iq2_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1062,7 +1102,8 @@ static void ggml_cpy_iq3_xxs_iq3_xxs(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1074,7 +1115,8 @@ static void ggml_cpy_iq1_s_iq1_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1086,7 +1128,8 @@ static void ggml_cpy_iq1_m_iq1_m(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1098,7 +1141,8 @@ static void ggml_cpy_iq4_nl_iq4_nl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK4_NL == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_NL, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1110,7 +1154,8 @@ static void ggml_cpy_iq3_s_iq3_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1122,7 +1167,8 @@ static void ggml_cpy_iq4_xs_iq4_xs(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
|
||||
@@ -261,6 +261,7 @@ class Keys:
|
||||
|
||||
class KDA:
|
||||
HEAD_DIM = "{arch}.kda.head_dim"
|
||||
SAFE_GATE = "{arch}.kda.safe_gate"
|
||||
GATE_LOWER_BOUND = "{arch}.kda.gate_lower_bound"
|
||||
|
||||
class WKV:
|
||||
@@ -552,6 +553,7 @@ class MODEL_ARCH(IntEnum):
|
||||
PLM = auto()
|
||||
BAILINGMOE = auto()
|
||||
BAILINGMOE2 = auto()
|
||||
BAILINGMOE3 = auto()
|
||||
DOTS1 = auto()
|
||||
ARCEE = auto()
|
||||
AFMOE = auto()
|
||||
@@ -1267,6 +1269,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
|
||||
MODEL_ARCH.PLM: "plm",
|
||||
MODEL_ARCH.BAILINGMOE: "bailingmoe",
|
||||
MODEL_ARCH.BAILINGMOE2: "bailingmoe2",
|
||||
MODEL_ARCH.BAILINGMOE3: "bailingmoe3",
|
||||
MODEL_ARCH.DOTS1: "dots1",
|
||||
MODEL_ARCH.ARCEE: "arcee",
|
||||
MODEL_ARCH.AFMOE: "afmoe",
|
||||
@@ -4234,6 +4237,50 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
|
||||
MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM,
|
||||
MODEL_TENSOR.LAYER_OUT_NORM,
|
||||
],
|
||||
MODEL_ARCH.BAILINGMOE3: [
|
||||
MODEL_TENSOR.TOKEN_EMBD,
|
||||
MODEL_TENSOR.OUTPUT_NORM,
|
||||
MODEL_TENSOR.OUTPUT,
|
||||
MODEL_TENSOR.ATTN_NORM,
|
||||
MODEL_TENSOR.ATTN_Q,
|
||||
MODEL_TENSOR.ATTN_Q_A,
|
||||
MODEL_TENSOR.ATTN_Q_B,
|
||||
MODEL_TENSOR.ATTN_Q_A_NORM,
|
||||
MODEL_TENSOR.ATTN_K,
|
||||
MODEL_TENSOR.ATTN_V,
|
||||
MODEL_TENSOR.ATTN_OUT,
|
||||
MODEL_TENSOR.ATTN_GATE,
|
||||
MODEL_TENSOR.ATTN_KV_A_MQA,
|
||||
MODEL_TENSOR.ATTN_KV_B,
|
||||
MODEL_TENSOR.ATTN_K_B,
|
||||
MODEL_TENSOR.ATTN_V_B,
|
||||
MODEL_TENSOR.ATTN_KV_A_NORM,
|
||||
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_GATE_SHEXP,
|
||||
MODEL_TENSOR.FFN_DOWN_SHEXP,
|
||||
MODEL_TENSOR.FFN_UP_SHEXP,
|
||||
MODEL_TENSOR.FFN_EXP_PROBS_B,
|
||||
MODEL_TENSOR.SSM_CONV1D_Q,
|
||||
MODEL_TENSOR.SSM_CONV1D_K,
|
||||
MODEL_TENSOR.SSM_CONV1D_V,
|
||||
MODEL_TENSOR.SSM_F_A,
|
||||
MODEL_TENSOR.SSM_BETA,
|
||||
MODEL_TENSOR.SSM_A,
|
||||
MODEL_TENSOR.SSM_G_A,
|
||||
MODEL_TENSOR.SSM_DT,
|
||||
MODEL_TENSOR.SSM_NORM,
|
||||
MODEL_TENSOR.NEXTN_EH_PROJ,
|
||||
MODEL_TENSOR.NEXTN_ENORM,
|
||||
MODEL_TENSOR.NEXTN_HNORM,
|
||||
MODEL_TENSOR.LAYER_OUT_NORM,
|
||||
],
|
||||
MODEL_ARCH.DOTS1: [
|
||||
MODEL_TENSOR.TOKEN_EMBD,
|
||||
MODEL_TENSOR.OUTPUT_NORM,
|
||||
@@ -5487,7 +5534,9 @@ KEY_SSM_GROUP_COUNT = Keys.SSM.GROUP_COUNT
|
||||
KEY_SSM_DT_B_C_RMS = Keys.SSM.DT_B_C_RMS
|
||||
|
||||
# KDA
|
||||
KEY_KDA_HEAD_DIM = Keys.KDA.HEAD_DIM
|
||||
KEY_KDA_HEAD_DIM = Keys.KDA.HEAD_DIM
|
||||
KEY_KDA_SAFE_GATE = Keys.KDA.SAFE_GATE
|
||||
KEY_KDA_GATE_LOWER_BOUND = Keys.KDA.GATE_LOWER_BOUND
|
||||
|
||||
# tokenization
|
||||
KEY_TOKENIZER_MODEL = Keys.Tokenizer.MODEL
|
||||
|
||||
@@ -1103,9 +1103,6 @@ class GGUFWriter:
|
||||
def add_ssm_dt_b_c_rms(self, value: bool) -> None:
|
||||
self.add_bool(Keys.SSM.DT_B_C_RMS.format(arch=self.arch), value)
|
||||
|
||||
def add_kda_gate_lower_bound(self, value: float) -> None:
|
||||
self.add_float32(Keys.KDA.GATE_LOWER_BOUND.format(arch=self.arch), value)
|
||||
|
||||
def add_expert_latent_length(self, value: int) -> None:
|
||||
self.add_uint32(Keys.LLM.EXPERT_LATENT_LENGTH.format(arch=self.arch), value)
|
||||
|
||||
@@ -1121,6 +1118,12 @@ class GGUFWriter:
|
||||
def add_kda_head_dim(self, value: int) -> None:
|
||||
self.add_uint32(Keys.KDA.HEAD_DIM.format(arch=self.arch), value)
|
||||
|
||||
def add_kda_safe_gate(self, value: bool) -> None:
|
||||
self.add_bool(Keys.KDA.SAFE_GATE.format(arch=self.arch), value)
|
||||
|
||||
def add_kda_gate_lower_bound(self, value: float) -> None:
|
||||
self.add_float32(Keys.KDA.GATE_LOWER_BOUND.format(arch=self.arch), value)
|
||||
|
||||
def add_tokenizer_model(self, model: str) -> None:
|
||||
self.add_string(Keys.Tokenizer.MODEL, model)
|
||||
|
||||
|
||||
@@ -255,6 +255,7 @@ class TensorNameMap:
|
||||
# Attention query
|
||||
MODEL_TENSOR.ATTN_Q: (
|
||||
"model.layers.{bid}.self_attn.q_proj", # llama-hf nemotron olmoe olmo2 phimoe
|
||||
"model.layers.{bid}.attention.q_proj", # bailingmoe3
|
||||
"layers.{bid}.self_attn.q_proj", # embeddinggemma
|
||||
"model.layers.{bid}.self_attn.q_proj_no_perm", # llama-custom
|
||||
"layers.{bid}.attention.wq", # llama-pth
|
||||
@@ -275,6 +276,7 @@ class TensorNameMap:
|
||||
# Attention key
|
||||
MODEL_TENSOR.ATTN_K: (
|
||||
"model.layers.{bid}.self_attn.k_proj", # llama-hf nemotron olmoe olmo2 phimoe
|
||||
"model.layers.{bid}.attention.k_proj", # bailingmoe3
|
||||
"layers.{bid}.self_attn.k_proj", # embeddinggemma
|
||||
"model.layers.{bid}.self_attn.k_proj_no_perm", # llama-custom
|
||||
"layers.{bid}.attention.wk", # llama-pth
|
||||
@@ -296,6 +298,7 @@ class TensorNameMap:
|
||||
# Attention value
|
||||
MODEL_TENSOR.ATTN_V: (
|
||||
"model.layers.{bid}.self_attn.v_proj", # llama-hf nemotron olmoe olmo2 phimoe
|
||||
"model.layers.{bid}.attention.v_proj", # bailingmoe3
|
||||
"layers.{bid}.self_attn.v_proj", # embeddinggemma
|
||||
"layers.{bid}.attention.wv", # llama-pth
|
||||
"encoder.layer.{bid}.attention.self.value", # bert
|
||||
@@ -321,6 +324,8 @@ class TensorNameMap:
|
||||
"transformer.h.{bid}.self_attention.dense", # falcon
|
||||
"h.{bid}.self_attention.dense", # bloom
|
||||
"model.layers.{bid}.self_attn.o_proj", # llama-hf nemotron olmoe olmo2 phimoe
|
||||
"model.layers.{bid}.attention.o_proj", # bailingmoe3
|
||||
"model.layers.{bid}.attention.dense", # bailingmoe3 MLA
|
||||
"layers.{bid}.self_attn.o_proj", # embeddinggemma
|
||||
"model.layers.{bid}.self_attn.out_proj", # lfm2 minimax-01
|
||||
"model.layers.{bid}.self_attn.linear_attn", # deci
|
||||
@@ -834,6 +839,7 @@ class TensorNameMap:
|
||||
"model.layers.{bid}.linear_attn.dt_proj", # qwen3next
|
||||
"backbone.layers.{bid}.mixer.dt", # nemotron-h-moe
|
||||
"model.layers.{bid}.self_attn.dt_proj", # kimi
|
||||
"model.layers.{bid}.attention.dt_proj", # bailingmoe3
|
||||
),
|
||||
|
||||
MODEL_TENSOR.SSM_DT_NORM: (
|
||||
@@ -848,6 +854,7 @@ class TensorNameMap:
|
||||
"model.layers.layers.{bid}.mixer.A_log", # plamo2
|
||||
"model.layers.{bid}.linear_attn.A_log", # qwen3next
|
||||
"model.layers.{bid}.self_attn.A_log", # kimi
|
||||
"model.layers.{bid}.attention.A_log", # bailingmoe3
|
||||
),
|
||||
|
||||
MODEL_TENSOR.SSM_B_NORM: (
|
||||
@@ -874,6 +881,7 @@ class TensorNameMap:
|
||||
"model.layers.{bid}.linear_attn.norm", # qwen3next
|
||||
"backbone.layers.{bid}.mixer.norm", # mamba2
|
||||
"model.layers.{bid}.self_attn.o_norm", # kimi
|
||||
"model.layers.{bid}.attention.o_norm", # bailingmoe3
|
||||
),
|
||||
|
||||
MODEL_TENSOR.SSM_OUT: (
|
||||
@@ -895,12 +903,15 @@ class TensorNameMap:
|
||||
# Kimi Linear KDA (using SSM_ prefix for consistency)
|
||||
MODEL_TENSOR.SSM_CONV1D_Q: (
|
||||
"model.layers.{bid}.self_attn.q_conv1d",
|
||||
"model.layers.{bid}.attention.q_conv1d",
|
||||
),
|
||||
MODEL_TENSOR.SSM_CONV1D_K: (
|
||||
"model.layers.{bid}.self_attn.k_conv1d",
|
||||
"model.layers.{bid}.attention.k_conv1d",
|
||||
),
|
||||
MODEL_TENSOR.SSM_CONV1D_V: (
|
||||
"model.layers.{bid}.self_attn.v_conv1d",
|
||||
"model.layers.{bid}.attention.v_conv1d",
|
||||
),
|
||||
MODEL_TENSOR.SSM_F_A: (
|
||||
"model.layers.{bid}.self_attn.f_a_proj",
|
||||
@@ -911,6 +922,7 @@ class TensorNameMap:
|
||||
MODEL_TENSOR.SSM_BETA: (
|
||||
"model.layers.{bid}.linear_attn.in_proj_b", # qwen3.5
|
||||
"model.layers.{bid}.self_attn.b_proj", # Kimi Linear
|
||||
"model.layers.{bid}.attention.b_proj", # bailingmoe3
|
||||
),
|
||||
# Kimi K3 latent MoE: routed experts operate in a down-projected space
|
||||
MODEL_TENSOR.FFN_ROUTED_DOWN: (
|
||||
@@ -1103,40 +1115,48 @@ class TensorNameMap:
|
||||
|
||||
MODEL_TENSOR.ATTN_Q_A: (
|
||||
"model.layers.{bid}.self_attn.q_a_proj", # deepseek2
|
||||
"model.layers.{bid}.attention.q_a_proj", # bailingmoe3 (Ling-3.0-tiny)
|
||||
"layers.{bid}.attention.wq_a", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_Q_B: (
|
||||
"model.layers.{bid}.self_attn.q_b_proj", # deepseek2
|
||||
"model.layers.{bid}.attention.q_b_proj", # bailingmoe3 (Ling-3.0-tiny)
|
||||
"layers.{bid}.attention.wq_b", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_KV_A_MQA: (
|
||||
"model.layers.{bid}.self_attn.kv_a_proj_with_mqa", # deepseek2
|
||||
"model.layers.{bid}.attention.kv_a_proj_with_mqa", # bailingmoe3
|
||||
"layers.{bid}.attention.wkv_a_with_mqa", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_KV_B: (
|
||||
"model.layers.{bid}.self_attn.kv_b_proj", # deepseek2
|
||||
"model.layers.{bid}.attention.kv_b_proj", # bailingmoe3
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_K_B: (
|
||||
"model.layers.{bid}.self_attn.k_b_proj", # deepseek2
|
||||
"model.layers.{bid}.attention.k_b_proj", # bailingmoe3
|
||||
"layers.{bid}.attention.k_b_proj", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_V_B: (
|
||||
"model.layers.{bid}.self_attn.v_b_proj", # deepseek2
|
||||
"model.layers.{bid}.attention.v_b_proj", # bailingmoe3
|
||||
"layers.{bid}.attention.v_b_proj", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_Q_A_NORM: (
|
||||
"model.layers.{bid}.self_attn.q_a_layernorm", # deepseek2
|
||||
"model.layers.{bid}.attention.q_a_layernorm", # bailingmoe3 (Ling-3.0-tiny)
|
||||
"layers.{bid}.attention.q_a_norm", # mistral-large
|
||||
),
|
||||
|
||||
MODEL_TENSOR.ATTN_KV_A_NORM: (
|
||||
"model.layers.{bid}.self_attn.kv_a_layernorm", # deepseek2
|
||||
"model.layers.{bid}.attention.kv_a_layernorm", # bailingmoe3
|
||||
"layers.{bid}.attention.kv_a_norm", # mistral-large
|
||||
),
|
||||
|
||||
|
||||
@@ -1 +1 @@
|
||||
2d191b5dee1a591c41ee8a653ce42bfcd9c8716d
|
||||
3834fd814e74e8af277939dabd69ecc780affd21
|
||||
|
||||
+5
-1
@@ -107,6 +107,7 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
|
||||
{ LLM_ARCH_PLM, "plm" },
|
||||
{ LLM_ARCH_BAILINGMOE, "bailingmoe" },
|
||||
{ LLM_ARCH_BAILINGMOE2, "bailingmoe2" },
|
||||
{ LLM_ARCH_BAILINGMOE3, "bailingmoe3" },
|
||||
{ LLM_ARCH_DOTS1, "dots1" },
|
||||
{ LLM_ARCH_ARCEE, "arcee" },
|
||||
{ LLM_ARCH_AFMOE, "afmoe" },
|
||||
@@ -317,7 +318,8 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
|
||||
{ LLM_KV_SSM_GROUP_COUNT, "%s.ssm.group_count" },
|
||||
{ LLM_KV_SSM_DT_B_C_RMS, "%s.ssm.dt_b_c_rms" },
|
||||
|
||||
{ LLM_KV_KDA_HEAD_DIM, "%s.kda.head_dim" },
|
||||
{ LLM_KV_KDA_HEAD_DIM, "%s.kda.head_dim" },
|
||||
{ LLM_KV_KDA_SAFE_GATE, "%s.kda.safe_gate" },
|
||||
{ LLM_KV_KDA_GATE_LOWER_BOUND, "%s.kda.gate_lower_bound" },
|
||||
|
||||
{ LLM_KV_WKV_HEAD_SIZE, "%s.wkv.head_size" },
|
||||
@@ -996,6 +998,7 @@ bool llm_arch_is_hybrid(const llm_arch & arch) {
|
||||
case LLM_ARCH_NEMOTRON_H_MOE:
|
||||
case LLM_ARCH_QWEN3NEXT:
|
||||
case LLM_ARCH_KIMI_LINEAR:
|
||||
case LLM_ARCH_BAILINGMOE3:
|
||||
case LLM_ARCH_KIMI_K3:
|
||||
case LLM_ARCH_QWEN35:
|
||||
case LLM_ARCH_QWEN35MOE:
|
||||
@@ -1061,6 +1064,7 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) {
|
||||
case LLM_ARCH_MINIMAX_M3:
|
||||
case LLM_ARCH_MISTRAL4:
|
||||
case LLM_ARCH_KIMI_LINEAR:
|
||||
case LLM_ARCH_BAILINGMOE3:
|
||||
case LLM_ARCH_KIMI_K3:
|
||||
case LLM_ARCH_QWEN3TTS:
|
||||
return false;
|
||||
|
||||
@@ -112,6 +112,7 @@ enum llm_arch {
|
||||
LLM_ARCH_PLM,
|
||||
LLM_ARCH_BAILINGMOE,
|
||||
LLM_ARCH_BAILINGMOE2,
|
||||
LLM_ARCH_BAILINGMOE3,
|
||||
LLM_ARCH_DOTS1,
|
||||
LLM_ARCH_ARCEE,
|
||||
LLM_ARCH_AFMOE,
|
||||
@@ -323,6 +324,7 @@ enum llm_kv {
|
||||
LLM_KV_SSM_DT_B_C_RMS,
|
||||
|
||||
LLM_KV_KDA_HEAD_DIM,
|
||||
LLM_KV_KDA_SAFE_GATE,
|
||||
LLM_KV_KDA_GATE_LOWER_BOUND,
|
||||
|
||||
LLM_KV_WKV_HEAD_SIZE,
|
||||
|
||||
@@ -2298,6 +2298,7 @@ uint32_t llama_context::graph_max_nodes(uint32_t n_tokens) const {
|
||||
res = std::max<uint32_t>(n_tokens * 160, 64u * model.n_tensors());
|
||||
} else if (model.arch == LLM_ARCH_QWEN3NEXT ||
|
||||
model.arch == LLM_ARCH_KIMI_LINEAR ||
|
||||
model.arch == LLM_ARCH_BAILINGMOE3 ||
|
||||
model.arch == LLM_ARCH_QWEN35 ||
|
||||
model.arch == LLM_ARCH_QWEN35MOE ||
|
||||
model.arch == LLM_ARCH_DEEPSEEK4 ||
|
||||
|
||||
@@ -170,6 +170,7 @@ struct llama_hparams {
|
||||
|
||||
// for Kimi Linear KDA
|
||||
uint32_t n_embd_head_kda = 0;
|
||||
bool kda_safe_gate = false;
|
||||
|
||||
// kimi-k3
|
||||
uint32_t n_expert_latent = 0; // routed_expert_hidden_size (0 = experts run at n_embd)
|
||||
|
||||
@@ -121,6 +121,7 @@ void llama_model_saver::add_kv(const enum llm_kv key, const Container & value, c
|
||||
}
|
||||
// instantiate for external usage:
|
||||
template void llama_model_saver::add_kv<std::vector<uint32_t>>(const enum llm_kv, const std::vector<uint32_t> &, const bool);
|
||||
template void llama_model_saver::add_kv<std::vector<float>>(const enum llm_kv, const std::vector<float> &, const bool);
|
||||
|
||||
void llama_model_saver::add_kv(const enum llm_kv key, const std::vector<std::string> & value) {
|
||||
std::vector<const char *> tmp(value.size());
|
||||
@@ -216,8 +217,10 @@ void llama_model_saver::add_kv_from_model() {
|
||||
add_kv(LLM_KV_EXPERT_LATENT_LENGTH, hparams.n_expert_latent);
|
||||
add_kv(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp);
|
||||
add_kv(LLM_KV_EXPERT_CHUNK_FEED_FORWARD_LENGTH, hparams.n_ff_chexp);
|
||||
add_kv(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp);
|
||||
add_kv(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp);
|
||||
add_kv(LLM_KV_SWIGLU_CLAMP_EXP, std::vector<float>(
|
||||
hparams.swiglu_clamp_exp.begin(), hparams.swiglu_clamp_exp.begin() + hparams.n_layer_all));
|
||||
add_kv(LLM_KV_SWIGLU_CLAMP_SHEXP, std::vector<float>(
|
||||
hparams.swiglu_clamp_shexp.begin(), hparams.swiglu_clamp_shexp.begin() + hparams.n_layer_all));
|
||||
add_kv(LLM_KV_USE_PARALLEL_RESIDUAL, hparams.use_par_res);
|
||||
// add_kv(LLM_KV_TENSOR_DATA_LAYOUT, ???);
|
||||
add_kv(LLM_KV_EXPERT_COUNT, hparams.n_expert);
|
||||
@@ -320,6 +323,7 @@ void llama_model_saver::add_kv_from_model() {
|
||||
add_kv(LLM_KV_SSM_DT_B_C_RMS, hparams.ssm_dt_b_c_rms);
|
||||
|
||||
add_kv(LLM_KV_KDA_HEAD_DIM, hparams.n_embd_head_kda);
|
||||
add_kv(LLM_KV_KDA_SAFE_GATE, hparams.kda_safe_gate);
|
||||
add_kv(LLM_KV_KDA_GATE_LOWER_BOUND, hparams.kda_gate_lower_bound);
|
||||
|
||||
add_kv(LLM_KV_WKV_HEAD_SIZE, hparams.wkv_head_size);
|
||||
|
||||
+9
-4
@@ -256,6 +256,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
|
||||
return new llama_model_bailingmoe(params);
|
||||
case LLM_ARCH_BAILINGMOE2:
|
||||
return new llama_model_bailingmoe2(params);
|
||||
case LLM_ARCH_BAILINGMOE3:
|
||||
return new llama_model_bailingmoe3(params);
|
||||
case LLM_ARCH_SEED_OSS:
|
||||
return new llama_model_seed_oss(params);
|
||||
case LLM_ARCH_DOTS1:
|
||||
@@ -821,6 +823,7 @@ const char * llm_type_name(llm_type type) {
|
||||
case LLM_TYPE_A13B: return "A13B";
|
||||
case LLM_TYPE_7B_A1B: return "7B.A1B";
|
||||
case LLM_TYPE_8B_A1B: return "8B.A1B";
|
||||
case LLM_TYPE_7_9B_A1_3B: return "7.9B.A1.3B";
|
||||
case LLM_TYPE_12B_A2_5B: return "12B.A2.5B";
|
||||
case LLM_TYPE_16B_A1B: return "16B.A1B";
|
||||
case LLM_TYPE_21B_A3B: return "21B.A3B";
|
||||
@@ -837,6 +840,7 @@ const char * llm_type_name(llm_type type) {
|
||||
case LLM_TYPE_118B_A8B: return "118B.A8B";
|
||||
case LLM_TYPE_120B_A12B: return "120B.A12B";
|
||||
case LLM_TYPE_122B_A10B: return "122B.A10B";
|
||||
case LLM_TYPE_124B_A5_1B: return "124B.A5.1B";
|
||||
case LLM_TYPE_196B_A11B: return "196B.A11B";
|
||||
case LLM_TYPE_230B_A10B: return "230B.A10B";
|
||||
case LLM_TYPE_428B_A23B: return "428B.A23B";
|
||||
@@ -1960,7 +1964,7 @@ void llama_model::print_info() const {
|
||||
LLAMA_LOG_INFO("%s: expert_weights_norm = %d\n", __func__, hparams.expert_weights_norm);
|
||||
}
|
||||
|
||||
if (arch == LLM_ARCH_BAILINGMOE2) {
|
||||
if (arch == LLM_ARCH_BAILINGMOE2 || arch == LLM_ARCH_BAILINGMOE3) {
|
||||
LLAMA_LOG_INFO("%s: n_layer_dense_lead = %d\n", __func__, hparams.n_layer_dense_lead);
|
||||
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
|
||||
LLAMA_LOG_INFO("%s: n_ff_shexp = %d\n", __func__, hparams.n_ff_shexp);
|
||||
@@ -2255,11 +2259,11 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params,
|
||||
// checks
|
||||
default:
|
||||
{
|
||||
// The MTP head is dense-attention only on hybrid Qwen3-Next/3.5/3.6, so use a plain
|
||||
// attention KV cache for the MTP context instead of the hybrid wrapper.
|
||||
// Dense MTP heads use a plain attention KV cache instead of the hybrid wrapper.
|
||||
const bool mtp_on_hybrid_qwen =
|
||||
params.ctx_type == LLAMA_CONTEXT_TYPE_MTP &&
|
||||
(arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE);
|
||||
(arch == LLM_ARCH_QWEN3NEXT || arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE ||
|
||||
arch == LLM_ARCH_BAILINGMOE3);
|
||||
|
||||
const bool mtp_on_hybrid_nemotron =
|
||||
params.ctx_type == LLAMA_CONTEXT_TYPE_MTP && arch == LLM_ARCH_NEMOTRON_H_MOE;
|
||||
@@ -2637,6 +2641,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
|
||||
case LLM_ARCH_GRANITE_SWITCH:
|
||||
case LLM_ARCH_CHAMELEON:
|
||||
case LLM_ARCH_BAILINGMOE:
|
||||
case LLM_ARCH_BAILINGMOE3:
|
||||
case LLM_ARCH_NEO_BERT:
|
||||
case LLM_ARCH_SMOLLM3:
|
||||
case LLM_ARCH_ARCEE:
|
||||
|
||||
@@ -118,6 +118,7 @@ enum llm_type {
|
||||
LLM_TYPE_A13B,
|
||||
LLM_TYPE_7B_A1B,
|
||||
LLM_TYPE_8B_A1B, // lfm2moe
|
||||
LLM_TYPE_7_9B_A1_3B, // Ling-3.0-tiny
|
||||
LLM_TYPE_12B_A2_5B,
|
||||
LLM_TYPE_16B_A1B,
|
||||
LLM_TYPE_21B_A3B, // Ernie MoE small
|
||||
@@ -134,6 +135,7 @@ enum llm_type {
|
||||
LLM_TYPE_118B_A8B, // Laguna-S-2
|
||||
LLM_TYPE_120B_A12B, // Nemotron 3 Super
|
||||
LLM_TYPE_122B_A10B, // Qwen3.5
|
||||
LLM_TYPE_124B_A5_1B, // Ling-3.0-flash
|
||||
LLM_TYPE_196B_A11B, // Step3.5-Flash
|
||||
LLM_TYPE_230B_A10B, // Minimax M2
|
||||
LLM_TYPE_428B_A23B, // Minimax M3
|
||||
|
||||
+6
-1
@@ -474,7 +474,12 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type
|
||||
} else if (ftype == LLAMA_FTYPE_MOSTLY_MXFP4_MOE) {
|
||||
// MoE tensors -> MXFP4
|
||||
// other tensors -> Q8_0
|
||||
if (tensor->ne[2] > 1) {
|
||||
// MLA projection tensors are also 3D, so match expert tensor roles explicitly.
|
||||
const bool is_bailingmoe3_expert = arch == LLM_ARCH_BAILINGMOE3 &&
|
||||
(category == tensor_category::FFN_UP ||
|
||||
category == tensor_category::FFN_GATE ||
|
||||
category == tensor_category::FFN_DOWN);
|
||||
if (tensor->ne[2] > 1 && (arch != LLM_ARCH_BAILINGMOE3 || is_bailingmoe3_expert)) {
|
||||
new_type = GGML_TYPE_MXFP4;
|
||||
} else {
|
||||
new_type = GGML_TYPE_Q8_0;
|
||||
|
||||
@@ -0,0 +1,532 @@
|
||||
#include "models.h"
|
||||
#include "llama-memory-recurrent.h"
|
||||
|
||||
void llama_model_bailingmoe3::load_arch_hparams(llama_model_loader & ml) {
|
||||
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
|
||||
ml.get_key(LLM_KV_ATTENTION_KEY_LENGTH_MLA, hparams.n_embd_head_k_mla_impl);
|
||||
ml.get_key(LLM_KV_ATTENTION_VALUE_LENGTH_MLA, hparams.n_embd_head_v_mla_impl);
|
||||
ml.get_key(LLM_KV_ATTENTION_KV_LORA_RANK, hparams.n_lora_kv);
|
||||
ml.get_key(LLM_KV_ATTENTION_Q_LORA_RANK, hparams.n_lora_q, false);
|
||||
ml.get_key(LLM_KV_SSM_CONV_KERNEL, hparams.ssm_d_conv);
|
||||
ml.get_key(LLM_KV_KDA_HEAD_DIM, hparams.n_embd_head_kda);
|
||||
if (!ml.get_key(LLM_KV_KDA_SAFE_GATE, hparams.kda_safe_gate, false)) {
|
||||
hparams.kda_safe_gate = true;
|
||||
}
|
||||
ml.get_key(LLM_KV_KDA_GATE_LOWER_BOUND, hparams.kda_gate_lower_bound);
|
||||
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
|
||||
ml.get_key(LLM_KV_EXPERT_SHARED_FEED_FORWARD_LENGTH, hparams.n_ff_shexp, false);
|
||||
ml.get_key(LLM_KV_EXPERT_SHARED_COUNT, hparams.n_expert_shared);
|
||||
ml.get_key(LLM_KV_LEADING_DENSE_BLOCK_COUNT, hparams.n_layer_dense_lead);
|
||||
ml.get_key(LLM_KV_EXPERT_WEIGHTS_SCALE, hparams.expert_weights_scale, false);
|
||||
ml.get_key(LLM_KV_EXPERT_WEIGHTS_NORM, hparams.expert_weights_norm, false);
|
||||
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);
|
||||
ml.get_key(LLM_KV_NEXTN_PREDICT_LAYERS, hparams.n_layer_nextn, false);
|
||||
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer_all, false);
|
||||
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer_all, false);
|
||||
|
||||
if (hparams.n_ff_shexp == 0) {
|
||||
hparams.n_ff_shexp = hparams.n_ff_exp * std::max(1u, hparams.n_expert_shared);
|
||||
}
|
||||
|
||||
GGML_ASSERT(hparams.kda_safe_gate);
|
||||
GGML_ASSERT(hparams.kda_gate_lower_bound < 0.0f);
|
||||
|
||||
for (uint32_t il = 0; il < hparams.n_layer(); ++il) {
|
||||
hparams.is_recr_impl[il] = hparams.n_head_kv(il) == 0;
|
||||
}
|
||||
|
||||
switch (hparams.n_layer()) {
|
||||
case 24: type = hparams.n_embd == 1536 && hparams.n_expert == 128 ? LLM_TYPE_7_9B_A1_3B : LLM_TYPE_UNKNOWN; break;
|
||||
case 42: type = hparams.n_embd == 2560 && hparams.n_expert == 512 ? LLM_TYPE_124B_A5_1B : LLM_TYPE_UNKNOWN; break;
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
}
|
||||
|
||||
void llama_model_bailingmoe3::load_arch_tensors(llama_model_loader & ml) {
|
||||
LLAMA_LOAD_LOCALS;
|
||||
|
||||
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);
|
||||
|
||||
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 == nullptr) {
|
||||
output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, TENSOR_DUPLICATED);
|
||||
}
|
||||
|
||||
const int64_t head_dim = hparams.n_embd_head_kda;
|
||||
const int64_t d_inner = head_dim * n_head;
|
||||
const int64_t d_conv = hparams.ssm_d_conv;
|
||||
const int64_t kv_lora_rank = hparams.n_lora_kv;
|
||||
const int64_t q_lora_rank = hparams.n_lora_q;
|
||||
const int64_t qk_rope_head_dim = hparams.n_rot();
|
||||
const int64_t qk_head_dim = hparams.n_embd_head_k_mla();
|
||||
const int64_t v_head_dim = hparams.n_embd_head_v_mla();
|
||||
|
||||
const bool mtp_only = (hparams.n_layer_nextn > 0) && (ml.get_weight("blk.0.attn_norm.weight") == nullptr);
|
||||
const std::string mtp_probe = "blk." + std::to_string(n_layer) + ".nextn.eh_proj.weight";
|
||||
const bool trunk_only = (hparams.n_layer_nextn > 0) && (ml.get_weight(mtp_probe.c_str()) == nullptr);
|
||||
const int trunk_flags = mtp_only ? TENSOR_NOT_REQUIRED : 0;
|
||||
int mtp_flags = trunk_only ? TENSOR_NOT_REQUIRED : 0;
|
||||
|
||||
if (!ml.load_mtp) {
|
||||
mtp_flags |= TENSOR_SKIP;
|
||||
}
|
||||
|
||||
for (int il = 0; il < n_layer; ++il) {
|
||||
auto & layer = layers[il];
|
||||
|
||||
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", il), { n_embd }, trunk_flags);
|
||||
|
||||
if (hparams.is_recr(il)) {
|
||||
layer.ssm_q_conv = create_tensor(tn(LLM_TENSOR_SSM_CONV1D_Q, "weight", il), { d_conv, 1, d_inner, 1 }, trunk_flags);
|
||||
layer.ssm_k_conv = create_tensor(tn(LLM_TENSOR_SSM_CONV1D_K, "weight", il), { d_conv, 1, d_inner, 1 }, trunk_flags);
|
||||
layer.ssm_v_conv = create_tensor(tn(LLM_TENSOR_SSM_CONV1D_V, "weight", il), { d_conv, 1, d_inner, 1 }, trunk_flags);
|
||||
|
||||
create_tensor_qkv(layer, il, n_embd, d_inner, d_inner, d_inner, trunk_flags);
|
||||
layer.ssm_f_a = create_tensor(tn(LLM_TENSOR_SSM_F_A, "weight", il), { n_embd, d_inner }, trunk_flags);
|
||||
layer.ssm_beta = create_tensor(tn(LLM_TENSOR_SSM_BETA, "weight", il), { n_embd, n_head }, trunk_flags);
|
||||
layer.ssm_a = create_tensor(tn(LLM_TENSOR_SSM_A, il), { 1, n_head }, trunk_flags);
|
||||
layer.ssm_dt_b = create_tensor(tn(LLM_TENSOR_SSM_DT, "bias", il), { d_inner }, trunk_flags);
|
||||
layer.ssm_g_a = create_tensor(tn(LLM_TENSOR_SSM_G_A, "weight", il), { n_embd, d_inner }, trunk_flags);
|
||||
layer.ssm_o_norm = create_tensor(tn(LLM_TENSOR_SSM_NORM, "weight", il), { head_dim }, trunk_flags);
|
||||
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", il), { d_inner, n_embd }, trunk_flags);
|
||||
} else {
|
||||
if (q_lora_rank > 0) {
|
||||
layer.wq_a = create_tensor(tn(LLM_TENSOR_ATTN_Q_A, "weight", il), { n_embd, q_lora_rank }, trunk_flags);
|
||||
layer.attn_q_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_A_NORM, "weight", il), { q_lora_rank }, trunk_flags);
|
||||
layer.wq_b = create_tensor(tn(LLM_TENSOR_ATTN_Q_B, "weight", il), { q_lora_rank, n_head * qk_head_dim }, trunk_flags);
|
||||
} else {
|
||||
layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "weight", il), { n_embd, n_head * qk_head_dim }, trunk_flags);
|
||||
}
|
||||
layer.wkv_a_mqa = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_MQA, "weight", il), { n_embd, kv_lora_rank + qk_rope_head_dim }, trunk_flags);
|
||||
layer.attn_kv_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_NORM, "weight", il), { kv_lora_rank }, trunk_flags);
|
||||
layer.wk_b = create_tensor(tn(LLM_TENSOR_ATTN_K_B, "weight", il), { qk_head_dim - qk_rope_head_dim, kv_lora_rank, n_head }, trunk_flags);
|
||||
layer.wv_b = create_tensor(tn(LLM_TENSOR_ATTN_V_B, "weight", il), { kv_lora_rank, v_head_dim, n_head }, trunk_flags);
|
||||
layer.wqkv_gate = create_tensor(tn(LLM_TENSOR_ATTN_GATE, "weight", il), { n_embd, n_head }, trunk_flags);
|
||||
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", il), { n_head * v_head_dim, n_embd }, trunk_flags);
|
||||
}
|
||||
|
||||
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", il), { n_embd }, trunk_flags);
|
||||
if ((uint32_t) il < hparams.n_layer_dense_lead) {
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", il), { n_embd, n_ff }, trunk_flags);
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", il), { n_embd, n_ff }, trunk_flags);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", il), { n_ff, n_embd }, trunk_flags);
|
||||
} else {
|
||||
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", il), { n_embd, n_expert }, trunk_flags);
|
||||
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", il), { n_expert }, trunk_flags);
|
||||
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, trunk_flags);
|
||||
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, trunk_flags);
|
||||
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp, n_embd, n_expert }, trunk_flags);
|
||||
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, trunk_flags);
|
||||
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, trunk_flags);
|
||||
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", il), { hparams.n_ff_shexp, n_embd }, trunk_flags);
|
||||
}
|
||||
}
|
||||
|
||||
for (int il = n_layer; il < n_layer_all; ++il) {
|
||||
auto & layer = layers[il];
|
||||
const int flags = mtp_flags;
|
||||
|
||||
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", il), { n_embd }, flags);
|
||||
if (q_lora_rank > 0) {
|
||||
layer.wq_a = create_tensor(tn(LLM_TENSOR_ATTN_Q_A, "weight", il), { n_embd, q_lora_rank }, flags);
|
||||
layer.attn_q_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_A_NORM, "weight", il), { q_lora_rank }, flags);
|
||||
layer.wq_b = create_tensor(tn(LLM_TENSOR_ATTN_Q_B, "weight", il), { q_lora_rank, n_head * qk_head_dim }, flags);
|
||||
} else {
|
||||
layer.wq = create_tensor(tn(LLM_TENSOR_ATTN_Q, "weight", il), { n_embd, n_head * qk_head_dim }, flags);
|
||||
}
|
||||
layer.wkv_a_mqa = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_MQA, "weight", il), { n_embd, kv_lora_rank + qk_rope_head_dim }, flags);
|
||||
layer.attn_kv_a_norm = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_NORM, "weight", il), { kv_lora_rank }, flags);
|
||||
layer.wk_b = create_tensor(tn(LLM_TENSOR_ATTN_K_B, "weight", il), { qk_head_dim - qk_rope_head_dim, kv_lora_rank, n_head }, flags);
|
||||
layer.wv_b = create_tensor(tn(LLM_TENSOR_ATTN_V_B, "weight", il), { kv_lora_rank, v_head_dim, n_head }, flags);
|
||||
layer.wqkv_gate = create_tensor(tn(LLM_TENSOR_ATTN_GATE, "weight", il), { n_embd, n_head }, flags);
|
||||
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", il), { n_head * v_head_dim, n_embd }, flags);
|
||||
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", il), { n_embd }, flags);
|
||||
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", il), { n_embd, n_expert }, flags);
|
||||
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", il), { n_expert }, flags);
|
||||
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, flags);
|
||||
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", il), { n_embd, hparams.n_ff_exp, n_expert }, flags);
|
||||
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", il), { hparams.n_ff_exp, n_embd, n_expert }, flags);
|
||||
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, flags);
|
||||
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", il), { n_embd, hparams.n_ff_shexp }, flags);
|
||||
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", il), { hparams.n_ff_shexp, n_embd }, flags);
|
||||
layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", il), { 2 * n_embd, n_embd }, flags);
|
||||
layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", il), { n_embd }, flags);
|
||||
layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", il), { n_embd }, flags);
|
||||
layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", il), { n_embd }, flags);
|
||||
}
|
||||
}
|
||||
|
||||
std::unique_ptr<llm_graph_context> llama_model_bailingmoe3::build_arch_graph(const llm_graph_params & params) const {
|
||||
if (params.gtype == LLM_GRAPH_TYPE_DECODER_MTP) {
|
||||
return std::make_unique<graph_mtp>(*this, params);
|
||||
}
|
||||
return std::make_unique<graph>(*this, params);
|
||||
}
|
||||
|
||||
static ggml_tensor * bailingmoe3_causal_conv1d(
|
||||
ggml_cgraph * gf,
|
||||
ggml_context * ctx0,
|
||||
ggml_tensor * conv_states_all,
|
||||
ggml_tensor * conv_state_all,
|
||||
int64_t qkv,
|
||||
ggml_tensor * x,
|
||||
ggml_tensor * proj_w,
|
||||
ggml_tensor * conv_w,
|
||||
int64_t d_conv,
|
||||
int64_t head_dim,
|
||||
int64_t n_head,
|
||||
int64_t n_seq_tokens,
|
||||
int64_t n_seqs,
|
||||
int64_t n_tokens,
|
||||
int64_t cache_head) {
|
||||
const int64_t d_inner = head_dim * n_head;
|
||||
const int64_t conv_state_size = (d_conv - 1) * d_inner;
|
||||
const int64_t total_state_size = 3 * conv_state_size;
|
||||
|
||||
ggml_tensor * conv_state = ggml_view_3d(ctx0, conv_state_all, d_conv - 1, d_inner, n_seqs,
|
||||
(d_conv - 1) * ggml_element_size(conv_state_all),
|
||||
total_state_size * ggml_element_size(conv_state_all),
|
||||
qkv * conv_state_size * ggml_element_size(conv_state_all));
|
||||
|
||||
ggml_tensor * x_proj = ggml_mul_mat(ctx0, proj_w, x);
|
||||
x_proj = ggml_reshape_3d(ctx0, x_proj, d_inner, n_seq_tokens, n_seqs);
|
||||
ggml_tensor * conv_x = ggml_concat(ctx0, conv_state, ggml_transpose(ctx0, x_proj), 0);
|
||||
|
||||
ggml_tensor * last_conv_x = ggml_view_3d(ctx0, conv_x, d_conv - 1, d_inner, n_seqs,
|
||||
conv_x->nb[1], conv_x->nb[2], n_seq_tokens * conv_x->nb[0]);
|
||||
ggml_build_forward_expand(gf, ggml_cpy(ctx0, last_conv_x,
|
||||
ggml_view_3d(ctx0, conv_states_all, d_conv - 1, d_inner, n_seqs,
|
||||
(d_conv - 1) * ggml_element_size(conv_states_all),
|
||||
total_state_size * ggml_element_size(conv_states_all),
|
||||
(cache_head * total_state_size + qkv * conv_state_size) * ggml_element_size(conv_states_all))));
|
||||
|
||||
ggml_tensor * conv_weight = ggml_reshape_2d(ctx0, conv_w, d_conv, d_inner);
|
||||
ggml_tensor * out = ggml_ssm_conv(ctx0, conv_x, conv_weight);
|
||||
out = ggml_silu(ctx0, ggml_reshape_2d(ctx0, out, d_inner, n_tokens));
|
||||
return ggml_reshape_4d(ctx0, out, head_dim, n_head, n_seq_tokens, n_seqs);
|
||||
}
|
||||
|
||||
llama_model_bailingmoe3::graph::graph(const llama_model & model, const llm_graph_params & params) :
|
||||
llm_build_delta_net_base(params), model(model) {
|
||||
ggml_tensor * inpL = build_inp_embd(model.tok_embd);
|
||||
cb(inpL, "model.input_embed", -1);
|
||||
|
||||
auto * inp = build_inp_mem_hybrid_k();
|
||||
auto * inp_rs = inp->get_recr();
|
||||
auto * inp_attn = inp->get_attn();
|
||||
|
||||
ggml_tensor * inp_pos = build_inp_pos();
|
||||
ggml_tensor * inp_out_ids = build_inp_out_ids();
|
||||
|
||||
const int64_t n_head = hparams.n_head();
|
||||
const int64_t head_dim = hparams.n_embd_head_kda;
|
||||
const int64_t d_inner = n_head * head_dim;
|
||||
const int64_t d_conv = hparams.ssm_d_conv;
|
||||
const int64_t n_seqs = ubatch.n_seqs;
|
||||
const int64_t n_seq_tokens = ubatch.n_seq_tokens;
|
||||
const int64_t qk_head_dim = hparams.n_embd_head_k_mla();
|
||||
const int64_t v_head_dim = hparams.n_embd_head_v_mla();
|
||||
const int64_t qk_rope_head_dim = hparams.n_rot();
|
||||
const int64_t qk_nope_head_dim = qk_head_dim - qk_rope_head_dim;
|
||||
const int64_t kv_lora_rank = hparams.n_lora_kv;
|
||||
const float kq_scale = 1.0f / sqrtf((float) qk_head_dim);
|
||||
|
||||
GGML_ASSERT(n_seqs > 0);
|
||||
GGML_ASSERT(ubatch.equal_seqs());
|
||||
GGML_ASSERT(ubatch.n_tokens == n_seq_tokens * n_seqs);
|
||||
|
||||
for (int il = 0; il < n_layer; ++il) {
|
||||
const auto & layer = model.layers[il];
|
||||
ggml_tensor * inpSA = inpL;
|
||||
ggml_tensor * cur = build_norm(inpL, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
|
||||
cb(cur, "attn_norm", il);
|
||||
|
||||
if (hparams.is_recr(il)) {
|
||||
const auto * mctx_cur = inp_rs->mctx;
|
||||
const auto cache_head = mctx_cur->get_head();
|
||||
ggml_tensor * conv_states_all = mctx_cur->get_r_l(il);
|
||||
ggml_tensor * conv_state_all = build_rs(inp_rs, conv_states_all, hparams.n_embd_r(), n_seqs);
|
||||
|
||||
ggml_tensor * q = bailingmoe3_causal_conv1d(
|
||||
gf, ctx0, conv_states_all, conv_state_all, 0, cur, layer.wq, layer.ssm_q_conv,
|
||||
d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
|
||||
ggml_tensor * k = bailingmoe3_causal_conv1d(
|
||||
gf, ctx0, conv_states_all, conv_state_all, 1, cur, layer.wk, layer.ssm_k_conv,
|
||||
d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
|
||||
ggml_tensor * v = bailingmoe3_causal_conv1d(
|
||||
gf, ctx0, conv_states_all, conv_state_all, 2, cur, layer.wv, layer.ssm_v_conv,
|
||||
d_conv, head_dim, n_head, n_seq_tokens, n_seqs, n_tokens, cache_head);
|
||||
|
||||
ggml_tensor * gate = ggml_mul_mat(ctx0, layer.ssm_f_a, cur);
|
||||
gate = ggml_add(ctx0, gate, layer.ssm_dt_b);
|
||||
gate = ggml_reshape_3d(ctx0, gate, head_dim, n_head, n_tokens);
|
||||
ggml_tensor * a = ggml_reshape_3d(ctx0, layer.ssm_a, 1, n_head, 1);
|
||||
gate = ggml_scale(ctx0, ggml_sigmoid(ctx0, ggml_mul(ctx0, gate, a)), hparams.kda_gate_lower_bound);
|
||||
gate = ggml_reshape_4d(ctx0, gate, head_dim, n_head, n_seq_tokens, n_seqs);
|
||||
cb(gate, "kda_gate", il);
|
||||
|
||||
ggml_tensor * beta = ggml_mul_mat(ctx0, layer.ssm_beta, cur);
|
||||
beta = ggml_sigmoid(ctx0, ggml_reshape_4d(ctx0, beta, 1, n_head, n_seq_tokens, n_seqs));
|
||||
|
||||
q = ggml_l2_norm(ctx0, q, hparams.f_norm_rms_eps);
|
||||
k = ggml_l2_norm(ctx0, k, hparams.f_norm_rms_eps);
|
||||
|
||||
ggml_tensor * states_all = mctx_cur->get_s_l(il);
|
||||
ggml_tensor * state = build_rs(inp_rs, states_all, hparams.n_embd_s(), n_seqs);
|
||||
state = ggml_reshape_4d(ctx0, state, head_dim, head_dim, n_head, n_seqs);
|
||||
|
||||
auto result = build_delta_net(q, k, v, gate, beta, state, il);
|
||||
ggml_tensor * out = ggml_cont(ctx0, result.first);
|
||||
ggml_build_forward_expand(gf, ggml_cpy(ctx0, result.second,
|
||||
ggml_view_1d(ctx0, states_all, hparams.n_embd_s() * n_seqs,
|
||||
cache_head * hparams.n_embd_s() * ggml_element_size(states_all))));
|
||||
|
||||
ggml_tensor * out_gate = ggml_mul_mat(ctx0, layer.ssm_g_a, cur);
|
||||
out_gate = ggml_reshape_3d(ctx0, out_gate, head_dim, n_head, n_tokens);
|
||||
out = ggml_reshape_3d(ctx0, out, head_dim, n_head, n_tokens);
|
||||
out = build_norm(out, layer.ssm_o_norm, nullptr, LLM_NORM_RMS, il);
|
||||
out = ggml_mul(ctx0, out, ggml_sigmoid(ctx0, out_gate));
|
||||
cur = ggml_mul_mat(ctx0, layer.wo, ggml_cont_2d(ctx0, out, d_inner, n_tokens));
|
||||
cb(cur, "kda_out", il);
|
||||
} else {
|
||||
ggml_tensor * attn_input = cur;
|
||||
ggml_tensor * q_all;
|
||||
if (layer.wq_a) {
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq_a, cur);
|
||||
cb(q_all, "q_a", il);
|
||||
q_all = build_norm(q_all, layer.attn_q_a_norm, nullptr, LLM_NORM_RMS, il);
|
||||
cb(q_all, "q_a_norm", il);
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq_b, q_all);
|
||||
cb(q_all, "q_b", il);
|
||||
} else {
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq, cur);
|
||||
}
|
||||
ggml_tensor * q_nope = ggml_view_3d(ctx0, q_all, qk_nope_head_dim, n_head, n_tokens,
|
||||
ggml_row_size(q_all->type, qk_head_dim),
|
||||
ggml_row_size(q_all->type, qk_head_dim) * n_head, 0);
|
||||
ggml_tensor * q_pe = ggml_view_3d(ctx0, q_all, qk_rope_head_dim, n_head, n_tokens,
|
||||
ggml_row_size(q_all->type, qk_head_dim),
|
||||
ggml_row_size(q_all->type, qk_head_dim) * n_head,
|
||||
ggml_row_size(q_all->type, qk_nope_head_dim));
|
||||
|
||||
ggml_tensor * kv_all = ggml_mul_mat(ctx0, layer.wkv_a_mqa, cur);
|
||||
ggml_tensor * kv = ggml_view_2d(ctx0, kv_all, kv_lora_rank, n_tokens,
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim), 0);
|
||||
ggml_tensor * k_pe = ggml_view_3d(ctx0, kv_all, qk_rope_head_dim, 1, n_tokens,
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
|
||||
ggml_row_size(kv_all->type, kv_lora_rank));
|
||||
|
||||
q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
kv = build_norm(kv, layer.attn_kv_a_norm, nullptr, LLM_NORM_RMS, il);
|
||||
|
||||
q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);
|
||||
q_nope = ggml_mul_mat(ctx0, layer.wk_b, q_nope);
|
||||
q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);
|
||||
|
||||
ggml_tensor * q = ggml_concat(ctx0, q_nope, q_pe, 0);
|
||||
kv = ggml_reshape_3d(ctx0, kv, kv_lora_rank, 1, n_tokens);
|
||||
ggml_tensor * k = ggml_concat(ctx0, kv, k_pe, 0);
|
||||
|
||||
cur = build_attn(inp_attn, nullptr, nullptr, nullptr,
|
||||
q, k, kv, nullptr, nullptr, layer.wv_b, kq_scale, il);
|
||||
|
||||
ggml_tensor * attn_gate = ggml_mul_mat(ctx0, layer.wqkv_gate, attn_input);
|
||||
attn_gate = ggml_sigmoid(ctx0, ggml_reshape_3d(ctx0, attn_gate, 1, n_head, n_tokens));
|
||||
cur = ggml_reshape_3d(ctx0, cur, v_head_dim, n_head, n_tokens);
|
||||
cur = ggml_mul(ctx0, cur, attn_gate);
|
||||
cur = ggml_mul_mat(ctx0, layer.wo, ggml_cont_2d(ctx0, cur, v_head_dim * n_head, n_tokens));
|
||||
cb(cur, "mla_out", il);
|
||||
}
|
||||
|
||||
if (il == n_layer - 1 && inp_out_ids && cparams.embeddings_nextn_masked) {
|
||||
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);
|
||||
cur = build_norm(ffn_inp, layer.ffn_norm, nullptr, LLM_NORM_RMS, il);
|
||||
|
||||
if ((uint32_t) il < hparams.n_layer_dense_lead) {
|
||||
cur = build_ffn(cur,
|
||||
layer.ffn_up, nullptr, nullptr,
|
||||
layer.ffn_gate, nullptr, nullptr,
|
||||
layer.ffn_down, nullptr, nullptr,
|
||||
nullptr, LLM_FFN_SILU, LLM_FFN_PAR, il);
|
||||
} else {
|
||||
ggml_tensor * moe = build_moe_ffn(cur,
|
||||
layer.ffn_gate_inp,
|
||||
layer.ffn_up_exps,
|
||||
layer.ffn_gate_exps,
|
||||
layer.ffn_down_exps,
|
||||
layer.ffn_exp_probs_b,
|
||||
n_expert, n_expert_used,
|
||||
LLM_FFN_SILU,
|
||||
hparams.expert_weights_norm,
|
||||
hparams.expert_weights_scale,
|
||||
(llama_expert_gating_func_type) hparams.expert_gating_func,
|
||||
il);
|
||||
ggml_tensor * shared = build_ffn(cur,
|
||||
layer.ffn_up_shexp, nullptr, nullptr,
|
||||
layer.ffn_gate_shexp, nullptr, nullptr,
|
||||
layer.ffn_down_shexp, nullptr, nullptr,
|
||||
nullptr, LLM_FFN_SILU, LLM_FFN_PAR, il);
|
||||
cur = ggml_add(ctx0, moe, shared);
|
||||
}
|
||||
|
||||
cur = ggml_add(ctx0, cur, ffn_inp);
|
||||
cur = build_cvec(cur, il);
|
||||
cb(cur, "l_out", il);
|
||||
inpL = cur;
|
||||
}
|
||||
|
||||
ggml_tensor * cur = build_norm(inpL, model.output_norm, nullptr, LLM_NORM_RMS, -1);
|
||||
cb(cur, "h_nextn", -1);
|
||||
res->t_h_nextn = cur;
|
||||
|
||||
if (!cparams.embeddings_nextn_masked && inp_out_ids) {
|
||||
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
|
||||
}
|
||||
|
||||
cb(cur, "result_norm", -1);
|
||||
res->t_embd = cur;
|
||||
|
||||
cur = ggml_mul_mat(ctx0, model.output, cur);
|
||||
cb(cur, "result_output", -1);
|
||||
res->t_logits = cur;
|
||||
ggml_build_forward_expand(gf, cur);
|
||||
}
|
||||
|
||||
llama_model_bailingmoe3::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params) :
|
||||
llm_graph_context(params) {
|
||||
GGML_ASSERT(hparams.n_layer_nextn == 1 && "BailingMoE3 MTP requires one NextN layer");
|
||||
|
||||
const int il = hparams.n_layer() + cparams.nextn_layer_offset;
|
||||
GGML_ASSERT(cparams.nextn_layer_offset >= 0 &&
|
||||
cparams.nextn_layer_offset < (int) hparams.n_layer_nextn &&
|
||||
"nextn_layer_offset out of range");
|
||||
const auto & layer = model.layers[il];
|
||||
|
||||
GGML_ASSERT(layer.nextn.eh_proj && "MTP block missing nextn.eh_proj");
|
||||
GGML_ASSERT(layer.nextn.enorm && "MTP block missing nextn.enorm");
|
||||
GGML_ASSERT(layer.nextn.hnorm && "MTP block missing nextn.hnorm");
|
||||
GGML_ASSERT(layer.nextn.shared_head_norm && "MTP block missing final norm");
|
||||
|
||||
const int64_t n_head = hparams.n_head();
|
||||
const int64_t qk_head_dim = hparams.n_embd_head_k_mla();
|
||||
const int64_t v_head_dim = hparams.n_embd_head_v_mla();
|
||||
const int64_t qk_rope_head_dim = hparams.n_rot();
|
||||
const int64_t qk_nope_head_dim = qk_head_dim - qk_rope_head_dim;
|
||||
const int64_t kv_lora_rank = hparams.n_lora_kv;
|
||||
const float kq_scale = 1.0f / sqrtf((float) qk_head_dim);
|
||||
|
||||
auto inp = std::make_unique<llm_graph_input_embd>(hparams.n_embd);
|
||||
inp->tokens = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_tokens);
|
||||
ggml_set_input(inp->tokens);
|
||||
inp->embd = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, hparams.n_embd, n_tokens);
|
||||
ggml_set_input(inp->embd);
|
||||
ggml_set_name(inp->embd, "mtp_h_input");
|
||||
|
||||
ggml_tensor * tok_embd = ggml_get_rows(ctx0, model.tok_embd, inp->tokens);
|
||||
ggml_tensor * h_norm = build_norm(inp->embd, layer.nextn.hnorm, nullptr, LLM_NORM_RMS, il);
|
||||
ggml_tensor * e_norm = build_norm(tok_embd, layer.nextn.enorm, nullptr, LLM_NORM_RMS, il);
|
||||
ggml_tensor * cur = ggml_mul_mat(ctx0, layer.nextn.eh_proj, ggml_concat(ctx0, e_norm, h_norm, 0));
|
||||
cb(cur, "mtp_eh_proj", il);
|
||||
|
||||
res->add_input(std::move(inp));
|
||||
|
||||
ggml_tensor * inp_pos = build_inp_pos();
|
||||
ggml_tensor * inp_out_ids = build_inp_out_ids();
|
||||
auto * inp_attn = build_attn_inp_k();
|
||||
|
||||
ggml_tensor * inpSA = cur;
|
||||
cur = build_norm(cur, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
|
||||
ggml_tensor * attn_input = cur;
|
||||
|
||||
ggml_tensor * q_all;
|
||||
if (layer.wq_a) {
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq_a, cur);
|
||||
cb(q_all, "q_a", il);
|
||||
q_all = build_norm(q_all, layer.attn_q_a_norm, nullptr, LLM_NORM_RMS, il);
|
||||
cb(q_all, "q_a_norm", il);
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq_b, q_all);
|
||||
cb(q_all, "q_b", il);
|
||||
} else {
|
||||
q_all = ggml_mul_mat(ctx0, layer.wq, cur);
|
||||
}
|
||||
ggml_tensor * q_nope = ggml_view_3d(ctx0, q_all, qk_nope_head_dim, n_head, n_tokens,
|
||||
ggml_row_size(q_all->type, qk_head_dim),
|
||||
ggml_row_size(q_all->type, qk_head_dim) * n_head, 0);
|
||||
ggml_tensor * q_pe = ggml_view_3d(ctx0, q_all, qk_rope_head_dim, n_head, n_tokens,
|
||||
ggml_row_size(q_all->type, qk_head_dim),
|
||||
ggml_row_size(q_all->type, qk_head_dim) * n_head,
|
||||
ggml_row_size(q_all->type, qk_nope_head_dim));
|
||||
|
||||
ggml_tensor * kv_all = ggml_mul_mat(ctx0, layer.wkv_a_mqa, cur);
|
||||
ggml_tensor * kv = ggml_view_2d(ctx0, kv_all, kv_lora_rank, n_tokens,
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim), 0);
|
||||
ggml_tensor * k_pe = ggml_view_3d(ctx0, kv_all, qk_rope_head_dim, 1, n_tokens,
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
|
||||
ggml_row_size(kv_all->type, kv_lora_rank + qk_rope_head_dim),
|
||||
ggml_row_size(kv_all->type, kv_lora_rank));
|
||||
|
||||
q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
kv = build_norm(kv, layer.attn_kv_a_norm, nullptr, LLM_NORM_RMS, il);
|
||||
|
||||
q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);
|
||||
q_nope = ggml_mul_mat(ctx0, layer.wk_b, q_nope);
|
||||
q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3);
|
||||
|
||||
ggml_tensor * q = ggml_concat(ctx0, q_nope, q_pe, 0);
|
||||
kv = ggml_reshape_3d(ctx0, kv, kv_lora_rank, 1, n_tokens);
|
||||
ggml_tensor * k = ggml_concat(ctx0, kv, k_pe, 0);
|
||||
|
||||
cur = build_attn(inp_attn, nullptr, nullptr, nullptr,
|
||||
q, k, kv, nullptr, nullptr, layer.wv_b, kq_scale, il);
|
||||
|
||||
ggml_tensor * attn_gate = ggml_mul_mat(ctx0, layer.wqkv_gate, attn_input);
|
||||
attn_gate = ggml_sigmoid(ctx0, ggml_reshape_3d(ctx0, attn_gate, 1, n_head, n_tokens));
|
||||
cur = ggml_reshape_3d(ctx0, cur, v_head_dim, n_head, n_tokens);
|
||||
cur = ggml_mul(ctx0, cur, attn_gate);
|
||||
cur = ggml_mul_mat(ctx0, layer.wo, ggml_cont_2d(ctx0, cur, v_head_dim * n_head, n_tokens));
|
||||
|
||||
ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
|
||||
cur = build_norm(ffn_inp, layer.ffn_norm, nullptr, LLM_NORM_RMS, il);
|
||||
|
||||
ggml_tensor * moe = build_moe_ffn(cur,
|
||||
layer.ffn_gate_inp,
|
||||
layer.ffn_up_exps,
|
||||
layer.ffn_gate_exps,
|
||||
layer.ffn_down_exps,
|
||||
layer.ffn_exp_probs_b,
|
||||
n_expert, n_expert_used,
|
||||
LLM_FFN_SILU,
|
||||
hparams.expert_weights_norm,
|
||||
hparams.expert_weights_scale,
|
||||
(llama_expert_gating_func_type) hparams.expert_gating_func,
|
||||
il);
|
||||
ggml_tensor * shared = build_ffn(cur,
|
||||
layer.ffn_up_shexp, nullptr, nullptr,
|
||||
layer.ffn_gate_shexp, nullptr, nullptr,
|
||||
layer.ffn_down_shexp, nullptr, nullptr,
|
||||
nullptr, LLM_FFN_SILU, LLM_FFN_PAR, il);
|
||||
cur = ggml_add(ctx0, moe, shared);
|
||||
cur = ggml_add(ctx0, cur, ffn_inp);
|
||||
cur = build_norm(cur, layer.nextn.shared_head_norm, nullptr, LLM_NORM_RMS, -1);
|
||||
|
||||
cb(cur, "h_nextn", -1);
|
||||
res->t_h_nextn = cur;
|
||||
|
||||
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
|
||||
cur = ggml_mul_mat(ctx0, model.output, cur);
|
||||
cb(cur, "result_output", -1);
|
||||
res->t_logits = cur;
|
||||
ggml_build_forward_expand(gf, cur);
|
||||
}
|
||||
@@ -1784,6 +1784,25 @@ struct llama_model_bailingmoe2 : public llama_model_base {
|
||||
};
|
||||
|
||||
|
||||
struct llama_model_bailingmoe3 : public llama_model_base {
|
||||
llama_model_bailingmoe3(const struct llama_model_params & params) : llama_model_base(params) {}
|
||||
void load_arch_hparams(llama_model_loader & ml) override;
|
||||
void load_arch_tensors(llama_model_loader & ml) override;
|
||||
|
||||
struct graph : public llm_build_delta_net_base {
|
||||
graph(const llama_model & model, const llm_graph_params & params);
|
||||
|
||||
const llama_model & model;
|
||||
};
|
||||
|
||||
struct graph_mtp : public llm_graph_context {
|
||||
graph_mtp(const llama_model & model, const llm_graph_params & params);
|
||||
};
|
||||
|
||||
std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
|
||||
};
|
||||
|
||||
|
||||
struct llama_model_seed_oss : public llama_model_base {
|
||||
llama_model_seed_oss(const struct llama_model_params & params) : llama_model_base(params) {}
|
||||
void load_arch_hparams(llama_model_loader & ml) override;
|
||||
|
||||
@@ -90,6 +90,7 @@ static void test_normalize_quotes_with_embedded_quotes(testing & t);
|
||||
|
||||
// TAG_WITH_TAGGED argument parsing tests
|
||||
static void test_tagged_args_with_embedded_quotes(testing & t);
|
||||
static void test_bailing_v3_tool_format(testing & t);
|
||||
|
||||
static void test_role_markers_all_templates(testing & t);
|
||||
|
||||
@@ -118,6 +119,7 @@ int main(int argc, char * argv[]) {
|
||||
t.test("standard_json_tools", test_standard_json_tools_formats);
|
||||
t.test("normalize_quotes_to_json", test_normalize_quotes_to_json);
|
||||
t.test("tagged_args_embedded_quotes", test_tagged_args_with_embedded_quotes);
|
||||
t.test("bailing_v3", test_bailing_v3_tool_format);
|
||||
t.test("role_markers_all_templates", test_role_markers_all_templates);
|
||||
|
||||
return t.summary();
|
||||
@@ -2081,6 +2083,68 @@ static void test_role_markers_all_templates(testing & t) {
|
||||
}
|
||||
}
|
||||
|
||||
static void test_bailing_v3_tool_format(testing & t) {
|
||||
const std::string template_source = R"JINJA(
|
||||
{# Bailing V3 chat template #}
|
||||
{%- if tools %}{{ tools | tojson }}{%- endif %}
|
||||
{%- for message in messages %}
|
||||
{%- if message.role == "user" %}
|
||||
{{- '<role>HUMAN</role>' + message.content + '<|role_end|>' }}
|
||||
{%- elif message.role == "assistant" %}
|
||||
{{- '<role>ASSISTANT</role>' }}
|
||||
{%- if message.tool_calls %}
|
||||
{%- for tool_call in message.tool_calls %}
|
||||
{%- set tc = tool_call.function %}
|
||||
{{- '<tool_call>' + tc.name }}
|
||||
{%- for k, v in tc.arguments.items() %}
|
||||
{{- '<arg_key>' + k + '</arg_key>' }}
|
||||
{{- '\n<arg_value>' + v + '</arg_value>' }}
|
||||
{%- endfor %}
|
||||
{{- '\n</tool_call>' }}
|
||||
{%- endfor %}
|
||||
{%- endif %}
|
||||
{{- '<|role_end|>' }}
|
||||
{%- endif %}
|
||||
{%- endfor %}
|
||||
{%- if add_generation_prompt %}{{- '<role>ASSISTANT</role>' }}{%- endif %}
|
||||
)JINJA";
|
||||
|
||||
common_chat_template tmpl(template_source, "", "");
|
||||
struct autoparser analysis;
|
||||
analysis.analyze_template(tmpl);
|
||||
|
||||
t.assert_equal("arg_value_suffix", "</arg_value>", analysis.tools.arguments.value_suffix);
|
||||
t.assert_true("intertag whitespace", analysis.tools.arguments.tolerate_intertag_whitespace);
|
||||
|
||||
generation_params inputs;
|
||||
inputs.tools = json::array({
|
||||
{
|
||||
{ "type", "function" },
|
||||
{ "function", {
|
||||
{ "name", "test_function_name" },
|
||||
{ "parameters", {
|
||||
{ "type", "object" },
|
||||
{ "properties", {
|
||||
{ "param1", { { "type", "string" } } },
|
||||
{ "param2", { { "type", "string" } } },
|
||||
} },
|
||||
} },
|
||||
} },
|
||||
},
|
||||
});
|
||||
inputs.reasoning_format = COMMON_REASONING_FORMAT_NONE;
|
||||
auto parser = analysis.build_parser(inputs, "");
|
||||
const std::string output =
|
||||
"<tool_call>test_function_name\n"
|
||||
"<arg_key>param1</arg_key>\n"
|
||||
"<arg_value>value1</arg_value>"
|
||||
"<arg_key>param2</arg_key>\n"
|
||||
"<arg_value>value2</arg_value>\n"
|
||||
"</tool_call>";
|
||||
common_peg_parse_context ctx(output, COMMON_PEG_PARSE_FLAG_LENIENT);
|
||||
t.assert_true("multi-argument tool call", parser.parse(ctx).success());
|
||||
}
|
||||
|
||||
// Test that reproduces the Seed-OSS template issue with embedded quotes
|
||||
static void test_tagged_args_with_embedded_quotes(testing & t) {
|
||||
json tools = build_edit_tool();
|
||||
@@ -2198,4 +2262,3 @@ static void test_tagged_args_with_embedded_quotes(testing & t) {
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -105,6 +105,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
|| arch == LLM_ARCH_DEEPSEEK32
|
||||
|| arch == LLM_ARCH_GLM_DSA
|
||||
|| arch == LLM_ARCH_KIMI_LINEAR
|
||||
|| arch == LLM_ARCH_BAILINGMOE3
|
||||
|| arch == LLM_ARCH_KIMI_K3
|
||||
|| arch == LLM_ARCH_MISTRAL4) {
|
||||
n_embd = 128;
|
||||
@@ -146,7 +147,8 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
ms.add_kv(LLM_KV_FULL_ATTENTION_INTERVAL, uint32_t(2));
|
||||
|
||||
if (arch == LLM_ARCH_PLAMO2 || arch == LLM_ARCH_JAMBA || arch == LLM_ARCH_NEMOTRON_H || arch == LLM_ARCH_NEMOTRON_H_MOE ||
|
||||
arch == LLM_ARCH_GRANITE_HYBRID || arch == LLM_ARCH_LFM2 || arch == LLM_ARCH_LFM2MOE || arch == LLM_ARCH_KIMI_LINEAR || arch == LLM_ARCH_KIMI_K3) {
|
||||
arch == LLM_ARCH_GRANITE_HYBRID || arch == LLM_ARCH_LFM2 || arch == LLM_ARCH_LFM2MOE || arch == LLM_ARCH_KIMI_LINEAR ||
|
||||
arch == LLM_ARCH_BAILINGMOE3 || arch == LLM_ARCH_KIMI_K3) {
|
||||
GGML_ASSERT(n_layer >= 2);
|
||||
std::vector<uint32_t> n_head_per_layer;
|
||||
n_head_per_layer.reserve(n_layer);
|
||||
@@ -165,6 +167,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
|| arch == LLM_ARCH_DEEPSEEK32
|
||||
|| arch == LLM_ARCH_GLM_DSA
|
||||
|| arch == LLM_ARCH_KIMI_LINEAR
|
||||
|| arch == LLM_ARCH_BAILINGMOE3
|
||||
|| arch == LLM_ARCH_KIMI_K3
|
||||
|| arch == LLM_ARCH_MISTRAL4) {
|
||||
ms.add_kv(LLM_KV_ATTENTION_KEY_LENGTH, uint32_t(576));
|
||||
@@ -244,6 +247,12 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
|
||||
ms.add_kv(LLM_KV_SSM_TIME_STEP_RANK, n_head);
|
||||
ms.add_kv(LLM_KV_SSM_GROUP_COUNT, arch == LLM_ARCH_PLAMO2 ? 0 : uint32_t(2));
|
||||
ms.add_kv(LLM_KV_KDA_HEAD_DIM, uint32_t(128));
|
||||
ms.add_kv(LLM_KV_KDA_SAFE_GATE, true);
|
||||
ms.add_kv(LLM_KV_KDA_GATE_LOWER_BOUND, -5.0f);
|
||||
if (arch == LLM_ARCH_BAILINGMOE3) {
|
||||
ms.add_kv(LLM_KV_SWIGLU_CLAMP_EXP, std::vector<float>({0.0f, 4.0f}));
|
||||
ms.add_kv(LLM_KV_SWIGLU_CLAMP_SHEXP, std::vector<float>({0.0f, 5.0f}));
|
||||
}
|
||||
ms.add_kv(LLM_KV_WKV_HEAD_SIZE, n_embd/n_head);
|
||||
ms.add_kv(LLM_KV_SHORTCONV_L_CACHE, uint32_t(3));
|
||||
ms.add_kv(LLM_KV_RESIDUAL_SCALE, 3.5565588200778455f);
|
||||
@@ -361,6 +370,7 @@ static bool moe_mandatory(const llm_arch arch) {
|
||||
case LLM_ARCH_EXAONE_MOE:
|
||||
case LLM_ARCH_BAILINGMOE:
|
||||
case LLM_ARCH_BAILINGMOE2:
|
||||
case LLM_ARCH_BAILINGMOE3:
|
||||
case LLM_ARCH_DOTS1:
|
||||
case LLM_ARCH_AFMOE:
|
||||
case LLM_ARCH_ERNIE4_5:
|
||||
@@ -610,6 +620,9 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg
|
||||
}
|
||||
const std::string config_name = moe ? "MoE" : "Dense";
|
||||
gguf_context_ptr gguf_ctx = get_gguf_ctx(arch, moe);
|
||||
if (arch == LLM_ARCH_BAILINGMOE3) {
|
||||
GGML_ASSERT(gguf_remove_key(gguf_ctx.get(), "bailingmoe3.kda.safe_gate") >= 0);
|
||||
}
|
||||
std::pair<llama_model_ptr, llama_context_ptr> model_and_ctx_cpu;
|
||||
std::vector<float> logits_cpu;
|
||||
for (device_config & dc : dev_configs) {
|
||||
|
||||
@@ -846,7 +846,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
|
||||
invalid_param = true;
|
||||
break;
|
||||
}
|
||||
LOG_WRN("DEPRECATED: -mmp and --mmap are deprecated in favour of --load-mode. Please use --load-mode mmap instead.");
|
||||
LOG_WRN("DEPRECATED: -mmp and --mmap are deprecated in favour of --load-mode. Please use --load-mode mmap instead.\n");
|
||||
auto p = string_split<bool>(argv[i], split_delim);
|
||||
|
||||
std::vector<llama_load_mode> modes;
|
||||
@@ -865,7 +865,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
|
||||
invalid_param = true;
|
||||
break;
|
||||
}
|
||||
LOG_WRN("DEPRECATED: -dio and --direct-io are deprecated in favour of --load-mode. Please use --load-mode dio instead.");
|
||||
LOG_WRN("DEPRECATED: -dio and --direct-io are deprecated in favour of --load-mode. Please use --load-mode dio instead.\n");
|
||||
auto p = string_split<bool>(argv[i], split_delim);
|
||||
|
||||
std::vector<llama_load_mode> modes;
|
||||
|
||||
Reference in New Issue
Block a user