does not work for e4b

This commit is contained in:
Concedo
2026-06-04 10:44:40 +08:00
106 changed files with 4684 additions and 1205 deletions
+12 -5
View File
@@ -354,7 +354,6 @@ static handle_model_result common_params_handle_model(struct common_params_model
model.path = "";
}
common_download_opts hf_opts = opts;
hf_opts.download_mmproj = true; // also look for mmproj when downloading hf model
auto download_result = common_download_model(model, hf_opts);
if (download_result.model_path.empty()) {
@@ -442,10 +441,11 @@ bool common_params_handle_models(common_params & params, llama_example curr_ex)
COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params.speculative.types.end();
common_download_opts opts;
opts.bearer_token = params.hf_token;
opts.offline = params.offline;
opts.skip_download = params.skip_download;
opts.download_mtp = spec_type_draft_mtp;
opts.bearer_token = params.hf_token;
opts.offline = params.offline;
opts.skip_download = params.skip_download;
opts.download_mtp = spec_type_draft_mtp;
opts.download_mmproj = !params.no_mmproj;
try {
auto res = common_params_handle_model(params.model, opts);
@@ -3028,6 +3028,13 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
params.timeout_write = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_TIMEOUT"));
add_opt(common_arg(
{"--sse-ping-interval"}, "N",
string_format("server SSE ping interval in seconds (-1 = disabled, default: %d)", params.sse_ping_interval),
[](common_params & params, int value) {
params.sse_ping_interval = value;
}
).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_SSE_PING_INTERVAL"));
add_opt(common_arg(
{"--threads-http"}, "N",
string_format("number of threads used to process HTTP requests (default: %d)", params.n_threads_http),
+13 -12
View File
@@ -1988,36 +1988,37 @@ bool common_replay_last_token(struct llama_context * ctx, llama_token last_token
bool common_prompt_batch_decode(
struct llama_context * ctx,
const std::vector<llama_token> & tokens,
const std::vector<llama_token> & all_tokens,
int n_new,
int & n_past,
int n_batch,
std::string_view state_path,
bool save_state) {
const int n_eval = tokens.size();
if (n_eval == 0) {
if (n_new == 0) {
return true;
}
const int offset = all_tokens.size() - n_new;
if (save_state && n_eval > 1) {
const int n_tokens_before_last = n_eval - 1;
if (save_state && n_new > 1) {
const int n_tokens_before_last = n_new - 1;
GGML_ASSERT(n_eval <= n_batch);
GGML_ASSERT(n_new <= n_batch);
// Decode all but the last token so we can save the memory state before decoding the last token.
// This is done so we can restore the session state later and replay the last token.
// Memory implementations in recurrent/hybrid models don't support removing tokens from their
// memory, so we can't just remove the last token from the memory and replay the last token which
// is the reason for this logic.
if (llama_decode(ctx, llama_batch_get_one(const_cast<llama_token*>(tokens.data()), n_tokens_before_last))) {
if (llama_decode(ctx, llama_batch_get_one(const_cast<llama_token*>(all_tokens.data() + offset), n_tokens_before_last))) {
LOG_ERR("%s : failed to eval\n", __func__);
return false;
}
n_past += n_tokens_before_last;
llama_state_save_file(ctx, state_path.data(), tokens.data(), n_tokens_before_last);
LOG_INF("saved session before last token to %s, n_tokens = %d\n", state_path.data(), n_tokens_before_last);
llama_state_save_file(ctx, state_path.data(), all_tokens.data(), all_tokens.size());
LOG_INF("saved session before last token to %s, n_new = %zu\n", state_path.data(), all_tokens.size());
llama_token last_token = tokens.back();
llama_token last_token = all_tokens.back();
llama_batch batch = llama_batch_get_one(&last_token, 1);
int32_t pos = n_past;
batch.pos = &pos;
@@ -2028,11 +2029,11 @@ bool common_prompt_batch_decode(
}
n_past++;
} else {
if (llama_decode(ctx, llama_batch_get_one(const_cast<llama_token*>(tokens.data()), n_eval))) {
if (llama_decode(ctx, llama_batch_get_one(const_cast<llama_token*>(all_tokens.data() + offset), n_new))) {
LOG_ERR("%s : failed to eval\n", __func__);
return false;
}
n_past += n_eval;
n_past += n_new;
}
return true;
+3 -1
View File
@@ -593,6 +593,7 @@ struct common_params {
bool reuse_port = false; // allow multiple sockets to bind to the same port
int32_t timeout_read = 3600; // http read timeout in seconds
int32_t timeout_write = timeout_read; // http write timeout in seconds
int32_t sse_ping_interval = 30; // SSE ping interval in seconds
int32_t n_threads_http = -1; // number of threads to process HTTP requests (TODO: support threadpool)
int32_t n_cache_reuse = 0; // min chunk size to reuse from the cache via KV shifting
bool cache_prompt = true; // whether to enable prompt caching
@@ -930,7 +931,8 @@ void common_batch_add(
// tokens from memory, so this approach works across all model architectures.
bool common_prompt_batch_decode(
struct llama_context * ctx,
const std::vector<llama_token> & embd,
const std::vector<llama_token> & all_tokens,
int n_new,
int & n_past,
int n_batch,
std::string_view state_path,
+10 -10
View File
@@ -3,7 +3,7 @@
#include "common.h"
#include "ggml.h"
#include "llama.h"
#include "../src/llama-ext.h" // staging API: llama_set_embeddings_pre_norm / llama_get_embeddings_pre_norm_ith (used by MTP)
#include "../src/llama-ext.h" // staging API: llama_set_embeddings_nextn / llama_get_embeddings_nextn_ith (used by MTP)
#include "log.h"
#include "ngram-cache.cpp"
#include "ngram-map.cpp"
@@ -162,7 +162,7 @@ struct common_speculative_impl {
virtual bool need_embd() const = 0;
// true if this implementation requires the target context to extract pre-norm embeddings
virtual bool need_embd_pre_norm() const { return false; }
virtual bool need_embd_nextn() const { return false; }
};
struct common_speculative_impl_draft_simple : public common_speculative_impl {
@@ -487,8 +487,8 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
}
}
llama_set_embeddings_pre_norm(ctx_tgt, true, /*masked*/ false);
llama_set_embeddings_pre_norm(ctx_dft, true, /*masked*/ true);
llama_set_embeddings_nextn(ctx_tgt, true, /*masked*/ false);
llama_set_embeddings_nextn(ctx_dft, true, /*masked*/ true);
pending_h.assign(n_seq, std::vector<float>(n_embd, 0.0f));
@@ -583,7 +583,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
// ^--- this is a problem
// TODO:this is generally true, but would be nice to assert it
{
const float * h_tgt = llama_get_embeddings_pre_norm(ctx_tgt);
const float * h_tgt = llama_get_embeddings_nextn(ctx_tgt);
std::memcpy(batch.embd + (size_t) 1 * n_embd, h_tgt, row_bytes * (n_tokens-1));
//{
@@ -625,7 +625,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
verify_h[seq_id].resize((size_t) n_rows * n_embd);
for (int32_t i = 0; i < n_rows; ++i) {
const float * h = llama_get_embeddings_pre_norm_ith(ctx_tgt, i_batch_beg[seq_id] + i);
const float * h = llama_get_embeddings_nextn_ith(ctx_tgt, i_batch_beg[seq_id] + i);
std::memcpy(verify_h[seq_id].data() + (size_t) i * n_embd, h, row_bytes);
}
@@ -686,7 +686,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
auto * smpl = smpls[seq_id].get();
common_sampler_sample(smpl, ctx_dft, i_batch, true);
h_row = llama_get_embeddings_pre_norm_ith(ctx_dft, i_batch);
h_row = llama_get_embeddings_nextn_ith(ctx_dft, i_batch);
++i_batch;
const auto * cur_p = common_sampler_get_candidates(smpl, true);
@@ -772,7 +772,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
return false;
}
bool need_embd_pre_norm() const override {
bool need_embd_nextn() const override {
return true;
}
};
@@ -1539,13 +1539,13 @@ bool common_speculative_need_embd(common_speculative * spec) {
return false;
}
bool common_speculative_need_embd_pre_norm(common_speculative * spec) {
bool common_speculative_need_embd_nextn(common_speculative * spec) {
if (spec == nullptr) {
return false;
}
for (auto & impl : spec->impls) {
if (impl->need_embd_pre_norm()) {
if (impl->need_embd_nextn()) {
return true;
}
}
+2 -2
View File
@@ -59,8 +59,8 @@ bool common_speculative_process(common_speculative * spec, const llama_batch & b
// true if any implementation requires target post-norm embeddings to be extracted
bool common_speculative_need_embd(common_speculative * spec);
// true if any implementation requires target pre-norm embeddings to be extracted
bool common_speculative_need_embd_pre_norm(common_speculative * spec);
// true if any implementation requires target nextn embeddings to be extracted
bool common_speculative_need_embd_nextn(common_speculative * spec);
// generate drafts for the sequences specified with `common_speculative_get_draft_params`
void common_speculative_draft(common_speculative * spec);
+3
View File
@@ -77,6 +77,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Gemma3nForConditionalGeneration": "gemma",
"Gemma4ForConditionalGeneration": "gemma",
"Gemma4ForCausalLM": "gemma",
"Gemma4UnifiedForConditionalGeneration": "gemma",
"GemmaForCausalLM": "gemma",
"Glm4ForCausalLM": "glm",
"Glm4MoeForCausalLM": "glm",
@@ -135,6 +136,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
"Mamba2ForCausalLM": "mamba",
"MambaForCausalLM": "mamba",
"MambaLMHeadModel": "mamba",
"MellumForCausalLM": "mellum",
"MiMoV2FlashForCausalLM": "mimo",
"MiMoV2ForCausalLM": "mimo",
"MiniCPM3ForCausalLM": "minicpm",
@@ -246,6 +248,7 @@ MMPROJ_MODEL_MAP: dict[str, str] = {
"Gemma3ForConditionalGeneration": "gemma",
"Gemma3nForConditionalGeneration": "gemma",
"Gemma4ForConditionalGeneration": "gemma",
"Gemma4UnifiedForConditionalGeneration": "gemma",
"Glm4vForConditionalGeneration": "qwen3vl",
"Glm4vMoeForConditionalGeneration": "qwen3vl",
"GlmOcrForConditionalGeneration": "qwen3vl",
+9
View File
@@ -1657,6 +1657,15 @@ class TextModel(ModelBase):
if chkhsh == "36f3066e97b7f3994b379aaacde306c1444c6ae84e81a5ae3cd2b7ed3b8c42d4":
# ref: https://huggingface.co/openbmb/MiniCPM5-1B
res = "minicpm5"
if chkhsh == "f241072145675bf8322086f115aebad05e9f869557a238bf2150a2a417d1bf60":
# ref: https://huggingface.co/ibm-granite/granite-embedding-97m-multilingual-r2
res = "granite-embed-multi-97m"
if chkhsh == "789696f5946cc0fc59371f39f6097cafed196b3acded6140432f26bbb1ae1669":
# ref: https://huggingface.co/ibm-granite/granite-embedding-311m-multilingual-r2
res = "granite-embed-multi-311m"
if chkhsh == "9dcf830ee9990cdbf78cc523a5f7bd9ad8f3f9890c2d3581d2785ad10f07049d":
# ref: https://huggingface.co/JetBrains/Mellum2-12B-A2.5B-Base
res = "mellum2"
if res is None:
logger.warning("\n")
+6
View File
@@ -603,6 +603,12 @@ class ModernBertModel(BertModel):
self.gguf_writer.add_sliding_window_pattern(sliding_window_pattern)
self.gguf_writer.add_rope_scaling_type(gguf.RopeScalingType.NONE)
self.gguf_writer.add_vocab_size(self.hparams["vocab_size"])
# FFN activation: ModernBert uses a GLU pair (ffn_up output is 2*n_ff). The
# original ModernBERT uses GELU (-> GeGLU); some derivatives such as IBM
# Granite Embedding 97m R2 use SiLU (-> SwiGLU). Persist this so the
# llama.cpp graph can pick the matching activation.
if hidden_act := self.hparams.get("hidden_activation"):
self.gguf_writer.add_hidden_act(hidden_act)
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
+79 -1
View File
@@ -3,7 +3,7 @@ from __future__ import annotations
import json
import re
from typing import Callable, Iterable, TYPE_CHECKING
from typing import Callable, Iterable, TYPE_CHECKING, Sequence
import torch
@@ -765,6 +765,26 @@ class Gemma4Model(Gemma3Model):
yield from super().modify_tensors(data_torch, name, bid)
@ModelBase.register("Gemma4UnifiedForConditionalGeneration")
class Gemma4UnifiedModel(Gemma4Model):
model_arch = gguf.MODEL_ARCH.GEMMA4
def _get_suppress_tokens(self) -> Sequence[int] | None:
gen_cfg_path = self.dir_model / "generation_config.json"
if gen_cfg_path.is_file():
with open(gen_cfg_path, encoding="utf-8") as f:
gen_cfg = json.load(f)
return gen_cfg.get("suppress_tokens")
return None
def set_gguf_parameters(self):
super().set_gguf_parameters()
suppress_tokens = self._get_suppress_tokens()
if suppress_tokens is not None:
self.gguf_writer.add_suppress_tokens(suppress_tokens)
@ModelBase.register("Gemma4ForConditionalGeneration")
class Gemma4VisionAudioModel(MmprojModel):
has_audio_encoder = True
@@ -839,3 +859,61 @@ class Gemma4VisionAudioModel(MmprojModel):
data_torch = data_torch.permute(0, 3, 1, 2).contiguous()
mapped_name = self.map_tensor_name(name, (".weight", ".bias", ".input_max", ".input_min", ".output_max", ".output_min"))
yield (mapped_name, data_torch)
@ModelBase.register("Gemma4UnifiedForConditionalGeneration")
class Gemma4UnifiedVisionAudioModel(Gemma4VisionAudioModel):
has_audio_encoder = True
has_vision_encoder = True
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
assert self.hparams_vision is not None
assert self.hparams_audio is not None
text_embd_dim = self.hparams_vision["mm_embed_dim"]
self.hparams_vision["hidden_size"] = text_embd_dim
self.hparams_audio["hidden_size"] = text_embd_dim
# this is a transformer-less vision tower, the params below are redundant but set to avoid error
self.hparams_vision["intermediate_size"] = 0
self.hparams_vision["num_layers"] = 0
self.hparams_vision["num_attention_heads"] = 0
self.hparams_audio["intermediate_size"] = 0
self.hparams_audio["num_layers"] = 0
self.hparams_audio["num_attention_heads"] = 0
def set_gguf_parameters(self):
super().set_gguf_parameters()
self.gguf_writer.add_clip_vision_projector_type(gguf.VisionProjectorType.GEMMA4UV)
self.gguf_writer.add_clip_audio_projector_type(gguf.VisionProjectorType.GEMMA4UA)
def modify_tensors(self, data_torch, name, bid):
if name.endswith("pos_embedding"):
name += ".weight"
data_torch = data_torch.permute(1, 0, 2)
elif ".pos_norm." in name:
# rename to patch_ln3 to reuse the tensor name scheme
name = name.replace(".pos_norm.", ".patch_ln3.")
elif "patch_dense.weight" in name:
# ggml im2col outputs in RR..GG..BB.. (CHW) order, but weight expects RGBRGB.. (HWC).
# Permute columns so column i aligns with CHW input position i.
assert self.hparams_vision is not None
p = self.hparams_vision["model_patch_size"]
i = torch.arange(p * p * 3)
ch = i // (p * p)
row = (i % (p * p)) // p
col = i % p
# perm[i] = HWC column index for CHW position i
perm = row * p * 3 + col * 3 + ch
data_torch = data_torch[:, perm]
elif "patch_ln1.weight" in name or "patch_ln1.bias" in name:
# same permutation for patch_ln1 as patch_dense to align with CHW input order
assert self.hparams_vision is not None
p = self.hparams_vision["model_patch_size"]
i = torch.arange(p * p * 3)
ch = i // (p * p)
row = (i % (p * p)) // p
col = i % p
# perm[i] = HWC index for CHW position i
perm = row * p * 3 + col * 3 + ch
data_torch = data_torch[perm]
return super().modify_tensors(data_torch, name, bid)
+61
View File
@@ -0,0 +1,61 @@
from __future__ import annotations
from typing import Iterable, TYPE_CHECKING
import torch
if TYPE_CHECKING:
from torch import Tensor
from .base import ModelBase, TextModel, gguf, logger
@ModelBase.register("MellumForCausalLM")
class MellumModel(TextModel):
model_arch = gguf.MODEL_ARCH.MELLUM
def set_gguf_parameters(self):
super().set_gguf_parameters()
if (moe_intermediate_size := self.hparams.get("moe_intermediate_size")) is not None:
self.gguf_writer.add_expert_feed_forward_length(moe_intermediate_size)
logger.info(f"gguf: expert feed forward length = {moe_intermediate_size}")
use_sliding_window = self.hparams.get("use_sliding_window")
sliding_window = self.hparams.get("sliding_window")
if (use_sliding_window is True or use_sliding_window is None) and sliding_window is not None:
self.gguf_writer.add_sliding_window(sliding_window)
logger.info(f"gguf: sliding window = {sliding_window}")
self.gguf_writer.add_sliding_window_pattern([t == "sliding_attention" for t in self.hparams["layer_types"]])
logger.info(f"gguf: sliding window pattern length = {len(self.hparams['layer_types'])}")
_experts: list[dict[str, Tensor]] | None = None
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
if name.find("experts") != -1:
n_experts = self.find_hparam(["num_local_experts", "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 w_name in ["down_proj", "gate_proj", "up_proj"]:
datas: list[Tensor] = []
for xid in range(n_experts):
ename = f"model.layers.{bid}.mlp.experts.{xid}.{w_name}.weight"
datas.append(self._experts[bid][ename])
del self._experts[bid][ename]
data_torch = torch.stack(datas, dim=0)
merged_name = f"model.layers.{bid}.mlp.experts.{w_name}.weight"
yield from super().modify_tensors(data_torch, merged_name, bid)
return
else:
return
yield from super().modify_tensors(data_torch, name, bid)
+103 -13
View File
@@ -99,6 +99,34 @@ class Step3VLTextModel(Qwen3Model):
class Step35Model(TextModel):
model_arch = gguf.MODEL_ARCH.STEP35
# The --mtp / --no-mtp toggles are ModelBase.mtp_only / no_mtp (set in
# convert_hf_to_gguf.py main()). Unlike Qwen3.5, which stores MTP under a
# `mtp.*` namespace, Step3.5 appends MTP layers at
# `model.layers.{num_hidden_layers + i}`, so we filter them by layer index.
# The trunk layer count is captured before indexing so the classmethod
# filter_tensors can tell the appended MTP block(s) apart from the trunk.
_n_main_layers: int | None = None
def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs)
# NextN/MTP layers are appended past num_hidden_layers; extend the
# tensor map to cover them so the MTP block's tensors get correctly
# indexed names. When --no-mtp drops the MTP blocks, fall back to the
# base num_hidden_layers so we don't reserve unused slots.
n_nextn = int(self.hparams.get("num_nextn_predict_layers", 0))
if n_nextn > 0 and not self.no_mtp:
self.block_count += n_nextn
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):
# filter_tensors is a classmethod and can't reach self.hparams; stash
# the trunk layer count here (before indexing runs) so it can detect
# the appended MTP layers by index.
hparams = {**self.hparams, **self.hparams.get("text_config", {})}
key = next((k for k in ["n_layers", "num_hidden_layers", "n_layer", "num_layers"] if k in hparams), None)
type(self)._n_main_layers = hparams.get(key)
return super().index_tensors(remote_hf_model_id=remote_hf_model_id)
def set_gguf_parameters(self):
rope_theta = self.hparams.get("rope_theta")
if isinstance(rope_theta, list):
@@ -119,8 +147,25 @@ class Step35Model(TextModel):
n_head_swa = attn_other.get("num_attention_heads", n_head_base)
n_kv_swa = attn_other.get("num_attention_groups", n_kv_base)
layer_types = layer_types[: self.block_count]
partial_rotary_factors = partial_rotary_factors[: self.block_count]
n_nextn = int(self.hparams.get("num_nextn_predict_layers", 0))
# The Step3p5 HF checkpoint stores layer_types/partial_rotary_factors
# entries for the MTP blocks past num_hidden_layers; preserve them so
# the MTP layer's attention shape, SWA flag, and partial RoPE dim are
# set correctly. Pad with full-attention defaults if the checkpoint
# truncated them.
def _pad(arr, n, default):
arr = list(arr)
if len(arr) < n:
arr = arr + [default] * (n - len(arr))
return arr[:n]
layer_types = _pad(layer_types, self.block_count, "full_attention")
partial_rotary_factors = _pad(
partial_rotary_factors,
self.block_count,
0.5, # full_attention default for Step3p5
)
assert [1.0 if lt == "sliding_attention" else 0.5 for lt in layer_types] == partial_rotary_factors
head_arr = [n_head_swa if lt == "sliding_attention" else n_head_base for lt in layer_types]
kv_arr = [n_kv_swa if lt == "sliding_attention" else n_kv_base for lt in layer_types]
@@ -157,31 +202,61 @@ class Step35Model(TextModel):
self.gguf_writer.add_layer_norm_rms_eps(self.hparams.get("rms_norm_eps", 1e-5))
# Optional per-layer SwiGLU clamps.
# Optional per-layer SwiGLU clamps. MTP layers default to no clamping (0.0).
if (limits := self.hparams.get("swiglu_limits")) is not None:
limits_f = [0.0 if v is None else float(v) for v in limits[: self.block_count]]
limits_f = _pad(
[0.0 if v is None else float(v) for v in limits],
self.block_count,
0.0,
)
self.gguf_writer.add_swiglu_clamp_exp(limits_f)
if (limits_shared := self.hparams.get("swiglu_limits_shared")) is not None:
limits_shared_f = [0.0 if v is None else float(v) for v in limits_shared[: self.block_count]]
limits_shared_f = _pad(
[0.0 if v is None else float(v) for v in limits_shared],
self.block_count,
0.0,
)
self.gguf_writer.add_swiglu_clamp_shexp(limits_shared_f)
if n_nextn > 0 and not self.no_mtp:
self.gguf_writer.add_nextn_predict_layers(n_nextn)
@classmethod
def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None:
name, gen = item
if (titem := super().filter_tensors(item)) is None:
return None
name, gen = titem
# Map router bias (expert selection bias) to a GGUF bias tensor
if name.endswith(".moe.router_bias"):
name += ".bias"
return super().filter_tensors((name, gen))
# Step3.5 appends the MTP block(s) past num_hidden_layers.
assert cls._n_main_layers is not None
is_mtp = (m := re.match(r"model\.layers\.(\d+)\.", name)) is not None and int(m.group(1)) >= cls._n_main_layers
# --no-mtp: drop the appended MTP block(s) entirely.
if is_mtp and cls.no_mtp:
return None
# --mtp: keep ONLY MTP-block tensors plus the shared embeddings/norm/
# lm_head (so the resulting GGUF carries just the draft head).
if cls.mtp_only and not is_mtp and name not in (
"model.embed_tokens.weight", "model.norm.weight", "lm_head.weight",
):
return None
# The checkpoint nests the per-MTP-layer shared head under
# `model.layers.{N+i}.transformer.shared_head.{norm,output}.weight`;
# strip the `transformer.` infix and rename `output` → `head` so the
# existing NEXTN_SHARED_HEAD_{NORM,HEAD} tensor mapping picks them up.
# Mirrors vllm's `_rewrite_spec_layer_name` (step3p5_mtp.py).
if is_mtp:
name = name.replace(".transformer.", ".")
name = name.replace("shared_head.output", "shared_head.head")
return name, gen
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None):
# remove mtp layers
if (m := re.match(r"model\.layers\.(\d+)\.", name)) is not None:
il = int(m.group(1))
n_main = int(self.hparams.get("num_hidden_layers", self.block_count))
if il >= n_main:
return
if name.endswith("norm.weight"):
data_torch += 1.0
@@ -190,6 +265,21 @@ class Step35Model(TextModel):
yield from super().modify_tensors(data_torch, name, bid)
def prepare_metadata(self, vocab_only: bool):
from_dir = self.fname_out.is_dir()
super().prepare_metadata(vocab_only=vocab_only)
# Mirror Qwen3.5's behavior: when emitting a draft-only file into a
# directory, prefix with "mtp-" so it doesn't collide with the trunk.
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"
def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]:
# Step35 can optionally use Llama-3 style RoPE scaling (HF: rope_scaling.rope_type == "llama3").
# llama.cpp represents this via a single extra tensor: "rope_freqs.weight" (aka MODEL_TENSOR.ROPE_FREQS).
+3 -2
View File
@@ -251,8 +251,9 @@ def main() -> None:
if args.mtp or args.no_mtp:
from conversion.qwen import _Qwen35MtpMixin
if not issubclass(model_class, _Qwen35MtpMixin):
logger.error("--mtp / --no-mtp are only supported for Qwen3.5/3.6 text variants today")
from conversion.step3 import Step35Model
if not (issubclass(model_class, _Qwen35MtpMixin) or issubclass(model_class, Step35Model)):
logger.error("--mtp / --no-mtp are only supported for Qwen3.5/3.6 and Step3.5 text variants today")
sys.exit(1)
if args.no_mtp:
model_class.no_mtp = True
+3
View File
@@ -158,6 +158,9 @@ models = [
{"name": "sarvam-moe", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/sarvamai/sarvam-30b", },
{"name": "talkie", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/lewtun/talkie-1930-13b-it-hf", },
{"name": "minicpm5", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/openbmb/MiniCPM5-1B"},
{"name": "granite-embed-multi-97m", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/ibm-granite/granite-embedding-97m-multilingual-r2", },
{"name": "granite-embed-multi-311m", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/ibm-granite/granite-embedding-311m-multilingual-r2", },
{"name": "mellum2", "tokt": TOKENIZER_TYPE.BPE, "repo": "https://huggingface.co/JetBrains/Mellum2-12B-A2.5B-Base"},
]
# some models are known to be broken upstream, so we will skip them as exceptions
+10 -1
View File
@@ -8955,7 +8955,12 @@ static void ggml_compute_forward_flash_attn_ext_f16(
k->type == v->type &&
neq1 >= Q_TILE_SZ);
#ifdef GGML_SIMD
use_tiled &= (DV % GGML_F32_EPR == 0);
#if defined(__ARM_FEATURE_SVE)
const int64_t f32_epr = svcntw();
#else
const int64_t f32_epr = GGML_F32_EPR;
#endif
use_tiled &= (DV % f32_epr == 0);
#endif
int current_chunk = ith;
@@ -11358,7 +11363,11 @@ static void ggml_compute_forward_fwht_f32(const ggml_compute_params * params, gg
// Scalar passes
#if defined(GGML_SIMD)
#if defined(__ARM_FEATURE_SVE)
const int step = svcntw();
#else
const int step = GGML_F32_EPR;
#endif
#else
const int step = n;
#endif
+6
View File
@@ -1618,6 +1618,12 @@ static bool ggml_cuda_kernel_can_use_pdl(const void * kernel) {
#endif //defined(GGML_CUDA_USE_PDL)
// PDL and __restrict__ need to be mutually exclusive, see https://github.com/ggml-org/llama.cpp/pull/24030
# if (defined(GGML_CUDA_USE_PDL) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_HOPPER)
# define GGML_CUDA_RESTRICT
# else
# define GGML_CUDA_RESTRICT __restrict__
# endif // defined(GGML_CUDA_USE_PDL) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= GGML_CUDA_CC_HOPPER
template<typename Kernel, typename... Args>
static __inline__ void ggml_cuda_kernel_launch(Kernel kernel, const ggml_cuda_kernel_launch_params & launch_params, Args&&... args) {
+70 -20
View File
@@ -44,6 +44,46 @@ typedef void (* fattn_kernel_t)(
typedef float (*vec_dot_KQ_t)(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8 , const void * __restrict__ Q_ds);
struct ggml_cuda_flash_attn_ext_f16_extra_data {
uintptr_t K;
uintptr_t V;
uintptr_t end;
};
static inline ggml_cuda_flash_attn_ext_f16_extra_data ggml_cuda_flash_attn_ext_get_f16_extra_data(
const ggml_tensor * dst, const bool need_f16_K, const bool need_f16_V) {
GGML_ASSERT(dst->op == GGML_OP_FLASH_ATTN_EXT);
const ggml_tensor * K = dst->src[1];
const ggml_tensor * V = dst->src[2];
GGML_ASSERT(K != nullptr);
GGML_ASSERT(V != nullptr);
const bool V_is_K_view = V->view_src && (V->view_src == K || (V->view_src == K->view_src && V->view_offs == K->view_offs));
ggml_cuda_flash_attn_ext_f16_extra_data data = {};
data.end = (uintptr_t) dst->data + ggml_nbytes(dst);
if (need_f16_K && K->type != GGML_TYPE_F16) {
data.end = GGML_PAD(data.end, 128);
data.K = data.end;
data.end += ggml_nelements(K)*ggml_type_size(GGML_TYPE_F16);
}
if (need_f16_V && V->type != GGML_TYPE_F16) {
if (V_is_K_view) {
data.V = data.K;
} else {
data.end = GGML_PAD(data.end, 128);
data.V = data.end;
data.end += ggml_nelements(V)*ggml_type_size(GGML_TYPE_F16);
}
}
return data;
}
template <int D, int nthreads>
static __device__ __forceinline__ float vec_dot_fattn_vec_KQ_f16(
const char * __restrict__ K_c, const void * __restrict__ Q_v, const int * __restrict__ Q_q8 , const void * __restrict__ Q_ds_v) {
@@ -678,8 +718,8 @@ static __global__ void flash_attn_mask_to_KV_max(
template<int D, int ncols1, int ncols2> // D == head size
__launch_bounds__(D, 1)
static __global__ void flash_attn_stream_k_fixup_uniform(
float * __restrict__ dst,
const float2 * __restrict__ dst_fixup,
float * dst_ptr,
const float2 * dst_fixup_ptr,
const int ne01, const int ne02,
const int ne12, const int nblocks_stream_k,
const int gqa_ratio,
@@ -689,6 +729,8 @@ static __global__ void flash_attn_stream_k_fixup_uniform(
const uint3 fd_iter_j) {
constexpr int ncols = ncols1*ncols2;
ggml_cuda_pdl_lc();
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const float2 * GGML_CUDA_RESTRICT dst_fixup = dst_fixup_ptr;
const int tile_idx = blockIdx.x; // One block per output tile.
const int j = blockIdx.y;
@@ -760,8 +802,8 @@ static __global__ void flash_attn_stream_k_fixup_uniform(
template <int D, int ncols1, int ncols2> // D == head size
__launch_bounds__(D, 1)
static __global__ void flash_attn_stream_k_fixup_general(
float * __restrict__ dst,
const float2 * __restrict__ dst_fixup,
float * dst_ptr,
const float2 * dst_fixup_ptr,
const int ne01, const int ne02,
const int gqa_ratio,
const int total_work,
@@ -769,6 +811,8 @@ static __global__ void flash_attn_stream_k_fixup_general(
const uint3 fd_iter_k_j_z,
const uint3 fd_iter_k_j,
const uint3 fd_iter_k) {
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const float2 * GGML_CUDA_RESTRICT dst_fixup = dst_fixup_ptr;
constexpr int ncols = ncols1*ncols2;
const int bidx0 = blockIdx.x;
@@ -867,11 +911,14 @@ static __global__ void flash_attn_stream_k_fixup_general(
template<int D> // D == head size
__launch_bounds__(D, 1)
static __global__ void flash_attn_combine_results(
const float * __restrict__ VKQ_parts,
const float2 * __restrict__ VKQ_meta,
float * __restrict__ dst,
const float * VKQ_parts_ptr,
const float2 * VKQ_meta_ptr,
float * dst_ptr,
const int parallel_blocks) {
ggml_cuda_pdl_lc();
const float * GGML_CUDA_RESTRICT VKQ_parts = VKQ_parts_ptr;
const float2 * GGML_CUDA_RESTRICT VKQ_meta = VKQ_meta_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
// Dimension 0: threadIdx.x
// Dimension 1: blockIdx.x
// Dimension 2: blockIdx.y
@@ -952,8 +999,9 @@ void launch_fattn(
const int cc = ggml_cuda_info().devices[id].cc;
const int nsm = ggml_cuda_info().devices[id].nsm;
ggml_cuda_pool_alloc<half> K_f16(pool);
ggml_cuda_pool_alloc<half> V_f16(pool);
const ggml_cuda_flash_attn_ext_f16_extra_data f16_extra =
ggml_cuda_flash_attn_ext_get_f16_extra_data(KQV, need_f16_K, need_f16_V);
ggml_cuda_pool_alloc<int> KV_max(pool);
ggml_cuda_pool_alloc<float> dst_tmp(pool);
ggml_cuda_pool_alloc<float2> dst_tmp_meta(pool);
@@ -972,10 +1020,11 @@ void launch_fattn(
const size_t bs = ggml_blck_size(K->type);
const size_t ts = ggml_type_size(K->type);
K_f16.alloc(ggml_nelements(K));
GGML_ASSERT(f16_extra.K != 0);
half * K_f16 = (half *) f16_extra.K;
if (ggml_is_contiguously_allocated(K)) {
to_fp16_cuda_t to_fp16 = ggml_get_to_fp16_cuda(K->type);
to_fp16(K_data, K_f16.ptr, ggml_nelements(K), main_stream);
to_fp16(K_data, K_f16, ggml_nelements(K), main_stream);
nb11 = nb11*bs*sizeof(half)/ts;
nb12 = nb12*bs*sizeof(half)/ts;
@@ -986,13 +1035,13 @@ void launch_fattn(
const int64_t s01 = nb11 / ts;
const int64_t s02 = nb12 / ts;
const int64_t s03 = nb13 / ts;
to_fp16(K_data, K_f16.ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, main_stream);
to_fp16(K_data, K_f16, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, main_stream);
nb11 = K->ne[0] * sizeof(half);
nb12 = K->ne[1] * nb11;
nb13 = K->ne[2] * nb12;
}
K_data = (char *) K_f16.ptr;
K_data = (char *) K_f16;
}
if (need_f16_V && V->type != GGML_TYPE_F16) {
@@ -1005,11 +1054,12 @@ void launch_fattn(
const size_t bs = ggml_blck_size(V->type);
const size_t ts = ggml_type_size(V->type);
V_f16.alloc(ggml_nelements(V));
GGML_ASSERT(f16_extra.V != 0);
half * V_f16 = (half *) f16_extra.V;
if (ggml_is_contiguously_allocated(V)) {
to_fp16_cuda_t to_fp16 = ggml_get_to_fp16_cuda(V->type);
to_fp16(V_data, V_f16.ptr, ggml_nelements(V), main_stream);
V_data = (char *) V_f16.ptr;
to_fp16(V_data, V_f16, ggml_nelements(V), main_stream);
V_data = (char *) V_f16;
nb21 = nb21*bs*sizeof(half)/ts;
nb22 = nb22*bs*sizeof(half)/ts;
@@ -1020,13 +1070,13 @@ void launch_fattn(
const int64_t s01 = nb21 / ts;
const int64_t s02 = nb22 / ts;
const int64_t s03 = nb23 / ts;
to_fp16(V_data, V_f16.ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, main_stream);
to_fp16(V_data, V_f16, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, main_stream);
nb21 = V->ne[0] * sizeof(half);
nb22 = V->ne[1] * nb21;
nb23 = V->ne[2] * nb22;
}
V_data = (char *) V_f16.ptr;
V_data = (char *) V_f16;
}
}
@@ -1153,8 +1203,8 @@ void launch_fattn(
GGML_ASSERT(block_dim.x % warp_size == 0);
// disabled PDL enrollment for now due to a compiler bug.
fattn_kernel<<<blocks_num, block_dim, nbytes_shared, main_stream>>>(
ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks_num, block_dim, nbytes_shared, main_stream);
ggml_cuda_kernel_launch(fattn_kernel, launch_params,
(const char *) Q->data,
K_data,
V_data,
+17 -9
View File
@@ -1703,14 +1703,14 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
template<int DKQ, int DV, int ncols1, int ncols2, bool use_logit_softcap, bool V_is_K_view>
__launch_bounds__(ggml_cuda_fattn_mma_get_nthreads(DKQ, DV, ncols1*ncols2), ggml_cuda_fattn_mma_get_occupancy(DKQ, DV, ncols1*ncols2))
static __global__ void flash_attn_ext_f16(
const char * __restrict__ Q,
const char * __restrict__ K,
const char * __restrict__ V,
const char * __restrict__ mask,
const char * __restrict__ sinks,
const int * __restrict__ KV_max,
float * __restrict__ dst,
float2 * __restrict__ dst_meta,
const char * Q_ptr,
const char * K_ptr,
const char * V_ptr,
const char * mask_ptr,
const char * sinks_ptr,
const int * KV_max_ptr,
float * dst_ptr,
float2 * dst_meta_ptr,
const float scale,
const float max_bias,
const float m0,
@@ -1726,6 +1726,14 @@ static __global__ void flash_attn_ext_f16(
const int32_t nb31, const int32_t nb32, const int64_t nb33) {
ggml_cuda_pdl_sync(); // TODO optimize placement
#if defined(FLASH_ATTN_AVAILABLE) && (defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE))
const char * GGML_CUDA_RESTRICT Q = Q_ptr;
const char * GGML_CUDA_RESTRICT K = K_ptr;
const char * GGML_CUDA_RESTRICT V = V_ptr;
const char * GGML_CUDA_RESTRICT mask = mask_ptr;
const char * GGML_CUDA_RESTRICT sinks = sinks_ptr;
const int * GGML_CUDA_RESTRICT KV_max = KV_max_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
float2 * GGML_CUDA_RESTRICT dst_meta = dst_meta_ptr;
// Skip unused kernel variants for faster compilation:
if (use_logit_softcap && !(DKQ == 128 || DKQ == 256 || DKQ == 512)) {
@@ -1871,7 +1879,7 @@ static __global__ void flash_attn_ext_f16(
(Q_f2, K_h2, V_h2, mask_h, sinks_f, dstk, dst_meta, scale, slope, logit_softcap,
ne01, ne02, gqa_ratio, ne11, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, zt_gqa, kb0_start, kb0_stop);
#else
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
GGML_UNUSED_VARS(Q_ptr, K_ptr, V_ptr, mask_ptr, sinks_ptr, KV_max_ptr, dst_ptr, dst_meta_ptr, scale,
max_bias, m0, m1, n_head_log2, logit_softcap,
ne00, ne01, ne02, ne03,
nb01, nb02, nb03,
+17 -9
View File
@@ -788,14 +788,14 @@ static __device__ __forceinline__ void flash_attn_tile_iter(
template<int DKQ, int DV, int ncols1, int ncols2, bool use_logit_softcap> // D == head size
__launch_bounds__(ggml_cuda_fattn_tile_get_nthreads(DKQ, DV, ncols1*ncols2), ggml_cuda_fattn_tile_get_occupancy(DKQ, DV, ncols1*ncols2))
static __global__ void flash_attn_tile(
const char * __restrict__ Q,
const char * __restrict__ K,
const char * __restrict__ V,
const char * __restrict__ mask,
const char * __restrict__ sinks,
const int * __restrict__ KV_max,
float * __restrict__ dst,
float2 * __restrict__ dst_meta,
const char * Q_ptr,
const char * K_ptr,
const char * V_ptr,
const char * mask_ptr,
const char * sinks_ptr,
const int * KV_max_ptr,
float * dst_ptr,
float2 * dst_meta_ptr,
const float scale,
const float max_bias,
const float m0,
@@ -810,6 +810,14 @@ static __global__ void flash_attn_tile(
const int32_t ne31, const int32_t ne32, const int32_t ne33,
const int32_t nb31, const int32_t nb32, const int64_t nb33) {
#ifdef FLASH_ATTN_AVAILABLE
const char * GGML_CUDA_RESTRICT Q = Q_ptr;
const char * GGML_CUDA_RESTRICT K = K_ptr;
const char * GGML_CUDA_RESTRICT V = V_ptr;
const char * GGML_CUDA_RESTRICT mask = mask_ptr;
const char * GGML_CUDA_RESTRICT sinks = sinks_ptr;
const int * GGML_CUDA_RESTRICT KV_max = KV_max_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
float2 * GGML_CUDA_RESTRICT dst_meta = dst_meta_ptr;
// Skip unused kernel variants for faster compilation:
@@ -1126,7 +1134,7 @@ static __global__ void flash_attn_tile(
}
}
#else
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
GGML_UNUSED_VARS(Q_ptr, K_ptr, V_ptr, mask_ptr, sinks_ptr, KV_max_ptr, dst_ptr, dst_meta_ptr, scale,
max_bias, m0, m1, n_head_log2, logit_softcap,
ne00, ne01, ne02, ne03,
nb01, nb02, nb03,
+17 -9
View File
@@ -19,14 +19,14 @@ static constexpr __device__ int ggml_cuda_fattn_vec_get_nthreads_device() {
template<int D, int ncols, ggml_type type_K, ggml_type type_V, bool use_logit_softcap> // D == head size
__launch_bounds__(ggml_cuda_fattn_vec_get_nthreads_device(), 1)
static __global__ void flash_attn_ext_vec(
const char * __restrict__ Q,
const char * __restrict__ K,
const char * __restrict__ V,
const char * __restrict__ mask,
const char * __restrict__ sinks,
const int * __restrict__ KV_max,
float * __restrict__ dst,
float2 * __restrict__ dst_meta,
const char * Q_ptr,
const char * K_ptr,
const char * V_ptr,
const char * mask_ptr,
const char * sinks_ptr,
const int * KV_max_ptr,
float * dst_ptr,
float2 * dst_meta_ptr,
const float scale,
const float max_bias,
const float m0,
@@ -42,6 +42,14 @@ static __global__ void flash_attn_ext_vec(
const int32_t nb31, const int32_t nb32, const int64_t nb33) {
ggml_cuda_pdl_lc();
#ifdef FLASH_ATTN_AVAILABLE
const char * GGML_CUDA_RESTRICT Q = Q_ptr;
const char * GGML_CUDA_RESTRICT K = K_ptr;
const char * GGML_CUDA_RESTRICT V = V_ptr;
const char * GGML_CUDA_RESTRICT mask = mask_ptr;
const char * GGML_CUDA_RESTRICT sinks = sinks_ptr;
const int * GGML_CUDA_RESTRICT KV_max = KV_max_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
float2 * GGML_CUDA_RESTRICT dst_meta = dst_meta_ptr;
// Skip unused kernel variants for faster compilation:
if (use_logit_softcap && !(D == 128 || D == 256)) {
@@ -506,7 +514,7 @@ static __global__ void flash_attn_ext_vec(
dst_meta[((sequence*int(ne01.z) + ic0 + tid)*ne02 + head)*gridDim.y + blockIdx.y] = make_float2(KQ_max[tid], KQ_sum[tid]);
}
#else
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
GGML_UNUSED_VARS(Q_ptr, K_ptr, V_ptr, mask_ptr, sinks_ptr, KV_max_ptr, dst_ptr, dst_meta_ptr, scale,
max_bias, m0, m1, n_head_log2, logit_softcap,
ne00, ne01, ne02, ne03,
nb01, nb02, nb03,
+17 -9
View File
@@ -24,14 +24,14 @@ namespace wmma = rocwmma;
template<int D, int ncols, int nwarps, int VKQ_stride, typename KQ_acc_t, bool use_logit_softcap>
__launch_bounds__(nwarps*ggml_cuda_get_physical_warp_size(), 1)
static __global__ void flash_attn_ext_f16(
const char * __restrict__ Q,
const char * __restrict__ K,
const char * __restrict__ V,
const char * __restrict__ mask,
const char * __restrict__ sinks,
const int * __restrict__ KV_max,
float * __restrict__ dst,
float2 * __restrict__ dst_meta,
const char * Q_ptr,
const char * K_ptr,
const char * V_ptr,
const char * mask_ptr,
const char * sinks_ptr,
const int * KV_max_ptr,
float * dst_ptr,
float2 * dst_meta_ptr,
const float scale,
const float max_bias,
const float m0,
@@ -46,6 +46,14 @@ static __global__ void flash_attn_ext_f16(
const int32_t ne31, const int32_t ne32, const int32_t ne33,
const int32_t nb31, const int32_t nb32, const int64_t nb33) {
#if defined(FLASH_ATTN_AVAILABLE) && (defined(GGML_HIP_ROCWMMA_FATTN) && defined(GGML_USE_WMMA_FATTN))
const char * GGML_CUDA_RESTRICT Q = Q_ptr;
const char * GGML_CUDA_RESTRICT K = K_ptr;
const char * GGML_CUDA_RESTRICT V = V_ptr;
const char * GGML_CUDA_RESTRICT mask = mask_ptr;
const char * GGML_CUDA_RESTRICT sinks = sinks_ptr;
const int * GGML_CUDA_RESTRICT KV_max = KV_max_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
float2 * GGML_CUDA_RESTRICT dst_meta = dst_meta_ptr;
// Skip unused kernel variants for faster compilation:
if (use_logit_softcap && !(D == 128 || D == 256)) {
NO_DEVICE_CODE;
@@ -494,7 +502,7 @@ static __global__ void flash_attn_ext_f16(
dst_meta[j_dst_unrolled] = dst_meta_val;
}
#else
GGML_UNUSED_VARS(Q, K, V, mask, sinks, KV_max, dst, dst_meta, scale,
GGML_UNUSED_VARS(Q_ptr, K_ptr, V_ptr, mask_ptr, sinks_ptr, KV_max_ptr, dst_ptr, dst_meta_ptr, scale,
max_bias, m0, m1, n_head_log2, logit_softcap,
ne00, ne01, ne02, ne03,
nb01, nb02, nb03,
+35
View File
@@ -530,6 +530,41 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
return BEST_FATTN_KERNEL_TILE;
}
size_t ggml_cuda_flash_attn_ext_get_alloc_size(int device, const ggml_tensor * dst) {
GGML_ASSERT(dst->op == GGML_OP_FLASH_ATTN_EXT);
const ggml_tensor * K = dst->src[1];
const ggml_tensor * V = dst->src[2];
GGML_ASSERT(K != nullptr);
GGML_ASSERT(V != nullptr);
const best_fattn_kernel kernel = ggml_cuda_get_best_fattn_kernel(device, dst);
bool need_f16_K = false;
bool need_f16_V = false;
switch (kernel) {
case BEST_FATTN_KERNEL_TILE:
case BEST_FATTN_KERNEL_WMMA_F16:
case BEST_FATTN_KERNEL_MMA_F16:
need_f16_K = true;
need_f16_V = true;
break;
case BEST_FATTN_KERNEL_VEC:
need_f16_K = K->type == GGML_TYPE_F32;
need_f16_V = V->type == GGML_TYPE_F32;
break;
case BEST_FATTN_KERNEL_NONE:
break;
}
const ggml_cuda_flash_attn_ext_f16_extra_data f16_extra =
ggml_cuda_flash_attn_ext_get_f16_extra_data(dst, need_f16_K, need_f16_V);
return f16_extra.end - (uintptr_t) dst->data;
}
void ggml_cuda_flash_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
ggml_cuda_set_device(ctx.device);
switch (ggml_cuda_get_best_fattn_kernel(ggml_cuda_get_device(), dst)) {
+2
View File
@@ -3,3 +3,5 @@
void ggml_cuda_flash_attn_ext(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
bool ggml_cuda_flash_attn_ext_supported(int device, const ggml_tensor * dst);
size_t ggml_cuda_flash_attn_ext_get_alloc_size(int device, const ggml_tensor * dst);
+4 -1
View File
@@ -42,7 +42,7 @@ static __global__ void k_get_rows(
template<typename src0_t, typename dst_t>
static __global__ void k_get_rows_float(
const src0_t * __restrict__ src0, const int32_t * __restrict__ src1, dst_t * __restrict__ dst,
const src0_t * src0_ptr, const int32_t * src1_ptr, dst_t * dst_ptr,
const int64_t ne00, /*const int64_t ne01, const int64_t ne02, const int64_t ne03,*/
/*const int64_t ne10,*/ const int64_t ne11, const uint3 ne12_fdv, /*const int64_t ne13,*/
/*const size_t s0,*/ const size_t s1, const size_t s2, const size_t s3,
@@ -50,6 +50,9 @@ static __global__ void k_get_rows_float(
const size_t s10, const size_t s11, const size_t s12/*, const size_t s13*/) {
ggml_cuda_pdl_lc();
const src0_t * GGML_CUDA_RESTRICT src0 = src0_ptr;
const int32_t * GGML_CUDA_RESTRICT src1 = src1_ptr;
dst_t * GGML_CUDA_RESTRICT dst = dst_ptr;
ggml_cuda_pdl_sync();
for (int64_t z = blockIdx.z; z < ne11*(int64_t)ne12_fdv.z; z += gridDim.z) {
for (int64_t i00 = blockIdx.y*blockDim.x + threadIdx.x; i00 < ne00; i00 += gridDim.y*blockDim.x) {
+5 -3
View File
@@ -801,7 +801,11 @@ static size_t ggml_backend_cuda_buffer_type_get_alignment(ggml_backend_buffer_ty
}
static size_t ggml_backend_cuda_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor) {
size_t size = ggml_nbytes(tensor);
ggml_backend_cuda_buffer_type_context * buft_ctx = (ggml_backend_cuda_buffer_type_context *) buft->context;
size_t size = tensor->op == GGML_OP_FLASH_ATTN_EXT
? ggml_cuda_flash_attn_ext_get_alloc_size(buft_ctx->device, tensor)
: ggml_nbytes(tensor);
int64_t ne0 = tensor->ne[0];
if (ggml_is_quantized(tensor->type)) {
@@ -812,8 +816,6 @@ static size_t ggml_backend_cuda_buffer_type_get_alloc_size(ggml_backend_buffer_t
}
return size;
GGML_UNUSED(buft);
}
static const ggml_backend_buffer_type_i ggml_backend_cuda_buffer_type_interface = {
+5 -1
View File
@@ -6,11 +6,15 @@
template <typename T, typename type_acc, int ncols_dst, int block_size, bool has_fusion = false, bool is_multi_token_id = false>
static __global__ void mul_mat_vec_f(
const T * __restrict__ x, const float * __restrict__ y, const int32_t * __restrict__ ids, const ggml_cuda_mm_fusion_args_device fusion, float * __restrict__ dst,
const T * x_ptr, const float * y_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
const int ncols2, const uint3 nchannels_y, const int stride_row, const int stride_col_y2, const int stride_col_dst,
const uint3 channel_ratio, const int stride_channel_x, const int stride_channel_y, const int stride_channel_dst,
const uint3 sample_ratio, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst,
const int ids_stride) {
const T * GGML_CUDA_RESTRICT x = x_ptr;
const float * GGML_CUDA_RESTRICT y = y_ptr;
const int32_t * GGML_CUDA_RESTRICT ids = ids_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const int row = blockIdx.x;
// for MUL_MAT_ID - blockIdx.y = n_expert_used, blockIdx.z = ncols_dst (tokens)
const int channel_dst = blockIdx.y;
+5 -1
View File
@@ -476,12 +476,16 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
__launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
static __global__ void mul_mat_vec_q(
const void * __restrict__ vx, const void * __restrict__ vy, const int32_t * __restrict__ ids, const ggml_cuda_mm_fusion_args_device fusion, float * __restrict__ dst,
const void * vx_ptr, const void * vy_ptr, const int32_t * ids_ptr, const ggml_cuda_mm_fusion_args_device fusion, float * dst_ptr,
const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
const uint32_t stride_col_dst, const uint3 channel_ratio, const uint32_t stride_channel_x,
const uint32_t stride_channel_y, const uint32_t stride_channel_dst, const uint3 sample_ratio,
const uint32_t stride_sample_x, const uint32_t stride_sample_y, const uint32_t stride_sample_dst,
const uint32_t ids_stride) {
const void * GGML_CUDA_RESTRICT vx = vx_ptr;
const void * GGML_CUDA_RESTRICT vy = vy_ptr;
const int32_t * GGML_CUDA_RESTRICT ids = ids_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
constexpr int qk = ggml_cuda_type_traits<type>::qk;
constexpr int qi = ggml_cuda_type_traits<type>::qi;
+3 -1
View File
@@ -3,10 +3,12 @@
__launch_bounds__(CUDA_QUANTIZE_BLOCK_SIZE, 1)
static __global__ void quantize_q8_1(
const float * __restrict__ x, void * __restrict__ vy,
const float * x_ptr, void * vy_ptr,
const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t ne0, const uint32_t ne1, const uint3 ne2) {
ggml_cuda_pdl_lc();
const float * GGML_CUDA_RESTRICT x = x_ptr;
void * GGML_CUDA_RESTRICT vy = vy_ptr;
const int64_t i0 = (int64_t)blockDim.x*blockIdx.x + threadIdx.x;
if (i0 >= ne0) {
+3 -1
View File
@@ -2,7 +2,9 @@
// Row reduction kernel template - compute sum (norm=false) or mean (norm=true)
template <bool norm>
static __global__ void reduce_rows_f32(const float * __restrict__ x, float * __restrict__ dst, const int ncols) {
static __global__ void reduce_rows_f32(const float * x_ptr, float * dst_ptr, const int ncols) {
const float * GGML_CUDA_RESTRICT x = x_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const int row = blockIdx.x;
const int col = threadIdx.x;
+6 -3
View File
@@ -111,9 +111,9 @@ static void set_rows_cuda_quant(
}
template <typename src_t, typename idx_t, typename dst_t>
static __global__ void k_set_rows(const src_t * __restrict__ src0,
const idx_t * __restrict__ src1,
dst_t * __restrict__ dst,
static __global__ void k_set_rows(const src_t * src0_ptr,
const idx_t * src1_ptr,
dst_t * dst_ptr,
const int64_t ne_total,
const int64_t ne10,
const int64_t ne11,
@@ -133,6 +133,9 @@ static __global__ void k_set_rows(const src_t * __restrict__ src0,
const uint3 ne02,
const uint3 ne11_fd,
const uint3 ne12_fd) {
const src_t * GGML_CUDA_RESTRICT src0 = src0_ptr;
const idx_t * GGML_CUDA_RESTRICT src1 = src1_ptr;
dst_t * GGML_CUDA_RESTRICT dst = dst_ptr;
const int64_t i = int64_t(blockDim.x) * blockIdx.x + threadIdx.x;
if (i >= ne_total) {
+7 -3
View File
@@ -3,12 +3,16 @@
#include "unary.cuh"
template <bool apply_silu, size_t split_d_inner, size_t d_conv>
static __global__ void ssm_conv_f32(const float * __restrict__ src0, const float * __restrict__ src1,
const float * __restrict__ bias,
static __global__ void ssm_conv_f32(const float * src0_ptr, const float * src1_ptr,
const float * bias_ptr,
const int src0_nb0, const int src0_nb1, const int src0_nb2, const int src1_nb1,
float * __restrict__ dst, const int dst_nb0, const int dst_nb1, const int dst_nb2,
float * dst_ptr, const int dst_nb0, const int dst_nb1, const int dst_nb2,
const int64_t n_t) {
ggml_cuda_pdl_lc();
const float * GGML_CUDA_RESTRICT src0 = src0_ptr;
const float * GGML_CUDA_RESTRICT src1 = src1_ptr;
const float * GGML_CUDA_RESTRICT bias = bias_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
GGML_UNUSED(src0_nb0);
const int tid = threadIdx.x;
const int bidx = blockIdx.x;
+22 -6
View File
@@ -17,14 +17,22 @@ using namespace cub;
#endif // __clang__
template <size_t splitD, size_t N, size_t L_template>
__global__ void __launch_bounds__(splitD, 1)
ssm_scan_f32(const float *__restrict__ src0, const float *__restrict__ src1, const float *__restrict__ src2,
const float *__restrict__ src3, const float *__restrict__ src4, const float *__restrict__ src5,
const int32_t * __restrict__ src6, float * __restrict__ dst,
ssm_scan_f32(const float * src0_ptr, const float * src1_ptr, const float * src2_ptr,
const float * src3_ptr, const float * src4_ptr, const float * src5_ptr,
const int32_t * src6_ptr, float * dst_ptr,
const int src0_nb2, const int src0_nb3, const int src1_nb2, const int src1_nb3,
const int src2_nb1, const int src2_nb2, const int src3_nb1,
const int src4_nb2, const int src4_nb3, const int src5_nb2, const int src5_nb3,
const int64_t s_off, const int64_t d_inner, const int64_t L_param)
{
const float * GGML_CUDA_RESTRICT src0 = src0_ptr;
const float * GGML_CUDA_RESTRICT src1 = src1_ptr;
const float * GGML_CUDA_RESTRICT src2 = src2_ptr;
const float * GGML_CUDA_RESTRICT src3 = src3_ptr;
const float * GGML_CUDA_RESTRICT src4 = src4_ptr;
const float * GGML_CUDA_RESTRICT src5 = src5_ptr;
const int32_t * GGML_CUDA_RESTRICT src6 = src6_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const size_t L = L_template == 0 ? L_param : L_template;
ggml_cuda_pdl_sync();
const float *s0_block = (const float *)((const char *)src0 + src6[blockIdx.x] * src0_nb3 + blockIdx.y * splitD * src0_nb2);
@@ -118,13 +126,21 @@ __global__ void __launch_bounds__(splitD, 1)
template <int c_factor, int d_state>
__global__ void __launch_bounds__(d_state, 1)
ssm_scan_f32_group(
const float * __restrict__ src0, const float * __restrict__ src1, const float * __restrict__ src2,
const float * __restrict__ src3, const float * __restrict__ src4, const float * __restrict__ src5,
const int32_t * __restrict__ src6, float * __restrict__ dst,
const float * src0_ptr, const float * src1_ptr, const float * src2_ptr,
const float * src3_ptr, const float * src4_ptr, const float * src5_ptr,
const int32_t * src6_ptr, float * dst_ptr,
const int src0_nb2, const int src0_nb3, const int src1_nb2, const int src1_nb3,
const int src2_nb1, const int src2_nb2, const int src3_nb1,
const int src4_nb2, const int src4_nb3, const int src5_nb2, const int src5_nb3,
const int64_t s_off, const int64_t n_head, const int64_t d_head, const int64_t n_group, const int64_t n_tok) {
const float * GGML_CUDA_RESTRICT src0 = src0_ptr;
const float * GGML_CUDA_RESTRICT src1 = src1_ptr;
const float * GGML_CUDA_RESTRICT src2 = src2_ptr;
const float * GGML_CUDA_RESTRICT src3 = src3_ptr;
const float * GGML_CUDA_RESTRICT src4 = src4_ptr;
const float * GGML_CUDA_RESTRICT src5 = src5_ptr;
const int32_t * GGML_CUDA_RESTRICT src6 = src6_ptr;
float * GGML_CUDA_RESTRICT dst = dst_ptr;
const int warp = threadIdx.x / WARP_SIZE;
const int lane = threadIdx.x % WARP_SIZE;
+17 -10
View File
@@ -56,6 +56,20 @@ struct htp_opnode {
}
std::vector<const ggml_tensor *> get_inputs() const {
if (fused.empty()) {
int last_non_null = -1;
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (node->src[i]) {
last_non_null = i;
}
}
std::vector<const ggml_tensor *> inputs(last_non_null + 1, nullptr);
for (int i = 0; i <= last_non_null; i++) {
inputs[i] = node->src[i];
}
return inputs;
}
std::vector<const ggml_tensor *> inputs(GGML_MAX_SRC, nullptr);
std::vector<const ggml_tensor *> outputs;
outputs.push_back(node);
@@ -82,12 +96,8 @@ struct htp_opnode {
};
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (fused.empty()) {
inputs[i] = node->src[i];
} else {
if (node->src[i]) {
add_input(node->src[i]);
}
if (node->src[i]) {
add_input(node->src[i]);
}
}
for (const auto * f : fused) {
@@ -98,10 +108,7 @@ struct htp_opnode {
}
}
if (!fused.empty()) {
inputs.resize(count);
}
inputs.resize(count);
return inputs;
}
+33
View File
@@ -150,6 +150,7 @@ class Keys:
EMBD_LENGTH_PER_LAYER_INP = "{arch}.embedding_length_per_layer_input"
SWIGLU_CLAMP_EXP = "{arch}.swiglu_clamp_exp"
SWIGLU_CLAMP_SHEXP = "{arch}.swiglu_clamp_shexp"
HIDDEN_ACT = "{arch}.hidden_activation"
DENSE_FEAT_IN_SIZE = "{arch}.{dense}_feat_in"
DENSE_FEAT_OUT_SIZE = "{arch}.{dense}_feat_out"
@@ -263,6 +264,7 @@ class Keys:
ADD_PREFIX = "tokenizer.ggml.add_space_prefix"
REMOVE_EXTRA_WS = "tokenizer.ggml.remove_extra_whitespaces"
PRECOMPILED_CHARSMAP = "tokenizer.ggml.precompiled_charsmap"
SUPPRESS_TOKENS = "tokenizer.ggml.suppress_tokens"
HF_JSON = "tokenizer.huggingface.json"
RWKV = "tokenizer.rwkv.world"
CHAT_TEMPLATE = "tokenizer.chat_template"
@@ -509,6 +511,7 @@ class MODEL_ARCH(IntEnum):
MAINCODER = auto()
KIMI_LINEAR = auto()
TALKIE = auto()
MELLUM = auto()
class VISION_PROJECTOR_TYPE(IntEnum):
@@ -729,6 +732,7 @@ class MODEL_TENSOR(IntEnum):
V_ENC_EMBD_CLS = auto()
V_ENC_EMBD_PATCH = auto()
V_ENC_EMBD_NORM = auto()
V_ENC_EMBD_PATCH_NORM = auto() # allow multiple norms in the same embd, e.g. for gemma4u
V_ENC_EMBD_POS = auto()
V_ENC_INPUT_NORM = auto()
V_ENC_ATTN_QKV = auto()
@@ -1029,6 +1033,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
MODEL_ARCH.MAINCODER: "maincoder",
MODEL_ARCH.KIMI_LINEAR: "kimi-linear",
MODEL_ARCH.TALKIE: "talkie",
MODEL_ARCH.MELLUM: "mellum",
}
VISION_PROJECTOR_TYPE_NAMES: dict[VISION_PROJECTOR_TYPE, str] = {
@@ -1247,6 +1252,7 @@ TENSOR_NAMES: dict[MODEL_TENSOR, str] = {
MODEL_TENSOR.V_ENC_EMBD_CLS: "v.class_embd",
MODEL_TENSOR.V_ENC_EMBD_PATCH: "v.patch_embd",
MODEL_TENSOR.V_ENC_EMBD_NORM: "v.norm_embd",
MODEL_TENSOR.V_ENC_EMBD_PATCH_NORM: "v.patch_norm.{bid}",
MODEL_TENSOR.V_ENC_EMBD_POS: "v.position_embd",
MODEL_TENSOR.V_ENC_ATTN_QKV: "v.blk.{bid}.attn_qkv",
MODEL_TENSOR.V_ENC_ATTN_Q: "v.blk.{bid}.attn_q",
@@ -1428,6 +1434,7 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.V_ENC_EMBD_CLS,
MODEL_TENSOR.V_ENC_EMBD_PATCH,
MODEL_TENSOR.V_ENC_EMBD_NORM,
MODEL_TENSOR.V_ENC_EMBD_PATCH_NORM,
MODEL_TENSOR.V_ENC_EMBD_POS,
MODEL_TENSOR.V_ENC_EMBD_IMGNL,
MODEL_TENSOR.V_ENC_EMBD_VSEP,
@@ -3994,6 +4001,13 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_GATE_SHEXP,
MODEL_TENSOR.FFN_DOWN_SHEXP,
MODEL_TENSOR.FFN_EXP_PROBS_B,
# NextN/MTP tensors (Step3p5 draft head)
MODEL_TENSOR.NEXTN_EH_PROJ,
MODEL_TENSOR.NEXTN_EMBED_TOKENS,
MODEL_TENSOR.NEXTN_ENORM,
MODEL_TENSOR.NEXTN_HNORM,
MODEL_TENSOR.NEXTN_SHARED_HEAD_HEAD,
MODEL_TENSOR.NEXTN_SHARED_HEAD_NORM,
],
MODEL_ARCH.LLAMA_EMBED: [
MODEL_TENSOR.TOKEN_EMBD,
@@ -4085,6 +4099,23 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
MODEL_TENSOR.FFN_UP,
MODEL_TENSOR.LAYER_OUT_SCALE,
],
MODEL_ARCH.MELLUM: [
MODEL_TENSOR.TOKEN_EMBD,
MODEL_TENSOR.OUTPUT_NORM,
MODEL_TENSOR.OUTPUT,
MODEL_TENSOR.ATTN_NORM,
MODEL_TENSOR.ATTN_Q,
MODEL_TENSOR.ATTN_Q_NORM,
MODEL_TENSOR.ATTN_K,
MODEL_TENSOR.ATTN_K_NORM,
MODEL_TENSOR.ATTN_V,
MODEL_TENSOR.ATTN_OUT,
MODEL_TENSOR.FFN_NORM,
MODEL_TENSOR.FFN_GATE_INP,
MODEL_TENSOR.FFN_GATE_EXP,
MODEL_TENSOR.FFN_DOWN_EXP,
MODEL_TENSOR.FFN_UP_EXP,
],
# TODO
}
@@ -4319,6 +4350,8 @@ class VisionProjectorType:
GEMMA3NA = "gemma3na"
GEMMA4V = "gemma4v"
GEMMA4A = "gemma4a"
GEMMA4UV = "gemma4uv" # "unified" variant
GEMMA4UA = "gemma4ua" # "unified" variant
PHI4 = "phi4"
IDEFICS3 = "idefics3"
PIXTRAL = "pixtral"
+6
View File
@@ -853,6 +853,9 @@ class GGUFWriter:
def add_swiglu_clamp_shexp(self, values: Sequence[float]) -> None:
self.add_array(Keys.LLM.SWIGLU_CLAMP_SHEXP.format(arch=self.arch), values)
def add_hidden_act(self, value: str) -> None:
self.add_string(Keys.LLM.HIDDEN_ACT.format(arch=self.arch), value)
def add_expert_group_scale(self, value: float) -> None:
self.add_float32(Keys.LLM.EXPERT_GROUP_SCALE.format(arch=self.arch), value)
@@ -1110,6 +1113,9 @@ class GGUFWriter:
self.add_string(Keys.Tokenizer.CHAT_TEMPLATE, value)
def add_suppress_tokens(self, tokens: Sequence[int]) -> None:
self.add_array(Keys.Tokenizer.SUPPRESS_TOKENS, tokens)
def add_normalizer_lowercase(self, value: bool) -> None:
self.add_bool(Keys.Tokenizer.NORMALIZER_LOWERCASE, value)
+6
View File
@@ -1426,6 +1426,7 @@ class TensorNameMap:
"model.vision_tower.patch_embedder.input_proj", # gemma4
"vision_tower.patch_embed.patchifier.proj", # dots.ocr
"vision_model.conv1", # Step3-VL
"model.vision_embedder.patch_dense", # gemma4 unified
),
MODEL_TENSOR.V_ENC_EMBD_NORM: (
@@ -1433,6 +1434,10 @@ class TensorNameMap:
"vision_tower.patch_embed.patchifier.norm", # dots.ocr
),
MODEL_TENSOR.V_ENC_EMBD_PATCH_NORM: (
"model.vision_embedder.patch_ln{bid}", # gemma4 unified
),
MODEL_TENSOR.V_ENC_EMBD_POS: (
"vision_tower.vision_model.embeddings.position_embedding",
"model.vision_tower.embeddings.position_embedding", # minicpmv4_6
@@ -1448,6 +1453,7 @@ class TensorNameMap:
"vision_model.radio_model.model.patch_generator.pos_embed", # Nemotron Nano v2 VL
"model.vision_tower.patch_embedder.position_embedding_table", # gemma4
"vision_model.positional_embedding", # Step3-VL
"model.vision_embedder.pos_embedding", # gemma4 unified
),
MODEL_TENSOR.V_ENC_EMBD_IMGNL: (
+3
View File
@@ -2279,6 +2279,9 @@ void kcpp_init_audio_proj(clip_ctx * ctx_a)
case PROJECTOR_TYPE_GEMMA4A:
audio_preproc = std::make_unique<mtmd_audio_preprocessor_gemma4a>(ctx_a);
break;
case PROJECTOR_TYPE_GEMMA4UA:
audio_preproc = std::make_unique<mtmd_audio_preprocessor_gemma4ua>(ctx_a);
break;
default:
GGML_ABORT("unsupported audio projector type");
}
+4
View File
@@ -135,6 +135,7 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
{ LLM_ARCH_MAINCODER, "maincoder" },
{ LLM_ARCH_KIMI_LINEAR, "kimi-linear" },
{ LLM_ARCH_TALKIE, "talkie" },
{ LLM_ARCH_MELLUM, "mellum" },
{ LLM_ARCH_UNKNOWN, "(unknown)" },
};
@@ -195,6 +196,7 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
{ LLM_KV_MOE_LATENT_SIZE, "%s.moe_latent_size" },
{ LLM_KV_NEXTN_PREDICT_LAYERS, "%s.nextn_predict_layers" },
{ LLM_KV_NUM_DEEPSTACK_LAYERS, "%s.n_deepstack_layers" },
{ LLM_KV_HIDDEN_ACT, "%s.hidden_activation" },
{ LLM_KV_POOLING_TYPE, "%s.pooling_type" },
{ LLM_KV_LOGIT_SCALE, "%s.logit_scale" },
{ LLM_KV_DECODER_START_TOKEN_ID, "%s.decoder_start_token_id" },
@@ -245,6 +247,7 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
{ LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, "%s.attention.indexer.key_length" },
{ LLM_KV_ATTENTION_INDEXER_TOP_K, "%s.attention.indexer.top_k" },
{ LLM_KV_ATTENTION_SHARED_KV_LAYERS, "%s.attention.shared_kv_layers" },
{ LLM_KV_ATTENTION_RECURRENT_LAYERS, "%s.attention.recurrent_layers" },
{ LLM_KV_ROPE_DIMENSION_COUNT, "%s.rope.dimension_count" },
{ LLM_KV_ROPE_DIMENSION_COUNT_SWA, "%s.rope.dimension_count_swa" },
@@ -326,6 +329,7 @@ static const std::map<llm_kv, const char *> LLM_KV_NAMES = {
{ LLM_KV_TOKENIZER_FIM_PAD_ID, "tokenizer.ggml.fim_pad_token_id" },
{ LLM_KV_TOKENIZER_FIM_REP_ID, "tokenizer.ggml.fim_rep_token_id" },
{ LLM_KV_TOKENIZER_FIM_SEP_ID, "tokenizer.ggml.fim_sep_token_id" },
{ LLM_KV_TOKENIZER_SUPPRESS_TOKENS, "tokenizer.ggml.suppress_tokens" },
{ LLM_KV_ADAPTER_TYPE, "adapter.type" },
{ LLM_KV_ADAPTER_LORA_ALPHA, "adapter.lora.alpha" },
+4
View File
@@ -139,6 +139,7 @@ enum llm_arch {
LLM_ARCH_MAINCODER,
LLM_ARCH_KIMI_LINEAR,
LLM_ARCH_TALKIE,
LLM_ARCH_MELLUM,
LLM_ARCH_UNKNOWN,
};
@@ -199,6 +200,7 @@ enum llm_kv {
LLM_KV_MOE_LATENT_SIZE,
LLM_KV_NEXTN_PREDICT_LAYERS,
LLM_KV_NUM_DEEPSTACK_LAYERS,
LLM_KV_HIDDEN_ACT,
LLM_KV_POOLING_TYPE,
LLM_KV_LOGIT_SCALE,
LLM_KV_DECODER_START_TOKEN_ID,
@@ -249,6 +251,7 @@ enum llm_kv {
LLM_KV_ATTENTION_INDEXER_KEY_LENGTH,
LLM_KV_ATTENTION_INDEXER_TOP_K,
LLM_KV_ATTENTION_SHARED_KV_LAYERS,
LLM_KV_ATTENTION_RECURRENT_LAYERS,
LLM_KV_ROPE_DIMENSION_COUNT,
LLM_KV_ROPE_DIMENSION_COUNT_SWA,
@@ -315,6 +318,7 @@ enum llm_kv {
LLM_KV_TOKENIZER_FIM_PAD_ID,
LLM_KV_TOKENIZER_FIM_REP_ID,
LLM_KV_TOKENIZER_FIM_SEP_ID,
LLM_KV_TOKENIZER_SUPPRESS_TOKENS,
LLM_KV_ADAPTER_TYPE,
LLM_KV_ADAPTER_LORA_ALPHA,
+71 -70
View File
@@ -61,19 +61,20 @@ llama_context::llama_context(
cparams.n_rs_seq = 0;
}
cparams.n_threads = params.n_threads;
cparams.n_threads_batch = params.n_threads_batch;
cparams.yarn_ext_factor = params.yarn_ext_factor >= 0.0f ? params.yarn_ext_factor : hparams.yarn_ext_factor;
cparams.yarn_attn_factor = params.yarn_attn_factor >= 0.0f ? params.yarn_attn_factor : hparams.yarn_attn_factor;
cparams.yarn_beta_fast = params.yarn_beta_fast >= 0.0f ? params.yarn_beta_fast : hparams.yarn_beta_fast;
cparams.yarn_beta_slow = params.yarn_beta_slow >= 0.0f ? params.yarn_beta_slow : hparams.yarn_beta_slow;
cparams.embeddings = params.embeddings;
cparams.embeddings_pre_norm = false;
cparams.embeddings_pre_norm_masked = false;
cparams.offload_kqv = params.offload_kqv;
cparams.no_perf = params.no_perf;
cparams.pooling_type = params.pooling_type;
cparams.warmup = false;
cparams.n_threads = params.n_threads;
cparams.n_threads_batch = params.n_threads_batch;
cparams.yarn_ext_factor = params.yarn_ext_factor >= 0.0f ? params.yarn_ext_factor : hparams.yarn_ext_factor;
cparams.yarn_attn_factor = params.yarn_attn_factor >= 0.0f ? params.yarn_attn_factor : hparams.yarn_attn_factor;
cparams.yarn_beta_fast = params.yarn_beta_fast >= 0.0f ? params.yarn_beta_fast : hparams.yarn_beta_fast;
cparams.yarn_beta_slow = params.yarn_beta_slow >= 0.0f ? params.yarn_beta_slow : hparams.yarn_beta_slow;
cparams.embeddings = params.embeddings;
cparams.embeddings_nextn = false;
cparams.embeddings_nextn_masked = false;
cparams.offload_kqv = params.offload_kqv;
cparams.no_perf = params.no_perf;
cparams.pooling_type = params.pooling_type;
cparams.warmup = false;
cparams.n_ctx = params.n_ctx == 0 ? hparams.n_ctx_train : params.n_ctx;
cparams.rope_freq_base = params.rope_freq_base == 0.0f ? hparams.rope_freq_base_train : params.rope_freq_base;
@@ -899,34 +900,34 @@ float * llama_context::get_embeddings_seq(llama_seq_id seq_id) {
return it->second.data();
}
float * llama_context::get_embeddings_pre_norm() {
float * llama_context::get_embeddings_nextn() {
output_reorder();
return embd_pre_norm.data;
return embd_nextn.data;
}
float * llama_context::get_embeddings_pre_norm_ith(int32_t i) {
float * llama_context::get_embeddings_nextn_ith(int32_t i) {
output_reorder();
try {
if (embd_pre_norm.data == nullptr) {
throw std::runtime_error("no pre-norm embeddings");
if (embd_nextn.data == nullptr) {
throw std::runtime_error("no nextn embeddings");
}
const uint32_t n_embd = model.hparams.n_embd;
if (!cparams.embeddings_pre_norm_masked) {
// unmasked: pre-norm rows are stored densely, indexed by raw token position.
if (i < 0 || (size_t)(i + 1) * n_embd > embd_pre_norm.size) {
throw std::runtime_error(format("out of range [0, %zu)", embd_pre_norm.size / n_embd));
if (!cparams.embeddings_nextn_masked) {
// unmasked: nextn rows are stored densely, indexed by raw token position.
if (i < 0 || (size_t)(i + 1) * n_embd > embd_nextn.size) {
throw std::runtime_error(format("out of range [0, %zu)", embd_nextn.size / n_embd));
}
return embd_pre_norm.data + (size_t) i * n_embd;
return embd_nextn.data + (size_t) i * n_embd;
}
const int64_t j = output_resolve_row(i);
return embd_pre_norm.data + j*n_embd;
return embd_nextn.data + j*n_embd;
} catch (const std::exception & err) {
LLAMA_LOG_ERROR("%s: invalid pre-norm embeddings id %d, reason: %s\n", __func__, i, err.what());
LLAMA_LOG_ERROR("%s: invalid nextn embeddings id %d, reason: %s\n", __func__, i, err.what());
#ifndef NDEBUG
GGML_ABORT("fatal error");
#else
@@ -1115,11 +1116,11 @@ void llama_context::set_embeddings(bool value) {
//sched_need_reserve = true;
}
void llama_context::set_embeddings_pre_norm(bool value, bool masked) {
void llama_context::set_embeddings_nextn(bool value, bool masked) {
LLAMA_LOG_DEBUG("%s: value = %d, masked = %d\n", __func__, value, masked);
cparams.embeddings_pre_norm = value;
cparams.embeddings_pre_norm_masked = masked;
cparams.embeddings_nextn = value;
cparams.embeddings_nextn_masked = masked;
}
void llama_context::set_causal_attn(bool value) {
@@ -1336,7 +1337,7 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll
}
int llama_context::encode(const llama_batch & batch_inp) {
// MTP hook batches carry both token (next-token id) and embd (h_pre_norm row),
// MTP hook batches carry both token (next-token id) and embd (h_nextn row),
// so accept either present rather than requiring exactly one.
GGML_ASSERT(batch_inp.token || batch_inp.embd);
@@ -1409,9 +1410,9 @@ int llama_context::encode(const llama_batch & batch_inp) {
}
}
auto * t_logits = res->get_logits();
auto * t_embd = res->get_embd_pooled() ? res->get_embd_pooled() : res->get_embd();
auto * t_h_pre_norm = cparams.embeddings_pre_norm ? res->get_h_pre_norm() : nullptr;
auto * t_logits = res->get_logits();
auto * t_embd = res->get_embd_pooled() ? res->get_embd_pooled() : res->get_embd();
auto * t_h_nextn = cparams.embeddings_nextn ? res->get_h_nextn() : nullptr;
// extract logits
if (logits.data && t_logits) {
@@ -1477,14 +1478,14 @@ int llama_context::encode(const llama_batch & batch_inp) {
}
}
// extract pre-norm embeddings (hidden state before the final output norm)
if (embd_pre_norm.data && t_h_pre_norm && cparams.pooling_type == LLAMA_POOLING_TYPE_NONE) {
ggml_backend_t backend_h = ggml_backend_sched_get_tensor_backend(sched.get(), t_h_pre_norm);
// extract nextn embeddings (hidden state before the final output norm)
if (embd_nextn.data && t_h_nextn && cparams.pooling_type == LLAMA_POOLING_TYPE_NONE) {
ggml_backend_t backend_h = ggml_backend_sched_get_tensor_backend(sched.get(), t_h_nextn);
GGML_ASSERT(backend_h != nullptr);
const uint32_t n_embd = hparams.n_embd;
GGML_ASSERT(n_tokens*n_embd <= (int64_t) embd_pre_norm.size);
ggml_backend_tensor_get_async(backend_h, t_h_pre_norm, embd_pre_norm.data, 0, n_tokens*n_embd*sizeof(float));
GGML_ASSERT(n_tokens*n_embd <= (int64_t) embd_nextn.size);
ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn.data, 0, n_tokens*n_embd*sizeof(float));
}
// TODO: hacky solution
@@ -1639,7 +1640,7 @@ static bool needs_raw_logits(const llama_ubatch & ubatch, const std::map<llama_s
}
int llama_context::decode(const llama_batch & batch_inp) {
// MTP hook batches carry both token (next-token id) and embd (h_pre_norm row),
// MTP hook batches carry both token (next-token id) and embd (h_nextn row),
// so accept either present rather than requiring exactly one.
GGML_ASSERT(batch_inp.token || batch_inp.embd);
@@ -1839,9 +1840,9 @@ int llama_context::decode(const llama_batch & batch_inp) {
// ggml_graph_dump_dot(gf, NULL, "llama.dot");
//}
auto * t_logits = res->get_logits();
auto * t_embd = cparams.embeddings ? res->get_embd() : nullptr;
auto * t_h_pre_norm = cparams.embeddings_pre_norm ? res->get_h_pre_norm() : nullptr;
auto * t_logits = res->get_logits();
auto * t_embd = cparams.embeddings ? res->get_embd() : nullptr;
auto * t_h_nextn = cparams.embeddings_nextn ? res->get_h_nextn() : nullptr;
if (t_embd && res->get_embd_pooled()) {
t_embd = res->get_embd_pooled();
@@ -1922,22 +1923,22 @@ int llama_context::decode(const llama_batch & batch_inp) {
}
}
// extract pre-norm embeddings (hidden state before the final output norm)
// extract nextn embeddings before
// only meaningful in LLAMA_POOLING_TYPE_NONE (per-token); other pooling modes are ignored.
{
const bool masked = cparams.embeddings_pre_norm_masked;
const bool masked = cparams.embeddings_nextn_masked;
const int64_t n_rows = masked ? n_outputs : (int64_t) ubatch.n_tokens;
const int64_t offset = masked ? n_outputs_prev : n_tokens_prev;
if (embd_pre_norm.data && t_h_pre_norm && n_rows > 0 && cparams.pooling_type == LLAMA_POOLING_TYPE_NONE) {
ggml_backend_t backend_h = ggml_backend_sched_get_tensor_backend(sched.get(), t_h_pre_norm);
if (embd_nextn.data && t_h_nextn && n_rows > 0 && cparams.pooling_type == LLAMA_POOLING_TYPE_NONE) {
ggml_backend_t backend_h = ggml_backend_sched_get_tensor_backend(sched.get(), t_h_nextn);
GGML_ASSERT(backend_h != nullptr);
const uint32_t n_embd = hparams.n_embd;
float * embd_pre_norm_out = embd_pre_norm.data + offset*n_embd;
const uint32_t n_embd = hparams.n_embd;
float * embd_nextn_out = embd_nextn.data + offset*n_embd;
GGML_ASSERT((offset + n_rows)*n_embd <= (int64_t) embd_pre_norm.size);
ggml_backend_tensor_get_async(backend_h, t_h_pre_norm, embd_pre_norm_out, 0, n_rows*n_embd*sizeof(float));
GGML_ASSERT((offset + n_rows)*n_embd <= (int64_t) embd_nextn.size);
ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn_out, 0, n_rows*n_embd*sizeof(float));
}
}
@@ -2029,9 +2030,9 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) {
const auto n_embd = hparams.n_embd;
const auto n_embd_out = hparams.n_embd_out();
bool has_logits = true;
bool has_embd = cparams.embeddings;
bool has_embd_pre_norm = cparams.embeddings_pre_norm;
bool has_logits = true;
bool has_embd = cparams.embeddings;
bool has_embd_nextn = cparams.embeddings_nextn;
// TODO: hacky enc-dec support
if (model.arch == LLM_ARCH_T5) {
@@ -2043,14 +2044,14 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) {
size_t backend_float_count = 0;
size_t backend_token_count = 0;
logits.size = has_logits ? n_vocab*n_outputs_max : 0;
embd.size = has_embd ? n_embd_out*n_outputs_max : 0;
embd_pre_norm.size = has_embd_pre_norm ? n_embd*n_outputs_max : 0;
logits.size = has_logits ? n_vocab*n_outputs_max : 0;
embd.size = has_embd ? n_embd_out*n_outputs_max : 0;
embd_nextn.size = has_embd_nextn ? n_embd*n_outputs_max : 0;
if (has_embd_pre_norm && !cparams.embeddings_pre_norm_masked) {
// unmasked: pre-norm row exists for every token in the batch, not just
if (has_embd_nextn && !cparams.embeddings_nextn_masked) {
// unmasked: nextn row exists for every token in the batch, not just
// those flagged via batch.logits[i] -> size by token count instead.
embd_pre_norm.size = (size_t) n_embd * n_batch;
embd_nextn.size = (size_t) n_embd * n_batch;
}
// Allocate backend sampling output buffers if there are backend samplers configured.
@@ -2067,7 +2068,7 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) {
const size_t prev_size = buf_output ? ggml_backend_buffer_get_size(buf_output.get()) : 0;
const size_t new_size =
(logits.size + embd.size + embd_pre_norm.size + backend_float_count) * sizeof(float) +
(logits.size + embd.size + embd_nextn.size + backend_float_count) * sizeof(float) +
( backend_token_count) * sizeof(llama_token);
// alloc only when more than the current capacity is required
@@ -2084,7 +2085,7 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) {
buf_output = nullptr;
logits.data = nullptr;
embd.data = nullptr;
embd_pre_norm.data = nullptr;
embd_nextn.data = nullptr;
}
auto * buft = ggml_backend_cpu_buffer_type();
@@ -2113,8 +2114,8 @@ uint32_t llama_context::output_reserve(int32_t n_outputs) {
embd = has_embd ? buffer_view<float>{(float *) (base + offset), embd.size} : buffer_view<float>{nullptr, 0};
offset += embd.size * sizeof(float);
embd_pre_norm = has_embd_pre_norm ? buffer_view<float>{(float *) (base + offset), embd_pre_norm.size} : buffer_view<float>{nullptr, 0};
offset += embd_pre_norm.size * sizeof(float);
embd_nextn = has_embd_nextn ? buffer_view<float>{(float *) (base + offset), embd_nextn.size} : buffer_view<float>{nullptr, 0};
offset += embd_nextn.size * sizeof(float);
if (has_sampling) {
sampling.logits = {(float *) (base + offset), (size_t)(n_vocab*n_outputs_max)};
@@ -2182,9 +2183,9 @@ void llama_context::output_reorder() {
}
}
if (embd_pre_norm.size > 0) {
if (embd_nextn.size > 0) {
for (uint64_t k = 0; k < n_embd; k++) {
std::swap(embd_pre_norm.data[i0*n_embd + k], embd_pre_norm.data[i1*n_embd + k]);
std::swap(embd_nextn.data[i0*n_embd + k], embd_nextn.data[i1*n_embd + k]);
}
}
@@ -3598,20 +3599,20 @@ float * llama_get_embeddings_seq(llama_context * ctx, llama_seq_id seq_id) {
return ctx->get_embeddings_seq(seq_id);
}
void llama_set_embeddings_pre_norm(llama_context * ctx, bool value, bool masked) {
ctx->set_embeddings_pre_norm(value, masked);
void llama_set_embeddings_nextn(llama_context * ctx, bool value, bool masked) {
ctx->set_embeddings_nextn(value, masked);
}
float * llama_get_embeddings_pre_norm(llama_context * ctx) {
float * llama_get_embeddings_nextn(llama_context * ctx) {
ctx->synchronize();
return ctx->get_embeddings_pre_norm();
return ctx->get_embeddings_nextn();
}
float * llama_get_embeddings_pre_norm_ith(llama_context * ctx, int32_t i) {
float * llama_get_embeddings_nextn_ith(llama_context * ctx, int32_t i) {
ctx->synchronize();
return ctx->get_embeddings_pre_norm_ith(i);
return ctx->get_embeddings_nextn_ith(i);
}
bool llama_set_sampler(llama_context * ctx, llama_seq_id seq_id, llama_sampler * smpl) {
+7 -7
View File
@@ -84,8 +84,8 @@ struct llama_context {
float * get_embeddings_ith(int32_t i);
float * get_embeddings_seq(llama_seq_id seq_id);
float * get_embeddings_pre_norm();
float * get_embeddings_pre_norm_ith(int32_t i);
float * get_embeddings_nextn();
float * get_embeddings_nextn_ith(int32_t i);
llama_token * get_sampled_tokens() const;
llama_token get_sampled_token_ith(int32_t idx);
@@ -110,7 +110,7 @@ struct llama_context {
void set_abort_callback(bool (*abort_callback)(void * data), void * abort_callback_data);
void set_embeddings (bool value);
void set_embeddings_pre_norm(bool value, bool masked);
void set_embeddings_nextn(bool value, bool masked);
void set_causal_attn(bool value);
void set_warmup(bool value);
@@ -282,10 +282,10 @@ private:
// populated only when pooling_type == LLAMA_POOLING_TYPE_NONE
buffer_view<float> embd = {nullptr, 0};
// hidden state before the final output norm (2-dimensional array: [n_outputs][n_embd])
// populated only when cparams.embeddings_pre_norm is enabled and the model graph
// sets llm_graph_result::t_h_pre_norm
buffer_view<float> embd_pre_norm = {nullptr, 0};
// hidden state required by the nextn layers (2-dimensional array: [n_outputs][n_embd])
// populated only when cparams.embeddings_nextn is enabled and the model graph
// sets llm_graph_result::t_h_nextn
buffer_view<float> embd_nextn = {nullptr, 0};
struct sampling_info {
// !samplers.empty() to check if any samplers are active
+2 -2
View File
@@ -29,8 +29,8 @@ struct llama_cparams {
float yarn_beta_slow;
bool embeddings;
bool embeddings_pre_norm; // also extract the hidden state before the final output norm
bool embeddings_pre_norm_masked; // extract for only rows where batch.logits != 0
bool embeddings_nextn; // also extract the hidden state before the final output norm
bool embeddings_nextn_masked; // extract for only rows where batch.logits != 0
bool causal_attn;
bool offload_kqv;
bool flash_attn;
+4 -8
View File
@@ -89,18 +89,14 @@ LLAMA_API ggml_backend_dev_t llama_model_get_device(const struct llama_model * m
LLAMA_API llama_memory_breakdown llama_get_memory_breakdown(const struct llama_context * ctx);
//
// pre-norm embeddings (hidden state before the final output norm)
//
// Set whether the context outputs pre-norm embeddings or not
// Set whether the context outputs nextn embeddings or not
// If masked == true, output the embeddings only for the tokens with batch.logits != 0
// If masked == false, output the embeddings for all tokens in the batch regardless of batch.logits
LLAMA_API void llama_set_embeddings_pre_norm(struct llama_context * ctx, bool value, bool masked);
LLAMA_API void llama_set_embeddings_nextn(struct llama_context * ctx, bool value, bool masked);
// mirrors:
// LLAMA_API float * llama_get_embeddings(struct llama_context * ctx);
LLAMA_API float * llama_get_embeddings_pre_norm (struct llama_context * ctx);
LLAMA_API float * llama_get_embeddings_nextn(struct llama_context * ctx);
// LLAMA_API float * llama_get_embeddings_ith(struct llama_context * ctx, int32_t i);
LLAMA_API float * llama_get_embeddings_pre_norm_ith(struct llama_context * ctx, int32_t i);
LLAMA_API float * llama_get_embeddings_nextn_ith(struct llama_context * ctx, int32_t i);
+2 -2
View File
@@ -929,8 +929,8 @@ void llm_graph_result::set_outputs() {
if (t_embd_pooled != nullptr) {
ggml_set_output(t_embd_pooled);
}
if (t_h_pre_norm != nullptr) {
ggml_set_output(t_h_pre_norm);
if (t_h_nextn != nullptr) {
ggml_set_output(t_h_nextn);
}
for (auto & [seq_id, t] : t_sampled) {
if (t != nullptr) {
+4 -3
View File
@@ -36,7 +36,8 @@ enum llm_graph_type {
LLM_GRAPH_TYPE_DECODER_MTP,
};
enum llm_ffn_op_type {
enum llm_ffn_op_type : int {
LLM_FFN_NONE = 0, // sentinel: unset; archs must assign before use
LLM_FFN_SILU,
LLM_FFN_GELU,
LLM_FFN_RELU,
@@ -702,7 +703,7 @@ public:
ggml_tensor * get_logits() const { return t_logits; }
ggml_tensor * get_embd() const { return t_embd; }
ggml_tensor * get_embd_pooled() const { return t_embd_pooled; }
ggml_tensor * get_h_pre_norm() const { return t_h_pre_norm; }
ggml_tensor * get_h_nextn() const { return t_h_nextn; }
ggml_cgraph * get_gf() const { return gf; }
ggml_context * get_ctx() const { return ctx_compute.get(); }
@@ -731,7 +732,7 @@ public:
ggml_tensor * t_logits = nullptr;
ggml_tensor * t_embd = nullptr;
ggml_tensor * t_embd_pooled = nullptr;
ggml_tensor * t_h_pre_norm = nullptr; // [n_embd, n_outputs] hidden state before final output norm
ggml_tensor * t_h_nextn = nullptr; // [n_embd, n_outputs] hidden state before final output norm
std::map<llama_seq_id, ggml_tensor*> t_sampled_logits;
std::map<llama_seq_id, ggml_tensor*> t_candidates;
+19 -6
View File
@@ -8,18 +8,31 @@
void llama_hparams::set_swa_pattern(uint32_t n_pattern, bool dense_first) {
if (dense_first) {
for (uint32_t il = 0; il < n_layer; ++il) {
swa_layers[il] = n_pattern == 0 || (il % n_pattern != 0);
is_swa_impl[il] = n_pattern == 0 || (il % n_pattern != 0);
}
} else {
for (uint32_t il = 0; il < n_layer; ++il) {
swa_layers[il] = n_pattern == 0 || (il % n_pattern < (n_pattern - 1));
is_swa_impl[il] = n_pattern == 0 || (il % n_pattern < (n_pattern - 1));
}
}
}
// TODO: implement
//void llama_hparams::set_recr_pattern(uint32_t n_pattern, bool dense_first) {
// if (dense_first) {
// for (uint32_t il = 0; il < n_layer; ++il) {
// is_recr_impl[il] = n_pattern == 0 || (il % n_pattern != 0);
// }
// } else {
// for (uint32_t il = 0; il < n_layer; ++il) {
// is_recr_impl[il] = n_pattern == 0 || (il % n_pattern < (n_pattern - 1));
// }
// }
//}
bool llama_hparams::is_swa_any() const {
for (uint32_t il = 0; il < n_layer; ++il) {
if (swa_layers[il]) {
if (is_swa_impl[il]) {
return true;
}
}
@@ -193,9 +206,9 @@ uint32_t llama_hparams::n_embd_s() const {
return ssm_d_state * ssm_d_inner;
}
bool llama_hparams::is_recurrent(uint32_t il) const {
bool llama_hparams::is_recr(uint32_t il) const {
if (il < n_layer) {
return recurrent_layer_arr[il];
return is_recr_impl[il];
}
GGML_ABORT("%s: il (%u) out of bounds (n_layer: %u)\n", __func__, il, n_layer);
@@ -207,7 +220,7 @@ uint32_t llama_hparams::n_pos_per_embd() const {
bool llama_hparams::is_swa(uint32_t il) const {
if (il < n_layer) {
return swa_layers[il];
return is_swa_impl[il];
}
GGML_ABORT("fatal error");
+30 -12
View File
@@ -23,6 +23,9 @@ enum llama_swa_type {
LLAMA_SWA_TYPE_SYMMETRIC = 3,
};
// forward declaration; full definition in llama-graph.h
enum llm_ffn_op_type : int;
struct llama_hparams_posnet {
uint32_t n_embd;
uint32_t n_layer;
@@ -34,6 +37,9 @@ struct llama_hparams_convnext {
};
struct llama_hparams {
// note: use the `_impl` suffix to avoid name conflict between members and getters
// for example: n_embd_out() vs n_embd_out_impl
bool vocab_only;
bool no_alloc;
bool rope_finetuned;
@@ -43,7 +49,7 @@ struct llama_hparams {
uint32_t n_ctx_train; // context size the model was trained on
uint32_t n_embd;
uint32_t n_layer;
int32_t n_layer_kv_from_start = -1; // if non-negative, the first n_layer_kv_from_start layers have KV cache
int32_t n_layer_kv_from_start = -1; // if non-negative, the first n_layer_kv_from_start layers have KV cache
uint32_t n_expert = 0;
uint32_t n_expert_used = 0;
uint32_t n_rel_attn_bkts = 0;
@@ -134,11 +140,15 @@ struct llama_hparams {
llama_swa_type swa_type = LLAMA_SWA_TYPE_NONE;
// the size of the sliding window (0 - no SWA)
uint32_t n_swa = 0;
// if swa_layers[il] == 1, then layer il is SWA
// if swa_layers[il] == 0, then layer il is dense (i.e. non-SWA)
// if is_swa_impl[il] == 1, then layer il is SWA
// if is_swa_impl[il] == 0, then layer il is dense (i.e. non-SWA)
// by default, all layers are dense
// note: using uint32_t type for compatibility reason
std::array<uint32_t, LLAMA_MAX_LAYERS> swa_layers;
std::array<uint32_t, LLAMA_MAX_LAYERS> is_swa_impl;
// for hybrid state space models
std::array<uint32_t, LLAMA_MAX_LAYERS> is_recr_impl;
// for State Space Models
uint32_t ssm_d_conv = 0;
@@ -150,9 +160,6 @@ struct llama_hparams {
// for Kimi Linear KDA
uint32_t n_embd_head_kda = 0;
// for hybrid state space models
std::array<bool, LLAMA_MAX_LAYERS> recurrent_layer_arr;
bool ssm_dt_b_c_rms = false;
float f_clamp_kqv = 0.0f;
@@ -227,6 +234,14 @@ struct llama_hparams {
enum llama_rope_scaling_type rope_scaling_type_train = LLAMA_ROPE_SCALING_TYPE_NONE;
// Resolved FFN gated activation flavor for archs that read
// `<arch>.hidden_activation` from the GGUF (e.g. ModernBert derivatives).
// Defaults to LLM_FFN_NONE (sentinel = 0); the mapping from the GGUF
// string to a real op is done at hparam-load time via
// llm_ffn_op_type_from_string() in llama-model.cpp, mirroring how
// rope_scaling_type_train is handled.
enum llm_ffn_op_type llm_ffn_op;
// Step35: optional per-layer clamps for (Swi)GLU
std::array<float, LLAMA_MAX_LAYERS> swiglu_clamp_exp; // clamping for expert FFN
std::array<float, LLAMA_MAX_LAYERS> swiglu_clamp_shexp; // shared expert
@@ -255,6 +270,14 @@ struct llama_hparams {
// return true if one of the layers is SWA
bool is_swa_any() const;
bool is_swa(uint32_t il) const;
// TODO: implement
//void set_recr_pattern(uint32_t n_pattern, bool dense_first = false);
// whether or not the given layer is recurrent (for hybrid models)
bool is_recr(uint32_t il) const;
uint32_t n_head(uint32_t il = 0) const;
uint32_t n_head_kv(uint32_t il = 0) const;
@@ -296,13 +319,8 @@ struct llama_hparams {
// dimension of the recurrent state embeddings
uint32_t n_embd_s() const;
// whether or not the given layer is recurrent (for hybrid models)
bool is_recurrent(uint32_t il) const;
uint32_t n_pos_per_embd() const;
bool is_swa(uint32_t il) const;
// note: currently only support if either all or none of the layers are MLA
bool is_mla() const;
+2 -2
View File
@@ -44,7 +44,7 @@ llama_memory_hybrid_iswa::llama_memory_hybrid_iswa(
n_ubatch,
n_pad,
filter_attn == nullptr ?
[&](int32_t il) { return !hparams.is_recurrent(il); }
[&](int32_t il) { return !hparams.is_recr(il); }
: filter_attn,
nullptr
)),
@@ -57,7 +57,7 @@ llama_memory_hybrid_iswa::llama_memory_hybrid_iswa(
n_seq_max,
n_rs_seq,
filter_recr == nullptr ?
[&](int32_t il) { return hparams.is_recurrent(il); }
[&](int32_t il) { return hparams.is_recr(il); }
: filter_recr
)) {}
+2 -2
View File
@@ -45,7 +45,7 @@ llama_memory_hybrid::llama_memory_hybrid(
n_swa,
swa_type,
filter_attn == nullptr ?
[&](int32_t il) { return !hparams.is_recurrent(il); }
[&](int32_t il) { return !hparams.is_recr(il); }
: filter_attn,
nullptr
)),
@@ -58,7 +58,7 @@ llama_memory_hybrid::llama_memory_hybrid(
n_seq_max,
n_rs_seq,
filter_recr == nullptr ?
[&](int32_t il) { return hparams.is_recurrent(il); }
[&](int32_t il) { return hparams.is_recr(il); }
: filter_recr
)) {}
+4 -4
View File
@@ -146,7 +146,7 @@ namespace GGUFMeta {
const enum gguf_type arr_type = gguf_get_arr_type(ctx, k);
return ArrayInfo {
arr_type,
size_t(gguf_get_arr_n(ctx, k)),
gguf_get_arr_n(ctx, k),
arr_type == GGUF_TYPE_STRING ? nullptr : gguf_get_arr_data(ctx, k),
};
}
@@ -445,7 +445,7 @@ namespace GGUFMeta {
}
if (n > N_MAX) {
throw std::runtime_error(format("n > N_MAX: %u > %u for key %s", (uint32_t) n, (uint32_t) N_MAX, key.c_str()));
throw std::runtime_error(format("n > N_MAX: %u > %u for key %s", n, (uint32_t) N_MAX, key.c_str()));
}
if (gguf_get_kv_type(metadata, kid) == GGUF_TYPE_ARRAY) {
@@ -502,9 +502,9 @@ namespace GGUFMeta {
}
// TODO: this is not very clever - figure out something better
template bool llama_model_loader::get_key_or_arr<std::array<int, 4>>(enum llm_kv kid, std::array<int, 4> & result, uint32_t n, bool required);
template bool llama_model_loader::get_key_or_arr<std::array<int, 4>> (enum llm_kv kid, std::array<int, 4> & result, uint32_t n, bool required);
template bool llama_model_loader::get_key_or_arr<std::array<uint32_t, 512>>(enum llm_kv kid, std::array<uint32_t, 512> & result, uint32_t n, bool required);
template bool llama_model_loader::get_key_or_arr<std::array<float, 512>>(enum llm_kv kid, std::array<float, 512> & result, uint32_t n, bool required);
template bool llama_model_loader::get_key_or_arr<std::array<float, 512>>(enum llm_kv kid, std::array<float, 512> & result, uint32_t n, bool required);
llama_model_loader::llama_model_loader(
+5 -4
View File
@@ -14,9 +14,6 @@
bool llama_model_saver_supports_arch(llm_arch arch) {
switch (arch) {
case LLM_ARCH_QWEN3NEXT:
case LLM_ARCH_QWEN35:
case LLM_ARCH_QWEN35MOE:
case LLM_ARCH_PLAMO3:
case LLM_ARCH_GEMMA3:
case LLM_ARCH_GEMMA3N:
@@ -29,6 +26,7 @@ bool llama_model_saver_supports_arch(llm_arch arch) {
case LLM_ARCH_APERTUS:
case LLM_ARCH_MIMO2:
case LLM_ARCH_STEP35:
case LLM_ARCH_MELLUM:
return false;
default:
return true;
@@ -106,6 +104,8 @@ void llama_model_saver::add_kv(const enum llm_kv key, const Container & value, c
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_INT8, value.data(), n_values);
} else if (std::is_same<typename Container::value_type, uint32_t>::value) {
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_UINT32, value.data(), n_values);
} else if (std::is_same<typename Container::value_type, bool>::value) {
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_BOOL, value.data(), n_values);
} else if (std::is_same<typename Container::value_type, int32_t>::value) {
gguf_set_arr_data(gguf_ctx, llm_kv(key).c_str(), GGUF_TYPE_INT32, value.data(), n_values);
} else if (std::is_same<typename Container::value_type, float>::value) {
@@ -244,7 +244,7 @@ void llama_model_saver::add_kv_from_model() {
add_kv(LLM_KV_EMBEDDING_SCALE, hparams.f_embedding_scale);
add_kv(LLM_KV_TOKEN_SHIFT_COUNT, hparams.token_shift_count);
add_kv(LLM_KV_INTERLEAVE_MOE_LAYER_STEP, hparams.n_moe_layer_step);
// add_kv(LLM_KV_FULL_ATTENTION_INTERVAL, ???);
// add_kv(LLM_KV_FULL_ATTENTION_INTERVAL, ???); // saved as LLM_KV_ATTENTION_RECURRENT_LAYERS instead
add_kv(LLM_KV_ATTENTION_HEAD_COUNT, hparams.n_head_arr, true);
add_kv(LLM_KV_ATTENTION_HEAD_COUNT_KV, hparams.n_head_kv_arr, true);
@@ -278,6 +278,7 @@ void llama_model_saver::add_kv_from_model() {
add_kv(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, hparams.indexer_n_head);
add_kv(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, hparams.indexer_head_size);
add_kv(LLM_KV_ATTENTION_INDEXER_TOP_K, hparams.indexer_top_k);
add_kv(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, true);
const float rope_scaling_factor = hparams.rope_freq_scale_train == 1.0f ? 0.0f : 1.0f/hparams.rope_freq_scale_train;
+45 -16
View File
@@ -107,6 +107,7 @@
#include "models/mamba-base.cpp"
#include "models/mamba.cpp"
#include "models/mamba2.cpp"
#include "models/mellum.cpp"
#include "models/mimo2.cpp"
#include "models/minicpm.cpp"
#include "models/minicpm3.cpp"
@@ -213,6 +214,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
return new llama_model_mpt(params);
case LLM_ARCH_STABLELM:
return new llama_model_stablelm(params);
case LLM_ARCH_MELLUM:
return new llama_model_mellum(params);
case LLM_ARCH_QWEN:
return new llama_model_qwen(params);
case LLM_ARCH_QWEN2:
@@ -503,10 +506,10 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
// count only the same type of previous layers to avoid this
auto get_il_eff = [&](const size_t il){
size_t ret = 0;
const bool il_is_recurrent = hparams.is_recurrent(il);
const bool il_is_swa = hparams.is_swa(il);
const bool il_is_recr = hparams.is_recr(il);
const bool il_is_swa = hparams.is_swa(il);
for (size_t il_prev = 0; il_prev < il; il_prev++) {
ret += hparams.is_recurrent(il_prev) == il_is_recurrent && hparams.is_swa(il_prev) == il_is_swa;
ret += hparams.is_recr(il_prev) == il_is_recr && hparams.is_swa(il_prev) == il_is_swa;
}
return ret;
};
@@ -683,7 +686,7 @@ struct ggml_backend_meta_split_state llama_meta_device_get_split_state(const str
};
auto get_split_granularity = [&](int64_t blck_size, uint32_t il, const std::vector<std::pair<int64_t, uint32_t>> & segments) -> std::vector<int64_t> {
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// linear attention
const int64_t head_dim = hparams.ssm_d_state;
const int64_t granularity_qkv = std::lcm(blck_size, head_dim);
@@ -896,6 +899,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_12B_A2_5B: return "12B.A2.5B";
case LLM_TYPE_16B_A1B: return "16B.A1B";
case LLM_TYPE_21B_A3B: return "21B.A3B";
case LLM_TYPE_24B_A2B: return "24B.A2B";
@@ -954,6 +958,28 @@ static llama_rope_scaling_type llama_rope_scaling_type_from_string(const std::st
return LLAMA_ROPE_SCALING_TYPE_UNSPECIFIED;
}
// Maps the GGUF `<arch>.hidden_activation` string to the FFN op type used by the
// graph builders. Only gated activations that map cleanly to llm_ffn_op_type are
// listed; unrecognized values fall back to GeGLU, which matches the historical
// default for ModernBert-style architectures.
static const std::map<std::string, llm_ffn_op_type> LLM_FFN_OP_TYPES_FROM_STRING = {
{ "gelu", LLM_FFN_GEGLU },
{ "geglu", LLM_FFN_GEGLU },
{ "silu", LLM_FFN_SWIGLU },
{ "swish", LLM_FFN_SWIGLU },
{ "swiglu", LLM_FFN_SWIGLU },
{ "relu", LLM_FFN_RELU },
{ "reglu", LLM_FFN_REGLU },
};
llm_ffn_op_type llm_ffn_op_type_from_string(const std::string & name, llm_ffn_op_type fallback) {
const auto it = LLM_FFN_OP_TYPES_FROM_STRING.find(name);
if (it != LLM_FFN_OP_TYPES_FROM_STRING.end()) {
return it->second;
}
return fallback;
}
// CPU: ACCEL -> GPU host -> CPU extra -> CPU
static buft_list_t make_cpu_buft_list(const std::vector<llama_device> & devices, bool use_extra_bufts, bool no_host) {
buft_list_t buft_list;
@@ -1183,18 +1209,16 @@ void llama_model_base::load_hparams(llama_model_loader & ml) {
std::fill(hparams.n_head_arr.begin(), hparams.n_head_arr.end(), 0);
std::fill(hparams.n_head_kv_arr.begin(), hparams.n_head_kv_arr.end(), 0);
std::fill(hparams.n_ff_arr.begin(), hparams.n_ff_arr.end(), 0);
std::fill(
hparams.recurrent_layer_arr.begin(),
hparams.recurrent_layer_arr.end(),
llm_arch_is_recurrent(ml.get_arch()));
std::fill(hparams.rope_sections.begin(), hparams.rope_sections.end(), 0);
std::fill(hparams.swa_layers.begin(), hparams.swa_layers.end(), 0);
std::fill(hparams.is_swa_impl.begin(), hparams.is_swa_impl.end(), 0);
std::fill(hparams.is_recr_impl.begin(), hparams.is_recr_impl.end(), llm_arch_is_recurrent(ml.get_arch()) ? 1 : 0);
std::fill(hparams.xielu_alpha_n.begin(), hparams.xielu_alpha_n.end(), 0.0f);
std::fill(hparams.xielu_alpha_p.begin(), hparams.xielu_alpha_p.end(), 0.0f);
std::fill(hparams.xielu_beta.begin(), hparams.xielu_beta.end(), 0.0f);
std::fill(hparams.xielu_eps.begin(), hparams.xielu_eps.end(), 0.0f);
std::fill(hparams.xielu_beta.begin(), hparams.xielu_beta.end(), 0.0f);
std::fill(hparams.xielu_eps.begin(), hparams.xielu_eps.end(), 0.0f);
std::fill(hparams.swiglu_clamp_exp.begin(), hparams.swiglu_clamp_exp.end(), 0.0f);
std::fill(hparams.swiglu_clamp_shexp.begin(), hparams.swiglu_clamp_shexp.end(), 0.0f);
@@ -1927,7 +1951,11 @@ void llama_model::print_info() const {
LLAMA_LOG_INFO("%s: n_ff_shexp = %d\n", __func__, hparams.n_ff_shexp);
}
if (arch == LLM_ARCH_QWEN3MOE || arch == LLM_ARCH_OPENAI_MOE || arch == LLM_ARCH_QWEN3VLMOE || arch == LLM_ARCH_RND1) {
if (arch == LLM_ARCH_MELLUM ||
arch == LLM_ARCH_QWEN3MOE ||
arch == LLM_ARCH_OPENAI_MOE ||
arch == LLM_ARCH_QWEN3VLMOE ||
arch == LLM_ARCH_RND1) {
LLAMA_LOG_INFO("%s: n_ff_exp = %d\n", __func__, hparams.n_ff_exp);
}
@@ -2144,18 +2172,18 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params,
filter_recr = [&](int32_t) { return true; };
} else if (arch == LLM_ARCH_NEMOTRON_H || arch == LLM_ARCH_NEMOTRON_H_MOE) {
filter_attn = [&](int32_t il) {
return !hparams.is_recurrent(il) && hparams.n_ff(il) == 0;
return !hparams.is_recr(il) && hparams.n_ff(il) == 0;
};
filter_recr = [&](int32_t il) {
return hparams.is_recurrent(il) && hparams.n_ff(il) == 0;
return hparams.is_recr(il) && hparams.n_ff(il) == 0;
};
} else if (arch == LLM_ARCH_QWEN35 || arch == LLM_ARCH_QWEN35MOE) {
const uint32_t n_main = hparams.n_layer - hparams.nextn_predict_layers;
filter_attn = [&, n_main](int32_t il) {
return (uint32_t)il < n_main && !hparams.is_recurrent(il);
return (uint32_t)il < n_main && !hparams.is_recr(il);
};
filter_recr = [&, n_main](int32_t il) {
return (uint32_t)il < n_main && hparams.is_recurrent(il);
return (uint32_t)il < n_main && hparams.is_recr(il);
};
}
@@ -2515,6 +2543,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
case LLM_ARCH_MIMO2:
case LLM_ARCH_STEP35:
case LLM_ARCH_TALKIE:
case LLM_ARCH_MELLUM:
return LLAMA_ROPE_TYPE_NEOX;
case LLM_ARCH_QWEN2VL:
+5
View File
@@ -116,6 +116,7 @@ enum llm_type {
LLM_TYPE_A13B,
LLM_TYPE_7B_A1B,
LLM_TYPE_8B_A1B, // lfm2moe
LLM_TYPE_12B_A2_5B,
LLM_TYPE_16B_A1B,
LLM_TYPE_21B_A3B, // Ernie MoE small
LLM_TYPE_24B_A2B, // lfm2moe
@@ -145,6 +146,10 @@ enum llm_type {
std::string llama_rope_scaling_type_name(llama_rope_scaling_type rope_scaling_type);
// Map a GGUF activation-name string to llm_ffn_op_type. Returns `fallback` if
// the string is empty or not recognized.
llm_ffn_op_type llm_ffn_op_type_from_string(const std::string & name, llm_ffn_op_type fallback);
struct llama_layer_posnet {
// resnet
struct ggml_tensor * norm1 = nullptr;
+36 -1
View File
@@ -578,6 +578,7 @@ struct llm_tokenizer_bpe : llm_tokenizer {
case LLAMA_VOCAB_PRE_TYPE_CODESHELL:
case LLAMA_VOCAB_PRE_TYPE_EXAONE:
case LLAMA_VOCAB_PRE_TYPE_MINERVA:
case LLAMA_VOCAB_PRE_TYPE_MELLUM2:
regex_exprs = {
"\\p{N}",
"'s|'t|'re|'ve|'m|'ll|'d| ?\\p{L}+| ?\\p{N}+| ?[^\\s\\p{L}\\p{N}]+|\\s+(?!\\S)",
@@ -657,6 +658,15 @@ struct llm_tokenizer_bpe : llm_tokenizer {
"[^\\r\\n\\p{L}\\p{N}]?((?=[\\p{L}])([^a-z]))*((?=[\\p{L}])([^A-Z]))+(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])?|[^\\r\\n\\p{L}\\p{N}]?((?=[\\p{L}])([^a-z]))+((?=[\\p{L}])([^A-Z]))*(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])?|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n/]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+",
};
break;
case LLAMA_VOCAB_PRE_TYPE_GRANITE_EMB_MULTI:
// Same lookaheads as GPT4O but with \p{M} added so combining marks
// (diacritics) attach to their base letters. Avoids excessive
// backtracking on scripts that use them heavily (Bengali, Hindi,
// Telugu, Thai, ...). See PR #22716 for benchmarks.
regex_exprs = {
"[^\\r\\n\\p{L}\\p{N}]?((?=[\\p{L}\\p{M}])([^a-z]))*((?=[\\p{L}\\p{M}])([^A-Z]))+(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])?|[^\\r\\n\\p{L}\\p{N}]?((?=[\\p{L}\\p{M}])([^a-z]))+((?=[\\p{L}\\p{M}])([^A-Z]))*(?:'[sS]|'[tT]|'[rR][eE]|'[vV][eE]|'[mM]|'[lL][lL]|'[dD])?|\\p{N}{1,3}| ?[^\\s\\p{L}\\p{N}]+[\\r\\n/]*|\\s*[\\r\\n]+|\\s+(?!\\S)|\\s+",
};
break;
case LLAMA_VOCAB_PRE_TYPE_TINY_AYA:
regex_exprs = {
// original regex from tokenizer.json: "\\d{1,3}(?=(?:\\d{3})*\\b)"
@@ -2030,6 +2040,8 @@ struct llama_vocab::impl {
// set of all tokens that cause "end of generation"
std::set<llama_token> special_eog_ids;
std::vector<llama_token> suppress_tokens;
std::unique_ptr<llm_tokenizer> tokenizer;
std::vector<char> precompiled_charsmap;
@@ -2378,7 +2390,8 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
tokenizer_pre == "jais-2") {
pre_type = LLAMA_VOCAB_PRE_TYPE_JAIS2;
} else if (
tokenizer_pre == "gemma4") {
tokenizer_pre == "gemma4" ||
tokenizer_pre == "granite-embed-multi-311m") {
pre_type = LLAMA_VOCAB_PRE_TYPE_GEMMA4;
escape_whitespaces = true;
} else if (
@@ -2488,6 +2501,11 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
tokenizer_pre == "talkie") {
pre_type = LLAMA_VOCAB_PRE_TYPE_GPT4O;
clean_spaces = false;
} else if (
tokenizer_pre == "granite-embed-multi-97m") {
pre_type = LLAMA_VOCAB_PRE_TYPE_GRANITE_EMB_MULTI;
clean_spaces = false;
ignore_merges = true;
} else if (
tokenizer_pre == "tiny_aya") {
pre_type = LLAMA_VOCAB_PRE_TYPE_TINY_AYA;
@@ -2546,6 +2564,9 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
tokenizer_pre == "solar-open") {
pre_type = LLAMA_VOCAB_PRE_TYPE_SOLAR_OPEN;
clean_spaces = false;
} else if (
tokenizer_pre == "mellum2") {
pre_type = LLAMA_VOCAB_PRE_TYPE_MELLUM2;
} else {
throw std::runtime_error(format("unknown pre-tokenizer type: '%s'", tokenizer_pre.c_str()));
}
@@ -2753,6 +2774,16 @@ void llama_vocab::impl::load(llama_model_loader & ml, const LLM_KV & kv) {
// Lowercase normalizer flag (consulted by WPM / whitespace BPE)
ml.get_key(LLM_KV_TOKENIZER_NORMALIZER_LOWERCASE, normalizer_lowercase, false);
// suppress tokens
{
const int suppress_idx = gguf_find_key(ctx, kv(LLM_KV_TOKENIZER_SUPPRESS_TOKENS).c_str());
if (suppress_idx != -1) {
const int n = gguf_get_arr_n(ctx, suppress_idx);
const int32_t * data = (const int32_t *) gguf_get_arr_data(ctx, suppress_idx);
suppress_tokens.assign(data, data + n);
}
}
// auto-detect special tokens by text
// TODO: convert scripts should provide these tokens through the KV metadata LLM_KV_TOKENIZER_...
// for now, we apply this workaround to find the tokens based on their text
@@ -4228,6 +4259,10 @@ bool llama_vocab::get_normalizer_lowercase() const {
return pimpl->normalizer_lowercase;
}
const std::vector<llama_token> & llama_vocab::get_suppress_tokens() const {
return pimpl->suppress_tokens;
}
int llama_vocab::max_token_len() const {
return pimpl->max_token_len;
}
+58 -54
View File
@@ -9,60 +9,62 @@
// pre-tokenization types
enum llama_vocab_pre_type {
LLAMA_VOCAB_PRE_TYPE_DEFAULT = 0,
LLAMA_VOCAB_PRE_TYPE_LLAMA3 = 1,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_LLM = 2,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_CODER = 3,
LLAMA_VOCAB_PRE_TYPE_FALCON = 4,
LLAMA_VOCAB_PRE_TYPE_MPT = 5,
LLAMA_VOCAB_PRE_TYPE_STARCODER = 6,
LLAMA_VOCAB_PRE_TYPE_GPT2 = 7,
LLAMA_VOCAB_PRE_TYPE_REFACT = 8,
LLAMA_VOCAB_PRE_TYPE_COMMAND_R = 9,
LLAMA_VOCAB_PRE_TYPE_STABLELM2 = 10,
LLAMA_VOCAB_PRE_TYPE_QWEN2 = 11,
LLAMA_VOCAB_PRE_TYPE_OLMO = 12,
LLAMA_VOCAB_PRE_TYPE_DBRX = 13,
LLAMA_VOCAB_PRE_TYPE_SMAUG = 14,
LLAMA_VOCAB_PRE_TYPE_PORO = 15,
LLAMA_VOCAB_PRE_TYPE_CHATGLM3 = 16,
LLAMA_VOCAB_PRE_TYPE_CHATGLM4 = 17,
LLAMA_VOCAB_PRE_TYPE_VIKING = 18,
LLAMA_VOCAB_PRE_TYPE_JAIS = 19,
LLAMA_VOCAB_PRE_TYPE_TEKKEN = 20,
LLAMA_VOCAB_PRE_TYPE_SMOLLM = 21,
LLAMA_VOCAB_PRE_TYPE_CODESHELL = 22,
LLAMA_VOCAB_PRE_TYPE_BLOOM = 23,
LLAMA_VOCAB_PRE_TYPE_GPT3_FINNISH = 24,
LLAMA_VOCAB_PRE_TYPE_EXAONE = 25,
LLAMA_VOCAB_PRE_TYPE_CHAMELEON = 26,
LLAMA_VOCAB_PRE_TYPE_MINERVA = 27,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK3_LLM = 28,
LLAMA_VOCAB_PRE_TYPE_GPT4O = 29,
LLAMA_VOCAB_PRE_TYPE_SUPERBPE = 30,
LLAMA_VOCAB_PRE_TYPE_TRILLION = 31,
LLAMA_VOCAB_PRE_TYPE_BAILINGMOE = 32,
LLAMA_VOCAB_PRE_TYPE_LLAMA4 = 33,
LLAMA_VOCAB_PRE_TYPE_PIXTRAL = 34,
LLAMA_VOCAB_PRE_TYPE_SEED_CODER = 35,
LLAMA_VOCAB_PRE_TYPE_HUNYUAN = 36,
LLAMA_VOCAB_PRE_TYPE_KIMI_K2 = 37,
LLAMA_VOCAB_PRE_TYPE_HUNYUAN_DENSE = 38,
LLAMA_VOCAB_PRE_TYPE_GROK_2 = 39,
LLAMA_VOCAB_PRE_TYPE_GRANITE_DOCLING = 40,
LLAMA_VOCAB_PRE_TYPE_MINIMAX_M2 = 41,
LLAMA_VOCAB_PRE_TYPE_AFMOE = 42,
LLAMA_VOCAB_PRE_TYPE_SOLAR_OPEN = 43,
LLAMA_VOCAB_PRE_TYPE_YOUTU = 44,
LLAMA_VOCAB_PRE_TYPE_EXAONE_MOE = 45,
LLAMA_VOCAB_PRE_TYPE_QWEN35 = 46,
LLAMA_VOCAB_PRE_TYPE_TINY_AYA = 47,
LLAMA_VOCAB_PRE_TYPE_JOYAI_LLM = 48,
LLAMA_VOCAB_PRE_TYPE_JAIS2 = 49,
LLAMA_VOCAB_PRE_TYPE_GEMMA4 = 50,
LLAMA_VOCAB_PRE_TYPE_SARVAM_MOE = 51,
LLAMA_VOCAB_PRE_TYPE_MINICPM5 = 52,
LLAMA_VOCAB_PRE_TYPE_WHITESPACE = 53,
LLAMA_VOCAB_PRE_TYPE_DEFAULT = 0,
LLAMA_VOCAB_PRE_TYPE_LLAMA3 = 1,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_LLM = 2,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK_CODER = 3,
LLAMA_VOCAB_PRE_TYPE_FALCON = 4,
LLAMA_VOCAB_PRE_TYPE_MPT = 5,
LLAMA_VOCAB_PRE_TYPE_STARCODER = 6,
LLAMA_VOCAB_PRE_TYPE_GPT2 = 7,
LLAMA_VOCAB_PRE_TYPE_REFACT = 8,
LLAMA_VOCAB_PRE_TYPE_COMMAND_R = 9,
LLAMA_VOCAB_PRE_TYPE_STABLELM2 = 10,
LLAMA_VOCAB_PRE_TYPE_QWEN2 = 11,
LLAMA_VOCAB_PRE_TYPE_OLMO = 12,
LLAMA_VOCAB_PRE_TYPE_DBRX = 13,
LLAMA_VOCAB_PRE_TYPE_SMAUG = 14,
LLAMA_VOCAB_PRE_TYPE_PORO = 15,
LLAMA_VOCAB_PRE_TYPE_CHATGLM3 = 16,
LLAMA_VOCAB_PRE_TYPE_CHATGLM4 = 17,
LLAMA_VOCAB_PRE_TYPE_VIKING = 18,
LLAMA_VOCAB_PRE_TYPE_JAIS = 19,
LLAMA_VOCAB_PRE_TYPE_TEKKEN = 20,
LLAMA_VOCAB_PRE_TYPE_SMOLLM = 21,
LLAMA_VOCAB_PRE_TYPE_CODESHELL = 22,
LLAMA_VOCAB_PRE_TYPE_BLOOM = 23,
LLAMA_VOCAB_PRE_TYPE_GPT3_FINNISH = 24,
LLAMA_VOCAB_PRE_TYPE_EXAONE = 25,
LLAMA_VOCAB_PRE_TYPE_CHAMELEON = 26,
LLAMA_VOCAB_PRE_TYPE_MINERVA = 27,
LLAMA_VOCAB_PRE_TYPE_DEEPSEEK3_LLM = 28,
LLAMA_VOCAB_PRE_TYPE_GPT4O = 29,
LLAMA_VOCAB_PRE_TYPE_SUPERBPE = 30,
LLAMA_VOCAB_PRE_TYPE_TRILLION = 31,
LLAMA_VOCAB_PRE_TYPE_BAILINGMOE = 32,
LLAMA_VOCAB_PRE_TYPE_LLAMA4 = 33,
LLAMA_VOCAB_PRE_TYPE_PIXTRAL = 34,
LLAMA_VOCAB_PRE_TYPE_SEED_CODER = 35,
LLAMA_VOCAB_PRE_TYPE_HUNYUAN = 36,
LLAMA_VOCAB_PRE_TYPE_KIMI_K2 = 37,
LLAMA_VOCAB_PRE_TYPE_HUNYUAN_DENSE = 38,
LLAMA_VOCAB_PRE_TYPE_GROK_2 = 39,
LLAMA_VOCAB_PRE_TYPE_GRANITE_DOCLING = 40,
LLAMA_VOCAB_PRE_TYPE_MINIMAX_M2 = 41,
LLAMA_VOCAB_PRE_TYPE_AFMOE = 42,
LLAMA_VOCAB_PRE_TYPE_SOLAR_OPEN = 43,
LLAMA_VOCAB_PRE_TYPE_YOUTU = 44,
LLAMA_VOCAB_PRE_TYPE_EXAONE_MOE = 45,
LLAMA_VOCAB_PRE_TYPE_QWEN35 = 46,
LLAMA_VOCAB_PRE_TYPE_TINY_AYA = 47,
LLAMA_VOCAB_PRE_TYPE_JOYAI_LLM = 48,
LLAMA_VOCAB_PRE_TYPE_JAIS2 = 49,
LLAMA_VOCAB_PRE_TYPE_GEMMA4 = 50,
LLAMA_VOCAB_PRE_TYPE_SARVAM_MOE = 51,
LLAMA_VOCAB_PRE_TYPE_MINICPM5 = 52,
LLAMA_VOCAB_PRE_TYPE_WHITESPACE = 53,
LLAMA_VOCAB_PRE_TYPE_GRANITE_EMB_MULTI = 54,
LLAMA_VOCAB_PRE_TYPE_MELLUM2 = 55,
};
struct LLM_KV;
@@ -143,6 +145,8 @@ struct llama_vocab {
bool get_treat_whitespace_as_suffix() const;
bool get_normalizer_lowercase () const;
const std::vector<llama_token> & get_suppress_tokens() const;
int max_token_len() const;
int find_bpe_rank(const std::string & token_left, const std::string & token_right) const;
+1 -1
View File
@@ -11,7 +11,7 @@ void llama_model_falcon_h1::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SSM_TIME_STEP_RANK, hparams.ssm_dt_rank);
ml.get_key(LLM_KV_SSM_GROUP_COUNT, hparams.ssm_n_group);
std::fill(hparams.recurrent_layer_arr.begin(), hparams.recurrent_layer_arr.end(), true);
std::fill(hparams.is_recr_impl.begin(), hparams.is_recr_impl.end(), true);
switch (hparams.n_layer) {
case 36:
+36 -1
View File
@@ -2,7 +2,7 @@
void llama_model_gemma4::load_arch_hparams(llama_model_loader & ml) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.swa_layers, hparams.n_layer);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer);
uint32_t n_kv_shared_layers = 0;
ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
@@ -142,6 +142,31 @@ std::unique_ptr<llm_graph_context> llama_model_gemma4::build_arch_graph(const ll
// idx * x->ne[0] * x->ne[1] * ggml_element_size(x));
// }
// TODO @ngxson : maybe improve this in the future
class llm_graph_input_logits_bias : public llm_graph_input_i {
public:
llm_graph_input_logits_bias(const llama_vocab & vocab) {
arr.resize(vocab.n_tokens(), 0.0f);
for (llama_token id : vocab.get_suppress_tokens()) {
if (0 <= id && id < (int32_t)vocab.n_tokens()) {
arr[id] = -INFINITY;
}
}
}
virtual ~llm_graph_input_logits_bias() = default;
void set_input(const llama_ubatch *) override {
const int64_t n_vocab = arr.size();
ggml_backend_tensor_set(logits_bias, arr.data(), 0, n_vocab*ggml_element_size(logits_bias));
}
// bool can_reuse(const llm_graph_params & params) override;
ggml_tensor * logits_bias = nullptr; // F32 [n_vocab]
std::vector<float> arr;
};
llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_params & params) :
llm_graph_context(params),
model(model),
@@ -388,6 +413,16 @@ llama_model_gemma4::graph::graph(const llama_model & model, const llm_graph_para
cur = ggml_scale(ctx0, cur, hparams.f_final_logit_softcapping);
}
// apply logits bias if needed (e.g. for gemma4_unified patch)
// this is to mirror the suppress_tokens patch on transformers, to avoid model from outputing <image|> and <audio|> tokens (which is a known issue related to the checkpoint)
// TODO: maybe handle this inside the sampling system in the future
if (!model.vocab.get_suppress_tokens().empty()) {
auto inp_bias = std::make_unique<llm_graph_input_logits_bias>(model.vocab);
inp_bias->logits_bias = ggml_new_tensor_1d(ctx0, GGML_TYPE_F32, inp_bias->arr.size());
cur = ggml_add(ctx0, cur, inp_bias->logits_bias);
res->add_input(std::move(inp_bias));
}
cb(cur, "result_output", -1);
res->t_logits = cur;
+3 -3
View File
@@ -20,7 +20,7 @@ void llama_model_granite_hybrid::load_arch_hparams(llama_model_loader & ml) {
// A layer is recurrent IFF the n_head_kv value is set to 0
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = hparams.n_head_kv(i) == 0;
hparams.is_recr_impl[i] = hparams.n_head_kv(i) == 0;
}
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -71,7 +71,7 @@ void llama_model_granite_hybrid::load_arch_tensors(llama_model_loader &) {
// norm
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
if (hparams.is_recurrent(i)) {
if (hparams.is_recr(i)) {
// ssm layers
layer.ssm_in = create_tensor(tn(LLM_TENSOR_SSM_IN, "weight", i), {n_embd, d_in_proj}, 0);
@@ -158,7 +158,7 @@ llama_model_granite_hybrid::graph::graph(const llama_model & model, const llm_gr
cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il);
cb(cur, "attn_norm", il);
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// ssm layer //
cur = build_mamba2_layer(inp->get_recr(), cur, model, ubatch, il);
} else {
+1 -1
View File
@@ -9,7 +9,7 @@ void llama_model_jamba::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = hparams.n_head_kv(i) == 0;
hparams.is_recr_impl[i] = hparams.n_head_kv(i) == 0;
}
switch (hparams.n_layer) {
+3 -3
View File
@@ -15,7 +15,7 @@ void llama_model_kimi_linear::load_arch_hparams(llama_model_loader & ml) {
// Mark KDA layers as recurrent using n_head_kv pattern (like Jamba)
// Set n_head_kv = 0 for KDA layers (recurrent), n_head_kv = n_head for MLA layers (attention)
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = hparams.n_head_kv(i) == 0; // KDA layers are recurrent
hparams.is_recr_impl[i] = hparams.n_head_kv(i) == 0; // KDA layers are recurrent
}
// MoE parameters - Kimi uses moe_intermediate_size = 1024
@@ -53,7 +53,7 @@ void llama_model_kimi_linear::load_arch_tensors(llama_model_loader &) {
const int64_t n_embd_head_v_kda = hparams.n_embd_head_kda;
const int64_t ssm_d_conv = hparams.ssm_d_conv;
if (hparams.is_recurrent(i)) {
if (hparams.is_recr(i)) {
// Conv1d weights: try 4D first, then 3D (quantization may remove trailing 1)
// 4D: [d_conv, 1, d_inner, 1], 3D: [d_conv, 1, d_inner]
layer.ssm_q_conv = create_tensor(tn(LLM_TENSOR_SSM_CONV1D_Q, "weight", i), {ssm_d_conv, 1, n_embd_head_k_kda * n_head, 1}, TENSOR_NOT_REQUIRED);
@@ -285,7 +285,7 @@ llama_model_kimi_linear::graph::graph(const llama_model & model, const llm_graph
ggml_build_forward_expand(gf, cur);
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// === KDA Layer (Kimi Delta Attention) with Recurrent State ===
// Reference: vLLM kda.py
const auto * mctx_cur = inp_rs->mctx;
+5 -5
View File
@@ -6,7 +6,7 @@ void llama_model_lfm2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SHORTCONV_L_CACHE, hparams.n_shortconv_l_cache);
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
for (uint32_t il = 0; il < hparams.n_layer; ++il) {
hparams.recurrent_layer_arr[il] = hparams.n_head_kv(il) == 0;
hparams.is_recr_impl[il] = hparams.n_head_kv(il) == 0;
}
hparams.n_layer_dense_lead = hparams.n_layer;
switch (hparams.n_ff()) {
@@ -19,7 +19,7 @@ void llama_model_lfm2::load_arch_hparams(llama_model_loader & ml) {
if (const auto is_swa = ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false); is_swa && hparams.n_swa > 0) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
for (uint32_t il = 0; il < hparams.n_layer; ++il) {
hparams.swa_layers[il] = !hparams.recurrent_layer_arr[il];
hparams.is_swa_impl[il] = !hparams.is_recr_impl[il];
}
}
}
@@ -59,7 +59,7 @@ void llama_model_lfm2::load_arch_tensors(llama_model_loader &) {
// for operator_norm
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
if (!hparams.is_recurrent(i)) {
if (!hparams.is_recr(i)) {
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {n_embd_head_k}, 0);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_k}, 0);
GGML_ASSERT(n_embd_v_gqa == n_embd_k_gqa);
@@ -235,8 +235,8 @@ llama_model_lfm2::graph<iswa>::graph(const llama_model & model, const llm_graph_
cur = build_norm(cur, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il);
cb(cur, "model.layers.{}.operator_norm", il);
cur = hparams.is_recurrent(il) ? build_shortconv_block(cur, inp_hybrid->get_recr(), il) :
build_attn_block(cur, inp_pos, inp_hybrid->get_attn(), il);
cur = hparams.is_recr(il) ? build_shortconv_block(cur, inp_hybrid->get_recr(), il) :
build_attn_block(cur, inp_pos, inp_hybrid->get_attn(), il);
if (il == n_layer - 1 && inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
+2 -2
View File
@@ -10,7 +10,7 @@ void llama_model_lfm2moe::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_GATING_FUNC, hparams.expert_gating_func);
for (uint32_t il = 0; il < hparams.n_layer; ++il) {
hparams.recurrent_layer_arr[il] = hparams.n_head_kv(il) == 0;
hparams.is_recr_impl[il] = hparams.n_head_kv(il) == 0;
}
switch (hparams.n_layer) {
@@ -55,7 +55,7 @@ void llama_model_lfm2moe::load_arch_tensors(llama_model_loader &) {
// for operator_norm
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
if (!hparams.is_recurrent(i)) {
if (!hparams.is_recr(i)) {
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {n_embd_head_k}, 0);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_k}, 0);
GGML_ASSERT(n_embd_v_gqa == n_embd_k_gqa);
+2 -1
View File
@@ -15,7 +15,8 @@ void llama_model_llama4::load_arch_hparams(llama_model_loader & ml) {
hparams.n_attn_temp_floor_scale = 8192;
hparams.f_attn_temp_scale = 0.1f;
hparams.f_attn_temp_offset = 1.0f;
uint32_t swa_period = 4; // pattern: 3 chunked - 1 full
uint32_t swa_period = 4; // pattern: 3 chunked - 1 full
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, swa_period, false);
hparams.set_swa_pattern(swa_period);
+225
View File
@@ -0,0 +1,225 @@
#include "models.h"
void llama_model_mellum::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false);
if (hparams.n_swa > 0) {
hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
uint32_t swa_period = 4;
const auto res = ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, swa_period, false);
if (res) {
hparams.set_swa_pattern(swa_period);
} else {
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer);
}
hparams.rope_freq_base_train_swa = hparams.rope_freq_base_train;
hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train;
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
} else {
hparams.swa_type = LLAMA_SWA_TYPE_NONE;
}
switch (hparams.n_layer) {
case 28: type = LLM_TYPE_12B_A2_5B; break;
default: type = LLM_TYPE_UNKNOWN;
}
}
void llama_model_mellum::load_arch_tensors(llama_model_loader &) {
LLAMA_LOAD_LOCALS;
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
// output
output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0);
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, 0);
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head, n_embd_gqa, n_embd_gqa, 0);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_k * n_head, n_embd}, 0);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_k}, 0);
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {n_embd_head_k}, 0);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0);
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
if (n_expert == 0) {
throw std::runtime_error("n_expert must be > 0 for Mellum");
}
if (n_expert_used == 0) {
throw std::runtime_error("n_expert_used must be > 0 for Mellum");
}
const int64_t n_ff_exp = hparams.n_ff_exp ? hparams.n_ff_exp : n_ff / n_expert_used;
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff_exp, n_expert}, 0);
}
}
std::unique_ptr<llm_graph_context> llama_model_mellum::build_arch_graph(const llm_graph_params & params) const {
if (hparams.swa_type == LLAMA_SWA_TYPE_STANDARD) {
return std::make_unique<graph<true>>(*this, params);
}
return std::make_unique<graph<false>>(*this, params);
}
template <bool iswa>
llama_model_mellum::graph<iswa>::graph(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
const int64_t n_embd_head = hparams.n_embd_head_v();
GGML_ASSERT(n_embd_head == hparams.n_embd_head_k());
GGML_ASSERT(n_embd_head == n_rot);
ggml_tensor * cur;
ggml_tensor * inpL;
inpL = build_inp_embd(model.tok_embd);
// inp_pos - contains the positions
ggml_tensor * inp_pos = build_inp_pos();
using inp_attn_type = std::conditional_t<iswa, llm_graph_input_attn_kv_iswa, llm_graph_input_attn_kv>;
inp_attn_type * inp_attn = nullptr;
if constexpr (iswa) {
inp_attn = build_attn_inp_kv_iswa();
} else {
inp_attn = build_attn_inp_kv();
}
ggml_tensor * inp_out_ids = build_inp_out_ids();
for (int il = 0; il < n_layer; ++il) {
ggml_tensor * inpSA = inpL;
// norm
cur = build_norm(inpL,
model.layers[il].attn_norm, nullptr,
LLM_NORM_RMS, il);
cb(cur, "attn_norm", il);
// self_attention
{
// compute Q and K and RoPE them
auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur,
n_embd_head, n_head, n_head_kv, il);
Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, nullptr, LLM_NORM_RMS, il);
cb(Qcur, "Qcur_normed", il);
Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, nullptr, LLM_NORM_RMS, il);
cb(Kcur, "Kcur_normed", il);
const bool is_swa = hparams.is_swa(il);
if (is_swa) {
// For sliding window layers, use regular rope with no yarn rope scaling.
// This is achieved here by setting freq_scale and attn_factor to 1.
// We also set ext_factor to 0 to avoid a few unnecessary computations.
Qcur = ggml_rope_ext(
ctx0, Qcur, inp_pos, nullptr,
n_rot, rope_type, n_ctx_orig, freq_base, 1.0,
0.0, 1.0, beta_fast, beta_slow
);
Kcur = ggml_rope_ext(
ctx0, Kcur, inp_pos, nullptr,
n_rot, rope_type, n_ctx_orig, freq_base, 1.0,
0.0, 1.0, beta_fast, beta_slow
);
} else {
Qcur = ggml_rope_ext(
ctx0, Qcur, inp_pos, nullptr,
n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
ext_factor, attn_factor, beta_fast, beta_slow
);
Kcur = ggml_rope_ext(
ctx0, Kcur, inp_pos, nullptr,
n_rot, rope_type, n_ctx_orig, freq_base, freq_scale,
ext_factor, attn_factor, beta_fast, beta_slow
);
}
cb(Qcur, "Qcur", il);
cb(Kcur, "Kcur", il);
cb(Vcur, "Vcur", il);
cur = build_attn(inp_attn,
model.layers[il].wo, model.layers[il].wo_b, model.layers[il].wo_s,
Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f/sqrtf(float(n_embd_head)), il);
}
if (il == n_layer - 1 && inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
}
ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
cb(ffn_inp, "ffn_inp", il);
// MoE
cur = build_norm(ffn_inp,
model.layers[il].ffn_norm, nullptr,
LLM_NORM_RMS, il);
cb(cur, "ffn_norm", il);
ggml_tensor * moe_out =
build_moe_ffn(cur,
model.layers[il].ffn_gate_inp,
model.layers[il].ffn_up_exps,
model.layers[il].ffn_gate_exps,
model.layers[il].ffn_down_exps,
nullptr,
n_expert, n_expert_used,
LLM_FFN_SILU, true,
hparams.expert_weights_scale,
LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
il,
nullptr, nullptr,
model.layers[il].ffn_up_exps_s,
model.layers[il].ffn_gate_exps_s,
model.layers[il].ffn_down_exps_s);
cb(moe_out, "ffn_moe_out", il);
cur = moe_out;
cur = ggml_add(ctx0, cur, ffn_inp);
cb(cur, "ffn_out", il);
cur = build_cvec(cur, il);
cb(cur, "l_out", il);
// input for next layer
inpL = cur;
}
cur = inpL;
cur = build_norm(cur,
model.output_norm, nullptr,
LLM_NORM_RMS, -1);
cb(cur, "result_norm", -1);
res->t_embd = cur;
// lm_head
cur = build_lora_mm(model.output, cur, model.output_s);
cb(cur, "result_output", -1);
res->t_logits = cur;
ggml_build_forward_expand(gf, cur);
}
template struct llama_model_mellum::graph<false>;
template struct llama_model_mellum::graph<true>;
+2 -1
View File
@@ -8,7 +8,8 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp);
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.swa_layers, hparams.n_layer);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer);
float value_scale = 0.0f;
if (ml.get_key(LLM_KV_ATTENTION_VALUE_SCALE, value_scale, false) && value_scale != 1.0f) {
+16
View File
@@ -411,6 +411,18 @@ struct llama_model_stablelm : public llama_model_base {
std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
};
struct llama_model_mellum : public llama_model_base {
llama_model_mellum(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;
template <bool iswa>
struct graph : public llm_graph_context {
graph(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_qwen : public llama_model_base {
llama_model_qwen(const struct llama_model_params & params) : llama_model_base(params) {}
@@ -1913,5 +1925,9 @@ struct llama_model_step35 : public llama_model_base {
graph(const llama_model & model, const llm_graph_params & params);
};
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;
};
+10 -1
View File
@@ -14,6 +14,14 @@ void llama_model_modern_bert::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
// Some ModernBert derivatives (e.g. IBM Granite Embedding 97m R2) use
// SiLU/SwiGLU in the FFN instead of the default GELU/GeGLU.
hparams.llm_ffn_op = LLM_FFN_GEGLU;
std::string hidden_act;
if (ml.get_key(LLM_KV_HIDDEN_ACT, hidden_act, false)) {
hparams.llm_ffn_op = llm_ffn_op_type_from_string(hidden_act, LLM_FFN_GEGLU);
}
switch (hparams.n_layer) {
case 12:
type = LLM_TYPE_47M; break; // granite-embedding-small
@@ -144,7 +152,8 @@ llama_model_modern_bert::graph::graph(const llama_model & model, const llm_graph
NULL, NULL, NULL,
model.layers[il].ffn_down, NULL, NULL,
NULL,
LLM_FFN_GEGLU, LLM_FFN_SEQ, il);
hparams.llm_ffn_op,
LLM_FFN_SEQ, il);
// attentions bypass the intermediate layer
cur = ggml_add(ctx0, cur, ffn_inp);
+3 -3
View File
@@ -10,7 +10,7 @@ void llama_model_nemotron_h::load_arch_hparams(llama_model_loader & ml) {
// A layer is recurrent IFF the n_head_kv value is set to 0 and
// the n_ff value is set to 0
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = (hparams.n_head_kv(i) == 0 && hparams.n_ff(i) == 0);
hparams.is_recr_impl[i] = (hparams.n_head_kv(i) == 0 && hparams.n_ff(i) == 0);
}
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
@@ -62,7 +62,7 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader &) {
// all blocks use the attn norm
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
if (hparams.is_recurrent(i)) {
if (hparams.is_recr(i)) {
// ssm layers
layer.ssm_in = create_tensor(tn(LLM_TENSOR_SSM_IN, "weight", i), {n_embd, d_in_proj}, 0);
@@ -143,7 +143,7 @@ llama_model_nemotron_h::graph::graph(const llama_model & model, const llm_graph_
cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il);
cb(cur, "attn_norm", il);
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// ssm layer //
cur = build_mamba2_layer(inp->get_recr(), cur, model, ubatch, il);
} else if (hparams.n_ff(il) == 0) {
+3 -3
View File
@@ -12,7 +12,7 @@ void llama_model_plamo2::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SSM_GROUP_COUNT, hparams.ssm_n_group);
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = hparams.n_head_kv(i) == 0;
hparams.is_recr_impl[i] = hparams.n_head_kv(i) == 0;
}
switch (hparams.n_layer) {
@@ -54,7 +54,7 @@ void llama_model_plamo2::load_arch_tensors(llama_model_loader &) {
for (int i = 0; i < n_layer; ++i) {
auto & layer = layers[i];
bool is_mamba_layer = hparams.is_recurrent(i);
bool is_mamba_layer = hparams.is_recr(i);
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
@@ -128,7 +128,7 @@ llama_model_plamo2::graph::graph(const llama_model & model, const llm_graph_para
cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il);
// check if this layer is Mamba or Attention
const bool is_mamba_layer = hparams.is_recurrent(il);
const bool is_mamba_layer = hparams.is_recr(il);
if (is_mamba_layer) {
// PLaMo-2 Mamba layer
+16 -18
View File
@@ -18,12 +18,13 @@ void llama_model_qwen35::load_arch_hparams(llama_model_loader & ml) {
// Mark recurrent layers (linear attention layers). MTP layers are dense
// attention-only and must be flagged non-recurrent.
{
if (!ml.get_key_or_arr(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, hparams.n_layer, false)) {
const uint32_t n_main = hparams.n_layer - hparams.nextn_predict_layers;
uint32_t full_attn_interval = 4;
ml.get_key(LLM_KV_FULL_ATTENTION_INTERVAL, full_attn_interval, false);
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = (i < n_main) && ((i + 1) % full_attn_interval != 0);
hparams.is_recr_impl[i] = (i < n_main) && ((i + 1) % full_attn_interval != 0);
}
}
@@ -69,7 +70,7 @@ void llama_model_qwen35::load_arch_tensors(llama_model_loader & ml) {
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", il), { n_embd }, flags);
layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, flags);
if (!hparams.is_recurrent(il)) {
if (!hparams.is_recr(il)) {
// Attention layers
create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, flags);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", il), { n_embd_head_k * n_head, n_embd }, flags);
@@ -168,7 +169,7 @@ llama_model_qwen35::graph::graph(const llama_model & model, const llm_graph_para
ggml_build_forward_expand(gf, cur);
// Determine layer type and build appropriate attention mechanism
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// Linear attention layer (gated delta net)
cur = build_layer_attn_linear(inp->get_recr(), cur, il);
} else {
@@ -176,7 +177,7 @@ llama_model_qwen35::graph::graph(const llama_model & model, const llm_graph_para
cur = build_layer_attn(inp->get_attn(), cur, inp_pos, sections, il);
}
if (il == n_transformer_layers - 1 && inp_out_ids && cparams.embeddings_pre_norm_masked) {
if (il == n_transformer_layers - 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);
}
@@ -208,16 +209,15 @@ llama_model_qwen35::graph::graph(const llama_model & model, const llm_graph_para
}
cur = inpL;
cb(cur, "h_pre_norm", -1);
res->t_h_pre_norm = cur;
cur = build_norm(cur, model.output_norm, nullptr, LLM_NORM_RMS, -1);
if (!cparams.embeddings_pre_norm_masked && inp_out_ids) {
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);
}
// Final norm
cur = build_norm(cur, model.output_norm, nullptr, LLM_NORM_RMS, -1);
cb(cur, "result_norm", -1);
res->t_embd = cur;
@@ -624,18 +624,16 @@ llama_model_qwen35::graph_mtp::graph_mtp(const llama_model & model, const llm_gr
cur = ggml_add(ctx0, cur, ffn_residual);
cb(cur, "mtp_post_ffn", il);
// Pre-norm hidden state: used by the AR draft loop to seed the next MTP step.
// (In the trunk graph this is `t_h_pre_norm`; the MTP head reuses the same slot.)
cb(cur, "h_pre_norm", -1);
res->t_h_pre_norm = cur;
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
ggml_tensor * head_norm_w = layer.nextn.shared_head_norm
? layer.nextn.shared_head_norm
: model.output_norm;
GGML_ASSERT(head_norm_w && "QWEN35 MTP: missing both nextn.shared_head_norm and output_norm");
cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1);
cb(cur, "h_nextn", -1);
res->t_h_nextn = cur;
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
cb(cur, "mtp_shared_head_norm", -1);
ggml_tensor * head_w = layer.nextn.shared_head_head ? layer.nextn.shared_head_head : model.output;
+17 -17
View File
@@ -21,12 +21,13 @@ void llama_model_qwen35moe::load_arch_hparams(llama_model_loader & ml) {
// Mark recurrent layers (linear attention layers). MTP layers are dense
// attention-only and must be flagged non-recurrent.
{
if (!ml.get_key_or_arr(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, hparams.n_layer, false)) {
const uint32_t n_main = hparams.n_layer - hparams.nextn_predict_layers;
uint32_t full_attn_interval = 4;
ml.get_key(LLM_KV_FULL_ATTENTION_INTERVAL, full_attn_interval, false);
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = (i < n_main) && ((i + 1) % full_attn_interval != 0);
hparams.is_recr_impl[i] = (i < n_main) && ((i + 1) % full_attn_interval != 0);
}
}
@@ -75,7 +76,7 @@ void llama_model_qwen35moe::load_arch_tensors(llama_model_loader & ml) {
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", il), { n_embd }, flags);
layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", il), { n_embd }, flags);
if (!hparams.is_recurrent(il)) {
if (!hparams.is_recr(il)) {
// Attention layers
create_tensor_qkv(layer, il, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, flags);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", il), { n_embd_head_k * n_head, n_embd }, flags);
@@ -191,7 +192,7 @@ llama_model_qwen35moe::graph::graph(const llama_model & model, const llm_graph_p
ggml_build_forward_expand(gf, cur);
// Determine layer type and build appropriate attention mechanism
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// Linear attention layer (gated delta net)
cur = build_layer_attn_linear(inp->get_recr(), cur, il);
} else {
@@ -199,7 +200,7 @@ llama_model_qwen35moe::graph::graph(const llama_model & model, const llm_graph_p
cur = build_layer_attn(inp->get_attn(), cur, inp_pos, sections, il);
}
if (il == n_transformer_layers - 1 && inp_out_ids && cparams.embeddings_pre_norm_masked) {
if (il == n_transformer_layers - 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);
}
@@ -231,16 +232,16 @@ llama_model_qwen35moe::graph::graph(const llama_model & model, const llm_graph_p
}
cur = inpL;
cb(cur, "h_pre_norm", -1);
res->t_h_pre_norm = cur;
// post-norm hidden state feeds both the LM head and the MTP seed below
cur = build_norm(cur, model.output_norm, nullptr, LLM_NORM_RMS, -1);
if (!cparams.embeddings_pre_norm_masked && inp_out_ids) {
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);
}
// Final norm
cur = build_norm(cur, model.output_norm, nullptr, LLM_NORM_RMS, -1);
cb(cur, "result_norm", -1);
res->t_embd = cur;
@@ -720,17 +721,16 @@ llama_model_qwen35moe::graph_mtp::graph_mtp(const llama_model & model, const llm
cur = ggml_add(ctx0, cur, ffn_residual);
cb(cur, "mtp_post_ffn", il);
// Pre-norm hidden state: used by the AR draft loop to seed the next MTP step.
cb(cur, "h_pre_norm", -1);
res->t_h_pre_norm = cur;
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
ggml_tensor * head_norm_w = layer.nextn.shared_head_norm
? layer.nextn.shared_head_norm
: model.output_norm;
GGML_ASSERT(head_norm_w && "QWEN35MOE MTP: missing both nextn.shared_head_norm and output_norm");
cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1);
cb(cur, "h_nextn", -1);
res->t_h_nextn= cur;
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
cb(cur, "mtp_shared_head_norm", -1);
ggml_tensor * head_w = layer.nextn.shared_head_head ? layer.nextn.shared_head_head : model.output;
+4 -4
View File
@@ -14,11 +14,11 @@ void llama_model_qwen3next::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_SSM_GROUP_COUNT, hparams.ssm_n_group);
// Mark recurrent layers (linear attention layers)
{
if (!ml.get_key_or_arr(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, hparams.n_layer, false)) {
uint32_t full_attn_interval = 4;
ml.get_key(LLM_KV_FULL_ATTENTION_INTERVAL, full_attn_interval, false);
for (uint32_t i = 0; i < hparams.n_layer; ++i) {
hparams.recurrent_layer_arr[i] = ((i + 1) % full_attn_interval != 0);
hparams.is_recr_impl[i] = ((i + 1) % full_attn_interval != 0);
}
}
@@ -68,7 +68,7 @@ void llama_model_qwen3next::load_arch_tensors(llama_model_loader &) {
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), { n_embd }, 0);
layer.attn_post_norm = create_tensor(tn(LLM_TENSOR_ATTN_POST_NORM, "weight", i), { n_embd }, 0);
if (!hparams.is_recurrent(i)) {
if (!hparams.is_recr(i)) {
// Attention layers
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head * 2, n_embd_k_gqa, n_embd_v_gqa, 0);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), { n_embd_head_k * n_head, n_embd }, 0);
@@ -129,7 +129,7 @@ llama_model_qwen3next::graph::graph(const llama_model & model, const llm_graph_p
ggml_build_forward_expand(gf, cur);
// Determine layer type and build appropriate attention mechanism
if (hparams.is_recurrent(il)) {
if (hparams.is_recr(il)) {
// Linear attention layer (gated delta net)
cur = build_layer_attn_linear(inp->get_recr(), cur, il);
} else {
+304 -12
View File
@@ -22,24 +22,42 @@ void llama_model_step35::load_arch_hparams(llama_model_loader & ml) {
ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa);
ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.swa_layers, hparams.n_layer);
ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer);
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer, false);
ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer, false);
switch (hparams.n_layer) {
// NextN/MTP (Step3p5): extra decoder block appended beyond the main stack.
ml.get_key(LLM_KV_NEXTN_PREDICT_LAYERS, hparams.nextn_predict_layers, false);
GGML_ASSERT(hparams.nextn_predict_layers < hparams.n_layer && "nextn_predict_layers must be < n_layer");
switch (hparams.n_layer - hparams.nextn_predict_layers) {
case 45: type = LLM_TYPE_196B_A11B; break;
default: type = LLM_TYPE_UNKNOWN;
}
}
void llama_model_step35::load_arch_tensors(llama_model_loader &) {
void llama_model_step35::load_arch_tensors(llama_model_loader & ml) {
LLAMA_LOAD_LOCALS;
const uint32_t n_main = n_layer - hparams.nextn_predict_layers;
const bool mtp_only = (hparams.nextn_predict_layers > 0) &&
(ml.get_weight("blk.0.attn_norm.weight") == nullptr);
// Trunk-only: the GGUF declares MTP layers in metadata but the actual MTP
// tensors live in a separate file (e.g. user split target/draft). Mark
// MTP tensors NOT_REQUIRED so the trunk loads cleanly.
const std::string mtp_probe = "blk." + std::to_string(n_main) + ".nextn.eh_proj.weight";
const bool trunk_only = (hparams.nextn_predict_layers > 0) &&
(ml.get_weight(mtp_probe.c_str()) == nullptr);
const int trunk_flags = mtp_only ? TENSOR_NOT_REQUIRED : 0;
const int mtp_flags = trunk_only ? TENSOR_NOT_REQUIRED : 0;
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
// output
output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0);
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, 0);
output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, trunk_flags);
// STEP35 supports per-layer partial RoPE dims; rope factors are stored as a single shared tensor
// ("rope_freqs.weight") and ggml uses only the first (n_rot_l/2) entries per layer.
@@ -51,14 +69,14 @@ void llama_model_step35::load_arch_tensors(llama_model_loader &) {
n_rot_max = n_rot;
}
for (int i = 0; i < n_layer; ++i) {
auto load_block_trunk = [&](int i, int flags) {
auto & layer = layers[i];
const uint32_t n_head_l = hparams.n_head(i);
const uint32_t n_embd_k_gqa = hparams.n_embd_k_gqa(i);
const uint32_t n_embd_v_gqa = hparams.n_embd_v_gqa(i);
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, flags);
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {n_embd_head_k}, TENSOR_NOT_REQUIRED);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_k}, TENSOR_NOT_REQUIRED);
@@ -70,13 +88,13 @@ void llama_model_step35::load_arch_tensors(llama_model_loader &) {
layer.rope_freqs = create_tensor(tn(LLM_TENSOR_ROPE_FREQS, "weight", i), {n_rot_max/2}, TENSOR_NOT_REQUIRED | (i != 0 ? TENSOR_DUPLICATED : 0));
}
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head_l, n_embd_k_gqa, n_embd_v_gqa, 0);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_v * n_head_l, n_embd}, 0);
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head_l, n_embd_k_gqa, n_embd_v_gqa, flags);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_v * n_head_l, n_embd}, flags);
// head-wise attention gate (Step35 self_attn.g_proj)
layer.wqkv_gate = create_tensor(tn(LLM_TENSOR_ATTN_GATE, "weight", i), {n_embd, n_head_l}, TENSOR_NOT_REQUIRED);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, flags);
// dense MLP (leading dense blocks)
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED);
@@ -95,10 +113,86 @@ void llama_model_step35::load_arch_tensors(llama_model_loader &) {
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", i), {n_embd, hparams.n_ff_shexp}, TENSOR_NOT_REQUIRED);
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", i), {n_embd, hparams.n_ff_shexp}, TENSOR_NOT_REQUIRED);
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", i), {hparams.n_ff_shexp, n_embd}, TENSOR_NOT_REQUIRED);
};
auto load_block_mtp = [&](int i, bool is_first_mtp) {
auto & layer = layers[i];
const uint32_t n_head_l = hparams.n_head(i);
const uint32_t n_embd_k_gqa = hparams.n_embd_k_gqa(i);
const uint32_t n_embd_v_gqa = hparams.n_embd_v_gqa(i);
// The MTP block is a full Step3p5 decoder layer (mtp_block) plus the
// NextN-specific wiring (enorm/hnorm/eh_proj + optional shared head).
// `mtp_flags` becomes NOT_REQUIRED when the GGUF is trunk-only.
//
// Only the FIRST MTP block (i == n_main) is required for the
// single-block MTP runtime; trailing MTP blocks are always tolerated
// as missing so pruned GGUFs (block 0 only) load cleanly. Override
// mtp_flags to NOT_REQUIRED for those.
const int eff_mtp_flags = is_first_mtp ? mtp_flags : (mtp_flags | TENSOR_NOT_REQUIRED);
layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, eff_mtp_flags);
layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {n_embd_head_k}, TENSOR_NOT_REQUIRED);
layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {n_embd_head_k}, TENSOR_NOT_REQUIRED);
if (hparams.rope_scaling_type_train == LLAMA_ROPE_SCALING_TYPE_LONGROPE) {
layer.rope_long = create_tensor(tn(LLM_TENSOR_ROPE_FACTORS_LONG, "weight", i), {n_rot_max/2}, TENSOR_NOT_REQUIRED | TENSOR_DUPLICATED);
layer.rope_short = create_tensor(tn(LLM_TENSOR_ROPE_FACTORS_SHORT, "weight", i), {n_rot_max/2}, TENSOR_NOT_REQUIRED | TENSOR_DUPLICATED);
} else {
layer.rope_freqs = create_tensor(tn(LLM_TENSOR_ROPE_FREQS, "weight", i), {n_rot_max/2}, TENSOR_NOT_REQUIRED | TENSOR_DUPLICATED);
}
create_tensor_qkv(layer, i, n_embd, n_embd_head_k * n_head_l, n_embd_k_gqa, n_embd_v_gqa, eff_mtp_flags);
layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_embd_head_v * n_head_l, n_embd}, eff_mtp_flags);
layer.wqkv_gate = create_tensor(tn(LLM_TENSOR_ATTN_GATE, "weight", i), {n_embd, n_head_l}, TENSOR_NOT_REQUIRED);
layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, eff_mtp_flags);
// dense MLP (leading dense blocks) — present if the MTP block isn't MoE
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED);
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, TENSOR_NOT_REQUIRED);
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, TENSOR_NOT_REQUIRED);
// MoE routed experts + selection bias (router_bias)
const int64_t n_ff_exp = hparams.n_ff_exp;
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_exp_probs_b = create_tensor(tn(LLM_TENSOR_FFN_EXP_PROBS_B, "bias", i), {n_expert}, TENSOR_NOT_REQUIRED);
layer.ffn_gate_shexp = create_tensor(tn(LLM_TENSOR_FFN_GATE_SHEXP, "weight", i), {n_embd, hparams.n_ff_shexp}, TENSOR_NOT_REQUIRED);
layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", i), {n_embd, hparams.n_ff_shexp}, TENSOR_NOT_REQUIRED);
layer.ffn_down_shexp = create_tensor(tn(LLM_TENSOR_FFN_DOWN_SHEXP, "weight", i), {hparams.n_ff_shexp, n_embd}, TENSOR_NOT_REQUIRED);
// NextN-specific tensors that define the MTP block.
layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), { 2 * n_embd, n_embd }, eff_mtp_flags);
layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), { n_embd }, eff_mtp_flags);
layer.nextn.hnorm = create_tensor(tn(LLM_TENSOR_NEXTN_HNORM, "weight", i), { n_embd }, eff_mtp_flags);
layer.nextn.embed_tokens = create_tensor(tn(LLM_TENSOR_NEXTN_EMBED_TOKENS, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED);
layer.nextn.shared_head_head = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_HEAD, "weight", i), { n_embd, n_vocab }, TENSOR_NOT_REQUIRED);
layer.nextn.shared_head_norm = create_tensor(tn(LLM_TENSOR_NEXTN_SHARED_HEAD_NORM, "weight", i), { n_embd }, TENSOR_NOT_REQUIRED);
};
for (int i = 0; i < (int) n_main; ++i) {
load_block_trunk(i, trunk_flags);
}
// Only the first MTP block (i == n_main) is required at runtime — the
// single-block-MTP graph in build_arch_graph always uses that one.
// Trailing MTP blocks are loaded if present (so an un-pruned GGUF with
// all MTP layers still works) but tolerated when absent via the pruning
// path. See scripts/prune_step35_extra_mtp.py for the pruner.
for (int i = (int) n_main; i < n_layer; ++i) {
load_block_mtp(i, /*is_first_mtp=*/ i == (int) n_main);
}
}
std::unique_ptr<llm_graph_context> llama_model_step35::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);
}
@@ -111,7 +205,9 @@ llama_model_step35::graph::graph(const llama_model & model, const llm_graph_para
auto * inp_attn = build_attn_inp_kv_iswa();
ggml_tensor * inp_out_ids = build_inp_out_ids();
for (int il = 0; il < n_layer; ++il) {
// MTP/NextN layers are loaded as extra decoder blocks but not executed in the main pass.
const int n_transformer_layers = n_layer - (int) hparams.nextn_predict_layers;
for (int il = 0; il < n_transformer_layers; ++il) {
ggml_tensor * inpSA = inpL;
const uint32_t n_head_l = hparams.n_head(il);
@@ -198,8 +294,8 @@ llama_model_step35::graph::graph(const llama_model & model, const llm_graph_para
cb(cur, "attn_proj", il);
}
if (il == n_layer - 1 && inp_out_ids) {
cur = ggml_get_rows(ctx0, cur, inp_out_ids);
if (il == n_transformer_layers - 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);
}
@@ -257,6 +353,13 @@ llama_model_step35::graph::graph(const llama_model & model, const llm_graph_para
cur = inpL;
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);
}
cur = build_norm(cur, model.output_norm, nullptr, LLM_NORM_RMS, -1);
cb(cur, "result_norm", -1);
res->t_embd = cur;
@@ -267,3 +370,192 @@ llama_model_step35::graph::graph(const llama_model & model, const llm_graph_para
ggml_build_forward_expand(gf, cur);
}
// LLM_GRAPH_TYPE_DECODER_MTP draft head for Step3p5 (MoE)
llama_model_step35::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params)
: llm_graph_context(params) {
GGML_ASSERT(hparams.nextn_predict_layers > 0 && "STEP35 MTP requires nextn_predict_layers > 0");
// Single-block MTP only: always run the first trained MTP block (Qwen
// MTP / vLLM single-MTP-layer style). Multi-block round-robin proved to
// be a much deeper refactor than this PR justifies; the trailing MTP
// blocks are loaded with TENSOR_NOT_REQUIRED so pruned GGUFs (with just
// block 0) also work — see load_arch_tensors below and
// scripts/prune_step35_extra_mtp.py.
const int il = (int) hparams.n_layer - (int) hparams.nextn_predict_layers;
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");
const uint32_t n_head_l = hparams.n_head(il);
const uint32_t n_head_kv_l = hparams.n_head_kv(il);
const float freq_base_l = model.get_rope_freq_base(cparams, il);
const float freq_scale_l = model.get_rope_freq_scale(cparams, il);
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_w = layer.nextn.embed_tokens ? layer.nextn.embed_tokens : model.tok_embd;
ggml_tensor * h_input = inp->embd;
ggml_tensor * tok_embd = ggml_get_rows(ctx0, tok_embd_w, inp->tokens);
cb(tok_embd, "mtp_tok_embd", il);
res->add_input(std::move(inp));
ggml_tensor * inp_pos = build_inp_pos();
auto * inp_attn = build_attn_inp_kv_iswa();
ggml_tensor * h_norm = build_norm(h_input, layer.nextn.hnorm, nullptr, LLM_NORM_RMS, il);
cb(h_norm, "mtp_hnorm", il);
ggml_tensor * e_norm = build_norm(tok_embd, layer.nextn.enorm, nullptr, LLM_NORM_RMS, il);
cb(e_norm, "mtp_enorm", il);
ggml_tensor * concat = ggml_concat(ctx0, e_norm, h_norm, /*dim=*/ 0);
cb(concat, "mtp_concat", il);
ggml_tensor * cur = build_lora_mm(layer.nextn.eh_proj, concat);
cb(cur, "mtp_eh_proj", il);
ggml_tensor * inpSA = cur;
// mtp_block: full Step3p5 decoder layer (attention with optional head-wise gate, then MoE/dense FFN)
cur = build_norm(cur, layer.attn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "mtp_attn_norm", il);
ggml_tensor * Qcur = build_lora_mm(layer.wq, cur, layer.wq_s);
ggml_tensor * Kcur = build_lora_mm(layer.wk, cur, layer.wk_s);
ggml_tensor * Vcur = build_lora_mm(layer.wv, cur, layer.wv_s);
cb(Qcur, "mtp_Qcur", il);
cb(Kcur, "mtp_Kcur", il);
cb(Vcur, "mtp_Vcur", il);
Qcur = ggml_reshape_3d(ctx0, Qcur, n_embd_head_k, n_head_l, n_tokens);
Kcur = ggml_reshape_3d(ctx0, Kcur, n_embd_head_k, n_head_kv_l, n_tokens);
Vcur = ggml_reshape_3d(ctx0, Vcur, n_embd_head_v, n_head_kv_l, n_tokens);
if (layer.attn_q_norm) {
Qcur = build_norm(Qcur, layer.attn_q_norm, nullptr, LLM_NORM_RMS, il);
cb(Qcur, "mtp_Qcur_normed", il);
}
if (layer.attn_k_norm) {
Kcur = build_norm(Kcur, layer.attn_k_norm, nullptr, LLM_NORM_RMS, il);
cb(Kcur, "mtp_Kcur_normed", il);
}
const bool is_swa = hparams.is_swa(il);
ggml_tensor * rope_factors = is_swa ? nullptr : model.get_rope_factors(cparams, il);
const int64_t n_rot_l = hparams.n_rot(il);
Qcur = ggml_rope_ext(
ctx0, Qcur, inp_pos, rope_factors,
n_rot_l, rope_type, n_ctx_orig, freq_base_l, freq_scale_l,
ext_factor, attn_factor, beta_fast, beta_slow);
Kcur = ggml_rope_ext(
ctx0, Kcur, inp_pos, rope_factors,
n_rot_l, rope_type, n_ctx_orig, freq_base_l, freq_scale_l,
ext_factor, attn_factor, beta_fast, beta_slow);
cb(Qcur, "mtp_Qcur_pos", il);
cb(Kcur, "mtp_Kcur_pos", il);
const float kq_scale = 1.0f / sqrtf(float(n_embd_head_k));
ggml_tensor * attn_out = build_attn(inp_attn,
nullptr, nullptr, nullptr,
Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
cb(attn_out, "mtp_attn_out", il);
// head-wise attention gate: sigmoid(g_proj(x))
if (layer.wqkv_gate) {
ggml_tensor * gate = build_lora_mm(layer.wqkv_gate, cur); // [n_head_l, n_tokens]
cb(gate, "mtp_attn_gate", il);
gate = ggml_sigmoid(ctx0, gate);
cb(gate, "mtp_attn_gate_sigmoid", il);
ggml_tensor * attn_3d = ggml_reshape_3d(ctx0, attn_out, n_embd_head_v, n_head_l, n_tokens);
ggml_tensor * gate_3d = ggml_reshape_3d(ctx0, gate, 1, n_head_l, n_tokens);
cb(gate_3d, "mtp_attn_gate_3d", il);
attn_3d = ggml_mul(ctx0, attn_3d, gate_3d);
cb(attn_3d, "mtp_attn_gated_3d", il);
attn_out = ggml_reshape_2d(ctx0, attn_3d, n_embd_head_v * n_head_l, n_tokens);
cb(attn_out, "mtp_attn_gated", il);
}
cur = build_lora_mm(layer.wo, attn_out, layer.wo_s);
cb(cur, "mtp_attn_proj", il);
cur = ggml_add(ctx0, cur, inpSA);
cb(cur, "mtp_attn_residual", il);
ggml_tensor * ffn_inp = cur;
cur = build_norm(cur, layer.ffn_norm, nullptr, LLM_NORM_RMS, il);
cb(cur, "mtp_ffn_norm", il);
// FFN: dense MLP or MoE (mirrors trunk path)
if (layer.ffn_gate_inp == nullptr) {
cur = build_ffn(cur,
layer.ffn_up, layer.ffn_up_b, nullptr,
layer.ffn_gate, layer.ffn_gate_b, nullptr,
layer.ffn_down, layer.ffn_down_b, nullptr,
nullptr,
LLM_FFN_SILU, LLM_FFN_PAR, il);
cb(cur, "mtp_ffn_out", il);
} else {
ggml_tensor * moe_out = 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);
cb(moe_out, "mtp_ffn_moe_out", il);
ggml_tensor * sh_out = 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);
cb(sh_out, "mtp_ffn_shared_out", il);
cur = ggml_add(ctx0, moe_out, sh_out);
cb(cur, "mtp_ffn_out", il);
}
cur = ggml_add(ctx0, cur, ffn_inp);
cb(cur, "mtp_post_ffn", il);
// Pre-norm hidden state: used by the AR draft loop to seed the next MTP step.
cb(cur, "h_nextn", -1);
res->t_h_nextn = cur;
ggml_tensor * head_norm_w = layer.nextn.shared_head_norm
? layer.nextn.shared_head_norm
: model.output_norm;
GGML_ASSERT(head_norm_w && "STEP35 MTP: missing both nextn.shared_head_norm and output_norm");
cur = build_norm(cur, head_norm_w, nullptr, LLM_NORM_RMS, -1);
cb(cur, "mtp_shared_head_norm", -1);
ggml_tensor * head_w = layer.nextn.shared_head_head ? layer.nextn.shared_head_head : model.output;
GGML_ASSERT(head_w && "STEP35 MTP: missing LM head (nextn.shared_head_head or model.output)");
cur = build_lora_mm(head_w, cur);
cb(cur, "result_output", -1);
res->t_logits = cur;
ggml_build_forward_expand(gf, cur);
}
+12 -13
View File
@@ -374,16 +374,10 @@ int llama_completion(int argc, char ** argv) {
__func__, n_match, embd_inp.size());
}
if (session_tokens.size() == n_match) {
// [TAG_CONTEXT_STATE_LOGITS]
// in this case, we are going to reuse the logits from the session
// if we ever decide to remove the logits from the session, we need to handle this somehow
// ref: https://github.com/ggml-org/llama.cpp/pull/18862#issuecomment-3756330941
}
// remove any "future" tokens that we might have inherited from the previous session
if (session_tokens.size() > n_match) {
if (!llama_memory_seq_rm(mem, -1, n_match, -1)) {
llama_pos pos = n_match > 0 ? (llama_pos)(n_match - 1) : 0;
if (!llama_memory_seq_rm(mem, -1, pos, -1)) {
LOG_WRN("%s: unable to reuse common prefix (for example, when the memory is recurrent)\n", __func__);
llama_memory_clear(mem, true);
session_tokens.clear();
@@ -399,7 +393,7 @@ int llama_completion(int argc, char ** argv) {
// Logits are not stored as part of the session state so we need to
// "replay" the last token to get logits for sampling.
if (!session_tokens.empty() && n_match > 0 && n_match == session_tokens.size()) {
if (!common_replay_last_token(ctx, session_tokens.back(), n_match)) {
if (!common_replay_last_token(ctx, session_tokens.back(), n_match - 1)) {
return 1;
}
@@ -696,12 +690,14 @@ int llama_completion(int argc, char ** argv) {
if (!embd.empty()) {
const bool is_last_batch = (n_consumed >= (int) embd_inp.size());
const bool save_now = session_do_save && is_last_batch;
if (!common_prompt_batch_decode(ctx, embd, n_past, params.n_batch, path_session, save_now)) {
session_tokens.insert(session_tokens.end(), embd.begin(), embd.end());
if (!common_prompt_batch_decode(ctx, session_tokens, embd.size(), n_past, params.n_batch, path_session, save_now)) {
return 1;
}
session_tokens.insert(session_tokens.end(), embd.begin(), embd.end());
n_session_consumed = session_tokens.size();
session_do_save = false;
n_session_consumed += embd.size();
if (save_now) {
session_do_save = false;
}
LOG_DBG("n_past = %d\n", n_past);
@@ -992,7 +988,10 @@ int llama_completion(int argc, char ** argv) {
if (!path_session.empty() && params.prompt_cache_all && !params.prompt_cache_ro) {
LOG("\n%s: saving final output to session file '%s'\n", __func__, path_session.c_str());
session_tokens.insert(session_tokens.end(), embd.begin(), embd.end());
llama_state_save_file(ctx, path_session.c_str(), session_tokens.data(), session_tokens.size());
LOG_INF("saved final session to %s, n_tokens = %ld\n", path_session.data(), session_tokens.size());
}
LOG("\n\n");
+5
View File
@@ -83,6 +83,7 @@
#define TN_PATCH_EMBD_1 "v.patch_embd.weight.1"
#define TN_PATCH_BIAS "v.patch_embd.bias"
#define TN_NORM_EMBD "v.norm_embd.%s"
#define TN_PATCH_NORM "v.patch_norm.%d.%s"
#define TN_ATTN_QKV "%s.blk.%d.attn_qkv.%s"
#define TN_ATTN_K "%s.blk.%d.attn_k.%s"
#define TN_ATTN_Q "%s.blk.%d.attn_q.%s"
@@ -317,6 +318,8 @@ enum projector_type {
PROJECTOR_TYPE_GEMMA3NA,
PROJECTOR_TYPE_GEMMA4V,
PROJECTOR_TYPE_GEMMA4A,
PROJECTOR_TYPE_GEMMA4UV,
PROJECTOR_TYPE_GEMMA4UA,
PROJECTOR_TYPE_PHI4,
PROJECTOR_TYPE_IDEFICS3,
PROJECTOR_TYPE_PIXTRAL,
@@ -369,6 +372,8 @@ static std::map<projector_type, std::string> PROJECTOR_TYPE_NAMES = {
{ PROJECTOR_TYPE_GEMMA3NA, "gemma3na"},
{ PROJECTOR_TYPE_GEMMA4V, "gemma4v"},
{ PROJECTOR_TYPE_GEMMA4A, "gemma4a"},
{ PROJECTOR_TYPE_GEMMA4UV, "gemma4uv"},
{ PROJECTOR_TYPE_GEMMA4UA, "gemma4ua"},
{ PROJECTOR_TYPE_PHI4, "phi4"},
{ PROJECTOR_TYPE_IDEFICS3, "idefics3"},
{ PROJECTOR_TYPE_PIXTRAL, "pixtral"},
+8
View File
@@ -339,6 +339,14 @@ struct clip_model {
ggml_tensor * norm_embd_w = nullptr;
ggml_tensor * norm_embd_b = nullptr;
// "indexed" patch embedding norms
ggml_tensor * patch_norm_1_w = nullptr;
ggml_tensor * patch_norm_1_b = nullptr;
ggml_tensor * patch_norm_2_w = nullptr;
ggml_tensor * patch_norm_2_b = nullptr;
ggml_tensor * patch_norm_3_w = nullptr;
ggml_tensor * patch_norm_3_b = nullptr;
ggml_tensor * pre_ln_w = nullptr;
ggml_tensor * pre_ln_b = nullptr;
+51 -4
View File
@@ -57,6 +57,8 @@
#include "models/exaone4_5.cpp"
#include "models/gemma4a.cpp"
#include "models/gemma4v.cpp"
#include "models/gemma4ua.cpp"
#include "models/gemma4uv.cpp"
#include "models/glm4v.cpp"
#include "models/granite-speech.cpp"
#include "models/hunyuanvl.cpp"
@@ -309,11 +311,11 @@ clip_graph::clip_graph(clip_ctx * ctx, const clip_image_f32 & img) :
n_embd(hparams.n_embd),
n_head(hparams.n_head),
n_head_kv(hparams.n_head_kv),
d_head(n_embd / n_head),
d_head(n_head > 0 ? n_embd / n_head : 0),
n_layer(hparams.n_layer),
n_mmproj_embd(clip_n_mmproj_embd(ctx)),
eps(hparams.eps),
kq_scale(1.0f / sqrtf((float)d_head)),
kq_scale(d_head > 0 ? 1.0f / sqrtf((float)d_head) : 0.0f),
flash_attn_type(ctx->flash_attn_type) {
struct ggml_init_params params = {
/*.mem_size =*/ ctx->buf_compute_meta.size(),
@@ -928,6 +930,10 @@ static ggml_cgraph * clip_image_build_graph(clip_ctx * ctx, const clip_image_f32
{
builder = std::make_unique<clip_graph_gemma4v>(ctx, img);
} break;
case PROJECTOR_TYPE_GEMMA4UV:
{
builder = std::make_unique<clip_graph_gemma4uv>(ctx, img);
} break;
case PROJECTOR_TYPE_PIXTRAL:
case PROJECTOR_TYPE_LIGHTONOCR:
{
@@ -1031,6 +1037,10 @@ static ggml_cgraph * clip_image_build_graph(clip_ctx * ctx, const clip_image_f32
{
builder = std::make_unique<clip_graph_gemma4a>(ctx, img);
} break;
case PROJECTOR_TYPE_GEMMA4UA:
{
builder = std::make_unique<clip_graph_gemma4ua>(ctx, img);
} break;
case PROJECTOR_TYPE_GRANITE_SPEECH:
{
builder = std::make_unique<clip_graph_granite_speech>(ctx, img);
@@ -1471,13 +1481,19 @@ struct clip_model_loader {
} break;
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
{
hparams.rope_theta = 100.0f;
hparams.n_merge = 3; // pooling_kernel_size
hparams.image_resize_algo = RESIZE_ALGO_BILINEAR;
get_u32(KEY_PROJ_SCALE_FACTOR, hparams.n_merge, false);
if (model.proj_type == PROJECTOR_TYPE_GEMMA4UV) {
// for "unified" variant, we directly use a bigger patch size, because the "token merging" is done directly on conv layer
hparams.patch_size = hparams.patch_size * hparams.n_merge;
hparams.n_merge = 1;
}
// @ngxson : the model performs quite poor with small images, we need to bump minimum image tokens to 40 to avoid that
hparams.set_limit_image_tokens(252, 280);
hparams.set_limit_image_tokens(40, 280);
hparams.set_warmup_n_tokens(256); // avoid OOM on warmup
} break;
@@ -1676,6 +1692,14 @@ struct clip_model_loader {
// since all gemma4a models use 1e-6, we just hardcode it here to avoid re-conversion
hparams.eps = 1e-6f;
} break;
case PROJECTOR_TYPE_GEMMA4UA:
{
// Encoder-free: raw 16 kHz waveform chunked into 640-sample frames.
hparams.audio_chunk_len = 0;
hparams.audio_sample_rate = 16000;
hparams.eps = 1e-6f;
hparams.n_mel_bins = 640;
} break;
case PROJECTOR_TYPE_GRANITE_SPEECH:
{
hparams.audio_chunk_len = 0;
@@ -2192,6 +2216,16 @@ struct clip_model_loader {
}
}
} break;
case PROJECTOR_TYPE_GEMMA4UV:
{
model.mm_input_proj_w = get_tensor(TN_MM_INP_PROJ);
model.patch_norm_1_w = get_tensor(string_format(TN_PATCH_NORM, 1, "weight"));
model.patch_norm_1_b = get_tensor(string_format(TN_PATCH_NORM, 1, "bias"));
model.patch_norm_2_w = get_tensor(string_format(TN_PATCH_NORM, 2, "weight"));
model.patch_norm_2_b = get_tensor(string_format(TN_PATCH_NORM, 2, "bias"));
model.patch_norm_3_w = get_tensor(string_format(TN_PATCH_NORM, 3, "weight")); // pos_norm
model.patch_norm_3_b = get_tensor(string_format(TN_PATCH_NORM, 3, "bias")); // pos_norm
} break;
case PROJECTOR_TYPE_GEMMA3NV:
{
model.mobilenet_stem_conv_w = get_tensor(TN_MNV5_STEM_CONV, false);
@@ -2605,6 +2639,10 @@ struct clip_model_loader {
}
}
} break;
case PROJECTOR_TYPE_GEMMA4UA:
{
model.mm_input_proj_w = get_tensor(string_format(TN_A_MM_INP_PROJ, "weight"));
} break;
case PROJECTOR_TYPE_LFM2A:
{
for (int i : {0, 2, 3, 5, 6}) {
@@ -3522,6 +3560,7 @@ void setup_init_vision_shim_kcpp(struct clip_ctx * ctx_v) {
image_preproc = std::make_unique<mtmd_image_preprocessor_dyn_size>(ctx_v);
} break;
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
{
// <|image> ... (image embeddings) ... <image|>
img_beg = "<|image>";
@@ -3751,6 +3790,7 @@ int clip_n_output_tokens(const struct clip_ctx * ctx, struct clip_image_f32 * im
} break;
case PROJECTOR_TYPE_GEMMA3:
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
case PROJECTOR_TYPE_IDEFICS3:
case PROJECTOR_TYPE_INTERNVL:
case PROJECTOR_TYPE_NEMOTRON_V2_VL:
@@ -3883,6 +3923,10 @@ int clip_n_output_tokens(const struct clip_ctx * ctx, struct clip_image_f32 * im
}
n_patches = n;
} break;
case PROJECTOR_TYPE_GEMMA4UA:
{
n_patches = img->nx; // no downsampling: one token per raw waveform frame
} break;
case PROJECTOR_TYPE_GRANITE_SPEECH:
{
const int ws = ctx->model.hparams.audio_proj_window_size;
@@ -4450,6 +4494,7 @@ bool clip_image_batch_encode(clip_ctx * ctx, const int n_threads, const clip_ima
set_input_i32("patches", patches);
} break;
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
{
// set (col, row) patch positions for learned positional embedding
const int n_cols = image_size_width / patch_size;
@@ -4531,6 +4576,7 @@ bool clip_image_batch_encode(clip_ctx * ctx, const int n_threads, const clip_ima
case PROJECTOR_TYPE_PHI4:
case PROJECTOR_TYPE_COGVLM:
case PROJECTOR_TYPE_YASA2:
case PROJECTOR_TYPE_GEMMA4UA:
{
// do nothing
} break;
@@ -5033,6 +5079,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
case PROJECTOR_TYPE_GEMMA3NV:
return ctx->model.mm_input_proj_w->ne[0];
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
return ctx->model.mm_input_proj_w->ne[1];
case PROJECTOR_TYPE_IDEFICS3:
return ctx->model.mm_fc_w->ne[1];
@@ -5067,7 +5114,7 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) {
return ctx->model.mm_fc_w->ne[1];
case PROJECTOR_TYPE_LFM2A:
return ctx->model.position_embeddings->ne[0];
case PROJECTOR_TYPE_GEMMA4A:
case PROJECTOR_TYPE_GEMMA4UA:
return ctx->model.hparams.projection_dim;
case PROJECTOR_TYPE_GRANITE_SPEECH:
return ctx->model.qf_proj_linear_w->ne[1];
+19
View File
@@ -0,0 +1,19 @@
#include "models.h"
#include <cmath>
ggml_cgraph * clip_graph_gemma4ua::build() {
ggml_tensor * inp = build_inp_raw(1);
auto cur = ggml_cont(ctx0, ggml_permute(ctx0, inp, 1, 0, 2, 3));
// Gemma4UnifiedMultimodalEmbedder
{
// embedding_pre_projection_norm
cur = ggml_rms_norm(ctx0, cur, hparams.eps);
cur = build_mm(model.mm_input_proj_w, cur);
cb(cur, "projected", -1);
}
ggml_build_forward_expand(gf, cur);
return gf;
}
+71
View File
@@ -0,0 +1,71 @@
#include "models.h"
#include <cmath>
ggml_cgraph * clip_graph_gemma4uv::build() {
ggml_tensor * inp_raw = build_inp_raw();
// Gemma4UnifiedVisionEmbedder uses default pytorch LayerNorm, not RMSNorm
float eps = 1e-5f; // default eps for pytorch LayerNorm
ggml_tensor * inp = nullptr;
{
// note: we cannot use ggml_conv_2d here because we need to apply norm after im2col
auto c = inp_raw->ne[2];
ggml_tensor * kernel = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, patch_size, patch_size, c);
inp = ggml_im2col(ctx0, kernel, inp_raw, patch_size, patch_size, 0, 0, 1, 1, true, inp_raw->type);
// inp shape: [patch_size * patch_size * c, n_patches_w, n_patches_h]
inp = ggml_reshape_2d(ctx0, inp, inp->ne[0], inp->ne[1] * inp->ne[2] * inp->ne[3]);
inp = build_norm(inp, model.patch_norm_1_w, model.patch_norm_1_b, NORM_TYPE_NORMAL, eps, -1);
// inp shape: [patch_size * patch_size * c, n_patches]
inp = ggml_mul_mat(ctx0, model.patch_embeddings_0, inp);
inp = ggml_add(ctx0, inp, model.patch_bias);
// inp shape: [n_embd, n_patches]
inp = build_norm(inp, model.patch_norm_2_w, model.patch_norm_2_b, NORM_TYPE_NORMAL, eps, -1);
}
ggml_tensor * pos_x = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_patches);
ggml_set_name(pos_x, "pos_x");
ggml_set_input(pos_x);
ggml_tensor * pos_y = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_patches);
ggml_set_name(pos_y, "pos_y");
ggml_set_input(pos_y);
{
const int64_t pos_size = model.position_embeddings->ne[1];
const size_t nb1 = ggml_row_size(model.position_embeddings->type, n_embd);
// positional embeddings are stored as lookup tables (one for x, one for y)
ggml_tensor * tbl_x = ggml_view_2d(ctx0, model.position_embeddings,
n_embd, pos_size, nb1, 0);
ggml_tensor * tbl_y = ggml_view_2d(ctx0, model.position_embeddings,
n_embd, pos_size, nb1, pos_size * nb1);
// ggml_get_rows: [n_embd, n_patches]
ggml_tensor * emb_x = ggml_get_rows(ctx0, tbl_x, pos_x);
ggml_tensor * emb_y = ggml_get_rows(ctx0, tbl_y, pos_y);
inp = ggml_add(ctx0, inp, emb_x);
inp = ggml_add(ctx0, inp, emb_y);
cb(inp, "pos_embd", -1);
// pos_norm
inp = build_norm(inp, model.patch_norm_3_w, model.patch_norm_3_b, NORM_TYPE_NORMAL, eps, -1);
}
auto cur = inp;
// Gemma4UnifiedMultimodalEmbedder
{
// embedding_pre_projection_norm
cur = ggml_rms_norm(ctx0, cur, hparams.eps);
cur = build_mm(model.mm_input_proj_w, cur);
cb(cur, "projected", -1);
}
ggml_build_forward_expand(gf, cur);
return gf;
}
+10
View File
@@ -18,6 +18,11 @@ struct clip_graph_gemma4v : clip_graph {
ggml_tensor * build_mm(ggml_tensor * w, ggml_tensor * x) const override;
};
struct clip_graph_gemma4uv : clip_graph {
clip_graph_gemma4uv(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
};
struct clip_graph_pixtral : clip_graph {
clip_graph_pixtral(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
@@ -142,6 +147,11 @@ struct clip_graph_gemma4a : clip_graph {
ggml_tensor * build_mm(ggml_tensor * w, ggml_tensor * x) const override;
};
struct clip_graph_gemma4ua : clip_graph {
clip_graph_gemma4ua(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
};
struct clip_graph_glm4v : clip_graph {
clip_graph_glm4v(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {}
ggml_cgraph * build() override;
+38
View File
@@ -942,6 +942,44 @@ bool mtmd_audio_preprocessor_gemma4a::preprocess(const float * s
return true;
}
//
// mtmd_audio_preprocessor_gemma4ua
//
void mtmd_audio_preprocessor_gemma4ua::initialize() {
// no-op: no FFT or filterbank needed
}
bool mtmd_audio_preprocessor_gemma4ua::preprocess(const float * samples,
size_t n_samples,
std::vector<mtmd_audio_mel> & output) {
if (n_samples == 0) {
return false;
}
const int frame_size = hparams.n_mel_bins; // 640 samples per token @ 16 kHz = 40 ms
const int n_tokens = ((int)n_samples + frame_size - 1) / frame_size;
mtmd_audio_mel mel;
mel.n_len = n_tokens;
mel.n_len_org = n_tokens;
mel.n_mel = frame_size;
mel.data.assign((size_t)frame_size * n_tokens, 0.0f);
// Store mel-major (data[f * n_tokens + t]) so the ggml tensor loads as
// [n_tokens, frame_size] with ne[0]=n_tokens, ne[1]=frame_size.
// The graph builder transposes before RMSNorm so normalization is over frame_size.
for (int t = 0; t < n_tokens; t++) {
for (int f = 0; f < frame_size; f++) {
size_t src = (size_t)t * frame_size + f;
mel.data[(size_t)f * n_tokens + t] = (src < n_samples) ? samples[src] : 0.0f;
}
}
output.push_back(std::move(mel));
return true;
}
//
// mtmd_audio_streaming_istft implementation
//
+6
View File
@@ -96,6 +96,12 @@ struct mtmd_audio_preprocessor_gemma4a : mtmd_audio_preprocessor {
mtmd_audio_cache cache;
};
struct mtmd_audio_preprocessor_gemma4ua : mtmd_audio_preprocessor {
mtmd_audio_preprocessor_gemma4ua(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
void initialize() override;
bool preprocess(const float * samples, size_t n_samples, std::vector<mtmd_audio_mel> & output) override;
};
struct mtmd_audio_preprocessor_qwen3a : mtmd_audio_preprocessor {
mtmd_audio_preprocessor_qwen3a(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {}
void initialize() override;
+8
View File
@@ -482,6 +482,7 @@ struct mtmd_context {
image_preproc = std::make_unique<mtmd_image_preprocessor_dyn_size>(ctx_v);
} break;
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
{
// <|image> ... (image embeddings) ... <image|>
img_beg = "<|image>";
@@ -576,6 +577,12 @@ struct mtmd_context {
aud_end = "<audio|>";
audio_preproc = std::make_unique<mtmd_audio_preprocessor_gemma4a>(ctx_a);
} break;
case PROJECTOR_TYPE_GEMMA4UA:
{
aud_beg = "<|audio>";
aud_end = "<audio|>";
audio_preproc = std::make_unique<mtmd_audio_preprocessor_gemma4ua>(ctx_a);
} break;
default:
throw std::runtime_error(string_format("%s: unexpected audio projector type %d\n", __func__, proj));
}
@@ -1142,6 +1149,7 @@ bool mtmd_decode_use_non_causal(const mtmd_context * ctx, const mtmd_input_chunk
switch (proj_type) {
case PROJECTOR_TYPE_GEMMA3:
case PROJECTOR_TYPE_GEMMA4V:
case PROJECTOR_TYPE_GEMMA4UV:
return true;
default:
return false;
+24 -5
View File
@@ -259,9 +259,9 @@ struct server_slot {
return task->need_embd() || (spec && common_speculative_need_embd(spec));
}
bool need_embd_pre_norm() const {
bool need_embd_nextn() const {
GGML_ASSERT(task);
return spec && common_speculative_need_embd_pre_norm(spec);
return spec && common_speculative_need_embd_nextn(spec);
}
// if the context does not have a memory module then all embeddings have to be computed within a single ubatch
@@ -3013,7 +3013,7 @@ private:
// embedding requires all tokens in the batch to be output;
// MTP also wants logits at every prompt position so the
// streaming hook can mirror t_h_pre_norm into ctx_dft.
// streaming hook can mirror t_h_nextn into ctx_dft.
common_batch_add(batch,
cur_tok,
slot.prompt.tokens.pos_next(),
@@ -3693,6 +3693,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_completions_impl(
auto res = create_response();
auto completion_id = gen_chatcmplid();
auto & rd = res->rd;
auto & params = this->params;
try {
std::vector<server_task> tasks;
@@ -3828,7 +3829,7 @@ std::unique_ptr<server_res_generator> server_routes::handle_completions_impl(
}
res->status = 200;
res->content_type = "text/event-stream";
res->next = [res_this = res.get(), res_type, &req](std::string & output) -> bool {
res->next = [res_this = res.get(), res_type, &req, &params](std::string & output) -> bool {
static auto format_error = [](task_response_type res_type, const json & res_json) {
if (res_type == TASK_RESPONSE_TYPE_ANTHROPIC) {
return format_anthropic_sse({
@@ -3873,7 +3874,25 @@ std::unique_ptr<server_res_generator> server_routes::handle_completions_impl(
}
// receive subsequent results
auto result = rd.next(req.should_stop);
bool timeout = false;
int64_t start_time = ggml_time_ms();
auto result = rd.next([&timeout, &req, &start_time, &params]() {
if (req.should_stop()) {
return true; // should_stop condition met
} else if (params.sse_ping_interval > 0 && ggml_time_ms() - start_time > (int64_t)params.sse_ping_interval * 1000) {
timeout = true;
return true; // timeout
}
return false;
});
if (timeout) {
// some clients may time out (e.g. undici) will time out if no data is received for a while, so we need to send a ping to keep the connection alive
SRV_DBG("%s", "sending SSE ping\n");
output = ":\n\n";
return true;
}
if (result == nullptr) {
SRV_DBG("%s", "stopping streaming due to should_stop condition\n");
GGML_ASSERT(req.should_stop());
-4
View File
@@ -381,10 +381,6 @@ server_task_result_ptr server_response_reader::next(const std::function<bool()>
if (result == nullptr) {
// timeout, check stop condition
if (should_stop()) {
const int64_t time_elapsed_ms = ggml_time_ms() - time_start_ms;
if (time_elapsed_ms > 30000) {
SRV_WRN("%s", "request cancelled after 30s, potentially a client-side timeout; please check your client's code\n");
}
return nullptr;
}
} else {
-2
View File
@@ -169,8 +169,6 @@ struct server_response_reader {
bool cancelled = false;
int polling_interval_seconds;
const int64_t time_start_ms = ggml_time_ms();
// tracking generation state and partial tool calls
// only used by streaming completions
std::vector<task_result_state> states;
-1
View File
@@ -1,6 +1,5 @@
aiohttp~=3.9.3
pytest~=8.3.3
huggingface_hub>=1.5.0,<2.0
numpy~=1.26.4
openai~=2.14.0
prometheus-client~=0.20.0
+1015 -9
View File
File diff suppressed because it is too large Load Diff
+1
View File
@@ -80,6 +80,7 @@
"dependencies": {
"@modelcontextprotocol/sdk": "^1.25.1",
"highlight.js": "^11.11.1",
"mermaid": "^11.15.0",
"mode-watcher": "^1.1.0",
"pdfjs-dist": "^5.4.54",
"rehype-highlight": "^7.0.2",
+1
View File
@@ -141,6 +141,7 @@
@apply bg-background text-foreground;
scrollbar-width: thin;
scrollbar-gutter: stable;
overflow: hidden; /* Added due to Mermaid rendering somehow causing the double scrollbar */
}
/* Global scrollbar styling - visible only on hover */
@@ -16,10 +16,23 @@
import { rehypeRestoreTableHtml } from './plugins/rehype/table-html-restorer';
import { rehypeEnhanceLinks } from './plugins/rehype/enhance-links';
import { rehypeEnhanceCodeBlocks } from './plugins/rehype/enhance-code-blocks';
import { rehypeEnhanceMermaidBlocks } from './plugins/rehype/enhance-mermaid-blocks';
import { rehypeMermaidPre } from './plugins/rehype/mermaid-pre';
import { rehypeResolveAttachmentImages } from './plugins/rehype/resolve-attachment-images';
import { rehypeRtlSupport } from './plugins/rehype/rehype-rtl-support';
import { remarkLiteralHtml } from './plugins/remark/literal-html';
import { copyCodeToClipboard, preprocessLaTeX, getImageErrorFallbackHtml } from '$lib/utils';
import {
getHastNodeId,
getMdastNodeHash,
isAppendMode,
getCodeInfoFromTarget
} from './markdown-utils';
import {
preprocessLaTeX,
getImageErrorFallbackHtml,
copyCodeToClipboard,
copyToClipboard
} from '$lib/utils';
import {
IMAGE_NOT_ERROR_BOUND_SELECTOR,
DATA_ERROR_BOUND_ATTR,
@@ -34,7 +47,12 @@
import githubDarkCss from 'highlight.js/styles/github-dark.css?inline';
import githubLightCss from 'highlight.js/styles/github.css?inline';
import { mode } from 'mode-watcher';
import { CodeBlockActions, DialogCodePreview } from '$lib/components/app';
import {
CodeBlockActions,
DialogCodePreview,
DialogMermaidPreview,
ActionIconCopyToClipboard
} from '$lib/components/app';
import { createAutoScrollController } from '$lib/hooks/use-auto-scroll.svelte';
import type { DatabaseMessageExtra } from '$lib/types/database';
import { config } from '$lib/stores/settings.svelte';
@@ -62,6 +80,9 @@
let previewDialogOpen = $state(false);
let previewCode = $state('');
let previewLanguage = $state('text');
let mermaidPreviewOpen = $state(false);
let mermaidPreviewSvgHtml = $state('');
let streamingCodeScrollContainer = $state<HTMLDivElement>();
// Auto-scroll controller for streaming code block content
@@ -102,7 +123,9 @@
}) // Add syntax highlighting
.use(rehypeRestoreTableHtml) // Restore limited HTML (e.g., <br>, <ul>) inside Markdown tables
.use(rehypeEnhanceLinks) // Add target="_blank" to links
.use(rehypeMermaidPre) // Convert mermaid blocks to <pre class="mermaid">
.use(rehypeEnhanceCodeBlocks) // Wrap code blocks with header and actions
.use(rehypeEnhanceMermaidBlocks) // Wrap mermaid blocks with header and actions
.use(rehypeResolveAttachmentImages, { attachments })
.use(rehypeRtlSupport) // Add bidirectional text support
.use(rehypeStringify, { allowDangerousHtml: true }); // Convert to HTML string
@@ -157,73 +180,7 @@
}
/**
* Extracts code information from a button click target within a code block.
* @param target - The clicked button element
* @returns Object with rawCode and language, or null if extraction fails
*/
function getCodeInfoFromTarget(target: HTMLElement) {
const wrapper = target.closest('.code-block-wrapper');
if (!wrapper) {
console.error('No wrapper found');
return null;
}
const codeElement = wrapper.querySelector<HTMLElement>('code[data-code-id]');
if (!codeElement) {
console.error('No code element found in wrapper');
return null;
}
const rawCode = codeElement.textContent ?? '';
const languageLabel = wrapper.querySelector<HTMLElement>('.code-language');
const language = languageLabel?.textContent?.trim() || 'text';
return { rawCode, language };
}
/**
* Generates a unique identifier for a HAST node based on its position.
* Used for stable block identification during incremental rendering.
* @param node - The HAST root content node
* @param indexFallback - Fallback index if position is unavailable
* @returns Unique string identifier for the node
*/
function getHastNodeId(node: HastRootContent, indexFallback: number): string {
const position = node.position;
if (position?.start?.offset != null && position?.end?.offset != null) {
return `hast-${position.start.offset}-${position.end.offset}`;
}
return `${node.type}-${indexFallback}`;
}
/**
* Generates a hash for MDAST node based on its position.
* Used for cache lookup during incremental rendering.
*/
function getMdastNodeHash(node: unknown, index: number): string {
const n = node as {
type?: string;
position?: { start?: { offset?: number }; end?: { offset?: number } };
};
if (n.position?.start?.offset != null && n.position?.end?.offset != null) {
return `${n.type}-${n.position.start.offset}-${n.position.end.offset}`;
}
return `${n.type}-idx${index}`;
}
/**
* Check if we're in append-only mode (streaming).
*/
function isAppendMode(newContent: string): boolean {
return previousContent.length > 0 && newContent.startsWith(previousContent);
}
* Transforms a single MDAST node to HTML string with caching.
/**
* Transforms a single MDAST node to HTML string with caching.
@@ -359,7 +316,7 @@
const nextBlocks: MarkdownBlock[] = [];
// Check if we're in append mode for cache reuse
const appendMode = isAppendMode(prefixMarkdown);
const appendMode = isAppendMode(prefixMarkdown, previousContent);
const previousBlockCount = appendMode ? renderedBlocks.length : 0;
// All prefix blocks are now stable since code block is separate
@@ -411,7 +368,7 @@
const nextBlocks: MarkdownBlock[] = [];
// Check if we're in append mode for cache reuse
const appendMode = isAppendMode(markdown);
const appendMode = isAppendMode(markdown, previousContent);
const previousBlockCount = appendMode ? renderedBlocks.length : 0;
for (let index = 0; index < stableCount; index++) {
@@ -496,6 +453,118 @@
}
}
/**
* Opens the mermaid diagram in a full-screen preview dialog with zoom/pan support.
* Also handles copy and preview button clicks for mermaid blocks.
* Uses event delegation: a single handler on the container.
*/
async function handleMermaidClick(event: MouseEvent) {
const target = event.target as HTMLElement;
// Check if clicking on copy or preview button in mermaid block
const copyBtn = target.closest('.mermaid-block-wrapper .copy-code-btn');
const previewBtn = target.closest('.mermaid-block-wrapper .preview-code-btn');
if (copyBtn || previewBtn) {
const wrapper = target.closest('.mermaid-block-wrapper');
if (!wrapper) return;
const preElement = wrapper.querySelector<HTMLElement>('pre.mermaid[data-mermaid-syntax]');
if (!preElement) return;
const mermaidSyntax = preElement.dataset.mermaidSyntax ?? '';
if (copyBtn) {
event.preventDefault();
event.stopPropagation();
try {
await copyToClipboard(mermaidSyntax);
} catch (error) {
console.error('Failed to copy mermaid syntax:', error);
}
return;
}
if (previewBtn) {
event.preventDefault();
event.stopPropagation();
const svg = preElement.querySelector('svg');
if (!svg) return;
mermaidPreviewSvgHtml = svg.outerHTML;
mermaidPreviewOpen = true;
return;
}
}
// Otherwise, open preview when clicking on the mermaid diagram itself
const mermaidEl = target.closest('.mermaid');
if (!mermaidEl) return;
const svg = mermaidEl.querySelector('svg');
if (!svg) return;
mermaidPreviewSvgHtml = svg.outerHTML;
mermaidPreviewOpen = true;
}
/**
* Handles mermaid preview dialog open state changes.
* Cleans up SVG content when dialog is closed.
*/
function handleMermaidPreviewOpenChange(open: boolean) {
mermaidPreviewOpen = open;
if (!open) {
mermaidPreviewSvgHtml = '';
}
}
/**
* Renders mermaid diagrams that haven't been rendered yet.
* Called after each markdown content update.
* Marks nodes immediately to prevent duplicate renders during streaming.
* Reads mode.current before await to ensure reactive tracking.
*/
async function renderMermaidDiagrams() {
if (!containerRef) return;
const nodes = containerRef.querySelectorAll('pre.mermaid:not([data-mermaid-rendered])');
if (nodes.length === 0) return;
// Mark nodes immediately to prevent duplicate renders if called again during streaming.
// This avoids needing a guard that would block node discovery.
nodes.forEach((node) => node.setAttribute('data-mermaid-rendered', 'true'));
// Read mode before await so Svelte tracks it reactively.
const isDark = mode.current === ColorMode.DARK;
// lazy load the mermaid dependecy only when needed to reduce bundle size.
const { default: mermaid } = await import('mermaid');
mermaid.initialize({
startOnLoad: false,
theme: isDark ? 'dark' : 'default',
securityLevel: 'strict',
flowchart: {
useMaxWidth: false,
htmlLabels: true
},
sequence: {
useMaxWidth: false
},
gantt: {
useMaxWidth: false
}
});
try {
await mermaid.run({
nodes: Array.from(nodes) as unknown as NodeListOf<HTMLElement>
});
} catch (error) {
console.error('Failed to render mermaid diagram:', error);
}
}
/**
* Handles image load errors by replacing the image with a fallback UI.
* Shows a placeholder with a link to open the image in a new tab.
@@ -577,6 +646,7 @@
if ((hasRenderedBlocks || hasUnstableBlock) && containerRef) {
setupCodeBlockActions();
setupImageErrorHandlers();
renderMermaidDiagrams();
}
});
@@ -596,15 +666,17 @@
});
</script>
<!-- svelte-ignore a11y_click_events_have_key_events -->
<!-- svelte-ignore a11y_no_static_element_interactions -->
<div
bind:this={containerRef}
class="{className}{config()[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS]
onclick={handleMermaidClick}
class="markdown-content {className}{config()[SETTINGS_KEYS.FULL_HEIGHT_CODE_BLOCKS]
? ' full-height-code-blocks'
: ''}"
>
{#each renderedBlocks as block (block.id)}
<div class="markdown-block" data-block-id={block.id} use:fadeInView={{ skipIfVisible: true }}>
<!-- eslint-disable-next-line no-at-html-tags -->
{@html block.html}
</div>
{/each}
@@ -617,34 +689,53 @@
{/if}
{#if incompleteCodeBlock}
<div class="code-block-wrapper streaming-code-block relative">
<div class="code-block-header">
<span class="code-language">{incompleteCodeBlock.language || 'text'}</span>
<CodeBlockActions
code={incompleteCodeBlock.code}
language={incompleteCodeBlock.language || 'text'}
disabled
onPreview={(code, lang) => {
previewCode = code;
previewLanguage = lang;
previewDialogOpen = true;
}}
/>
{#if incompleteCodeBlock.language === 'mermaid'}
<div class="mermaid-block-wrapper streaming-mermaid-block">
<div class="code-block-header">
<span class="code-language">mermaid</span>
<div class="code-block-actions">
<ActionIconCopyToClipboard
text={incompleteCodeBlock.code}
canCopy={false}
ariaLabel="Diagram incomplete"
/>
</div>
</div>
<div class="mermaid-loading-placeholder">
<span class="mermaid-loading-text">Generating diagram...</span>
</div>
</div>
<div
bind:this={streamingCodeScrollContainer}
class="streaming-code-scroll-container"
onscroll={() => streamingAutoScroll.handleScroll()}
>
<pre class="streaming-code-pre"><code
class="hljs language-{incompleteCodeBlock.language || 'text'}"
>{@html highlightCode(
incompleteCodeBlock.code,
incompleteCodeBlock.language || 'text'
)}</code
></pre>
{:else}
<div class="code-block-wrapper streaming-code-block relative">
<div class="code-block-header">
<span class="code-language">{incompleteCodeBlock.language || 'text'}</span>
<CodeBlockActions
code={incompleteCodeBlock.code}
language={incompleteCodeBlock.language || 'text'}
disabled
onPreview={(code, lang) => {
previewCode = code;
previewLanguage = lang;
previewDialogOpen = true;
}}
/>
</div>
<div
bind:this={streamingCodeScrollContainer}
class="streaming-code-scroll-container"
onscroll={() => streamingAutoScroll.handleScroll()}
>
<pre class="streaming-code-pre"><code
class="hljs language-{incompleteCodeBlock.language || 'text'}"
>{@html highlightCode(
incompleteCodeBlock.code,
incompleteCodeBlock.language || 'text'
)}</code
></pre>
</div>
</div>
</div>
{/if}
{/if}
</div>
@@ -655,566 +746,12 @@
onOpenChange={handlePreviewDialogOpenChange}
/>
<DialogMermaidPreview
open={mermaidPreviewOpen}
svgHtml={mermaidPreviewSvgHtml}
onOpenChange={handleMermaidPreviewOpenChange}
/>
<style>
.markdown-block--unstable {
display: contents;
}
/* Streaming code block uses .code-block-wrapper styles */
.streaming-code-block .streaming-code-pre {
background: transparent;
padding: 0.5rem;
margin: 0;
overflow-x: visible;
border-radius: 0;
border: none;
font-size: 0.875rem;
}
/* Base typography styles */
div :global(p) {
margin-block: 1rem;
line-height: 1.75;
}
div :global(:is(h1, h2, h3, h4, h5, h6):first-child) {
margin-top: 0;
}
/* Headers with consistent spacing */
div :global(h1) {
font-size: 1.875rem;
font-weight: 700;
line-height: 1.2;
margin: 1.5rem 0 0.75rem 0;
}
div :global(h2) {
font-size: 1.5rem;
font-weight: 600;
line-height: 1.3;
margin: 1.25rem 0 0.5rem 0;
}
div :global(h3) {
font-size: 1.25rem;
font-weight: 600;
margin: 1.5rem 0 0.5rem 0;
line-height: 1.4;
}
div :global(h4) {
font-size: 1.125rem;
font-weight: 600;
margin: 0.75rem 0 0.25rem 0;
}
div :global(h5) {
font-size: 1rem;
font-weight: 600;
margin: 0.5rem 0 0.25rem 0;
}
div :global(h6) {
font-size: 0.875rem;
font-weight: 600;
margin: 0.5rem 0 0.25rem 0;
}
/* Text formatting */
div :global(strong) {
font-weight: 600;
}
div :global(em) {
font-style: italic;
}
div :global(del) {
text-decoration: line-through;
opacity: 0.7;
}
/* Inline code */
div :global(code:not(pre code)) {
background: var(--muted);
color: var(--muted-foreground);
padding: 0.125rem 0.375rem;
border-radius: 0.375rem;
font-size: 0.875rem;
}
div :global(pre) {
display: inline;
margin: 0 !important;
overflow: hidden !important;
background: var(--muted);
overflow-x: auto;
border-radius: 1rem;
border: none;
line-height: 1 !important;
}
div :global(pre code) {
padding: 0 !important;
display: inline !important;
}
div :global(code) {
background: transparent;
color: var(--code-foreground);
}
/* Links */
div :global(a) {
color: var(--primary);
text-decoration: underline;
text-underline-offset: 2px;
transition: color 0.2s ease;
overflow-wrap: anywhere;
word-break: break-all;
}
div :global(a:hover) {
color: var(--primary);
}
/* Lists */
div :global(ul) {
list-style-type: disc;
margin-inline-start: 1.5rem;
margin-bottom: 1rem;
}
div :global(ol) {
list-style-type: decimal;
margin-inline-start: 1.5rem;
margin-bottom: 1rem;
}
div :global(li) {
margin-bottom: 0.25rem;
padding-inline-start: 0.5rem;
}
div :global(li::marker) {
color: var(--muted-foreground);
}
/* Nested lists */
div :global(ul ul) {
list-style-type: circle;
margin-top: 0.25rem;
margin-bottom: 0.25rem;
}
div :global(ol ol) {
list-style-type: lower-alpha;
margin-top: 0.25rem;
margin-bottom: 0.25rem;
}
/* Task lists */
div :global(.task-list-item) {
list-style: none;
margin-inline-start: 0;
padding-inline-start: 0;
}
div :global(.task-list-item-checkbox) {
margin-right: 0.5rem;
margin-top: 0.125rem;
}
/* Blockquotes */
div :global(blockquote) {
border-left: 4px solid var(--border);
padding: 0.5rem 1rem;
margin: 1.5rem 0;
font-style: italic;
color: var(--muted-foreground);
background: var(--muted);
border-radius: 0 0.375rem 0.375rem 0;
}
/* Tables */
div :global(table) {
width: 100%;
margin: 1.5rem 0;
border-collapse: collapse;
border: 1px solid var(--border);
border-radius: 0.375rem;
overflow: hidden;
}
div :global(th) {
background: hsl(var(--muted) / 0.3);
border: 1px solid var(--border);
padding: 0.5rem 0.75rem;
text-align: left;
font-weight: 600;
}
div :global(td) {
border: 1px solid var(--border);
padding: 0.5rem 0.75rem;
}
div :global(tr:nth-child(even)) {
background: hsl(var(--muted) / 0.1);
}
/* User message markdown should keep table borders visible on light primary backgrounds */
div.markdown-user-content :global(table),
div.markdown-user-content :global(th),
div.markdown-user-content :global(td),
div.markdown-user-content :global(.table-wrapper) {
border-color: currentColor;
}
/* Horizontal rules */
div :global(hr) {
border: none;
border-top: 1px solid var(--border);
margin: 1.5rem 0;
}
/* Images */
div :global(img) {
border-radius: 0.5rem;
box-shadow:
0 1px 3px 0 rgb(0 0 0 / 0.1),
0 1px 2px -1px rgb(0 0 0 / 0.1);
margin: 1.5rem 0;
max-width: 100%;
height: auto;
}
/* Code blocks */
div :global(.code-block-wrapper) {
margin: 1.5rem 0;
border-radius: 0.75rem;
overflow: hidden;
border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);
background: var(--code-background);
box-shadow: 0 1px 2px 0 rgb(0 0 0 / 0.05);
min-height: var(--min-message-height);
max-height: var(--max-message-height);
}
:global(.dark) div :global(.code-block-wrapper) {
border-color: color-mix(in oklch, var(--border) 20%, transparent);
}
/* Scroll container for code blocks (both streaming and completed) */
div :global(.code-block-scroll-container),
.streaming-code-scroll-container {
min-height: var(--min-message-height);
max-height: var(--max-message-height);
overflow-y: auto;
overflow-x: auto;
padding: 3rem 1rem 1rem;
line-height: 1.3;
}
.full-height-code-blocks :global(.code-block-wrapper) {
max-height: none;
}
.full-height-code-blocks :global(.code-block-scroll-container),
.full-height-code-blocks .streaming-code-scroll-container {
max-height: none;
overflow-y: visible;
}
div :global(.code-block-header) {
display: flex;
justify-content: space-between;
align-items: center;
padding: 0.5rem 1rem 0;
font-size: 0.875rem;
position: absolute;
top: 0;
left: 0;
right: 0;
}
div :global(.code-language) {
color: var(--color-foreground);
font-weight: 500;
font-family:
ui-monospace, SFMono-Regular, 'SF Mono', Monaco, 'Cascadia Code', 'Roboto Mono', Consolas,
'Liberation Mono', Menlo, monospace;
text-transform: uppercase;
font-size: 0.75rem;
letter-spacing: 0.05em;
}
div :global(.code-block-actions) {
display: flex;
align-items: center;
gap: 0.5rem;
}
div :global(.copy-code-btn),
div :global(.preview-code-btn) {
display: flex;
align-items: center;
justify-content: center;
padding: 0;
background: transparent;
color: var(--code-foreground);
cursor: pointer;
transition: all 0.2s ease;
}
div :global(.copy-code-btn:hover),
div :global(.preview-code-btn:hover) {
transform: scale(1.05);
}
div :global(.copy-code-btn:active),
div :global(.preview-code-btn:active) {
transform: scale(0.95);
}
div :global(.code-block-wrapper pre) {
background: transparent;
margin: 0;
border-radius: 0;
border: none;
font-size: 0.875rem;
}
/* Mentions and hashtags */
div :global(.mention) {
color: hsl(var(--primary));
font-weight: 500;
text-decoration: none;
}
div :global(.mention:hover) {
text-decoration: underline;
}
div :global(.hashtag) {
color: hsl(var(--primary));
font-weight: 500;
text-decoration: none;
}
div :global(.hashtag:hover) {
text-decoration: underline;
}
/* Advanced table enhancements */
div :global(table) {
transition: all 0.2s ease;
}
div :global(table:hover) {
box-shadow:
0 4px 6px -1px rgb(0 0 0 / 0.1),
0 2px 4px -2px rgb(0 0 0 / 0.1);
}
div :global(th:hover),
div :global(td:hover) {
background: var(--muted);
}
/* Disable hover effects when rendering user messages */
.markdown-user-content :global(a),
.markdown-user-content :global(a:hover) {
color: inherit;
}
.markdown-user-content :global(table:hover) {
box-shadow: none;
}
.markdown-user-content :global(th:hover),
.markdown-user-content :global(td:hover) {
background: inherit;
}
/* Enhanced blockquotes */
div :global(blockquote) {
transition: all 0.2s ease;
position: relative;
}
div :global(blockquote:hover) {
border-left-width: 6px;
background: var(--muted);
transform: translateX(2px);
}
div :global(blockquote::before) {
content: '"';
position: absolute;
top: -0.5rem;
left: 0.5rem;
font-size: 3rem;
color: var(--muted-foreground);
font-family: serif;
line-height: 1;
}
/* Enhanced images */
div :global(img) {
transition: all 0.3s ease;
cursor: pointer;
}
div :global(img:hover) {
transform: scale(1.02);
box-shadow:
0 10px 15px -3px rgb(0 0 0 / 0.1),
0 4px 6px -4px rgb(0 0 0 / 0.1);
}
/* Image zoom overlay */
div :global(.image-zoom-overlay) {
position: fixed;
top: 0;
left: 0;
right: 0;
bottom: 0;
background: rgba(0, 0, 0, 0.8);
display: flex;
align-items: center;
justify-content: center;
z-index: 1000;
cursor: pointer;
}
div :global(.image-zoom-overlay img) {
max-width: 90vw;
max-height: 90vh;
border-radius: 0.5rem;
box-shadow: 0 25px 50px -12px rgb(0 0 0 / 0.25);
}
/* Enhanced horizontal rules */
div :global(hr) {
border: none;
height: 2px;
background: linear-gradient(to right, transparent, var(--border), transparent);
margin: 2rem 0;
position: relative;
}
div :global(hr::after) {
content: '';
position: absolute;
top: 50%;
left: 50%;
transform: translate(-50%, -50%);
width: 1rem;
height: 1rem;
background: var(--border);
border-radius: 50%;
}
/* Scrollable tables */
div :global(.table-wrapper) {
overflow-x: auto;
margin: 1.5rem 0;
border-radius: 0.5rem;
border: 1px solid var(--border);
}
div :global(.table-wrapper table) {
margin: 0;
border: none;
}
/* Responsive adjustments */
@media (max-width: 640px) {
div :global(h1) {
font-size: 1.5rem;
}
div :global(h2) {
font-size: 1.25rem;
}
div :global(h3) {
font-size: 1.125rem;
}
div :global(table) {
font-size: 0.875rem;
}
div :global(th),
div :global(td) {
padding: 0.375rem 0.5rem;
}
div :global(.table-wrapper) {
margin: 0.5rem -1rem;
border-radius: 0;
border-left: none;
border-right: none;
}
}
/* Dark mode adjustments */
@media (prefers-color-scheme: dark) {
div :global(blockquote:hover) {
background: var(--muted);
}
}
/* Image load error fallback */
div :global(.image-load-error) {
display: flex;
align-items: center;
justify-content: center;
margin: 1.5rem 0;
padding: 1.5rem;
border-radius: 0.5rem;
background: var(--muted);
border: 1px dashed var(--border);
}
div :global(.image-error-content) {
display: flex;
flex-direction: column;
align-items: center;
gap: 0.75rem;
color: var(--muted-foreground);
text-align: center;
}
div :global(.image-error-content svg) {
opacity: 0.5;
}
div :global(.image-error-text) {
font-size: 0.875rem;
}
div :global(.image-error-link) {
display: inline-flex;
align-items: center;
gap: 0.375rem;
padding: 0.5rem 1rem;
font-size: 0.875rem;
font-weight: 500;
color: var(--primary);
background: var(--background);
border: 1px solid var(--border);
border-radius: 0.375rem;
text-decoration: none;
transition: all 0.2s ease;
}
div :global(.image-error-link:hover) {
background: var(--muted);
border-color: var(--primary);
}
@import './markdown-content.css';
</style>
@@ -0,0 +1,685 @@
.markdown-block--unstable {
display: contents;
}
/* Streaming code block uses .code-block-wrapper styles */
.streaming-code-block .streaming-code-pre {
background: transparent;
padding: 0.5rem;
margin: 0;
overflow-x: visible;
border-radius: 0;
border: none;
font-size: 0.875rem;
}
/* Base typography styles */
.markdown-content :global(p) {
margin-block: 1rem;
line-height: 1.75;
}
.markdown-content :global(:is(h1, h2, h3, h4, h5, h6):first-child) {
margin-top: 0;
}
/* Headers with consistent spacing */
.markdown-content :global(h1) {
font-size: 1.875rem;
font-weight: 700;
line-height: 1.2;
margin: 1.5rem 0 0.75rem 0;
}
.markdown-content :global(h2) {
font-size: 1.5rem;
font-weight: 600;
line-height: 1.3;
margin: 1.25rem 0 0.5rem 0;
}
.markdown-content :global(h3) {
font-size: 1.25rem;
font-weight: 600;
margin: 1.5rem 0 0.5rem 0;
line-height: 1.4;
}
.markdown-content :global(h4) {
font-size: 1.125rem;
font-weight: 600;
margin: 0.75rem 0 0.25rem 0;
}
.markdown-content :global(h5) {
font-size: 1rem;
font-weight: 600;
margin: 0.5rem 0 0.25rem 0;
}
.markdown-content :global(h6) {
font-size: 0.875rem;
font-weight: 600;
margin: 0.5rem 0 0.25rem 0;
}
/* Text formatting */
.markdown-content :global(strong) {
font-weight: 600;
}
.markdown-content :global(em) {
font-style: italic;
}
.markdown-content :global(del) {
text-decoration: line-through;
opacity: 0.7;
}
/* Inline code */
.markdown-content :global(code:not(pre code)) {
background: var(--muted);
color: var(--muted-foreground);
padding: 0.125rem 0.375rem;
border-radius: 0.375rem;
font-size: 0.875rem;
}
.markdown-content :global(pre) {
display: inline;
margin: 0 !important;
overflow: hidden !important;
background: var(--muted);
overflow-x: auto;
border-radius: 1rem;
border: none;
line-height: 1 !important;
}
.markdown-content :global(pre code) {
padding: 0 !important;
display: inline !important;
}
.markdown-content :global(code) {
background: transparent;
color: var(--code-foreground);
}
/* Links */
.markdown-content :global(a) {
color: var(--primary);
text-decoration: underline;
text-underline-offset: 2px;
transition: color 0.2s ease;
overflow-wrap: anywhere;
word-break: break-all;
}
.markdown-content :global(a:hover) {
color: var(--primary);
}
/* Lists */
.markdown-content :global(ul) {
list-style-type: disc;
margin-inline-start: 1.5rem;
margin-bottom: 1rem;
}
.markdown-content :global(ol) {
list-style-type: decimal;
margin-inline-start: 1.5rem;
margin-bottom: 1rem;
}
.markdown-content :global(li) {
margin-bottom: 0.25rem;
padding-inline-start: 0.5rem;
}
.markdown-content :global(li::marker) {
color: var(--muted-foreground);
}
/* Nested lists */
.markdown-content :global(ul ul) {
list-style-type: circle;
margin-top: 0.25rem;
margin-bottom: 0.25rem;
}
.markdown-content :global(ol ol) {
list-style-type: lower-alpha;
margin-top: 0.25rem;
margin-bottom: 0.25rem;
}
/* Task lists */
.markdown-content :global(.task-list-item) {
list-style: none;
margin-inline-start: 0;
padding-inline-start: 0;
}
.markdown-content :global(.task-list-item-checkbox) {
margin-right: 0.5rem;
margin-top: 0.125rem;
}
/* Blockquotes */
.markdown-content :global(blockquote) {
border-left: 4px solid var(--border);
padding: 0.5rem 1rem;
margin: 1.5rem 0;
font-style: italic;
color: var(--muted-foreground);
background: var(--muted);
border-radius: 0 0.375rem 0.375rem 0;
}
/* Tables */
.markdown-content :global(table) {
width: 100%;
margin: 1.5rem 0;
border-collapse: collapse;
border: 1px solid var(--border);
border-radius: 0.375rem;
overflow: hidden;
}
.markdown-content :global(th) {
background: hsl(var(--muted) / 0.3);
border: 1px solid var(--border);
padding: 0.5rem 0.75rem;
text-align: left;
font-weight: 600;
}
.markdown-content :global(td) {
border: 1px solid var(--border);
padding: 0.5rem 0.75rem;
}
.markdown-content :global(tr:nth-child(even)) {
background: hsl(var(--muted) / 0.1);
}
/* User message markdown should keep table borders visible on light primary backgrounds */
div.markdown-user-content :global(table),
div.markdown-user-content :global(th),
div.markdown-user-content :global(td),
div.markdown-user-content :global(.table-wrapper) {
border-color: currentColor;
}
/* Horizontal rules */
.markdown-content :global(hr) {
border: none;
border-top: 1px solid var(--border);
margin: 1.5rem 0;
}
/* Images */
.markdown-content :global(img) {
border-radius: 0.5rem;
box-shadow:
0 1px 3px 0 rgb(0 0 0 / 0.1),
0 1px 2px -1px rgb(0 0 0 / 0.1);
margin: 1.5rem 0;
max-width: 100%;
height: auto;
}
/* Code blocks */
.markdown-content :global(.code-block-wrapper) {
margin: 1.5rem 0;
border-radius: 0.75rem;
overflow: hidden;
border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);
background: var(--code-background);
box-shadow: 0 1px 2px 0 rgb(0 0 0 / 0.05);
min-height: var(--min-message-height);
max-height: var(--max-message-height);
}
.markdown-content:global(.dark) :global(.code-block-wrapper) {
border-color: color-mix(in oklch, var(--border) 20%, transparent);
}
/* Scroll container for code blocks (both streaming and completed) */
.markdown-content :global(.code-block-scroll-container),
.streaming-code-scroll-container {
min-height: var(--min-message-height);
max-height: var(--max-message-height);
overflow-y: auto;
overflow-x: auto;
padding: 3rem 1rem 1rem;
line-height: 1.3;
}
.full-height-code-blocks :global(.code-block-wrapper) {
max-height: none;
}
.full-height-code-blocks :global(.code-block-scroll-container),
.full-height-code-blocks .streaming-code-scroll-container {
max-height: none;
overflow-y: visible;
}
.markdown-content :global(.code-block-header) {
display: flex;
justify-content: space-between;
align-items: center;
padding: 0.5rem 1rem 0;
font-size: 0.875rem;
position: absolute;
top: 0;
left: 0;
right: 0;
}
.markdown-content :global(.code-language) {
color: var(--color-foreground);
font-weight: 500;
font-family:
ui-monospace, SFMono-Regular, 'SF Mono', Monaco, 'Cascadia Code', 'Roboto Mono', Consolas,
'Liberation Mono', Menlo, monospace;
text-transform: uppercase;
font-size: 0.75rem;
letter-spacing: 0.05em;
}
.markdown-content :global(.code-block-actions) {
display: flex;
align-items: center;
gap: 0.5rem;
}
.markdown-content :global(.copy-code-btn),
.markdown-content :global(.preview-code-btn) {
display: flex;
align-items: center;
justify-content: center;
padding: 0;
background: transparent;
color: var(--code-foreground);
cursor: pointer;
transition: all 0.2s ease;
}
.markdown-content :global(.copy-code-btn:hover),
.markdown-content :global(.preview-code-btn:hover) {
transform: scale(1.05);
}
.markdown-content :global(.copy-code-btn:active),
.markdown-content :global(.preview-code-btn:active) {
transform: scale(0.95);
}
.markdown-content :global(.code-block-wrapper pre) {
background: transparent;
margin: 0;
border-radius: 0;
border: none;
font-size: 0.875rem;
}
/* Mentions and hashtags */
.markdown-content :global(.mention) {
color: hsl(var(--primary));
font-weight: 500;
text-decoration: none;
}
.markdown-content :global(.mention:hover) {
text-decoration: underline;
}
.markdown-content :global(.hashtag) {
color: hsl(var(--primary));
font-weight: 500;
text-decoration: none;
}
.markdown-content :global(.hashtag:hover) {
text-decoration: underline;
}
/* Advanced table enhancements */
.markdown-content :global(table) {
transition: all 0.2s ease;
}
.markdown-content :global(table:hover) {
box-shadow:
0 4px 6px -1px rgb(0 0 0 / 0.1),
0 2px 4px -2px rgb(0 0 0 / 0.1);
}
.markdown-content :global(th:hover),
.markdown-content :global(td:hover) {
background: var(--muted);
}
/* Disable hover effects when rendering user messages */
.markdown-user-content :global(a),
.markdown-user-content :global(a:hover) {
color: inherit;
}
.markdown-user-content :global(table:hover) {
box-shadow: none;
}
.markdown-user-content :global(th:hover),
.markdown-user-content :global(td:hover) {
background: inherit;
}
/* Enhanced blockquotes */
.markdown-content :global(blockquote) {
transition: all 0.2s ease;
position: relative;
}
.markdown-content :global(blockquote:hover) {
border-left-width: 6px;
background: var(--muted);
transform: translateX(2px);
}
.markdown-content :global(blockquote::before) {
content: '"';
position: absolute;
top: -0.5rem;
left: 0.5rem;
font-size: 3rem;
color: var(--muted-foreground);
font-family: serif;
line-height: 1;
}
/* Enhanced images */
.markdown-content :global(img) {
transition: all 0.3s ease;
cursor: pointer;
}
.markdown-content :global(img:hover) {
transform: scale(1.02);
box-shadow:
0 10px 15px -3px rgb(0 0 0 / 0.1),
0 4px 6px -4px rgb(0 0 0 / 0.1);
}
/* Image zoom overlay */
.markdown-content :global(.image-zoom-overlay) {
position: fixed;
top: 0;
left: 0;
right: 0;
bottom: 0;
background: rgba(0, 0, 0, 0.8);
display: flex;
align-items: center;
justify-content: center;
z-index: 1000;
cursor: pointer;
}
.markdown-content :global(.image-zoom-overlay img) {
max-width: 90vw;
max-height: 90vh;
border-radius: 0.5rem;
box-shadow: 0 25px 50px -12px rgb(0 0 0 / 0.25);
}
/* Enhanced horizontal rules */
.markdown-content :global(hr) {
border: none;
height: 2px;
background: linear-gradient(to right, transparent, var(--border), transparent);
margin: 2rem 0;
position: relative;
}
.markdown-content :global(hr::after) {
content: '';
position: absolute;
top: 50%;
left: 50%;
transform: translate(-50%, -50%);
width: 1rem;
height: 1rem;
background: var(--border);
border-radius: 50%;
}
/* Scrollable tables */
.markdown-content :global(.table-wrapper) {
overflow-x: auto;
margin: 1.5rem 0;
border-radius: 0.5rem;
border: 1px solid var(--border);
}
.markdown-content :global(.table-wrapper table) {
margin: 0;
border: none;
}
/* Responsive adjustments */
@media (max-width: 640px) {
.markdown-content :global(h1) {
font-size: 1.5rem;
}
.markdown-content :global(h2) {
font-size: 1.25rem;
}
.markdown-content :global(h3) {
font-size: 1.125rem;
}
.markdown-content :global(table) {
font-size: 0.875rem;
}
.markdown-content :global(th),
.markdown-content :global(td) {
padding: 0.375rem 0.5rem;
}
.markdown-content :global(.table-wrapper) {
margin: 0.5rem -1rem;
border-radius: 0;
border-left: none;
border-right: none;
}
}
/* Dark mode adjustments */
@media (prefers-color-scheme: dark) {
.markdown-content :global(blockquote:hover) {
background: var(--muted);
}
}
/* Image load error fallback */
.markdown-content :global(.image-load-error) {
display: flex;
align-items: center;
justify-content: center;
margin: 1.5rem 0;
padding: 1.5rem;
border-radius: 0.5rem;
background: var(--muted);
border: 1px dashed var(--border);
}
.markdown-content :global(.image-error-content) {
display: flex;
flex-direction: column;
align-items: center;
gap: 0.75rem;
color: var(--muted-foreground);
text-align: center;
}
.markdown-content :global(.image-error-content svg) {
opacity: 0.5;
}
.markdown-content :global(.image-error-text) {
font-size: 0.875rem;
}
.markdown-content :global(.image-error-link) {
display: inline-flex;
align-items: center;
gap: 0.375rem;
padding: 0.5rem 1rem;
font-size: 0.875rem;
font-weight: 500;
color: var(--primary);
background: var(--background);
border: 1px solid var(--border);
border-radius: 0.375rem;
text-decoration: none;
transition: all 0.2s ease;
}
.markdown-content :global(.image-error-link:hover) {
background: var(--muted);
border-color: var(--primary);
}
/* Mermaid diagrams */
.markdown-content :global(pre.mermaid) {
background: transparent;
border: none;
padding: 0;
text-align: center;
font-family: inherit;
cursor: pointer;
transition: opacity 0.15s ease;
position: relative;
}
/* Hide mermaid code text until rendered - prevents flash */
.markdown-content :global(pre.mermaid:not([data-mermaid-rendered])),
.markdown-content :global(pre.mermaid[data-mermaid-rendered]:not(:has(svg))) {
display: none;
}
.markdown-content :global(pre.mermaid:hover) {
opacity: 0.85;
}
.markdown-content :global(pre.mermaid svg) {
max-width: 90%;
margin: 0 auto;
height: auto;
display: block;
padding: 3rem 1rem;
}
/* Mermaid block wrapper - matches code block styling */
.markdown-content :global(.mermaid-block-wrapper) {
margin: 1.5rem 0;
border-radius: 0.75rem;
overflow: hidden;
border: 1px solid color-mix(in oklch, var(--border) 30%, transparent);
background: var(--code-background);
box-shadow: 0 1px 2px 0 rgb(0 0 0 / 0.05);
position: relative;
min-height: var(--min-message-height);
max-height: var(--max-message-height);
}
.markdown-content:global(.dark) :global(.mermaid-block-wrapper) {
border-color: color-mix(in oklch, var(--border) 20%, transparent);
}
.markdown-content :global(.mermaid-scroll-container) {
min-height: 350px;
max-height: var(--max-message-height);
overflow-y: auto;
overflow-x: auto;
display: flex;
align-items: center;
justify-content: center;
padding: 3rem 1rem 1rem;
}
.full-height-code-blocks :global(.mermaid-block-wrapper) {
max-height: none;
}
.full-height-code-blocks :global(.mermaid-scroll-container) {
max-height: none;
overflow-y: visible;
}
/* Mermaid block uses same header styling as code blocks */
.markdown-content :global(.mermaid-block-wrapper .code-block-header) {
display: flex;
justify-content: space-between;
align-items: center;
padding: 0.5rem 1rem 0;
font-size: 0.875rem;
position: absolute;
top: 0;
left: 0;
right: 0;
}
.markdown-content :global(.mermaid-block-wrapper .code-block-actions) {
display: flex;
align-items: center;
gap: 0.5rem;
}
/* Mermaid pre element - remove default margins */
.markdown-content :global(.mermaid-block-wrapper pre.mermaid) {
background: transparent;
border: none;
padding: 0;
margin: 0;
text-align: center;
}
/* Mermaid SVG should be bigger */
.markdown-content :global(.mermaid-block-wrapper pre.mermaid svg) {
width: unset !important;
height: auto;
display: block;
padding: 3rem 1rem;
}
/* Streaming mermaid block - empty preview box */
.mermaid-streaming-block {
min-height: 300px;
display: flex;
align-items: center;
justify-content: center;
}
.mermaid-loading-placeholder {
display: flex;
align-items: center;
justify-content: center;
padding: 3rem;
color: var(--muted-foreground);
}
.mermaid-loading-text {
font-size: 0.875rem;
font-style: italic;
}
@@ -0,0 +1,264 @@
/**
* Event handler factories for markdown content components.
* Uses dependency injection pattern to avoid direct component state access.
*/
import { copyCodeToClipboard, copyToClipboard } from '$lib/utils';
export interface PreviewState {
previewDialogOpen: boolean;
previewCode: string;
previewLanguage: string;
setPreviewDialogOpen: (open: boolean) => void;
setPreviewCode: (code: string) => void;
setPreviewLanguage: (lang: string) => void;
}
export interface MermaidPreviewState {
mermaidPreviewOpen: boolean;
mermaidPreviewSvgHtml: string;
setMermaidPreviewOpen: (open: boolean) => void;
setMermaidPreviewSvgHtml: (html: string) => void;
}
export interface RenderedBlocksState {
renderedBlocks: Array<{ id: string; html: string; contentHash?: string }>;
setRenderedBlocks: (blocks: Array<{ id: string; html: string; contentHash?: string }>) => void;
}
/**
* Creates a click handler for copy buttons in code blocks.
* Copies the code content to clipboard.
*/
export function createHandleCopyClick() {
return async function handleCopyClick(event: Event) {
event.preventDefault();
event.stopPropagation();
const target = event.currentTarget as HTMLButtonElement | null;
if (!target) return;
const wrapper = target.closest('.code-block-wrapper');
if (!wrapper) return;
const codeElement = wrapper.querySelector<HTMLElement>('code[data-code-id]');
if (!codeElement) return;
const rawCode = codeElement.textContent ?? '';
try {
await copyCodeToClipboard(rawCode);
} catch (error) {
console.error('Failed to copy code:', error);
}
};
}
/**
* Creates a handler for preview dialog open state changes.
* Clears preview content when dialog is closed.
*/
export function createHandlePreviewDialogOpenChange(previewState: PreviewState) {
return function handlePreviewDialogOpenChange(open: boolean) {
previewState.setPreviewDialogOpen(open);
if (!open) {
previewState.setPreviewCode('');
previewState.setPreviewLanguage('text');
}
};
}
/**
* Creates a click handler for preview buttons within HTML code blocks.
* Opens a preview dialog with the rendered HTML content.
*/
export function createHandlePreviewClick(previewState: PreviewState) {
return async function handlePreviewClick(event: Event) {
event.preventDefault();
event.stopPropagation();
const target = event.currentTarget as HTMLButtonElement | null;
if (!target) return;
const wrapper = target.closest('.code-block-wrapper');
if (!wrapper) return;
const codeElement = wrapper.querySelector<HTMLElement>('code[data-code-id]');
if (!codeElement) return;
const rawCode = codeElement.textContent ?? '';
const languageLabel = wrapper.querySelector<HTMLElement>('.code-language');
const language = languageLabel?.textContent?.trim() || 'text';
previewState.setPreviewCode(rawCode);
previewState.setPreviewLanguage(language);
previewState.setPreviewDialogOpen(true);
};
}
/**
* Creates a click handler for mermaid block interactions.
* Handles copy, preview, and diagram click events via event delegation.
*/
export function createHandleMermaidClick(mermaidState: MermaidPreviewState) {
return async function handleMermaidClick(event: MouseEvent) {
const target = event.target as HTMLElement;
// Check if clicking on copy or preview button in mermaid block
const copyBtn = target.closest('.mermaid-block-wrapper .copy-code-btn');
const previewBtn = target.closest('.mermaid-block-wrapper .preview-code-btn');
if (copyBtn || previewBtn) {
const wrapper = target.closest('.mermaid-block-wrapper');
if (!wrapper) return;
const preElement = wrapper.querySelector<HTMLElement>('pre.mermaid[data-mermaid-syntax]');
if (!preElement) return;
const mermaidSyntax = preElement.dataset.mermaidSyntax ?? '';
if (copyBtn) {
event.preventDefault();
event.stopPropagation();
try {
await copyToClipboard(mermaidSyntax);
} catch (error) {
console.error('Failed to copy mermaid syntax:', error);
}
return;
}
if (previewBtn) {
event.preventDefault();
event.stopPropagation();
const svg = preElement.querySelector('svg');
if (!svg) return;
mermaidState.setMermaidPreviewSvgHtml(svg.outerHTML);
mermaidState.setMermaidPreviewOpen(true);
return;
}
}
// Otherwise, open preview when clicking on the mermaid diagram itself
const mermaidEl = target.closest('.mermaid');
if (!mermaidEl) return;
const svg = mermaidEl.querySelector('svg');
if (!svg) return;
mermaidState.setMermaidPreviewSvgHtml(svg.outerHTML);
mermaidState.setMermaidPreviewOpen(true);
};
}
/**
* Creates a handler for mermaid preview dialog open state changes.
* Cleans up SVG content when dialog is closed.
*/
export function createHandleMermaidPreviewOpenChange(mermaidState: MermaidPreviewState) {
return function handleMermaidPreviewOpenChange(open: boolean) {
mermaidState.setMermaidPreviewOpen(open);
if (!open) {
mermaidState.setMermaidPreviewSvgHtml('');
}
};
}
/**
* Creates an error handler for images that fail to load (e.g., CORS issues).
* Shows fallback UI for broken images.
*/
export function createHandleImageError(
renderedBlocksState: RenderedBlocksState,
IMAGE_NOT_ERROR_BOUND_SELECTOR: string,
DATA_ERROR_BOUND_ATTR: string,
BOOL_TRUE_STRING: string
) {
return async function handleImageError(event: Event) {
const img = event.target as HTMLImageElement;
if (!img) return;
const blockId = img.closest('[data-block-id]')?.getAttribute('data-block-id');
if (!blockId) return;
const block = renderedBlocksState.renderedBlocks.find((b) => b.id === blockId);
if (!block) return;
// Skip if already handled
if (img.dataset[DATA_ERROR_BOUND_ATTR] === BOOL_TRUE_STRING) return;
img.dataset[DATA_ERROR_BOUND_ATTR] = BOOL_TRUE_STRING;
// Get the fallback HTML and replace the image
const fallbackHtml = `<div class="image-error-placeholder" data-original-src="${img.src}">
<span class="image-error-icon"></span>
<span class="image-error-text">Failed to load image</span>
</div>`;
// Replace the img element with fallback in the block's HTML
const newHtml = block.html.replace(/img[^>]*src=["']([^"']*)[^>]*>/g, (match, src) => {
if (src === img.src) {
return fallbackHtml.replace('data-original-src=""', `data-original-src="${src}"`);
}
return match;
});
// Update the block
const newBlocks = renderedBlocksState.renderedBlocks.map((b) =>
b.id === blockId ? { ...b, html: newHtml } : b
);
renderedBlocksState.setRenderedBlocks(newBlocks);
};
}
/**
* Creates a function to set up code block action event listeners.
* Binds click handlers to copy and preview buttons within code blocks.
*/
export function createSetupCodeBlockActions(
handleCopyClick: (event: Event) => void,
handlePreviewClick: (event: Event) => void
) {
return function setupCodeBlockActions(containerRef: HTMLElement | null) {
if (!containerRef) return;
const wrappers = containerRef.querySelectorAll<HTMLElement>('.code-block-wrapper');
for (const wrapper of wrappers) {
const copyButton = wrapper.querySelector<HTMLButtonElement>('.copy-code-btn');
const previewButton = wrapper.querySelector<HTMLButtonElement>('.preview-code-btn');
if (copyButton && copyButton.dataset.listenerBound !== 'true') {
copyButton.dataset.listenerBound = 'true';
copyButton.addEventListener('click', handleCopyClick);
}
if (previewButton && previewButton.dataset.listenerBound !== 'true') {
previewButton.dataset.listenerBound = 'true';
previewButton.addEventListener('click', handlePreviewClick);
}
}
};
}
/**
* Creates a function to set up image error handlers.
* Attaches error handlers to images to show fallback UI when loading fails.
*/
export function createSetupImageErrorHandlers(
handleImageError: (event: Event) => void,
IMAGE_NOT_ERROR_BOUND_SELECTOR: string,
DATA_ERROR_BOUND_ATTR: string,
BOOL_TRUE_STRING: string
) {
return function setupImageErrorHandlers(containerRef: HTMLElement | null) {
if (!containerRef) return;
const images = containerRef.querySelectorAll<HTMLImageElement>(IMAGE_NOT_ERROR_BOUND_SELECTOR);
for (const img of images) {
img.dataset[DATA_ERROR_BOUND_ATTR] = BOOL_TRUE_STRING;
img.addEventListener('error', handleImageError);
}
};
}
@@ -0,0 +1,84 @@
/**
* Utility functions for markdown processing in MarkdownContent component.
*/
import type { RootContent as HastRootContent } from 'hast';
/**
* Generates a unique identifier for a HAST node based on its position.
* Used for stable block identification during incremental rendering.
* @param node - The HAST root content node
* @param indexFallback - Fallback index if position is unavailable
* @returns Unique string identifier for the node
*/
export function getHastNodeId(node: HastRootContent, indexFallback: number): string {
const position = node.position;
if (position?.start?.offset != null && position?.end?.offset != null) {
return `hast-${position.start.offset}-${position.end.offset}`;
}
return `${node.type}-${indexFallback}`;
}
/**
* Generates a hash for MDAST node based on its position.
* Used for cache lookup during incremental rendering.
*/
export function getMdastNodeHash(node: unknown, index: number): string {
const n = node as {
type?: string;
position?: { start?: { offset?: number }; end?: { offset?: number } };
};
if (n.position?.start?.offset != null && n.position?.end?.offset != null) {
return `${n.type}-${n.position.start.offset}-${n.position.end.offset}`;
}
return `${n.type}-idx${index}`;
}
/**
* Determines if the new content is an append (new content added to existing blocks).
* This is used to optimize cache reuse during streaming updates.
*
* @param newContent - The new markdown content
* @param previousContent - The previous markdown content to check against
* @returns true if the content appears to be an append operation
*/
export function isAppendMode(newContent: string, previousContent: string): boolean {
return previousContent.length > 0 && newContent.startsWith(previousContent);
}
export interface CodeInfo {
rawCode: string;
language: string;
}
/**
* Extracts code information from a button click target within a code block.
* @param target - The clicked button element
* @returns Object with rawCode and language, or null if extraction fails
*/
export function getCodeInfoFromTarget(target: HTMLElement): CodeInfo | null {
const wrapper = target.closest('.code-block-wrapper');
if (!wrapper) {
console.error('No wrapper found');
return null;
}
const codeElement = wrapper.querySelector<HTMLElement>('code[data-code-id]');
if (!codeElement) {
console.error('No code element found in wrapper');
return null;
}
const rawCode = codeElement.textContent ?? '';
const languageLabel = wrapper.querySelector<HTMLElement>('.code-language');
const language = languageLabel?.textContent?.trim() || 'text';
return { rawCode, language };
}
@@ -0,0 +1,151 @@
/**
* Shared utilities for enhanced code blocks and mermaid diagram blocks.
* Contains common HAST element creation functions to avoid code duplication.
*/
import type { Element, ElementContent } from 'hast';
import {
CODE_BLOCK_HEADER_CLASS,
CODE_BLOCK_ACTIONS_CLASS,
CODE_LANGUAGE_CLASS,
COPY_CODE_BTN_CLASS,
PREVIEW_CODE_BTN_CLASS,
RELATIVE_CLASS,
COPY_ICON_SVG,
PREVIEW_ICON_SVG
} from '$lib/constants';
export interface BlockIdGenerator {
(id: number): string;
}
/**
* Creates an icon element with the given SVG content.
*/
export function createIconElement(svg: string): Element {
return {
type: 'element',
tagName: 'span',
properties: {},
children: [{ type: 'raw', value: svg } as unknown as ElementContent]
};
}
/**
* Creates a button element with icon.
*/
export function createButton(
className: string,
title: string,
iconSvg: string,
id: string,
idAttribute: string
): Element {
return {
type: 'element',
tagName: 'button',
properties: {
className: [className],
[idAttribute]: id,
title,
type: 'button'
},
children: [createIconElement(iconSvg)]
};
}
/**
* Creates a copy button element.
*/
export function createCopyButton(id: string, idAttribute: string, title: string = 'Copy'): Element {
return createButton(COPY_CODE_BTN_CLASS, title, COPY_ICON_SVG, id, idAttribute);
}
/**
* Creates a preview button element.
*/
export function createPreviewButton(
id: string,
idAttribute: string,
title: string = 'Preview'
): Element {
return createButton(PREVIEW_CODE_BTN_CLASS, title, PREVIEW_ICON_SVG, id, idAttribute);
}
/**
* Creates a block header with language label and action buttons.
*/
export function createBlockHeader(
language: string,
id: string,
idAttribute: string,
actions: Element[],
languageClassName: string = CODE_LANGUAGE_CLASS
): Element {
return {
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_HEADER_CLASS] },
children: [
{
type: 'element',
tagName: 'span',
properties: { className: [languageClassName] },
children: [{ type: 'text', value: language }]
},
{
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_ACTIONS_CLASS] },
children: actions
}
]
};
}
/**
* Creates a scroll container element.
*/
export function createScrollContainer(preElement: Element, scrollContainerClass: string): Element {
return {
type: 'element',
tagName: 'div',
properties: { className: [scrollContainerClass] },
children: [preElement]
};
}
/**
* Creates a wrapper element with header and scroll container.
*/
export function createWrapper(
header: Element,
preElement: Element,
wrapperClass: string,
scrollContainerClass: string,
additionalAttributes?: Record<string, string>
): Element {
return {
type: 'element',
tagName: 'div',
properties: {
className: [wrapperClass, RELATIVE_CLASS],
...additionalAttributes
} as Element['properties'],
children: [header, createScrollContainer(preElement, scrollContainerClass)]
};
}
/**
* Generates a unique block ID using a global counter.
*/
export function generateBlockId(prefix: string, windowKey: keyof Window): string {
if (typeof window !== 'undefined') {
const idx = window[windowKey] as number | undefined;
const next = (idx ?? 0) + 1;
(window as unknown as Record<string, number>)[windowKey] = next;
return `${prefix}-${next}`;
}
// Fallback for SSR - use timestamp + random
return `${prefix}-${Date.now()}-${Math.random().toString(36).slice(2, 7)}`;
}
@@ -13,16 +13,14 @@
import type { Plugin } from 'unified';
import type { Root, Element, ElementContent } from 'hast';
import { visit } from 'unist-util-visit';
import { CODE_BLOCK_SCROLL_CONTAINER_CLASS, CODE_BLOCK_WRAPPER_CLASS } from '$lib/constants';
import {
CODE_BLOCK_SCROLL_CONTAINER_CLASS,
CODE_BLOCK_WRAPPER_CLASS,
CODE_BLOCK_HEADER_CLASS,
CODE_BLOCK_ACTIONS_CLASS,
CODE_LANGUAGE_CLASS,
COPY_CODE_BTN_CLASS,
PREVIEW_CODE_BTN_CLASS,
RELATIVE_CLASS
} from '$lib/constants';
createBlockHeader,
createCopyButton,
createPreviewButton,
createWrapper,
generateBlockId
} from './code-block-utils';
declare global {
interface Window {
@@ -30,87 +28,6 @@ declare global {
}
}
const COPY_ICON_SVG = `<svg xmlns="http://www.w3.org/2000/svg" width="16" height="16" viewBox="0 0 24 24" fill="none" stroke="currentColor" stroke-width="2" stroke-linecap="round" stroke-linejoin="round" class="lucide lucide-copy-icon lucide-copy"><rect width="14" height="14" x="8" y="8" rx="2" ry="2"/><path d="M4 16c-1.1 0-2-.9-2-2V4c0-1.1.9-2 2-2h10c1.1 0 2 .9 2 2"/></svg>`;
const PREVIEW_ICON_SVG = `<svg xmlns="http://www.w3.org/2000/svg" width="16" height="16" viewBox="0 0 24 24" fill="none" stroke="currentColor" stroke-width="2" stroke-linecap="round" stroke-linejoin="round" class="lucide lucide-eye lucide-eye-icon"><path d="M2.062 12.345a1 1 0 0 1 0-.69C3.5 7.73 7.36 5 12 5s8.5 2.73 9.938 6.655a1 1 0 0 1 0 .69C20.5 16.27 16.64 19 12 19s-8.5-2.73-9.938-6.655"/><circle cx="12" cy="12" r="3"/></svg>`;
function createIconElement(svg: string): Element {
return {
type: 'element',
tagName: 'span',
properties: {},
children: [{ type: 'raw', value: svg } as unknown as ElementContent]
};
}
function createButton(className: string, title: string, iconSvg: string, codeId: string): Element {
return {
type: 'element',
tagName: 'button',
properties: {
className: [className],
'data-code-id': codeId,
title,
type: 'button'
},
children: [createIconElement(iconSvg)]
};
}
function createCopyButton(codeId: string): Element {
return createButton(COPY_CODE_BTN_CLASS, 'Copy code', COPY_ICON_SVG, codeId);
}
function createPreviewButton(codeId: string): Element {
return createButton(PREVIEW_CODE_BTN_CLASS, 'Preview code', PREVIEW_ICON_SVG, codeId);
}
function createHeader(language: string, codeId: string): Element {
const actions: Element[] = [createCopyButton(codeId)];
if (language.toLowerCase() === 'html') {
actions.push(createPreviewButton(codeId));
}
return {
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_HEADER_CLASS] },
children: [
{
type: 'element',
tagName: 'span',
properties: { className: [CODE_LANGUAGE_CLASS] },
children: [{ type: 'text', value: language }]
},
{
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_ACTIONS_CLASS] },
children: actions
}
]
};
}
function createScrollContainer(preElement: Element): Element {
return {
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_SCROLL_CONTAINER_CLASS] },
children: [preElement]
};
}
function createWrapper(header: Element, preElement: Element): Element {
return {
type: 'element',
tagName: 'div',
properties: { className: [CODE_BLOCK_WRAPPER_CLASS, RELATIVE_CLASS] },
children: [header, createScrollContainer(preElement)]
};
}
function extractLanguage(codeElement: Element): string {
const className = codeElement.properties?.className;
if (!Array.isArray(className)) return 'text';
@@ -124,17 +41,6 @@ function extractLanguage(codeElement: Element): string {
return 'text';
}
/**
* Generates a unique code block ID using a global counter.
*/
function generateCodeId(): string {
if (typeof window !== 'undefined') {
return `code-${(window.idxCodeBlock = (window.idxCodeBlock ?? 0) + 1)}`;
}
// Fallback for SSR - use timestamp + random
return `code-${Date.now()}-${Math.random().toString(36).slice(2, 7)}`;
}
/**
* Rehype plugin to enhance code blocks with wrapper, header, and action buttons.
* This plugin wraps <pre><code> elements with a container that includes:
@@ -154,15 +60,26 @@ export const rehypeEnhanceCodeBlocks: Plugin<[], Root> = () => {
if (!codeElement) return;
const language = extractLanguage(codeElement);
const codeId = generateCodeId();
const codeId = generateBlockId('code', 'idxCodeBlock');
codeElement.properties = {
...codeElement.properties,
'data-code-id': codeId
};
const header = createHeader(language, codeId);
const wrapper = createWrapper(header, node);
const actions: Element[] = [createCopyButton(codeId, 'data-code-id', 'Copy code')];
if (language.toLowerCase() === 'html') {
actions.push(createPreviewButton(codeId, 'data-code-id', 'Preview code'));
}
const header = createBlockHeader(language, codeId, 'data-code-id', actions);
const wrapper = createWrapper(
header,
node,
CODE_BLOCK_WRAPPER_CLASS,
CODE_BLOCK_SCROLL_CONTAINER_CLASS
);
// Replace pre with wrapper in parent
(parent.children as ElementContent[])[index] = wrapper;
@@ -0,0 +1,85 @@
/**
* Rehype plugin to enhance mermaid diagram blocks with wrapper, header, and action buttons.
*
* Wraps <pre class="mermaid"> elements with a container that includes:
* - Language label ("mermaid")
* - Copy button (copies mermaid syntax to clipboard)
* - Preview button (opens fullscreen preview dialog)
*
* This operates directly on the HAST tree for better performance,
* avoiding the need to stringify and re-parse HTML.
*/
import type { Plugin } from 'unified';
import type { Root, Element, ElementContent } from 'hast';
import { visit } from 'unist-util-visit';
import { MERMAID_WRAPPER_CLASS, MERMAID_SCROLL_CONTAINER_CLASS } from '$lib/constants';
import {
createBlockHeader,
createCopyButton,
createPreviewButton,
createWrapper,
generateBlockId
} from './code-block-utils';
declare global {
interface Window {
idxMermaidBlock?: number;
}
}
/**
* Rehype plugin to enhance mermaid diagram blocks with wrapper, header, and action buttons.
* This plugin wraps <pre class="mermaid"> elements with a container that includes:
* - Language label ("mermaid")
* - Copy button
* - Preview button
*/
export const rehypeEnhanceMermaidBlocks: Plugin<[], Root> = () => {
return (tree: Root) => {
visit(tree, 'element', (node: Element, index, parent) => {
if (node.tagName !== 'pre' || !parent || index === undefined) return;
const className = node.properties?.className;
if (!Array.isArray(className)) return;
const isMermaid = className.some((cls) => typeof cls === 'string' && cls === 'mermaid');
if (!isMermaid) return;
const mermaidId = generateBlockId('mermaid', 'idxMermaidBlock');
// Extract the mermaid syntax (text content of the pre element)
const diagramText = node.children
.map((child) => {
if (child.type === 'text') return child.value;
return '';
})
.join('');
// Store the mermaid syntax in data attribute for copy functionality
node.properties = {
...node.properties,
'data-mermaid-syntax': diagramText,
'data-mermaid-id': mermaidId
};
const actions = [
createCopyButton(mermaidId, 'data-mermaid-id', 'Copy mermaid syntax'),
createPreviewButton(mermaidId, 'data-mermaid-id', 'Preview diagram')
];
const header = createBlockHeader('mermaid', mermaidId, 'data-mermaid-id', actions);
const wrapper = createWrapper(
header,
node,
MERMAID_WRAPPER_CLASS,
MERMAID_SCROLL_CONTAINER_CLASS,
{ 'data-mermaid-id': mermaidId }
);
// Replace pre with wrapper in parent
(parent.children as ElementContent[])[index] = wrapper;
});
};
};
@@ -0,0 +1,67 @@
import type { Plugin } from 'unified';
import type { Root, Element, ElementContent, Text } from 'hast';
import { visit } from 'unist-util-visit';
/**
* Recursively extracts all text content from a HAST node.
* Handles nested elements (e.g., span wrappers from syntax highlighting).
*/
function extractText(node: ElementContent): string {
if (node.type === 'text') return node.value;
if (node.type === 'element') {
return (node.children ?? []).map(extractText).join('');
}
return '';
}
/**
* Rehype plugin to convert mermaid code blocks to <pre class="mermaid"> elements.
*
* Transforms:
* <pre><code class="language-mermaid">graph TD; A-->B</code></pre>
* into:
* <pre class="mermaid">graph TD; A-->B</pre>
*
* The mermaid library renders these client-side via mermaid.run().
*
* Must run BEFORE rehypeEnhanceCodeBlocks so mermaid blocks are not wrapped
* with code block headers/buttons (they have no <code> child, so they're skipped).
*/
export const rehypeMermaidPre: Plugin<[], Root> = () => {
return (tree: Root) => {
visit(tree, 'element', (node: Element, index, parent) => {
if (node.tagName !== 'pre' || !parent || index === undefined) return;
const codeElement = node.children.find(
(child): child is Element => child.type === 'element' && child.tagName === 'code'
);
if (!codeElement) return;
const className = codeElement.properties?.className;
if (!Array.isArray(className)) return;
const isMermaid = className.some(
(cls) => typeof cls === 'string' && cls === 'language-mermaid'
);
if (!isMermaid) return;
// Recursively extract text to handle nested spans from syntax highlighting
const diagramText = codeElement.children.map(extractText).join('').trim();
if (!diagramText) return;
const mermaidPre: Element = {
type: 'element',
tagName: 'pre',
properties: {
className: ['mermaid']
},
children: [{ type: 'text', value: diagramText } as Text]
};
(parent.children as ElementContent[])[index] = mermaidPre;
});
};
};
@@ -0,0 +1,126 @@
<script lang="ts">
import MermaidPreviewControls from './MermaidPreviewControls.svelte';
interface Props {
svgHtml: string;
}
let { svgHtml }: Props = $props();
// Zoom and pan state
let scale = $state(1);
let translateX = $state(0);
let translateY = $state(0);
let isDragging = $state(false);
const containerRef = { current: null as HTMLDivElement | null };
// Drag start position
let dragStartX = 0;
let dragStartY = 0;
let dragStartTranslateX = 0;
let dragStartTranslateY = 0;
const MIN_SCALE = 0.1;
const MAX_SCALE = 10;
const ZOOM_STEP = 0.15;
function resetView() {
scale = 1;
translateX = 0;
translateY = 0;
}
function zoomIn() {
scale = Math.min(scale + ZOOM_STEP, MAX_SCALE);
}
function zoomOut() {
scale = Math.max(scale - ZOOM_STEP, MIN_SCALE);
}
function handleWheel(event: WheelEvent) {
event.preventDefault();
const delta = event.deltaY > 0 ? -ZOOM_STEP : ZOOM_STEP;
scale = Math.min(Math.max(scale + delta, MIN_SCALE), MAX_SCALE);
}
// Imperatively attach a non-passive wheel listener so preventDefault() actually works
// (Svelte 5 wheel listeners are passive by default, making preventDefault() a no-op)
$effect(() => {
const el = containerRef.current;
if (!el) return;
function onWheel(e: WheelEvent) {
handleWheel(e);
}
el.addEventListener('wheel', onWheel, { passive: false });
return () => el.removeEventListener('wheel', onWheel);
});
function handlePointerDown(event: PointerEvent) {
if (event.button !== 0 && event.pointerType === 'mouse') return;
isDragging = true;
dragStartX = event.clientX;
dragStartY = event.clientY;
dragStartTranslateX = translateX;
dragStartTranslateY = translateY;
(event.currentTarget as HTMLElement).setPointerCapture(event.pointerId);
}
function handlePointerMove(event: PointerEvent) {
if (!isDragging) return;
translateX = dragStartTranslateX + (event.clientX - dragStartX);
translateY = dragStartTranslateY + (event.clientY - dragStartY);
}
function handlePointerUp() {
isDragging = false;
}
</script>
<div
bind:this={containerRef.current}
class="mermaid-preview relative flex items-center justify-center overflow-hidden bg-muted/20"
>
<!-- svelte-ignore a11y_no_static_element_interactions -->
<div
class="mermaid-preview-diagram transform-origin-center inline-block min-h-fit min-w-fit will-change-transform {isDragging &&
'select-none'}"
style="transform: translate({translateX}px, {translateY}px) scale({scale}); cursor: {isDragging
? 'grabbing'
: 'grab'};"
onpointerdown={handlePointerDown}
onpointermove={handlePointerMove}
onpointerup={handlePointerUp}
onpointerleave={handlePointerUp}
>
<!-- eslint-disable-next-line no-at-html-tags -->
{@html svgHtml}
</div>
<MermaidPreviewControls
{scale}
{svgHtml}
onZoomIn={zoomIn}
onZoomOut={zoomOut}
onResetView={resetView}
/>
</div>
<style lang="postcss" scoped>
/* Styles for SVGs rendered via {@html} — no Tailwind class can target child elements */
.mermaid-preview-diagram :global(svg) {
min-height: min(50vh, 12rem);
min-width: min(80vw, 20rem);
max-width: none !important;
max-height: none !important;
height: auto !important;
width: auto !important;
display: block;
}
</style>
@@ -0,0 +1,74 @@
<script lang="ts">
import { Download } from '@lucide/svelte';
import ZoomInIcon from '@lucide/svelte/icons/zoom-in';
import ZoomOutIcon from '@lucide/svelte/icons/zoom-out';
import RotateCcwIcon from '@lucide/svelte/icons/rotate-ccw';
interface Props {
scale: number;
svgHtml: string;
onZoomIn: () => void;
onZoomOut: () => void;
onResetView: () => void;
}
let { scale, svgHtml, onZoomIn, onZoomOut, onResetView }: Props = $props();
function downloadSvg() {
if (!svgHtml) return;
const blob = new Blob([svgHtml], { type: 'image/svg+xml' });
const url = URL.createObjectURL(blob);
const a = document.createElement('a');
a.href = url;
a.download = 'diagram.svg';
a.click();
URL.revokeObjectURL(url);
}
</script>
<div
class="mermaid-preview-controls absolute bottom-8 flex shrink-0 items-center justify-center p-3"
>
<div class="mermaid-preview-controls-inner flex items-center gap-1 rounded-lg bg-muted p-1">
<button
class="mermaid-preview-btn flex h-8 w-8 cursor-pointer items-center justify-center rounded-md border-0 bg-transparent text-foreground transition-colors hover:bg-muted-foreground/15 active:bg-muted-foreground/25"
onclick={onZoomOut}
title="Zoom out"
aria-label="Zoom out"
>
<ZoomOutIcon class="mermaid-preview-btn-icon h-4 w-4" />
</button>
<span
class="mermaid-preview-zoom-label min-w-[3.5rem] px-0.5 text-center text-xs font-medium text-muted-foreground tabular-nums select-none"
>{Math.round(scale * 100)}%</span
>
<button
class="mermaid-preview-btn flex h-8 w-8 cursor-pointer items-center justify-center rounded-md border-0 bg-transparent text-foreground transition-colors hover:bg-muted-foreground/15 active:bg-muted-foreground/25"
onclick={onZoomIn}
title="Zoom in"
aria-label="Zoom in"
>
<ZoomInIcon class="mermaid-preview-btn-icon h-4 w-4" />
</button>
<div class="mermaid-preview-controls-separator mx-1 h-5 w-px bg-border/50"></div>
<button
class="mermaid-preview-btn flex h-8 w-8 cursor-pointer items-center justify-center rounded-md border-0 bg-transparent text-foreground transition-colors hover:bg-muted-foreground/15 active:bg-muted-foreground/25"
onclick={onResetView}
title="Reset view"
aria-label="Reset view"
>
<RotateCcwIcon class="mermaid-preview-btn-icon h-4 w-4" />
</button>
<div class="mermaid-preview-controls-separator mx-1 h-5 w-px bg-border/50"></div>
<button
class="mermaid-preview-btn flex h-8 w-8 cursor-pointer items-center justify-center rounded-md border-0 bg-transparent text-foreground transition-colors hover:bg-muted-foreground/15 active:bg-muted-foreground/25"
onclick={downloadSvg}
title="Download SVG"
aria-label="Download SVG"
>
<Download class="mermaid-preview-btn-icon h-4 w-4" />
</button>
</div>
</div>
@@ -77,3 +77,22 @@ export { default as SyntaxHighlightedCode } from './SyntaxHighlightedCode.svelte
* ```
*/
export { default as CollapsibleContentBlock } from './CollapsibleContentBlock.svelte';
/**
* **MermaidPreview** - Interactive Mermaid diagram viewer
*
* Renders Mermaid-generated SVG diagrams with zoom, pan, and fit-to-view controls.
*
* **Features:**
* - Mouse wheel zoom in/out
* - Click-drag panning with pointer capture
* - Fit to view and reset view controls
* - Download as SVG
* - Responsive scaling with viewBox detection
*
* @example
* ```svelte
* <MermaidPreview svgHtml={diagramSvg} />
* ```
*/
export { default as MermaidPreview } from './MermaidPreview.svelte';

Some files were not shown because too many files have changed in this diff Show More