diff --git a/common/arg.cpp b/common/arg.cpp index 9753441313a7..79480e06f9d2 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -539,6 +539,13 @@ void common_models_handler_apply(common_models_handler & handler, common_params } }; + // an explicit draft file selection (e.g. -md with -hfd) disables the sidecar resolution of the draft repo + if (!params.speculative.draft.mparams.hf_file.empty()) { + plan_spec.mtp = {}; + plan_spec.dflash = {}; + plan_spec.eagle3 = {}; + } + // infer the speculative type from the sidecar shipped by the draft repo when none is requested if (spec_types_is_default(params)) { if (!plan_spec.mtp.local_path.empty()) { @@ -588,6 +595,11 @@ void common_models_handler_apply(common_models_handler & handler, common_params }); } + // a wired draft sidecar counts as an explicit draft for the main plan fallback below + if (spec_sidecar_found) { + had_spec_url = true; + } + // handle plan_spec (e.g. --spec-draft-hf) if (!plan_spec.model_files.empty() && !had_spec_url && !spec_sidecar_found) { add_tasks(plan_spec.model_files, plan_spec.primary, params.speculative.draft.mparams); @@ -1049,6 +1061,31 @@ static std::vector parse_device_list(const std::string & val return devices; } +void common_print_available_devices() { + constexpr size_t MiB = 1024 * 1024; + std::vector devices; + + ggml_backend_load_all(); + + for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { + auto * dev = ggml_backend_dev_get(i); + if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) { + devices.push_back(dev); + } + } + printf("Available devices:\n"); + + if (devices.empty()) { + printf(" (none)\n"); + return; + } + for (auto * dev : devices) { + size_t free, total; + ggml_backend_dev_memory(dev, &free, &total); + printf(" %s: %s (%zu MiB, %zu MiB free)\n", ggml_backend_dev_name(dev), ggml_backend_dev_description(dev), total / MiB, free / MiB); + } +} + static void add_rpc_devices(const std::string & servers) { auto rpc_servers = string_split(servers, ','); if (rpc_servers.empty()) { @@ -2508,7 +2545,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } add_opt(common_arg( {"--mlock"}, - "DEPRECATED in favor of `--load-mode`: mmap + force system to keep model in RAM rather than swapping or compressing", + "DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing", [](common_params & params) { LOG_WRN("DEPRECATED: --mlock is deprecated. use --load-mode mlock instead\n"); params.load_mode = LLAMA_LOAD_MODE_MLOCK; @@ -2537,13 +2574,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex "model loading mode (default: mmap)\n" "- none: no special loading mode\n" "- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)\n" - "- mlock: mmap + force system to keep model in RAM rather than swapping or compressing\n" + "- mlock: force system to keep model in RAM rather than swapping or compressing\n" + "- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing\n" "- dio: use DirectIO if available\n", [](common_params & params, const std::string & value) { - /**/ if (value == "none") { params.load_mode = LLAMA_LOAD_MODE_NONE; } - else if (value == "mmap") { params.load_mode = LLAMA_LOAD_MODE_MMAP; } - else if (value == "mlock") { params.load_mode = LLAMA_LOAD_MODE_MLOCK; } - else if (value == "dio") { params.load_mode = LLAMA_LOAD_MODE_DIRECT_IO; } + /**/ if (value == "none") { params.load_mode = LLAMA_LOAD_MODE_NONE; } + else if (value == "mmap") { params.load_mode = LLAMA_LOAD_MODE_MMAP; } + else if (value == "mlock") { params.load_mode = LLAMA_LOAD_MODE_MLOCK; } + else if (value == "mmap+mlock") { params.load_mode = LLAMA_LOAD_MODE_MMAP_MLOCK; } + else if (value == "dio") { params.load_mode = LLAMA_LOAD_MODE_DIRECT_IO; } else { throw std::invalid_argument("invalid value"); } } ).set_env("LLAMA_ARG_LOAD_MODE")); @@ -2574,20 +2613,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex {"--list-devices"}, "print list of available devices and exit", [](common_params &) { - ggml_backend_load_all(); - std::vector devices; - for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { - auto * dev = ggml_backend_dev_get(i); - if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) { - devices.push_back(dev); - } - } - printf("Available devices:\n"); - for (auto * dev : devices) { - size_t free, total; - ggml_backend_dev_memory(dev, &free, &total); - printf(" %s: %s (%zu MiB, %zu MiB free)\n", ggml_backend_dev_name(dev), ggml_backend_dev_description(dev), total / 1024 / 1024, free / 1024 / 1024); - } + common_print_available_devices(); exit(0); } )); diff --git a/common/arg.h b/common/arg.h index 54a38b9cce4a..8f609e356fe2 100644 --- a/common/arg.h +++ b/common/arg.h @@ -123,6 +123,9 @@ struct common_params_context { // if one argument has invalid value, it will automatically display usage of the specific argument (and not the full usage message) bool common_params_parse(int argc, char ** argv, common_params & params, llama_example ex, void(*print_usage)(int, char **) = nullptr); +// load all backends and print the list of available (non-CPU) devices to stdout +void common_print_available_devices(); + // parse input arguments from CLI into a map bool common_params_to_map(int argc, char ** argv, llama_example ex, std::map & out_map); diff --git a/common/download.cpp b/common/download.cpp index e8e938426f2a..3776c6c7eb68 100644 --- a/common/download.cpp +++ b/common/download.cpp @@ -568,16 +568,30 @@ static hf_cache::hf_files get_split_files(const hf_cache::hf_files & files, } // pick the best sibling GGUF whose filename contains `keyword` (e.g. "mmproj" / "mtp"), -// preferring deeper shared directory prefix with the model, then closest quantization +// preferring deeper shared directory prefix with the model, then exact `tag` match, +// then closest quantization to the tag when given, or to the model otherwise static hf_cache::hf_file find_best_sibling(const hf_cache::hf_files & files, const std::string & model, - const std::string & keyword) { + const std::string & keyword, + const std::string & tag = "") { hf_cache::hf_file best; size_t best_depth = 0; int best_diff = 0; + bool best_exact = false; bool found = false; - auto model_bits = extract_quant_bits(model); + std::string tag_upper = tag; + for (char & c : tag_upper) { + c = (char) std::toupper((unsigned char) c); + } + + int model_bits = 0; + if (!tag_upper.empty()) { + auto pos = tag_upper.find_first_of("0123456789"); + model_bits = pos == std::string::npos ? 0 : std::stoi(tag_upper.substr(pos)); + } else { + model_bits = extract_quant_bits(model); + } auto model_parts = string_split(model, '/'); auto model_dir = model_parts.end() - 1; @@ -600,10 +614,19 @@ static hf_cache::hf_file find_best_sibling(const hf_cache::hf_files & files, auto bits = extract_quant_bits(f.path); auto diff = std::abs(bits - model_bits); - if (!found || depth > best_depth || (depth == best_depth && diff < best_diff)) { + std::string path_upper = f.path; + for (char & c : path_upper) { + c = (char) std::toupper((unsigned char) c); + } + bool exact = !tag_upper.empty() && path_upper.find("-" + tag_upper + ".") != std::string::npos; + + if (!found || depth > best_depth || + (depth == best_depth && exact && !best_exact) || + (depth == best_depth && exact == best_exact && diff < best_diff)) { best = f; best_depth = depth; best_diff = diff; + best_exact = exact; found = true; } } @@ -616,18 +639,21 @@ static hf_cache::hf_file find_best_mmproj(const hf_cache::hf_files & files, } static hf_cache::hf_file find_best_mtp(const hf_cache::hf_files & files, - const std::string & model) { - return find_best_sibling(files, model, "mtp-"); + const std::string & model, + const std::string & tag = "") { + return find_best_sibling(files, model, "mtp-", tag); } static hf_cache::hf_file find_best_eagle3(const hf_cache::hf_files & files, - const std::string & model) { - return find_best_sibling(files, model, "eagle3-"); + const std::string & model, + const std::string & tag = "") { + return find_best_sibling(files, model, "eagle3-", tag); } static hf_cache::hf_file find_best_dflash(const hf_cache::hf_files & files, - const std::string & model) { - return find_best_sibling(files, model, "dflash-"); + const std::string & model, + const std::string & tag = "") { + return find_best_sibling(files, model, "dflash-", tag); } static bool gguf_filename_is_model(const std::string & filepath) { @@ -736,27 +762,36 @@ common_download_hf_plan common_download_get_hf_plan(const common_params_model & } } else { primary = find_best_model(all, tag); - if (primary.path.empty()) { + // a requested sidecar can resolve on its own, without a full model of the same tag + if (primary.path.empty() && !opts.download_mtp && !opts.download_dflash && !opts.download_eagle3) { LOG_ERR("%s: no GGUF files found in repository %s\n", __func__, repo.c_str()); list_available_gguf_files(all); return plan; } } - plan.primary = primary; - plan.model_files = get_split_files(all, primary); + if (!primary.path.empty()) { + plan.primary = primary; + plan.model_files = get_split_files(all, primary); + } - if (opts.download_mmproj) { + if (opts.download_mmproj && !primary.path.empty()) { plan.mmproj = find_best_mmproj(all, primary.path); } if (opts.download_mtp) { - plan.mtp = find_best_mtp(all, primary.path); + plan.mtp = find_best_mtp(all, primary.path, tag); } if (opts.download_dflash) { - plan.dflash = find_best_dflash(all, primary.path); + plan.dflash = find_best_dflash(all, primary.path, tag); } if (opts.download_eagle3) { - plan.eagle3 = find_best_eagle3(all, primary.path); + plan.eagle3 = find_best_eagle3(all, primary.path, tag); + } + + if (primary.path.empty() && + plan.mtp.local_path.empty() && plan.dflash.local_path.empty() && plan.eagle3.local_path.empty()) { + LOG_ERR("%s: no GGUF files found in repository %s\n", __func__, repo.c_str()); + list_available_gguf_files(all); } return plan; diff --git a/common/fit.cpp b/common/fit.cpp index c79221cb00fa..c82d066ad444 100644 --- a/common/fit.cpp +++ b/common/fit.cpp @@ -136,7 +136,7 @@ static std::vector common_get_device_memory_data_impl( devs.push_back(llama_model_get_device(model, i)); } - hp_ngl = llama_model_n_layer(model); + hp_ngl = llama_model_n_layer(model) + llama_model_n_layer_nextn(model); hp_n_ctx_train = llama_model_n_ctx_train(model); hp_n_expert = llama_model_n_expert(model); diff --git a/common/speculative.cpp b/common/speculative.cpp index 3cb08767bd46..ee94d7c37662 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -2284,7 +2284,7 @@ common_speculative_init_result::common_speculative_init_result( std::string model_path; if (has_draft) { model_path = params.speculative.draft.mparams.path; - LOG_TRC("%s: loading draft model '%s'\n", __func__, model_path.c_str()); + LOG_INF("%s: loading draft model '%s'\n", __func__, model_path.c_str()); llama_model * model_dft = llama_model_load_from_file(params.model.path.c_str(), mparams); if (model_dft == NULL) { @@ -2304,7 +2304,7 @@ common_speculative_init_result::common_speculative_init_result( } else if (spec_mtp) { model_path = params.model.path; - LOG_TRC("%s: creating MTP draft context against the target model '%s'\n", __func__, model_path.c_str()); + LOG_INF("%s: creating MTP draft context against the target model '%s'\n", __func__, model_path.c_str()); llama_context * ctx_dft = llama_init_from_model(model_tgt, cparams); if (ctx_dft == nullptr) { diff --git a/conversion/__init__.py b/conversion/__init__.py index b2bb7e5161eb..45c001b78f96 100644 --- a/conversion/__init__.py +++ b/conversion/__init__.py @@ -167,6 +167,7 @@ "ModernBertForMaskedLM": "bert", "ModernBertForSequenceClassification": "bert", "ModernBertModel": "bert", + "NanbeigeForCausalLM": "nanbeige", "NemotronForCausalLM": "nemotron", "NemotronHForCausalLM": "nemotron", "NeoBERT": "bert", diff --git a/conversion/mimo.py b/conversion/mimo.py index 11ec2867940a..ca2ed28ad391 100644 --- a/conversion/mimo.py +++ b/conversion/mimo.py @@ -1,8 +1,9 @@ from __future__ import annotations +import json import re -from typing import Callable, TYPE_CHECKING +from typing import Any, Callable, Iterable, TYPE_CHECKING import torch @@ -229,7 +230,13 @@ def prepare_tensors(self): @ModelBase.register("MiMoV2ForCausalLM") -class MiMoV2VisionModel(MmprojModel): +class MiMoV2VisionAudioModel(MmprojModel): + has_audio_encoder = True + + _audio_tok_hparams: dict[str, Any] | None = None + _rvq_codebook_sizes: list[int] | None = None + _code_embd: dict[int, Tensor] | None = None + def __init__(self, *args, **kwargs): super().__init__(*args, **kwargs) assert self.hparams_vision is not None @@ -253,10 +260,22 @@ def __init__(self, *args, **kwargs): self.visual_token_window_size = int(hp.get("visual_token_window_size", -1)) self.use_sink = bool(hp.get("use_sink", False)) + def get_audio_config(self) -> dict[str, Any] | None: + if self._audio_tok_hparams is None: + path = self.dir_model / "audio_tokenizer" / "config.json" + with open(path, "r", encoding="utf-8") as f: + cfg = json.load(f) + # aliases so MmprojModel.find_aparam() / n_block_keys can resolve them + cfg["hidden_size"] = cfg["d_model"] + cfg["intermediate_size"] = cfg["encoder_ffn_dim"] + cfg["num_attention_heads"] = cfg["encoder_attention_heads"] + self._audio_tok_hparams = cfg + return self._audio_tok_hparams + def set_gguf_parameters(self): super().set_gguf_parameters() - self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.MIMOVL) + self.gguf_writer.add_clip_vision_projector_type(gguf.VisionProjectorType.MIMOVL) self.gguf_writer.add_vision_use_silu(True) self.gguf_writer.add_vision_head_count_kv(self.num_kv_heads) self.gguf_writer.add_vision_spatial_merge_size(self.spatial_merge_size) @@ -266,19 +285,45 @@ def set_gguf_parameters(self): self.gguf_writer.add_vision_min_pixels(int(self.preprocessor_config["min_pixels"])) self.gguf_writer.add_vision_max_pixels(int(self.preprocessor_config["max_pixels"])) + assert self.hparams_audio is not None + self.gguf_writer.add_clip_audio_projector_type(gguf.VisionProjectorType.MIMO_AUDIO) + self.gguf_writer.add_audio_num_mel_bins(self.hparams_audio["n_mels"]) + self.gguf_writer.add_audio_attention_layernorm_eps(self.hparams_audio.get("layer_norm_eps", 1e-5)) + + assert self._rvq_codebook_sizes is not None + self.gguf_writer.add_audio_rvq_num_quantizers(len(self._rvq_codebook_sizes)) + self.gguf_writer.add_audio_rvq_codebook_size(self._rvq_codebook_sizes) + + n_layer = self.hparams_audio["encoder_layers"] + swa_per_block = self.hparams_audio.get("swa_per_block", 1) + if self.hparams_audio.get("hybrid_attention") and swa_per_block > 1: + wa_pattern = [0 if i % swa_per_block < swa_per_block - 1 else -1 for i in range(n_layer)] + else: + wa_pattern = [-1] * n_layer + self.gguf_writer.add_audio_wa_pattern_mode(wa_pattern) + self.gguf_writer.add_audio_window_size(int(self.hparams_audio["encoder_attn_window_size"][0])) + + audio_cfg = self.global_config["audio_config"] + self.gguf_writer.add_audio_local_block_count(int(audio_cfg["input_local_layers"])) + self.gguf_writer.add_audio_local_group_size(int(audio_cfg["group_size"])) + def tensor_force_quant(self, name, new_name, bid, n_dims): - # Sinks must be F32: any sink-style softmax/mask add in ggml requires - # F32, and we fold sinks into a host-built F32 mask at encode time. - if new_name.endswith(".attn_sinks"): + # for audio encoder: keep codebook in F32 + if new_name in ( + gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.A_ENC_RVQ_CODEBOOK] + ".weight", + gguf.TENSOR_NAMES[gguf.MODEL_TENSOR.A_MM_CODE_EMBD] + ".weight", + ): + return gguf.GGMLQuantizationType.F32 + if ("encoder.conv" in name or "encoder.down_sample_layer" in name) and name.endswith(".weight"): return gguf.GGMLQuantizationType.F32 return super().tensor_force_quant(name, new_name, bid, n_dims) @classmethod def filter_tensors(cls, item: tuple[str, Callable[[], Tensor]]) -> tuple[str, Callable[[], Tensor]] | None: name, _ = item - if not name.startswith("visual."): - return None - return super().filter_tensors(item) + if name.startswith("visual.") or name.startswith("speech_embeddings.") or name.startswith("audio_encoder."): + return super().filter_tensors(item) + return None def modify_tensors(self, data_torch, name, bid): # Conv3D patch embed: split along the temporal axis (kt=2) into two Conv2D @@ -292,4 +337,64 @@ def modify_tensors(self, data_torch, name, bid): yield (embd_name + ".weight.1", data_torch[:, :, 1, ...]) return + if m := re.match(r"^speech_embeddings\.(\d+)\.weight$", name): + if self._code_embd is None: + self._code_embd = {} + self._code_embd[int(m.group(1))] = data_torch + + n_channels = int(self.global_config["audio_config"]["audio_channels"]) + if len(self._code_embd) < n_channels: + return + merged = torch.stack([self._code_embd.pop(i) for i in range(n_channels)], dim=0) + yield (self.format_tensor_name(gguf.MODEL_TENSOR.A_MM_CODE_EMBD), merged) + return + + if "conv1.bias" in name or "conv2.bias" in name: + # transpose conv1/conv2 bias so it broadcasts against [n_frames, C_out, 1] + data_torch = data_torch.unsqueeze(-1) + + if name == "audio_encoder.projection.mlp.0.weight": + yield (self.format_tensor_name(gguf.MODEL_TENSOR.A_MMPROJ, 1), data_torch) + return + if name == "audio_encoder.projection.mlp.2.weight": + yield (self.format_tensor_name(gguf.MODEL_TENSOR.A_MMPROJ, 2), data_torch) + return + yield from super().modify_tensors(data_torch, name, bid) + + def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]: + # note: audio encoder is in its own subdir "audio_tokenizer" + from safetensors.torch import load_file + + tok_dir = self.dir_model / "audio_tokenizer" + state_dict = load_file(tok_dir / "model.safetensors") + + codebook_re = re.compile(r"^encoder\.quantizer\.vq\.layers\.(\d+)\._codebook\.embed$") + codebooks: dict[int, Tensor] = {} + + # EMA/training-only RVQ buffers - not needed for inference (nearest-codebook + # lookup only reads "_codebook.embed") + skip_suffixes = ( + "_codebook.cluster_size", + "_codebook.embed_avg", + "_codebook.inited", + ) + for name, tensor in state_dict.items(): + if name.endswith(skip_suffixes): + continue + if m := codebook_re.match(name): + codebooks[int(m.group(1))] = tensor + continue + yield name, tensor + + # gather codebooks and merge into 3D tensor, similar to MoE MLP tensors + n_q = len(codebooks) + ordered = [codebooks[i] for i in range(n_q)] + self._rvq_codebook_sizes = [int(cb.shape[0]) for cb in ordered] + max_bins = max(self._rvq_codebook_sizes) + dim = ordered[0].shape[1] + merged = ordered[0].new_zeros(n_q, max_bins, dim) + for i, cb in enumerate(ordered): + merged[i, : cb.shape[0], :] = cb + + yield (self.format_tensor_name(gguf.MODEL_TENSOR.A_ENC_RVQ_CODEBOOK), merged) diff --git a/conversion/nanbeige.py b/conversion/nanbeige.py new file mode 100644 index 000000000000..f1fc425b3a09 --- /dev/null +++ b/conversion/nanbeige.py @@ -0,0 +1,24 @@ +from __future__ import annotations + +from .base import ModelBase, gguf, logger +from .llama import LlamaModel + + +@ModelBase.register("NanbeigeForCausalLM") +class NanbeigeModel(LlamaModel): + model_arch = gguf.MODEL_ARCH.NANBEIGE + undo_permute = True + + def set_gguf_parameters(self): + super().set_gguf_parameters() + hparams = self.hparams + + n_loops = int(hparams.get("num_loops", 1) or 1) + if n_loops < 1: + n_loops = 1 + self.gguf_writer.add_num_loops(n_loops) + logger.info(f"gguf: num_loops = {n_loops}") + + skip_loop_final_norm = bool(hparams.get("skip_loop_final_norm", False)) + self.gguf_writer.add_skip_loop_final_norm(skip_loop_final_norm) + logger.info(f"gguf: skip_loop_final_norm = {skip_loop_final_norm}") diff --git a/docs/development/HOWTO-add-model.md b/docs/development/HOWTO-add-model.md index 632e79881a43..102f479eb02c 100644 --- a/docs/development/HOWTO-add-model.md +++ b/docs/development/HOWTO-add-model.md @@ -144,6 +144,8 @@ Examples: - Gemma 3 folds the `1 +` of its `norm(1 + weight)` normalization into the weights at conversion time, so the graph just does a plain RMS norm. - Qwen3-Next applies its tensor permutation during conversion (in `modify_tensors`), so the graph can consume the already-permuted weights directly. +Exception: a plain `weight * scale` with a constant scale is usually better left to inference time rather than folded into the weight at conversion. The scale conceptually applies to the activation, not the weight, so folding it into the weight can hurt numerical stability, and it shifts the weight's value range in a way that can make quantization worse. In this case, write the scale to GGUF as its own metadata key (e.g. `%s.attention.output_scale`, `%s.attention.value_scale`, `%s.embedding_scale`) and apply it in the graph, instead of pre-multiplying the weight tensor during conversion. + ### Working with ggml_rope_ext PyTorch implementations usually prefer explicitly calculating `freq_cis`/`sin`/`cos` components. However, in llama.cpp, most RoPE operations can be handled via `ggml_rope_ext`, which does not require a sin/cos matrix. This saves memory while allowing the GGML RoPE kernel to be fused with other ops. diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 87615921c09b..7f4e252dca39 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -906,26 +906,35 @@ static int ggml_backend_sched_backend_id_from_cur(ggml_backend_sched_t sched, st } // operations with weights are preferably run on the same backend as the weights - for (int i = 0; i < GGML_MAX_SRC; i++) { - const struct ggml_tensor * src = tensor->src[i]; - if (src == NULL) { - continue; - } - // skip ROPE since the rope freqs tensor is too small to choose a backend based on it - // not an ideal solution - if (tensor->op != GGML_OP_ROPE && src->buffer != NULL && src->buffer->usage == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) { - int src_backend_id = ggml_backend_sched_backend_from_buffer(sched, src, tensor); - // check if a backend with higher prio wants to offload the op - if (sched->op_offload && src_backend_id == sched->n_backends - 1 && ggml_backend_buffer_is_host(src->buffer)) { - for (int b = 0; b < src_backend_id; b++) { - if (ggml_backend_supports_op(sched->backends[b], tensor) && ggml_backend_offload_op(sched->backends[b], tensor)) { - SET_CAUSE(tensor, "1.off"); - return b; + // TODO: there are exceptions (see below) - not an ideal solution + bool allow = true; + + // skip ROPE since the rope freqs tensor is too small to choose a backend based on it + allow = allow && tensor->op != GGML_OP_ROPE; + + // skip FLASH_ATTN_EXT since the sinks tensor is too small to choose a based based on it + allow = allow && tensor->op != GGML_OP_FLASH_ATTN_EXT; + + if (allow) { + for (int i = 0; i < GGML_MAX_SRC; i++) { + const struct ggml_tensor * src = tensor->src[i]; + if (src == NULL) { + continue; + } + if (src->buffer != NULL && src->buffer->usage == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) { + int src_backend_id = ggml_backend_sched_backend_from_buffer(sched, src, tensor); + // check if a backend with higher prio wants to offload the op + if (sched->op_offload && src_backend_id == sched->n_backends - 1 && ggml_backend_buffer_is_host(src->buffer)) { + for (int b = 0; b < src_backend_id; b++) { + if (ggml_backend_supports_op(sched->backends[b], tensor) && ggml_backend_offload_op(sched->backends[b], tensor)) { + SET_CAUSE(tensor, "1.off"); + return b; + } } } + SET_CAUSE(tensor, "1.wgt%d", i); + return src_backend_id; } - SET_CAUSE(tensor, "1.wgt%d", i); - return src_backend_id; } } diff --git a/ggml/src/ggml-cpu/llamafile/sgemm.cpp b/ggml/src/ggml-cpu/llamafile/sgemm.cpp index 23bcd54c122a..99b7d5afa2f9 100644 --- a/ggml/src/ggml-cpu/llamafile/sgemm.cpp +++ b/ggml/src/ggml-cpu/llamafile/sgemm.cpp @@ -1797,14 +1797,6 @@ class tinyBLAS_Q0_AVX { //PPC Implementation #if defined(__MMA__) -#define SAVE_ACC(ACC, ii, jj) \ - __builtin_mma_disassemble_acc(vec_C, ACC); \ - for (int I = 0; I < 4; I++) { \ - for (int J = 0; J < 4; J++) { \ - *((float*)(C+ii+((jj+J)*ldc)+I)) = *((float*)&vec_C[I]+J); \ - } \ - } \ - template struct mma_instr; @@ -1834,10 +1826,49 @@ class tinyBLAS_HP16_PPC { } void matmul(int64_t m, int64_t n) { - mnpack(0, m, 0, n); + int64_t mc = 256; + int64_t nc = 256; + int64_t kc = 256; + #if defined(_AIX) || defined(__BIG_ENDIAN__) + mc = 128; + nc = 128; + kc = 128; + #endif + if (k < kc) { + kc = k; + } + bool can_use_tiled = (m % mc == 0) && (n % nc == 0) && (k % kc == 0); + if (can_use_tiled) { + matmul_tiled(m, n, mc, nc, kc); + } else { + mnpack(0, m, 0, n); + } } private: + __attribute__((always_inline)) + inline void save_acc(acc_t * ACC, int64_t ii, int64_t jj) { + vec_t vec_C[4]; + __builtin_mma_disassemble_acc(vec_C, ACC); + for (int I = 0; I < 4; I++) { + for (int J = 0; J < 4; J++) { + *((float *)(C+ii+((jj+J)*ldc)+I)) = *((float *)&vec_C[I]+J); + } + } + } + + __attribute__((always_inline)) + inline void add_save_acc(acc_t * ACC, int64_t ii, int64_t jj) { + vec_t vec_C[4]; + __builtin_mma_disassemble_acc(vec_C, ACC); + for (int I = 0; I < 4; I++) { + for (int J = 0; J < 4; J++) { + float * c_ptr = (float *)(C+ii+((jj+J)*ldc)+I); + *c_ptr += *((float *)&vec_C[I]+J); + } + } + } + void vector_permute_store(vec_t *c, int numVec, unsigned char *vecOffset) { vec_t t[8], s[8]; vec_t swiz1 = {0, 1, 2, 3, 16, 17, 18, 19, 4, 5, 6, 7, 20, 21, 22, 23}; @@ -1896,6 +1927,7 @@ class tinyBLAS_HP16_PPC { j = (rows >> 3); if (j > 0) { do { + aoffsets[0] = aoffset; if (cols == 4) { aoffsets[0] = aoffset; for (int it = 1; it < 4; ++it) @@ -1910,17 +1942,17 @@ class tinyBLAS_HP16_PPC { } i = (cols >> 3); if (i > 0) { - aoffsets[0] = aoffset; for (int it = 1; it < 8; ++it) { aoffsets[it] = aoffsets[it-1] + lda; } aoffset += 8 * lda; + do { for (int it = 0; it < 8; ++it) c_arr[it] = vec_xl(0, (vector unsigned char*)aoffsets[it]); vector_permute_store(c_arr, 8, vecOffset); for (int it = 0; it < 8; ++it) - aoffsets[it] = aoffsets[it] + 8*lda; + aoffsets[it] = aoffsets[it] + 8; vecOffset += 128; i--; } while(i > 0); @@ -2147,8 +2179,8 @@ class tinyBLAS_HP16_PPC { mma_instr::outer_product(&acc_1, vec_A[x], vec_B[x+4]); } } - SAVE_ACC(&acc_0, ii, jj); - SAVE_ACC(&acc_1, ii, jj+4); + save_acc(&acc_0, ii, jj); + save_acc(&acc_1, ii, jj+4); } void KERNEL_8x4(int64_t ii, int64_t jj) { @@ -2164,8 +2196,8 @@ class tinyBLAS_HP16_PPC { mma_instr::outer_product(&acc_1, vec_A[x+4], vec_B[x]); } } - SAVE_ACC(&acc_0, ii, jj); - SAVE_ACC(&acc_1, ii+4, jj); + save_acc(&acc_0, ii, jj); + save_acc(&acc_1, ii+4, jj); } @@ -2186,13 +2218,64 @@ class tinyBLAS_HP16_PPC { mma_instr::outer_product(&acc_3, vec_A[x+4], vec_B[x+4]); } } - - SAVE_ACC(&acc_0, ii, jj); - SAVE_ACC(&acc_1, ii, jj+4); - SAVE_ACC(&acc_2, ii+4, jj); - SAVE_ACC(&acc_3, ii+4, jj+4); + save_acc(&acc_0, ii, jj); + save_acc(&acc_1, ii, jj+4); + save_acc(&acc_2, ii+4, jj); + save_acc(&acc_3, ii+4, jj+4); } + inline void MMA_16x8(vec_t * vec_A0, vec_t * vec_A1, vec_t * vec_B, acc_t * acc) { + for (int x = 0; x < 4; x ++) { + mma_instr::outer_product(&acc[0], vec_A0[x], vec_B[x]); + mma_instr::outer_product(&acc[1], vec_A0[x], vec_B[x+4]); + mma_instr::outer_product(&acc[2], vec_A0[x+4], vec_B[x]); + mma_instr::outer_product(&acc[3], vec_A0[x+4], vec_B[x+4]); + mma_instr::outer_product(&acc[4], vec_A1[x], vec_B[x]); + mma_instr::outer_product(&acc[5], vec_A1[x], vec_B[x+4]); + mma_instr::outer_product(&acc[6], vec_A1[x+4], vec_B[x]); + mma_instr::outer_product(&acc[7], vec_A1[x+4], vec_B[x+4]); + } + } + void KERNEL(int64_t ii, int64_t jj, int64_t mc, int64_t nc, int64_t kc, vec_t * vec_A, vec_t * vec_B, int64_t kk) { + for (int64_t i = 0; i < mc; i += 16) { + int A_base_addr = (mc / 8) * (i / 8) * 8; + for (int64_t j = 0; j < nc; j += 8) { + int B_base_addr = (nc / 8) * (j / 8) * 8; + acc_t acc[8]; + vec_t A0_block[8]; vec_t A1_block[8]; + for (int x = 0; x < 8; x++) + __builtin_mma_xxsetaccz(&acc[x]); + for (int64_t l = 0; l < kc; l += 8) { + int A0_block_idx = A_base_addr + (l / 8) * 8; + int A1_block_idx = A0_block_idx + (mc / 8) * 8; + int B_block_idx = B_base_addr + (l / 8) * 8; + vec_t* A0_block = &vec_A[A0_block_idx]; + vec_t* A1_block = &vec_A[A1_block_idx]; + vec_t* B_block = &vec_B[B_block_idx]; + MMA_16x8(A0_block, A1_block, B_block, acc); + } + if (kk == 0) { + save_acc(&acc[0], ii + i, jj + j); + save_acc(&acc[1], ii + i, jj + j + 4); + save_acc(&acc[2], ii + i + 4, jj + j); + save_acc(&acc[3], ii + i + 4, jj + j + 4); + save_acc(&acc[4], ii + i + 8, jj + j); + save_acc(&acc[5], ii + i + 8, jj + j + 4); + save_acc(&acc[6], ii + i + 12, jj + j); + save_acc(&acc[7], ii + i + 12, jj + j + 4); + } else { + add_save_acc(&acc[0], ii + i, jj + j); + add_save_acc(&acc[1], ii + i, jj + j + 4); + add_save_acc(&acc[2], ii + i + 4, jj + j); + add_save_acc(&acc[3], ii + i + 4, jj + j + 4); + add_save_acc(&acc[4], ii + i + 8, jj + j); + add_save_acc(&acc[5], ii + i + 8, jj + j + 4); + add_save_acc(&acc[6], ii + i + 12, jj + j); + add_save_acc(&acc[7], ii + i + 12, jj + j + 4); + } + } + } + } template void gemm_small(int64_t m0, int64_t m, int64_t n0, int64_t n) { int64_t ytiles = (m - m0) / RM; @@ -2281,6 +2364,29 @@ class tinyBLAS_HP16_PPC { } } + void matmul_tiled(int64_t m, int64_t n, int64_t mc, int64_t nc, int64_t kc) { + int64_t ytiles = m / mc; + int64_t xtiles = n / nc; + int64_t tiles = xtiles * ytiles; + int64_t duty = (tiles + nth - 1) / nth; + int64_t start = duty * ith; + int64_t end = start + duty; + if (end > tiles) { + end = tiles; + } + for (int64_t job = start; job < end; ++job) { + int64_t ii = (job / xtiles) * mc; + int64_t jj = (job % xtiles) * nc; + for (int64_t kk = 0; kk < k; kk += kc) { + vec_t A_pack[kc * mc / 8]; + vec_t B_pack[kc * nc / 8]; + packNormal(A + (ii * lda) + kk, lda, kc, mc, (uint8_t *)A_pack); + packNormal(B + (jj * ldb) + kk, ldb, kc, nc, (uint8_t *)B_pack); + KERNEL(ii, jj, mc, nc, kc, A_pack, B_pack, kk); + } + } + } + template NOINLINE void gemm(int64_t m0, int64_t m, int64_t n0, int64_t n) { int64_t ytiles = (m - m0) / RM; diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt index 5351dcae12db..bbc51797c182 100644 --- a/ggml/src/ggml-hip/CMakeLists.txt +++ b/ggml/src/ggml-hip/CMakeLists.txt @@ -154,5 +154,3 @@ if (GGML_HIP_RCCL) endif() target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas) - -target_compile_options(ggml-hip PRIVATE "$<$:-ffast-math;-fno-finite-math-only>") diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index 1c17d20df12b..a8d9c0d804bf 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -199,9 +199,20 @@ if (GGML_SYCL_DEVICE_ARCH) -fsycl-targets=spir64_gen "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\"" ) + + # Pass through parallel job (process) count for parallelising the + # `llvm-foreach -- ocloc` invocation for compiling AOT device images. + include(ProcessorCount) + ProcessorCount(_ggml_sycl_nproc) + if (_ggml_sycl_nproc LESS 1) + set(_ggml_sycl_nproc 1) + endif() + set(GGML_SYCL_MAX_PARALLEL_LINK_JOBS ${_ggml_sycl_nproc} CACHE STRING + "Parallel ocloc jobs for spir64_gen AOT device-image lowering") target_link_options( ggml-sycl PRIVATE -fsycl-targets=spir64_gen "SHELL:-Xsycl-target-backend=spir64_gen \"-device ${GGML_SYCL_DEVICE_ARCH}\"" + -fsycl-max-parallel-link-jobs=${GGML_SYCL_MAX_PARALLEL_LINK_JOBS} ) endif() diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 2071e3eaa8a4..f9264425f6e1 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -145,6 +145,8 @@ class LLM: TOKEN_SHIFT_COUNT = "{arch}.token_shift_count" INTERLEAVE_MOE_LAYER_STEP = "{arch}.interleave_moe_layer_step" FULL_ATTENTION_INTERVAL = "{arch}.full_attention_interval" + NUM_LOOPS = "{arch}.num_loops" + SKIP_LOOP_FINAL_NORM = "{arch}.skip_loop_final_norm" HASH_LAYER_COUNT = "{arch}.hash_layer_count" ACTIVATION_SPARSITY_SCALE = "{arch}.activation_sparsity_scale" ALTUP_ACTIVE_IDX = "{arch}.altup.active_idx" @@ -374,6 +376,12 @@ class ClipAudio: CONV_KERNEL_SIZE = "clip.audio.conv_kernel_size" MAX_POS_EMB = "clip.audio.max_pos_emb" FEATURE_LAYERS = "clip.audio.feature_layer" # Granite Speech Plus + RVQ_NUM_QUANTIZERS = "clip.audio.rvq.num_quantizers" + RVQ_CODEBOOK_SIZE = "clip.audio.rvq.codebook_size" + WA_PATTERN_MODE = "clip.audio.wa_pattern_mode" # per-layer -1 (full) / 0 (windowed) + WINDOW_SIZE = "clip.audio.window_size" + LOCAL_BLOCK_COUNT = "clip.audio.local_block_count" # mimo-v2.5: input_local_transformer layer count + LOCAL_GROUP_SIZE = "clip.audio.local_group_size" # mimo-v2.5: input_local_transformer grouping size class Attention: HEAD_COUNT = "clip.audio.attention.head_count" @@ -545,6 +553,7 @@ class MODEL_ARCH(IntEnum): KIMI_LINEAR = auto() TALKIE = auto() MELLUM = auto() + NANBEIGE = auto() class VISION_PROJECTOR_TYPE(IntEnum): @@ -942,6 +951,9 @@ class MODEL_TENSOR(IntEnum): A_ENC_FFN_SCALE_1 = auto() # gemma3n A_ENC_FFN_GATE_1 = auto() # lfm2, gemma3n A_ENC_FFN_DOWN_1 = auto() # lfm2, gemma3n + A_ENC_DOWNSAMPLE_CONV = auto() # mimo-audio-tokenizer: post-transformer downsample conv + A_ENC_DOWNSAMPLE_NORM = auto() # mimo-audio-tokenizer: post-transformer downsample norm + A_ENC_RVQ_CODEBOOK = auto() # mimo-audio-tokenizer: residual vector quantizer codebook, per quantizer index A_MMPROJ = auto() A_MMPROJ_FC = auto() A_MM_NORM_PRE = auto() @@ -950,6 +962,17 @@ class MODEL_TENSOR(IntEnum): A_MM_HARD_EMB_NORM = auto() # gemma3n A_MM_SOFT_EMB_NORM = auto() # gemma3n A_MM_INP_PROJ = auto() # gemma3n + A_MM_CODE_EMBD = auto() # mimo: text-side RVQ code embedding table ("text codebook"), merged 3D [n_channels, vocab, dim] + A_MM_LOCAL_ATTN_Q = auto() # mimo: input_local_transformer (LLM-side connector) + A_MM_LOCAL_ATTN_K = auto() + A_MM_LOCAL_ATTN_V = auto() + A_MM_LOCAL_ATTN_OUT = auto() + A_MM_LOCAL_FFN_GATE = auto() + A_MM_LOCAL_FFN_UP = auto() + A_MM_LOCAL_FFN_DOWN = auto() + A_MM_LOCAL_LN1 = auto() + A_MM_LOCAL_LN2 = auto() + A_MM_LOCAL_NORM = auto() # final norm after all input_local_transformer layers A_PER_DIM_K_SCALE = auto() # gemma4 A_PER_DIM_SCALE = auto() # gemma4 # nextn/mtp @@ -1134,6 +1157,7 @@ class MODEL_TENSOR(IntEnum): MODEL_ARCH.KIMI_LINEAR: "kimi-linear", MODEL_ARCH.TALKIE: "talkie", MODEL_ARCH.MELLUM: "mellum", + MODEL_ARCH.NANBEIGE: "nanbeige", } VISION_PROJECTOR_TYPE_NAMES: dict[VISION_PROJECTOR_TYPE, str] = { @@ -1528,6 +1552,9 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.A_ENC_FFN_UP_1: "a.blk.{bid}.ffn_up_1", MODEL_TENSOR.A_ENC_FFN_GATE_1: "a.blk.{bid}.ffn_gate_1", MODEL_TENSOR.A_ENC_FFN_DOWN_1: "a.blk.{bid}.ffn_down_1", + MODEL_TENSOR.A_ENC_DOWNSAMPLE_CONV: "a.downsample.conv", + MODEL_TENSOR.A_ENC_DOWNSAMPLE_NORM: "a.downsample.norm", + MODEL_TENSOR.A_ENC_RVQ_CODEBOOK: "a.rvq.codebook", MODEL_TENSOR.A_MMPROJ: "mm.a.mlp.{bid}", MODEL_TENSOR.A_MMPROJ_FC: "mm.a.fc", MODEL_TENSOR.A_MM_NORM_PRE: "mm.a.norm_pre", @@ -1536,6 +1563,17 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.A_MM_SOFT_EMB_NORM: "mm.a.soft_emb_norm", # gemma3n MODEL_TENSOR.A_MM_EMBEDDING: "mm.a.embedding", # gemma3n MODEL_TENSOR.A_MM_HARD_EMB_NORM: "mm.a.hard_emb_norm", # gemma3n + MODEL_TENSOR.A_MM_CODE_EMBD: "mm.a.code_embd", + MODEL_TENSOR.A_MM_LOCAL_ATTN_Q: "mm.a.local_blk.{bid}.attn_q", + MODEL_TENSOR.A_MM_LOCAL_ATTN_K: "mm.a.local_blk.{bid}.attn_k", + MODEL_TENSOR.A_MM_LOCAL_ATTN_V: "mm.a.local_blk.{bid}.attn_v", + MODEL_TENSOR.A_MM_LOCAL_ATTN_OUT: "mm.a.local_blk.{bid}.attn_out", + MODEL_TENSOR.A_MM_LOCAL_FFN_GATE: "mm.a.local_blk.{bid}.ffn_gate", + MODEL_TENSOR.A_MM_LOCAL_FFN_UP: "mm.a.local_blk.{bid}.ffn_up", + MODEL_TENSOR.A_MM_LOCAL_FFN_DOWN: "mm.a.local_blk.{bid}.ffn_down", + MODEL_TENSOR.A_MM_LOCAL_LN1: "mm.a.local_blk.{bid}.ln1", + MODEL_TENSOR.A_MM_LOCAL_LN2: "mm.a.local_blk.{bid}.ln2", + MODEL_TENSOR.A_MM_LOCAL_NORM: "mm.a.local_norm", MODEL_TENSOR.A_PER_DIM_K_SCALE: "a.blk.{bid}.per_dim_k_scale", # gemma4 MODEL_TENSOR.A_PER_DIM_SCALE: "a.blk.{bid}.per_dim_scale", # gemma4 # lfm2 audio @@ -1737,10 +1775,24 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.A_ENC_FFN_UP_1, MODEL_TENSOR.A_ENC_FFN_GATE_1, MODEL_TENSOR.A_ENC_FFN_DOWN_1, + MODEL_TENSOR.A_ENC_DOWNSAMPLE_CONV, + MODEL_TENSOR.A_ENC_DOWNSAMPLE_NORM, + MODEL_TENSOR.A_ENC_RVQ_CODEBOOK, MODEL_TENSOR.A_MMPROJ, MODEL_TENSOR.A_MMPROJ_FC, MODEL_TENSOR.A_MM_NORM_PRE, MODEL_TENSOR.A_MM_NORM_MID, + MODEL_TENSOR.A_MM_CODE_EMBD, + MODEL_TENSOR.A_MM_LOCAL_ATTN_Q, + MODEL_TENSOR.A_MM_LOCAL_ATTN_K, + MODEL_TENSOR.A_MM_LOCAL_ATTN_V, + MODEL_TENSOR.A_MM_LOCAL_ATTN_OUT, + MODEL_TENSOR.A_MM_LOCAL_FFN_GATE, + MODEL_TENSOR.A_MM_LOCAL_FFN_UP, + MODEL_TENSOR.A_MM_LOCAL_FFN_DOWN, + MODEL_TENSOR.A_MM_LOCAL_LN1, + MODEL_TENSOR.A_MM_LOCAL_LN2, + MODEL_TENSOR.A_MM_LOCAL_NORM, MODEL_TENSOR.A_ENC_NORM_CONV, MODEL_TENSOR.A_ENC_LINEAR_POS, MODEL_TENSOR.A_ENC_POS_BIAS_U, @@ -4505,7 +4557,22 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.FFN_DOWN_EXP, MODEL_TENSOR.FFN_UP_EXP, ], - # TODO + MODEL_ARCH.NANBEIGE: [ + MODEL_TENSOR.TOKEN_EMBD, + MODEL_TENSOR.OUTPUT_NORM, + MODEL_TENSOR.OUTPUT, + MODEL_TENSOR.ROPE_FREQS, + MODEL_TENSOR.ATTN_NORM, + MODEL_TENSOR.ATTN_Q, + MODEL_TENSOR.ATTN_K, + MODEL_TENSOR.ATTN_V, + MODEL_TENSOR.ATTN_OUT, + MODEL_TENSOR.ATTN_ROT_EMBD, + MODEL_TENSOR.FFN_NORM, + MODEL_TENSOR.FFN_GATE, + MODEL_TENSOR.FFN_DOWN, + MODEL_TENSOR.FFN_UP, + ], } # tensors that will not be serialized @@ -4572,6 +4639,10 @@ class MODEL_TENSOR(IntEnum): MODEL_TENSOR.ROPE_FREQS, MODEL_TENSOR.ATTN_ROT_EMBD, ], + MODEL_ARCH.NANBEIGE: [ + MODEL_TENSOR.ROPE_FREQS, + MODEL_TENSOR.ATTN_ROT_EMBD, + ], } # @@ -4781,6 +4852,7 @@ class VisionProjectorType: MINICPMV4_6 = "minicpmv4_6" GRANITE_SPEECH = "granite_speech" # audio MIMOVL = "mimovl" + MIMO_AUDIO = "mimo_audio" GRANITE4_VISION = "granite4_vision" diff --git a/gguf-py/gguf/gguf_writer.py b/gguf-py/gguf/gguf_writer.py index ba08f8d65004..bd8629aa119e 100644 --- a/gguf-py/gguf/gguf_writer.py +++ b/gguf-py/gguf/gguf_writer.py @@ -908,6 +908,12 @@ def add_wkv_head_size(self, size: int) -> None: def add_token_shift_count(self, count: int) -> None: self.add_uint32(Keys.LLM.TOKEN_SHIFT_COUNT.format(arch=self.arch), count) + def add_num_loops(self, count: int) -> None: + self.add_uint32(Keys.LLM.NUM_LOOPS.format(arch=self.arch), count) + + def add_skip_loop_final_norm(self, value: bool) -> None: + self.add_bool(Keys.LLM.SKIP_LOOP_FINAL_NORM.format(arch=self.arch), value) + def add_interleave_moe_layer_step(self, value: int) -> None: self.add_uint32(Keys.LLM.INTERLEAVE_MOE_LAYER_STEP.format(arch=self.arch), value) @@ -1344,6 +1350,24 @@ def add_audio_attention_layernorm_eps(self, value: float) -> None: def add_audio_num_mel_bins(self, value: int) -> None: self.add_uint32(Keys.ClipAudio.NUM_MEL_BINS, value) + def add_audio_rvq_num_quantizers(self, value: int) -> None: + self.add_uint32(Keys.ClipAudio.RVQ_NUM_QUANTIZERS, value) + + def add_audio_rvq_codebook_size(self, values: Sequence[int]) -> None: + self.add_array(Keys.ClipAudio.RVQ_CODEBOOK_SIZE, values) + + def add_audio_wa_pattern_mode(self, modes: Sequence[int]) -> None: + self.add_array(Keys.ClipAudio.WA_PATTERN_MODE, modes) + + def add_audio_window_size(self, value: int) -> None: + self.add_uint32(Keys.ClipAudio.WINDOW_SIZE, value) + + def add_audio_local_block_count(self, value: int) -> None: + self.add_uint32(Keys.ClipAudio.LOCAL_BLOCK_COUNT, value) + + def add_audio_local_group_size(self, value: int) -> None: + self.add_uint32(Keys.ClipAudio.LOCAL_GROUP_SIZE, value) + def add_audio_stack_factor(self, value: int) -> None: self.add_uint32(Keys.ClipAudio.Projector.STACK_FACTOR, value) diff --git a/gguf-py/gguf/tensor_mapping.py b/gguf-py/gguf/tensor_mapping.py index 62d7a827e35c..8299ac25b432 100644 --- a/gguf-py/gguf/tensor_mapping.py +++ b/gguf-py/gguf/tensor_mapping.py @@ -2095,6 +2095,7 @@ class TensorNameMap: "conformer.pre_encode.conv.{bid}", # lfm2 "model.audio_tower.subsample_conv_projection.conv_{bid}.conv", # gemma3n "conformer.subsample_conv_projection.layer{bid}.conv", # gemma4 + "encoder.conv{bid}", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_CONV1D_NORM: ( @@ -2119,6 +2120,7 @@ class TensorNameMap: MODEL_TENSOR.A_POST_NORM: ( "audio_tower.layer_norm", # ultravox "audio_tower.ln_post", # qwen2omni + "encoder.layer_norm", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_ATTN_Q: ( @@ -2127,6 +2129,7 @@ class TensorNameMap: "conformer.layers.{bid}.attention.attn.q_proj", # gemma3n "conformer.layers.{bid}.self_attn.q_proj", # gemma4 "encoder.layers.{bid}.attn.to_q", # granite_speech + "encoder.layers.{bid}.self_attn.q_proj", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_ATTN_K: ( @@ -2135,6 +2138,7 @@ class TensorNameMap: "conformer.layers.{bid}.attention.attn.k_proj", # gemma3n "conformer.layers.{bid}.self_attn.k_proj", # gemma4 "encoder.layers.{bid}.attn.to_k", # granite_speech (split from to_kv) + "encoder.layers.{bid}.self_attn.k_proj", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_ATTN_V: ( @@ -2143,6 +2147,7 @@ class TensorNameMap: "conformer.layers.{bid}.attention.attn.v_proj", # gemma3n "conformer.layers.{bid}.self_attn.v_proj", # gemma4 "encoder.layers.{bid}.attn.to_v", # granite_speech (split from to_kv) + "encoder.layers.{bid}.self_attn.v_proj", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_ATTN_K_REL: ( @@ -2171,6 +2176,7 @@ class TensorNameMap: "conformer.layers.{bid}.norm_self_att", # lfm2 "conformer.layers.{bid}.attention.pre_attn_norm", # gemma3n "encoder.layers.{bid}.attn.pre_norm", # granite_speech + "encoder.layers.{bid}.self_attn_layer_norm", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_OUTPUT: ( @@ -2179,6 +2185,7 @@ class TensorNameMap: "conformer.layers.{bid}.attention.post", # gemma3n "conformer.layers.{bid}.self_attn.post", # gemma4 "encoder.layers.{bid}.attn.to_out", # granite_speech + "encoder.layers.{bid}.self_attn.out_proj", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_OUTPUT_NORM: ( @@ -2186,6 +2193,7 @@ class TensorNameMap: "conformer.layers.{bid}.norm_out", # lfm2 "conformer.layers.{bid}.attention.post_norm", # gemma3n "encoder.layers.{bid}.post_norm", # granite_speech + "encoder.layers.{bid}.final_layer_norm", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_FFN_NORM: ( @@ -2210,6 +2218,7 @@ class TensorNameMap: "conformer.layers.{bid}.ffw_layer_start.ffw_layer_1", # gemma3n "conformer.layers.{bid}.feed_forward1.ffw_layer_1", # gemma4 "encoder.layers.{bid}.ff1.up_proj", # granite_speech + "encoder.layers.{bid}.fc1", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_FFN_GATE: (), @@ -2220,6 +2229,7 @@ class TensorNameMap: "conformer.layers.{bid}.ffw_layer_start.ffw_layer_2", # gemma3n "conformer.layers.{bid}.feed_forward1.ffw_layer_2", # gemma4 "encoder.layers.{bid}.ff1.down_proj", # granite_speech + "encoder.layers.{bid}.fc2", # mimo-audio-tokenizer ), MODEL_TENSOR.A_ENC_FFN_UP_1: ( @@ -2243,6 +2253,19 @@ class TensorNameMap: "encoder.layers.{bid}.ff2.pre_norm", # granite_speech ), + MODEL_TENSOR.A_ENC_DOWNSAMPLE_CONV: ( + "encoder.down_sample_layer.0", # mimo-audio-tokenizer + ), + + MODEL_TENSOR.A_ENC_DOWNSAMPLE_NORM: ( + "encoder.down_sample_norm", # mimo-audio-tokenizer + ), + + # note: the raw per-quantizer "encoder.quantizer.vq.layers.{i}._codebook.embed" + # tensors are merged (padded + stacked, like MoE experts) into this single 3D + # tensor in conversion code, so no raw-name mapping is registered here. + MODEL_TENSOR.A_ENC_RVQ_CODEBOOK: (), + MODEL_TENSOR.A_ENC_FFN_POST_NORM_1: ( "conformer.layers.{bid}.ffw_layer_end.post_layer_norm", # gemma3n "conformer.layers.{bid}.feed_forward2.post_layer_norm", # gemma4 @@ -2294,6 +2317,42 @@ class TensorNameMap: "audio.multi_modal_projector.ln_mid", # ultravox ), + # note: the raw per-channel "speech_embeddings.{i}" tensors are merged + # (stacked, like MoE experts) into this single 3D tensor in conversion + # code, so no raw-name mapping is registered here. + MODEL_TENSOR.A_MM_CODE_EMBD: (), + + MODEL_TENSOR.A_MM_LOCAL_ATTN_Q: ( + "audio_encoder.input_local_transformer.layers.{bid}.self_attn.q_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_ATTN_K: ( + "audio_encoder.input_local_transformer.layers.{bid}.self_attn.k_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_ATTN_V: ( + "audio_encoder.input_local_transformer.layers.{bid}.self_attn.v_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_ATTN_OUT: ( + "audio_encoder.input_local_transformer.layers.{bid}.self_attn.o_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_FFN_GATE: ( + "audio_encoder.input_local_transformer.layers.{bid}.mlp.gate_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_FFN_UP: ( + "audio_encoder.input_local_transformer.layers.{bid}.mlp.up_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_FFN_DOWN: ( + "audio_encoder.input_local_transformer.layers.{bid}.mlp.down_proj", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_LN1: ( + "audio_encoder.input_local_transformer.layers.{bid}.input_layernorm", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_LN2: ( + "audio_encoder.input_local_transformer.layers.{bid}.post_attention_layernorm", # mimo-v2.5 + ), + MODEL_TENSOR.A_MM_LOCAL_NORM: ( + "audio_encoder.input_local_transformer.norm", # mimo-v2.5 + ), + MODEL_TENSOR.A_ENC_CONV_DW: ( "conformer.layers.{bid}.conv.depthwise_conv", # lfm2 "conformer.layers.{bid}.lconv1d.depthwise_conv1d", # gemma3n diff --git a/include/llama.h b/include/llama.h index 9fab69317006..3c6d22be8999 100644 --- a/include/llama.h +++ b/include/llama.h @@ -203,10 +203,11 @@ extern "C" { }; enum llama_load_mode { - LLAMA_LOAD_MODE_NONE = 0, // no special loading mode - LLAMA_LOAD_MODE_MMAP = 1, // memory map the model - LLAMA_LOAD_MODE_MLOCK = 2, // mmap + force system to keep model in RAM rather than swapping or compressing - LLAMA_LOAD_MODE_DIRECT_IO = 3, // use direct I/O if available + LLAMA_LOAD_MODE_NONE = 0, // no special loading mode + LLAMA_LOAD_MODE_MMAP = 1, // memory map the model + LLAMA_LOAD_MODE_MLOCK = 2, // force system to keep model in RAM rather than swapping or compressing + LLAMA_LOAD_MODE_MMAP_MLOCK = 3, // mmap + force system to keep model in RAM rather than swapping or compressing + LLAMA_LOAD_MODE_DIRECT_IO = 4, // use direct I/O if available }; LLAMA_API const char * llama_load_mode_name(enum llama_load_mode load_mode); diff --git a/skills/add-new-model/SKILL.md b/skills/add-new-model/SKILL.md index 68be866c7b8d..f76d1abfd768 100644 --- a/skills/add-new-model/SKILL.md +++ b/skills/add-new-model/SKILL.md @@ -76,6 +76,7 @@ These recur often enough in review comments on past add-model PRs that they're w - Don't ship unfinished or unverified speculative-decoding (e.g. MTP) scaffolding in the base model PR - if it hasn't actually been confirmed to work, pull it out and land it as its own follow-up. - Conversion code should call into the base class's existing hparam logic (e.g. `super().set_gguf_parameters()`) rather than re-deriving it - large blocks of code that duplicate what `TextModel`/`MmprojModel` already provide will get flagged as redundant. - Do constant tensor modifications (e.g. `norm(1 + weight)`) and permutations/chunking at conversion time, not in the graph - see HOWTO-add-model.md's "Prefer conversion-time tensor modifications" tip (Gemma 3 folds its `1 +` into the weights, Qwen3-Next permutes in `modify_tensors`). Doing these at runtime in the graph is very likely to be rejected as over-complicated; if you genuinely can't do it at conversion time, open a discussion first explaining why rather than implementing it in the graph. + - Exception: a plain `weight * scale` with a constant scale is usually better applied at inference time instead of being folded into the weight at conversion. The scale conceptually applies to the activation, not the weight, so folding it in can hurt numerical stability, and it shifts the weight's value range in a way that can make quantization worse. ## Validation checklist diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index 39bf2c79590b..c01706785070 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -143,6 +143,7 @@ static const std::map LLM_ARCH_NAMES = { { LLM_ARCH_KIMI_LINEAR, "kimi-linear" }, { LLM_ARCH_TALKIE, "talkie" }, { LLM_ARCH_MELLUM, "mellum" }, + { LLM_ARCH_NANBEIGE, "nanbeige" }, { LLM_ARCH_UNKNOWN, "(unknown)" }, }; @@ -221,6 +222,8 @@ static const std::map LLM_KV_NAMES = { { LLM_KV_TOKEN_SHIFT_COUNT, "%s.token_shift_count" }, { LLM_KV_INTERLEAVE_MOE_LAYER_STEP, "%s.interleave_moe_layer_step" }, { LLM_KV_FULL_ATTENTION_INTERVAL, "%s.full_attention_interval" }, + { LLM_KV_NUM_LOOPS, "%s.num_loops" }, + { LLM_KV_SKIP_LOOP_FINAL_NORM, "%s.skip_loop_final_norm" }, { LLM_KV_ATTENTION_HEAD_COUNT, "%s.attention.head_count" }, { LLM_KV_ATTENTION_HEAD_COUNT_KV, "%s.attention.head_count_kv" }, diff --git a/src/llama-arch.h b/src/llama-arch.h index 2e3916a0beee..1c9aebb0bbdc 100644 --- a/src/llama-arch.h +++ b/src/llama-arch.h @@ -148,6 +148,7 @@ enum llm_arch { LLM_ARCH_EAGLE3, LLM_ARCH_MINIMAX_M3, LLM_ARCH_DFLASH, + LLM_ARCH_NANBEIGE, LLM_ARCH_UNKNOWN, }; @@ -226,6 +227,8 @@ enum llm_kv { LLM_KV_TOKEN_SHIFT_COUNT, LLM_KV_INTERLEAVE_MOE_LAYER_STEP, LLM_KV_FULL_ATTENTION_INTERVAL, + LLM_KV_NUM_LOOPS, + LLM_KV_SKIP_LOOP_FINAL_NORM, LLM_KV_ATTENTION_HEAD_COUNT, LLM_KV_ATTENTION_HEAD_COUNT_KV, diff --git a/src/llama-context.cpp b/src/llama-context.cpp index c512477c0eab..9b399d6096b1 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -2339,6 +2339,7 @@ uint32_t llama_context::graph_max_nodes(uint32_t n_tokens) const { model.arch == LLM_ARCH_QWEN35 || model.arch == LLM_ARCH_QWEN35MOE || model.arch == LLM_ARCH_DEEPSEEK4 || + model.arch == LLM_ARCH_NANBEIGE || model.arch == LLM_ARCH_MINIMAX_M3) { return std::max(n_tokens * 40, 32u * model.n_tensors()); } @@ -2473,11 +2474,12 @@ llm_graph_cb llama_context::graph_get_cb() const { ggml_set_name(cur, name); } - // norm may be automatically assigned to the backend of the previous layer, increasing data transfer between backends + // - norm may be automatically assigned to the backend of the previous layer, increasing data transfer between backends + // - force the last op of the layer on the specified backend to avoid running it on the backend of the next layer due to scheduling // FIXME: fix in ggml_backend_sched const bool full_offload = model.n_gpu_layers() > model.hparams.n_layer_all; if (ubatch.n_tokens < 32 || full_offload) { - if (il != -1 && strcmp(name, "norm") == 0) { + if (il != -1 && (strcmp(name, "norm") == 0 || strcmp(name, "l_last") == 0)) { const auto & dev_layer = model.dev_layer(il); for (const auto & backend : backends) { if (ggml_backend_get_device(backend.get()) == dev_layer) { diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 43447f57d30b..510586e96c20 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -542,7 +542,7 @@ llama_model_loader::llama_model_loader( tensor_buft_overrides = param_tensor_buft_overrides_p; - this->use_mmap = load_mode == LLAMA_LOAD_MODE_MMAP || load_mode == LLAMA_LOAD_MODE_MLOCK; + this->use_mmap = load_mode == LLAMA_LOAD_MODE_MMAP || load_mode == LLAMA_LOAD_MODE_MMAP_MLOCK; this->use_direct_io = load_mode == LLAMA_LOAD_MODE_DIRECT_IO; if (!fname.empty()) { diff --git a/src/llama-model.cpp b/src/llama-model.cpp index 51796921081f..be0a0df55d62 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp @@ -85,6 +85,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params return new llama_model_stablelm(params); case LLM_ARCH_MELLUM: return new llama_model_mellum(params); + case LLM_ARCH_NANBEIGE: + return new llama_model_nanbeige(params); case LLM_ARCH_QWEN: return new llama_model_qwen(params); case LLM_ARCH_QWEN2: @@ -1249,7 +1251,7 @@ void llama_model_base::load_vocab(llama_model_loader & ml) { bool llama_model_base::load_tensors(llama_model_loader & ml) { const auto & split_mode = params.split_mode; - const bool use_mlock = params.load_mode == LLAMA_LOAD_MODE_MLOCK; + const bool use_mlock = params.load_mode == LLAMA_LOAD_MODE_MLOCK || params.load_mode == LLAMA_LOAD_MODE_MMAP_MLOCK; const auto & tensor_split = params.tensor_split; const int n_layer_all = hparams.n_layer_all; @@ -2491,6 +2493,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) { case LLM_ARCH_LLAMA_EMBED: case LLM_ARCH_MAINCODER: case LLM_ARCH_GLM_DSA: + case LLM_ARCH_NANBEIGE: return LLAMA_ROPE_TYPE_NORM; // the pairs of head values are offset by n_rot/2 diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 7c0bac07d096..92ebc11b99f3 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -359,6 +359,10 @@ static bool tensor_allows_quantization(const llama_model_quantize_params * param quantize &= name.find(".patch_embd") == std::string::npos; quantize &= name.find(".patch_merger") == std::string::npos; + // audio codebook + quantize &= name.find("a.rvq.codebook") == std::string::npos; + quantize &= name.find("mm.a.code_embd") == std::string::npos; + return quantize; } diff --git a/src/llama.cpp b/src/llama.cpp index 11ac9656d9f9..d22e4c81a9e3 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -54,6 +54,8 @@ const char * llama_load_mode_name(enum llama_load_mode load_mode) { return "mmap"; case LLAMA_LOAD_MODE_MLOCK: return "mlock"; + case LLAMA_LOAD_MODE_MMAP_MLOCK: + return "mmap+mlock"; case LLAMA_LOAD_MODE_DIRECT_IO: return "dio"; } @@ -61,10 +63,11 @@ const char * llama_load_mode_name(enum llama_load_mode load_mode) { } enum llama_load_mode llama_load_mode_from_str(const char * str) { - if (std::strcmp(str, "none") == 0) { return LLAMA_LOAD_MODE_NONE; } - if (std::strcmp(str, "mmap") == 0) { return LLAMA_LOAD_MODE_MMAP; } - if (std::strcmp(str, "mlock") == 0) { return LLAMA_LOAD_MODE_MLOCK; } - if (std::strcmp(str, "dio") == 0) { return LLAMA_LOAD_MODE_DIRECT_IO; } + if (std::strcmp(str, "none") == 0) { return LLAMA_LOAD_MODE_NONE; } + if (std::strcmp(str, "mmap") == 0) { return LLAMA_LOAD_MODE_MMAP; } + if (std::strcmp(str, "mlock") == 0) { return LLAMA_LOAD_MODE_MLOCK; } + if (std::strcmp(str, "mmap+mlock") == 0) { return LLAMA_LOAD_MODE_MMAP_MLOCK; } + if (std::strcmp(str, "dio") == 0) { return LLAMA_LOAD_MODE_DIRECT_IO; } throw std::invalid_argument(std::string("unknown load mode: ") + str); } diff --git a/src/models/deepseek4.cpp b/src/models/deepseek4.cpp index 5ad6473ce203..2d41dace0b28 100644 --- a/src/models/deepseek4.cpp +++ b/src/models/deepseek4.cpp @@ -1133,6 +1133,10 @@ llama_model_deepseek4::graph::graph(const llama_model & model, const llm_graph_p &post, &comb, il); cb(cur, "hc_ffn_pre", il); + ggml_build_forward_expand(gf, residual); + ggml_build_forward_expand(gf, post); + ggml_build_forward_expand(gf, comb); + cur = build_norm(cur, model.layers[il].ffn_norm, nullptr, LLM_NORM_RMS, il); cb(cur, "ffn_norm", il); @@ -1175,7 +1179,7 @@ llama_model_deepseek4::graph::graph(const llama_model & model, const llm_graph_p inpL = build_hc_post(cur, residual, post, comb, il); inpL = build_cvec(inpL, il); - cb(inpL, "l_out", il); + cb(inpL, "l_last", il); } if (inp_out_ids) { diff --git a/src/models/models.h b/src/models/models.h index 916459e12782..92ebfafa1e29 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -424,6 +424,22 @@ struct llama_model_mellum : public llama_model_base { std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; }; +struct llama_model_nanbeige : public llama_model_base { + llama_model_nanbeige(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; + + int n_loops = 1; + int n_layer_phys = 0; + bool skip_loop_final_norm = false; + + struct graph : public llm_graph_context { + graph(const llama_model & model, const llm_graph_params & params); + }; + + std::unique_ptr 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) {} void load_arch_hparams(llama_model_loader & ml) override; diff --git a/src/models/nanbeige.cpp b/src/models/nanbeige.cpp new file mode 100644 index 000000000000..3a546600fa27 --- /dev/null +++ b/src/models/nanbeige.cpp @@ -0,0 +1,184 @@ +#include "models.h" + +void llama_model_nanbeige::load_arch_hparams(llama_model_loader & ml) { + ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); + + uint32_t n_loops_u = 1; + ml.get_key(LLM_KV_NUM_LOOPS, n_loops_u, false); + GGML_ASSERT(n_loops_u >= 1); + + skip_loop_final_norm = false; + ml.get_key(LLM_KV_SKIP_LOOP_FINAL_NORM, skip_loop_final_norm, false); + + n_layer_phys = (int) hparams.n_layer(); + + // Bound-check before casting: signed int mul can overflow and bypass the guard. + GGML_ASSERT((size_t) n_layer_phys * (size_t) n_loops_u <= (size_t) LLAMA_MAX_LAYERS); + n_loops = (int) n_loops_u; + + // Expand logical layer count before load_tensors() allocates layers / KV. + if (n_loops > 1) { + for (int j = 1; j < n_loops; ++j) { + for (int i = 0; i < n_layer_phys; ++i) { + const int dst = i + j * n_layer_phys; + hparams.n_head_arr[dst] = hparams.n_head_arr[i]; + hparams.n_head_kv_arr[dst] = hparams.n_head_kv_arr[i]; + hparams.n_ff_arr[dst] = hparams.n_ff_arr[i]; + hparams.is_swa_impl[dst] = hparams.is_swa_impl[i]; + hparams.is_recr_impl[dst] = hparams.is_recr_impl[i]; + } + } + hparams.n_layer_all = (uint32_t) ((size_t) n_layer_phys * (size_t) n_loops); + } + + type = LLM_TYPE_UNKNOWN; +} + +void llama_model_nanbeige::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_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0); + output = create_tensor(tn(LLM_TENSOR_OUTPUT, "weight"), {n_embd, n_vocab}, TENSOR_NOT_REQUIRED); + if (output == NULL) { + output = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, TENSOR_DUPLICATED); + } + + const int n_phys = n_layer_phys > 0 ? n_layer_phys : n_layer; + for (int i = 0; i < n_phys; ++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_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); + + layer.rope_freqs = create_tensor(tn(LLM_TENSOR_ROPE_FREQS, "weight", i), {n_rot/2}, + TENSOR_NOT_REQUIRED | (i != 0 ? TENSOR_DUPLICATED : 0)); + + layer.ffn_norm = create_tensor(tn(LLM_TENSOR_FFN_NORM, "weight", i), {n_embd}, 0); + layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0); + layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), { n_ff, n_embd}, 0); + layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0); + } + + // Share physical weights across loops; each slot still has its own KV index. + if (n_loops > 1) { + for (int j = 1; j < n_loops; ++j) { + for (int i = 0; i < n_phys; ++i) { + layers[i + j * n_phys] = layers[i]; + } + } + } +} + +std::unique_ptr llama_model_nanbeige::build_arch_graph(const llm_graph_params & params) const { + return std::make_unique(*this, params); +} + +llama_model_nanbeige::graph::graph(const llama_model & model, const llm_graph_params & params) : + llm_graph_context(params) { + const auto & nb = static_cast(model); + + const int64_t n_embd_head = hparams.n_embd_head_v(); + GGML_ASSERT(n_embd_head == hparams.n_embd_head_k()); + + const int n_phys = nb.n_layer_phys > 0 ? nb.n_layer_phys : (int) n_layer; + const int n_loops = nb.n_loops > 0 ? nb.n_loops : 1; + + ggml_tensor * cur; + ggml_tensor * inpL; + + inpL = build_inp_embd(model.tok_embd); + + ggml_tensor * inp_pos = build_inp_pos(); + + auto * inp_attn = build_attn_inp_kv(); + + const float kq_scale = hparams.f_attention_scale == 0.0f + ? 1.0f / sqrtf(float(n_embd_head)) + : hparams.f_attention_scale; + + ggml_tensor * inp_out_ids = build_inp_out_ids(); + + for (int il = 0; il < n_layer; ++il) { + ggml_tensor * inpSA = inpL; + + cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "attn_norm", il); + + { + ggml_tensor * rope_factors = model.get_rope_factors(cparams, il); + + auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur, + n_embd_head, n_head, n_head_kv, il); + + Qcur = ggml_rope_ext( + ctx0, Qcur, inp_pos, rope_factors, + 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, rope_factors, + 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, kq_scale, il); + cb(cur, "attn_out", il); + } + + if (il == n_layer - 1 && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids); + } + + ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA); + cb(ffn_inp, "ffn_inp", il); + + cur = build_norm(ffn_inp, model.layers[il].ffn_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "ffn_norm", il); + + cur = build_ffn(cur, + model.layers[il].ffn_up, model.layers[il].ffn_up_b, model.layers[il].ffn_up_s, + model.layers[il].ffn_gate, model.layers[il].ffn_gate_b, model.layers[il].ffn_gate_s, + model.layers[il].ffn_down, model.layers[il].ffn_down_b, model.layers[il].ffn_down_s, + NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); + cb(cur, "ffn_out", il); + + cur = ggml_add(ctx0, cur, ffn_inp); + cb(cur, "ffn_out", il); + + cur = build_cvec(cur, il); + cb(cur, "l_out", il); + + inpL = cur; + + if (n_loops > 1 && + ((il + 1) % n_phys) == 0 && + (il + 1) < n_layer && + !nb.skip_loop_final_norm) { + cur = build_norm(inpL, model.output_norm, NULL, LLM_NORM_RMS, il); + cb(cur, "loop_norm", il); + inpL = cur; + } + } + + cur = inpL; + + cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); + cb(cur, "result_norm", -1); + res->t_embd = cur; + + 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); +} diff --git a/tests/test-arg-parser.cpp b/tests/test-arg-parser.cpp index 000ecd9aaa76..1d3584f903c4 100644 --- a/tests/test-arg-parser.cpp +++ b/tests/test-arg-parser.cpp @@ -143,6 +143,10 @@ static void test(void) { assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); assert(params.load_mode == LLAMA_LOAD_MODE_MLOCK); + argv = {"binary_name", "-lm", "mmap+mlock"}; + assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); + assert(params.load_mode == LLAMA_LOAD_MODE_MMAP_MLOCK); + argv = {"binary_name", "-lm", "dio"}; assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); assert(params.load_mode == LLAMA_LOAD_MODE_DIRECT_IO); @@ -187,6 +191,11 @@ static void test(void) { assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); assert(params.load_mode == LLAMA_LOAD_MODE_MLOCK); + setenv("LLAMA_ARG_LOAD_MODE", "mmap+mlock", true); + argv = {"binary_name"}; + assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); + assert(params.load_mode == LLAMA_LOAD_MODE_MMAP_MLOCK); + setenv("LLAMA_ARG_LOAD_MODE", "dio", true); argv = {"binary_name"}; assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON)); diff --git a/tests/test-save-load-state.cpp b/tests/test-save-load-state.cpp index bbb025617f6f..6e93ce6fb8da 100644 --- a/tests/test-save-load-state.cpp +++ b/tests/test-save-load-state.cpp @@ -44,8 +44,6 @@ static llama_tokens generate_tokens(llama_context * ctx, llama_sampler * smpl, i n_past++; } - llama_synchronize(ctx); - return result; } diff --git a/tools/cli/README.md b/tools/cli/README.md index 6ee447b07301..972ea04dc7d7 100644 --- a/tools/cli/README.md +++ b/tools/cli/README.md @@ -55,10 +55,10 @@ | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | | `-np, --parallel N` | number of parallel sequences to decode (default: 1)
(env: LLAMA_ARG_N_PARALLEL) | | `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | -| `--mlock` | DEPRECATED in favor of `--load-mode`: mmap + force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | +| `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | -| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | +| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: force system to keep model in RAM rather than swapping or compressing
- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | | `--numa TYPE` | attempt optimizations that help on some NUMA systems
- distribute: spread execution evenly over all nodes
- isolate: only spawn threads on CPUs on the node that execution started on
- numactl: use the CPU map provided by numactl
if run without this previously, it is recommended to drop the system page cache before using this
see https://github.com/ggml-org/llama.cpp/issues/1437
(env: LLAMA_ARG_NUMA) | | `-dev, --device ` | comma-separated list of devices to use for offloading (none = don't offload)
use --list-devices to see a list of available devices
(env: LLAMA_ARG_DEVICE) | | `--list-devices` | print list of available devices and exit | diff --git a/tools/completion/README.md b/tools/completion/README.md index 17f7cd765900..bce71d68d949 100644 --- a/tools/completion/README.md +++ b/tools/completion/README.md @@ -138,10 +138,10 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1 | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | | `-np, --parallel N` | number of parallel sequences to decode (default: 1)
(env: LLAMA_ARG_N_PARALLEL) | | `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | -| `--mlock` | DEPRECATED in favor of `--load-mode`: mmap + force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | +| `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | -| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | +| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: force system to keep model in RAM rather than swapping or compressing
- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | | `--numa TYPE` | attempt optimizations that help on some NUMA systems
- distribute: spread execution evenly over all nodes
- isolate: only spawn threads on CPUs on the node that execution started on
- numactl: use the CPU map provided by numactl
if run without this previously, it is recommended to drop the system page cache before using this
see https://github.com/ggml-org/llama.cpp/issues/1437
(env: LLAMA_ARG_NUMA) | | `-dev, --device ` | comma-separated list of devices to use for offloading (none = don't offload)
use --list-devices to see a list of available devices
(env: LLAMA_ARG_DEVICE) | | `--list-devices` | print list of available devices and exit | diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp index 29ad352d0cf3..c17a27b54019 100644 --- a/tools/llama-bench/llama-bench.cpp +++ b/tools/llama-bench/llama-bench.cpp @@ -429,45 +429,45 @@ static void print_usage(int /* argc */, char ** argv) { } printf("\n"); printf("test parameters:\n"); - printf(" -m, --model (default: %s)\n", join(cmd_params_defaults.model, ",").c_str()); - printf(" -hf, -hfr, --hf-repo /[:quant] Hugging Face model repository; quant is optional, case-insensitive\n"); - printf(" default to Q4_K_M, or falls back to the first file in the repo if Q4_K_M doesn't exist.\n"); - printf(" example: ggml-org/GLM-4.7-Flash-GGUF:Q4_K_M\n"); - printf(" (default: unused)\n"); - printf(" -hff, --hf-file Hugging Face model file. If specified, it will override the quant in --hf-repo\n"); - printf(" (default: unused)\n"); - printf(" -hft, --hf-token Hugging Face access token\n"); - printf(" (default: value from HF_TOKEN environment variable)\n"); - printf(" --offline Offline mode: forces use of cache, prevents network access\n"); - printf(" (default: disabled)\n"); - printf(" -p, --n-prompt (default: %s)\n", join(cmd_params_defaults.n_prompt, ",").c_str()); - printf(" -n, --n-gen (default: %s)\n", join(cmd_params_defaults.n_gen, ",").c_str()); - printf(" -pg (default: %s)\n", join(transform_to_str(cmd_params_defaults.n_pg, pair_str), ",").c_str()); - printf(" -d, --n-depth (default: %s)\n", join(cmd_params_defaults.n_depth, ",").c_str()); - printf(" -b, --batch-size (default: %s)\n", join(cmd_params_defaults.n_batch, ",").c_str()); - printf(" -ub, --ubatch-size (default: %s)\n", join(cmd_params_defaults.n_ubatch, ",").c_str()); - printf(" -ctk, --cache-type-k (default: %s)\n", join(transform_to_str(cmd_params_defaults.type_k, ggml_type_name), ",").c_str()); - printf(" -ctv, --cache-type-v (default: %s)\n", join(transform_to_str(cmd_params_defaults.type_v, ggml_type_name), ",").c_str()); - printf(" -t, --threads (default: %s)\n", join(cmd_params_defaults.n_threads, ",").c_str()); - printf(" -C, --cpu-mask (default: %s)\n", join(cmd_params_defaults.cpu_mask, ",").c_str()); - printf(" --cpu-strict <0|1> (default: %s)\n", join(cmd_params_defaults.cpu_strict, ",").c_str()); - printf(" --poll <0...100> (default: %s)\n", join(cmd_params_defaults.poll, ",").c_str()); - printf(" -ngl, --n-gpu-layers (default: %s)\n", join(cmd_params_defaults.n_gpu_layers, ",").c_str()); - printf(" -ncmoe, --n-cpu-moe (default: %s)\n", join(cmd_params_defaults.n_cpu_moe, ",").c_str()); - printf(" -sm, --split-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.split_mode, split_mode_str), ",").c_str()); - printf(" -mg, --main-gpu (default: %s)\n", join(cmd_params_defaults.main_gpu, ",").c_str()); - printf(" -nkvo, --no-kv-offload <0|1> (default: %s)\n", join(cmd_params_defaults.no_kv_offload, ",").c_str()); - printf(" -fa, --flash-attn (default: %s)\n", join(transform_to_str(cmd_params_defaults.flash_attn, llama_flash_attn_type_name), ",").c_str()); - printf(" -dev, --device (default: auto)\n"); - printf(" -lm, --load-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.load_mode, llama_load_mode_name), ",").c_str()); - printf(" -mmp, --mmap <0|1> (DEPRECATED IN FAVOUR OF --load-mode)\n"); - printf(" -dio, --direct-io <0|1> (DEPRECATED IN FAVOUR OF --load-mode)\n"); - printf(" -embd, --embeddings <0|1> (default: %s)\n", join(cmd_params_defaults.embeddings, ",").c_str()); - printf(" -ts, --tensor-split (default: 0)\n"); + printf(" -m, --model (default: %s)\n", join(cmd_params_defaults.model, ",").c_str()); + printf(" -hf, -hfr, --hf-repo /[:quant] Hugging Face model repository; quant is optional, case-insensitive\n"); + printf(" default to Q4_K_M, or falls back to the first file in the repo if Q4_K_M doesn't exist.\n"); + printf(" example: ggml-org/GLM-4.7-Flash-GGUF:Q4_K_M\n"); + printf(" (default: unused)\n"); + printf(" -hff, --hf-file Hugging Face model file. If specified, it will override the quant in --hf-repo\n"); + printf(" (default: unused)\n"); + printf(" -hft, --hf-token Hugging Face access token\n"); + printf(" (default: value from HF_TOKEN environment variable)\n"); + printf(" --offline Offline mode: forces use of cache, prevents network access\n"); + printf(" (default: disabled)\n"); + printf(" -p, --n-prompt (default: %s)\n", join(cmd_params_defaults.n_prompt, ",").c_str()); + printf(" -n, --n-gen (default: %s)\n", join(cmd_params_defaults.n_gen, ",").c_str()); + printf(" -pg (default: %s)\n", join(transform_to_str(cmd_params_defaults.n_pg, pair_str), ",").c_str()); + printf(" -d, --n-depth (default: %s)\n", join(cmd_params_defaults.n_depth, ",").c_str()); + printf(" -b, --batch-size (default: %s)\n", join(cmd_params_defaults.n_batch, ",").c_str()); + printf(" -ub, --ubatch-size (default: %s)\n", join(cmd_params_defaults.n_ubatch, ",").c_str()); + printf(" -ctk, --cache-type-k (default: %s)\n", join(transform_to_str(cmd_params_defaults.type_k, ggml_type_name), ",").c_str()); + printf(" -ctv, --cache-type-v (default: %s)\n", join(transform_to_str(cmd_params_defaults.type_v, ggml_type_name), ",").c_str()); + printf(" -t, --threads (default: %s)\n", join(cmd_params_defaults.n_threads, ",").c_str()); + printf(" -C, --cpu-mask (default: %s)\n", join(cmd_params_defaults.cpu_mask, ",").c_str()); + printf(" --cpu-strict <0|1> (default: %s)\n", join(cmd_params_defaults.cpu_strict, ",").c_str()); + printf(" --poll <0...100> (default: %s)\n", join(cmd_params_defaults.poll, ",").c_str()); + printf(" -ngl, --n-gpu-layers (default: %s)\n", join(cmd_params_defaults.n_gpu_layers, ",").c_str()); + printf(" -ncmoe, --n-cpu-moe (default: %s)\n", join(cmd_params_defaults.n_cpu_moe, ",").c_str()); + printf(" -sm, --split-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.split_mode, split_mode_str), ",").c_str()); + printf(" -mg, --main-gpu (default: %s)\n", join(cmd_params_defaults.main_gpu, ",").c_str()); + printf(" -nkvo, --no-kv-offload <0|1> (default: %s)\n", join(cmd_params_defaults.no_kv_offload, ",").c_str()); + printf(" -fa, --flash-attn (default: %s)\n", join(transform_to_str(cmd_params_defaults.flash_attn, llama_flash_attn_type_name), ",").c_str()); + printf(" -dev, --device (default: auto)\n"); + printf(" -lm, --load-mode (default: %s)\n", join(transform_to_str(cmd_params_defaults.load_mode, llama_load_mode_name), ",").c_str()); + printf(" -mmp, --mmap <0|1> (DEPRECATED IN FAVOUR OF --load-mode)\n"); + printf(" -dio, --direct-io <0|1> (DEPRECATED IN FAVOUR OF --load-mode)\n"); + printf(" -embd, --embeddings <0|1> (default: %s)\n", join(cmd_params_defaults.embeddings, ",").c_str()); + printf(" -ts, --tensor-split (default: 0)\n"); printf(" -ot --override-tensor =;...\n"); - printf(" (default: disabled)\n"); - printf(" -nopo, --no-op-offload <0|1> (default: 0)\n"); - printf(" --no-host <0|1> (default: %s)\n", join(cmd_params_defaults.no_host, ",").c_str()); + printf(" (default: disabled)\n"); + printf(" -nopo, --no-op-offload <0|1> (default: 0)\n"); + printf(" --no-host <0|1> (default: %s)\n", join(cmd_params_defaults.no_host, ",").c_str()); printf("\n"); printf( "Multiple values can be given for each parameter by separating them with ','\n" @@ -670,22 +670,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { break; } } else if (arg == "--list-devices") { - std::vector devices; - for (size_t i = 0; i < ggml_backend_dev_count(); ++i) { - auto * dev = ggml_backend_dev_get(i); - if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) { - devices.push_back(dev); - } - } - printf("Available devices:\n"); - if (devices.empty()) { - printf(" (none)\n"); - } - for (auto * dev : devices) { - size_t free, total; - ggml_backend_dev_memory(dev, &free, &total); - printf(" %s: %s (%zu MiB, %zu MiB free)\n", ggml_backend_dev_name(dev), ggml_backend_dev_description(dev), total / 1024 / 1024, free / 1024 / 1024); - } + common_print_available_devices(); exit(0); } else if (arg == "-t" || arg == "--threads") { if (++i >= argc) { @@ -785,6 +770,8 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { mode = LLAMA_LOAD_MODE_MMAP; } else if (m == "mlock") { mode = LLAMA_LOAD_MODE_MLOCK; + } else if (m == "mmap+mlock") { + mode = LLAMA_LOAD_MODE_MMAP_MLOCK; } else if (m == "dio") { mode = LLAMA_LOAD_MODE_DIRECT_IO; } else { diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt index fd7ddceb0bf0..18a8288ba048 100644 --- a/tools/mtmd/CMakeLists.txt +++ b/tools/mtmd/CMakeLists.txt @@ -51,6 +51,7 @@ add_library(mtmd models/qwen3vl.cpp models/mimovl.cpp models/qwen3a.cpp + models/mimo-audio.cpp models/step3vl.cpp models/siglip.cpp models/whisper-enc.cpp diff --git a/tools/mtmd/clip-graph.h b/tools/mtmd/clip-graph.h index a95de20a3122..29352abb4c0b 100644 --- a/tools/mtmd/clip-graph.h +++ b/tools/mtmd/clip-graph.h @@ -13,6 +13,14 @@ struct build_vit_opts { ggml_tensor * attn_mask = nullptr; + // TODO @ngxson : merge attn_mask and attn_mask_layers into one call + std::vector attn_mask_layers; // one per layer + + // hook at layer output embeddings + std::function callback_layer_out = nullptr; + + // whether to skip the automatic post-layernorm (model.post_ln_w) applied at the end + bool skip_post_ln = false; }; struct clip_graph { diff --git a/tools/mtmd/clip-impl.h b/tools/mtmd/clip-impl.h index 42374311ce7b..09204113801f 100644 --- a/tools/mtmd/clip-impl.h +++ b/tools/mtmd/clip-impl.h @@ -82,6 +82,12 @@ #define KEY_A_PROJ_WINDOW_SIZE "clip.audio.projector.window_size" #define KEY_A_PROJ_DOWNSAMPLE_RATE "clip.audio.projector.downsample_rate" #define KEY_A_PROJ_HEAD_COUNT "clip.audio.projector.head_count" +#define KEY_A_RVQ_NUM_QUANTIZERS "clip.audio.rvq.num_quantizers" // mimo-audio-tokenizer +#define KEY_A_RVQ_CODEBOOK_SIZE "clip.audio.rvq.codebook_size" // mimo-audio-tokenizer: per-quantizer bin count +#define KEY_A_WA_PATTERN_MODE "clip.audio.wa_pattern_mode" // mimo-audio-tokenizer, per-layer -1 (full) / 0 (windowed) +#define KEY_A_ATTN_WINDOW_SIZE "clip.audio.window_size" // mimo-audio-tokenizer: sliding-window radius +#define KEY_A_LOCAL_BLOCK_COUNT "clip.audio.local_block_count" // mimo-v2.5: input_local_transformer layer count +#define KEY_A_LOCAL_GROUP_SIZE "clip.audio.local_group_size" // mimo-v2.5: input_local_transformer grouping size // // tensor name constants @@ -175,6 +181,24 @@ #define TN_MM_NORM_PRE "mm.a.norm_pre.%s" #define TN_MM_NORM_MID "mm.a.norm_mid.%s" +// mimo-audio-tokenizer +#define TN_A_DOWNSAMPLE_CONV "a.downsample.conv.%s" +#define TN_A_DOWNSAMPLE_NORM "a.downsample.norm.%s" +#define TN_A_RVQ_CODEBOOK "a.rvq.codebook.%s" +// mimo-v2.5: text-side RVQ code embedding ("text codebook") +#define TN_MM_A_CODE_EMBD "mm.a.code_embd.%s" +// mimo-v2.5: LLM-side connector (input_local_transformer) +#define TN_MM_A_LOCAL_ATTN_Q "mm.a.local_blk.%d.attn_q.%s" +#define TN_MM_A_LOCAL_ATTN_K "mm.a.local_blk.%d.attn_k.%s" +#define TN_MM_A_LOCAL_ATTN_V "mm.a.local_blk.%d.attn_v.%s" +#define TN_MM_A_LOCAL_ATTN_OUT "mm.a.local_blk.%d.attn_out.%s" +#define TN_MM_A_LOCAL_FFN_GATE "mm.a.local_blk.%d.ffn_gate.%s" +#define TN_MM_A_LOCAL_FFN_UP "mm.a.local_blk.%d.ffn_up.%s" +#define TN_MM_A_LOCAL_FFN_DOWN "mm.a.local_blk.%d.ffn_down.%s" +#define TN_MM_A_LOCAL_LN1 "mm.a.local_blk.%d.ln1.%s" +#define TN_MM_A_LOCAL_LN2 "mm.a.local_blk.%d.ln2.%s" +#define TN_MM_A_LOCAL_NORM "mm.a.local_norm.%s" + // cogvlm #define TN_MM_POST_FC_NORM "mm.post_fc_norm.%s" #define TN_MM_H_TO_4H "mm.up.%s" @@ -374,6 +398,7 @@ enum projector_type { PROJECTOR_TYPE_MIMOVL, PROJECTOR_TYPE_MINIMAX_M3, PROJECTOR_TYPE_GRANITE4_VISION, + PROJECTOR_TYPE_MIMO_AUDIO, PROJECTOR_TYPE_UNKNOWN, }; @@ -429,6 +454,7 @@ static std::map PROJECTOR_TYPE_NAMES = { { PROJECTOR_TYPE_MIMOVL, "mimovl"}, { PROJECTOR_TYPE_MINIMAX_M3, "minimax_m3"}, { PROJECTOR_TYPE_GRANITE4_VISION, "granite4_vision"}, + { PROJECTOR_TYPE_MIMO_AUDIO, "mimo_audio"}, }; static projector_type clip_projector_type_from_string(const std::string & str) { diff --git a/tools/mtmd/clip-model.h b/tools/mtmd/clip-model.h index 850957d7de1c..8dc87549766e 100644 --- a/tools/mtmd/clip-model.h +++ b/tools/mtmd/clip-model.h @@ -124,6 +124,14 @@ struct clip_hparams { int32_t audio_window_len = -1; int32_t audio_hop_len = -1; + // mimo-audio-tokenizer: residual vector quantizer + int32_t rvq_num_quantizers = 0; + std::vector rvq_codebook_size; // per-quantizer bin count (ragged, e.g. 1024/1024/256/128x17) + + // mimo-v2.5: LLM-side connector (input_local_transformer) + int32_t audio_local_n_layer = 0; + int32_t audio_local_group_size = 0; + // legacy bool has_llava_projector = false; int minicpmv_version = 0; @@ -537,6 +545,20 @@ struct clip_model { ggml_tensor * mm_norm_pre_b = nullptr; ggml_tensor * mm_norm_mid_w = nullptr; + // mimo-audio-tokenizer: post-transformer downsample + RVQ codebook + ggml_tensor * downsample_conv_w = nullptr; // no bias + ggml_tensor * downsample_norm_w = nullptr; + ggml_tensor * downsample_norm_b = nullptr; + ggml_tensor * rvq_codebook = nullptr; // merged 3D [n_q, max_bins, dim] + + // mimo-v2.5: text-side RVQ code embedding ("text codebook") + ggml_tensor * mm_a_code_embd = nullptr; // merged 3D [n_channels, vocab, dim] + + // mimo-v2.5: LLM-side connector (input_local_transformer, separate from the + // audio_tokenizer's own encoder `layers`) + std::vector mm_a_local_layers; + ggml_tensor * mm_a_local_norm_w = nullptr; + // qwen3a ggml_tensor * conv2d_1_w = nullptr; ggml_tensor * conv2d_1_b = nullptr; diff --git a/tools/mtmd/clip.cpp b/tools/mtmd/clip.cpp index e0e2107a0be3..04614b93bd27 100644 --- a/tools/mtmd/clip.cpp +++ b/tools/mtmd/clip.cpp @@ -340,6 +340,11 @@ ggml_tensor * clip_graph::build_vit( auto & layer = model.layers[il]; ggml_tensor * cur = inpL; // inpL = residual, cur = hidden_states + ggml_tensor * attn_mask = opts.attn_mask; + if (opts.attn_mask_layers.size() > (size_t) il) { + attn_mask = opts.attn_mask_layers[il]; + } + // layernorm1 cur = build_norm(cur, layer.ln_1_w, layer.ln_1_b, norm_t, eps, il); cb(cur, "layer_inp_normed", il); @@ -452,7 +457,7 @@ ggml_tensor * clip_graph::build_vit( // build_attn returns a flat 2D [n_embd, n_pos*B] cur = build_attn(layer.o_w, layer.o_b, - Qcur, Kcur, Vcur, opts.attn_mask, kq_scale, il); + Qcur, Kcur, Vcur, attn_mask, kq_scale, il); cb(cur, "attn_out", il); } @@ -471,6 +476,10 @@ ggml_tensor * clip_graph::build_vit( inpL = cur; // inpL = residual, cur = hidden_states + if (opts.callback_layer_out) { + opts.callback_layer_out(cur, il); + } + cb(cur, "ffn_inp", il); // layernorm2 (pre-ffn norm) @@ -519,7 +528,7 @@ ggml_tensor * clip_graph::build_vit( } // post-layernorm - if (model.post_ln_w) { + if (model.post_ln_w && !opts.skip_post_ln) { inpL = build_norm(inpL, model.post_ln_w, model.post_ln_b, norm_t, eps, -1); } @@ -1012,6 +1021,10 @@ static std::unique_ptr clip_get_graph_builder(clip_ctx * ctx, const { builder = std::make_unique(ctx, img); } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + builder = std::make_unique(ctx, img); + } break; case PROJECTOR_TYPE_YOUTUVL: { builder = std::make_unique(ctx, img); @@ -1575,6 +1588,45 @@ struct clip_model_loader { hparams.audio_window_len = 400; hparams.audio_hop_len = 160; } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + get_u32(KEY_A_RVQ_NUM_QUANTIZERS, hparams.rvq_num_quantizers, false); + get_arr_int(KEY_A_RVQ_CODEBOOK_SIZE, hparams.rvq_codebook_size, false); + if (hparams.rvq_num_quantizers <= 0) { + throw std::runtime_error(string_format("%s: mimo_audio: missing %s\n", __func__, KEY_A_RVQ_NUM_QUANTIZERS)); + } + if ((int) hparams.rvq_codebook_size.size() != hparams.rvq_num_quantizers) { + throw std::runtime_error(string_format( + "%s: mimo_audio: %s length (%zu) must equal %s (%d)\n", __func__, + KEY_A_RVQ_CODEBOOK_SIZE, hparams.rvq_codebook_size.size(), + KEY_A_RVQ_NUM_QUANTIZERS, hparams.rvq_num_quantizers)); + } + hparams.ffn_op = FFN_GELU_ERF; // PyTorch F.gelu default (approximate="none") + hparams.rope_theta = 10000.0f; + + // audio preprocessing params (mel spectrogram) + hparams.audio_sample_rate = 24000; + hparams.audio_n_fft = 960; + hparams.audio_window_len = 960; + hparams.audio_hop_len = 240; + + get_u32(KEY_A_ATTN_WINDOW_SIZE, hparams.attn_window_size); + std::vector wa_pattern; + get_arr_int(KEY_A_WA_PATTERN_MODE, wa_pattern, true); + if ((int) wa_pattern.size() != hparams.n_layer) { + throw std::runtime_error(string_format( + "%s: mimo_audio: %s length (%zu) must equal n_layer (%d)\n", __func__, + KEY_A_WA_PATTERN_MODE, wa_pattern.size(), hparams.n_layer)); + } + hparams.wa_pattern_mode.assign(wa_pattern.begin(), wa_pattern.end()); + + get_u32(KEY_A_LOCAL_BLOCK_COUNT, hparams.audio_local_n_layer); + get_u32(KEY_A_LOCAL_GROUP_SIZE, hparams.audio_local_group_size); + if (hparams.audio_local_group_size <= 0) { + throw std::runtime_error(string_format( + "%s: mimo_audio: %s must be > 0\n", __func__, KEY_A_LOCAL_GROUP_SIZE)); + } + } break; case PROJECTOR_TYPE_PADDLEOCR: { hparams.n_merge = 2; @@ -2444,6 +2496,54 @@ struct clip_model_loader { model.mm_2_w = get_tensor(string_format(TN_MM_AUDIO_MLP, 2, "weight")); model.mm_2_b = get_tensor(string_format(TN_MM_AUDIO_MLP, 2, "bias")); } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + model.conv1d_1_w = get_tensor(string_format(TN_CONV1D, 1, "weight")); + model.conv1d_1_b = get_tensor(string_format(TN_CONV1D, 1, "bias")); + model.conv1d_2_w = get_tensor(string_format(TN_CONV1D, 2, "weight")); + model.conv1d_2_b = get_tensor(string_format(TN_CONV1D, 2, "bias")); + model.downsample_conv_w = get_tensor(string_format(TN_A_DOWNSAMPLE_CONV, "weight")); + model.downsample_norm_w = get_tensor(string_format(TN_A_DOWNSAMPLE_NORM, "weight")); + model.downsample_norm_b = get_tensor(string_format(TN_A_DOWNSAMPLE_NORM, "bias")); + model.rvq_codebook = get_tensor(string_format(TN_A_RVQ_CODEBOOK, "weight"), false); + model.mm_a_code_embd = get_tensor(string_format(TN_MM_A_CODE_EMBD, "weight"), false); + if (!model.rvq_codebook || !model.mm_a_code_embd) { + throw std::runtime_error(string_format("%s: mimo_audio: missing %s or %s\n", __func__, + TN_A_RVQ_CODEBOOK, TN_MM_A_CODE_EMBD)); + } + // hparams.rvq_codebook_size comes from GGUF metadata and is independent of the + // tensors' actual shapes - bound it so codebook/code_embd views built from it + // (mimo-audio.cpp) can never read past either tensor's allocated bins/vocab. + for (int32_t bins : hparams.rvq_codebook_size) { + if (bins <= 0 || bins > model.rvq_codebook->ne[1] || bins > model.mm_a_code_embd->ne[1]) { + throw std::runtime_error(string_format( + "%s: mimo_audio: %s entry (%d) out of range for codebook/code_embd tensors\n", + __func__, KEY_A_RVQ_CODEBOOK_SIZE, bins)); + } + } + + // LLM-side connector: input_local_transformer + projection + model.mm_a_local_layers.resize(hparams.audio_local_n_layer); + for (int il = 0; il < hparams.audio_local_n_layer; il++) { + auto & layer = model.mm_a_local_layers[il]; + layer.q_w = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_Q, il, "weight")); + layer.q_b = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_Q, il, "bias")); + layer.k_w = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_K, il, "weight")); + layer.k_b = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_K, il, "bias")); + layer.v_w = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_V, il, "weight")); + layer.v_b = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_V, il, "bias")); + layer.o_w = get_tensor(string_format(TN_MM_A_LOCAL_ATTN_OUT, il, "weight")); + layer.ff_gate_w = get_tensor(string_format(TN_MM_A_LOCAL_FFN_GATE, il, "weight")); + layer.ff_up_w = get_tensor(string_format(TN_MM_A_LOCAL_FFN_UP, il, "weight")); + layer.ff_down_w = get_tensor(string_format(TN_MM_A_LOCAL_FFN_DOWN, il, "weight")); + layer.ln_1_w = get_tensor(string_format(TN_MM_A_LOCAL_LN1, il, "weight")); + layer.ln_2_w = get_tensor(string_format(TN_MM_A_LOCAL_LN2, il, "weight")); + } + model.mm_a_local_norm_w = get_tensor(string_format(TN_MM_A_LOCAL_NORM, "weight")); + + model.mm_1_w = get_tensor(string_format(TN_MM_AUDIO_MLP, 1, "weight")); + model.mm_2_w = get_tensor(string_format(TN_MM_AUDIO_MLP, 2, "weight")); + } break; case PROJECTOR_TYPE_VOXTRAL: { model.conv1d_1_w = get_tensor(string_format(TN_CONV1D, 1, "weight")); @@ -3549,6 +3649,15 @@ int clip_n_output_tokens(const clip_ctx * ctx, const clip_image_f32 * img) { { n_patches = img->nx(); // no downsampling: one token per raw waveform frame } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + // conv1(s=1) + conv2(s=2) -> RVQ-encoder downsample conv(k=2,s=2) + int n = img->nx(); + n = (n - 1) / 2 + 1; // conv1 + conv2 + n = (n - 2) / 2 + 1; // downsample conv + const int group_size = params.audio_local_group_size; + n_patches = (n + group_size - 1) / group_size; + } break; case PROJECTOR_TYPE_GRANITE_SPEECH: { const int ws = ctx->model.hparams.audio_proj_window_size; @@ -4376,6 +4485,58 @@ bool clip_image_batch_encode(clip_ctx * ctx, int n_threads, const clip_image_f32 set_input_f32("pos_emb", pos_emb); } } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + GGML_ASSERT(imgs.entries.size() == 1); + const int n_frames = imgs.entries.front().nx(); + const int n_pos = (n_frames - 1) / 2 + 1; // matches conv1(s=1)+conv2(s=2) output length + + std::vector positions(n_pos); + for (int i = 0; i < n_pos; i++) { + positions[i] = i; + } + set_input_i32("mimo_audio_positions", positions); + + const int window = hparams.attn_window_size; + GGML_ASSERT(window > 0); + + const float neg_inf = std::numeric_limits::lowest(); + std::vector full_mask((size_t) n_pos * n_pos); + std::vector window_mask((size_t) n_pos * n_pos); + for (int q = 0; q < n_pos; q++) { + for (int k = 0; k < n_pos; k++) { + const bool causal_ok = k <= q; + full_mask[(size_t) q * n_pos + k] = causal_ok ? 0.0f : neg_inf; + window_mask[(size_t) q * n_pos + k] = (causal_ok && (q - k) <= window) ? 0.0f : neg_inf; + } + } + set_input_f32("mimo_audio_full_mask", full_mask); + set_input_f32("mimo_audio_window_mask", window_mask); + + // input_local_transformer: block-diagonal mask + in-group positions + { + const int n_pos_ds = (n_pos - 2) / 2 + 1; // matches downsample conv (k=2,s=2,p=0) + const int group_size = hparams.audio_local_group_size; + GGML_ASSERT(group_size > 0); + const int n_groups = (n_pos_ds + group_size - 1) / group_size; + const int n_padded = n_groups * group_size; + + std::vector local_positions(n_padded); + for (int i = 0; i < n_padded; i++) { + local_positions[i] = i % group_size; + } + set_input_i32("mimo_audio_local_positions", local_positions); + + std::vector local_mask((size_t) n_padded * n_padded); + for (int q = 0; q < n_padded; q++) { + for (int k = 0; k < n_padded; k++) { + const bool same_group = (q / group_size) == (k / group_size); + local_mask[(size_t) q * n_padded + k] = same_group ? 0.0f : neg_inf; + } + } + set_input_f32("mimo_audio_local_mask", local_mask); + } + } break; case PROJECTOR_TYPE_LFM2A: { GGML_ASSERT(imgs.entries.size() == 1); @@ -4678,6 +4839,8 @@ int clip_n_mmproj_embd(const struct clip_ctx * ctx) { return ctx->model.qf_proj_blocks.size() * ctx->model.hparams.projection_dim; case PROJECTOR_TYPE_GLM4V: return ctx->model.mm_ffn_down_w->ne[1]; + case PROJECTOR_TYPE_MIMO_AUDIO: + return ctx->model.mm_2_w->ne[1]; default: GGML_ABORT("Unknown projector type"); } diff --git a/tools/mtmd/models/mimo-audio.cpp b/tools/mtmd/models/mimo-audio.cpp new file mode 100644 index 000000000000..481b36cc8d60 --- /dev/null +++ b/tools/mtmd/models/mimo-audio.cpp @@ -0,0 +1,218 @@ +#include "models.h" + +ggml_cgraph * clip_graph_mimo_audio::build() { + ggml_tensor * inp = build_inp_raw(1); // [n_frames, n_mel, 1] + + ggml_tensor * cur = ggml_conv_1d_ph(ctx0, model.conv1d_1_w, inp, 1, 1); + cur = ggml_add(ctx0, cur, model.conv1d_1_b); + cur = ggml_gelu_erf(ctx0, cur); + + cur = ggml_conv_1d_ph(ctx0, model.conv1d_2_w, cur, 2, 1); + cur = ggml_add(ctx0, cur, model.conv1d_2_b); + cur = ggml_gelu_erf(ctx0, cur); + + ggml_tensor * inpL = ggml_cont(ctx0, ggml_transpose(ctx0, cur)); // [n_embd, n_pos] + const int64_t n_pos = inpL->ne[1]; + cb(inpL, "after_conv1d", -1); + + GGML_ASSERT((int) hparams.wa_pattern_mode.size() == n_layer); + + ggml_tensor * inp_pos = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_pos); + ggml_set_name(inp_pos, "mimo_audio_positions"); + ggml_set_input(inp_pos); + + ggml_tensor * full_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_pos, n_pos); + ggml_set_name(full_mask, "mimo_audio_full_mask"); + ggml_set_input(full_mask); + + ggml_tensor * window_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_pos, n_pos); + ggml_set_name(window_mask, "mimo_audio_window_mask"); + ggml_set_input(window_mask); + + build_vit_opts opts; + opts.attn_mask_layers.resize(n_layer); + for (int il = 0; il < n_layer; il++) { + opts.attn_mask_layers[il] = hparams.wa_pattern_mode[il] == -1 ? full_mask : window_mask; + } + // the skip connection below must be added before the post-transformer norm, + // so build_vit must not apply that norm itself + opts.skip_post_ln = true; + + // encoder_skip_layer_id=3 (1-indexed) -> capture output of layer index 2 + const int skip_capture_il = 2; + GGML_ASSERT(n_layer > skip_capture_il); + ggml_tensor * skip_hidden = nullptr; + opts.callback_layer_out = [&](ggml_tensor * layer_cur, int il) { + if (il == skip_capture_il) { + skip_hidden = layer_cur; + } + }; + + auto add_pos = [&](ggml_tensor * x, const clip_layer &) { + return ggml_rope_ext(ctx0, x, inp_pos, nullptr, d_head, + GGML_ROPE_TYPE_NEOX, 0, hparams.rope_theta, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f); + }; + + inpL = build_vit(inpL, n_pos, NORM_TYPE_NORMAL, hparams.ffn_op, nullptr, add_pos, opts); + inpL = ggml_reshape_2d(ctx0, inpL, n_embd, n_pos); // build_vit restores a (size-1) batch dim + + GGML_ASSERT(skip_hidden != nullptr); + inpL = ggml_add(ctx0, inpL, skip_hidden); + + inpL = build_norm(inpL, model.post_ln_w, model.post_ln_b, NORM_TYPE_NORMAL, eps, -1); + cb(inpL, "after_transformer", -1); + + // downsample: strided conv (no bias) + gelu + layernorm + { + ggml_tensor * ds = ggml_cont(ctx0, ggml_transpose(ctx0, inpL)); // [n_pos, n_embd] + ds = ggml_conv_1d(ctx0, model.downsample_conv_w, ds, 2, 0, 1); + ds = ggml_gelu_erf(ctx0, ds); + ds = ggml_cont(ctx0, ggml_transpose(ctx0, ds)); // [n_embd, n_pos/2] + ds = build_norm(ds, model.downsample_norm_w, model.downsample_norm_b, NORM_TYPE_NORMAL, eps, -1); + inpL = ds; + } + cb(inpL, "after_downsample", -1); + + // RVQ quantize: codebook ne=[dim, max_bins, n_q] + // quantize input vector to codes (type=I32) + std::vector codes; + { + GGML_ASSERT(model.rvq_codebook != nullptr); + const int64_t dim = model.rvq_codebook->ne[0]; + GGML_ASSERT(dim == inpL->ne[0]); + GGML_ASSERT((int64_t) hparams.rvq_codebook_size.size() == model.rvq_codebook->ne[2]); + + ggml_tensor * residual = inpL; // [dim, n_pos_ds] + + for (size_t q = 0; q < hparams.rvq_codebook_size.size(); q++) { + const int64_t bins = hparams.rvq_codebook_size[q]; + ggml_tensor * codebook_q = ggml_view_2d(ctx0, model.rvq_codebook, dim, bins, + model.rvq_codebook->nb[1], q * model.rvq_codebook->nb[2]); + codebook_q = ggml_cont(ctx0, codebook_q); + + ggml_tensor * codebook_norm = ggml_sum_rows(ctx0, ggml_sqr(ctx0, codebook_q)); // [1, bins] + codebook_norm = ggml_cont(ctx0, ggml_transpose(ctx0, codebook_norm)); // [bins, 1] + + ggml_tensor * dot = ggml_mul_mat(ctx0, codebook_q, residual); // [bins, n_pos_ds] + ggml_tensor * scores = ggml_sub(ctx0, ggml_scale(ctx0, dot, 2.0f), codebook_norm); + + ggml_tensor * idx = ggml_argmax(ctx0, scores); // [n_pos_ds] + codes.push_back(idx); + + ggml_tensor * quant = ggml_get_rows(ctx0, codebook_q, idx); // [dim, n_pos_ds] + residual = ggml_sub(ctx0, residual, quant); + cb(idx, "rvq_code", (int) q); + } + } + + // convert codes to LLM embeddings + ggml_tensor * code_embd_sum = nullptr; + { + GGML_ASSERT(model.mm_a_code_embd != nullptr); + const int64_t dim = model.mm_a_code_embd->ne[0]; + const int64_t vocab = model.mm_a_code_embd->ne[1]; + GGML_ASSERT((int64_t) codes.size() == model.mm_a_code_embd->ne[2]); + GGML_ASSERT(dim == inpL->ne[0]); + + for (size_t i = 0; i < codes.size(); i++) { + ggml_tensor * table_i = ggml_view_2d(ctx0, model.mm_a_code_embd, dim, vocab, + model.mm_a_code_embd->nb[1], i * model.mm_a_code_embd->nb[2]); + table_i = ggml_cont(ctx0, table_i); + + ggml_tensor * embd_i = ggml_get_rows(ctx0, table_i, codes[i]); // [dim, n_pos_ds] + code_embd_sum = code_embd_sum ? ggml_add(ctx0, code_embd_sum, embd_i) : embd_i; + } + cb(code_embd_sum, "code_embd_sum", -1); + } + + // input_local_transformer + // groups of `group_size` consecutive downsampled frames are processed together, attending only within their own group. + // Implemented as a block-diagonal mask + in-group-repeating positions + // (rather than a real batch dim) - same technique as the encoder's masks above, and as gemma4a's / deepseekocr2's chunked attention. + + // note: hand-rolled here instead of build_vit() because this is a second, independent layer stack + // (own layer array/count, RMSNorm instead of LN, SiLU FFN, own RoPE theta) + + ggml_tensor * projected; + { + const int group_size = hparams.audio_local_group_size; + GGML_ASSERT(group_size > 0); + const int64_t n_pos_ds = code_embd_sum->ne[1]; + const int64_t n_groups = (n_pos_ds + group_size - 1) / group_size; + const int64_t n_padded = n_groups * group_size; + + ggml_tensor * cur_local = code_embd_sum; + if (n_padded != n_pos_ds) { + cur_local = ggml_pad(ctx0, cur_local, 0, (int) (n_padded - n_pos_ds), 0, 0); + } + + ggml_tensor * local_pos = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, n_padded); + ggml_set_name(local_pos, "mimo_audio_local_positions"); + ggml_set_input(local_pos); + + ggml_tensor * local_mask = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_padded, n_padded); + ggml_set_name(local_mask, "mimo_audio_local_mask"); + ggml_set_input(local_mask); + + const float local_rope_theta = 640000.0f; // audio_config.rope_theta (differs from the encoder's) + auto apply_local_rope = [&](ggml_tensor * x) { + return ggml_rope_ext(ctx0, x, local_pos, nullptr, d_head, + GGML_ROPE_TYPE_NEOX, 0, local_rope_theta, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f); + }; + + for (int il = 0; il < hparams.audio_local_n_layer; il++) { + auto & layer = model.mm_a_local_layers[il]; + + ggml_tensor * attn_in = build_norm(cur_local, layer.ln_1_w, nullptr, NORM_TYPE_RMS, eps, il); + + ggml_tensor * Qcur = build_mm(layer.q_w, attn_in); + if (layer.q_b) { + Qcur = ggml_add(ctx0, Qcur, layer.q_b); + } + ggml_tensor * Kcur = build_mm(layer.k_w, attn_in); + if (layer.k_b) { + Kcur = ggml_add(ctx0, Kcur, layer.k_b); + } + ggml_tensor * Vcur = build_mm(layer.v_w, attn_in); + if (layer.v_b) { + Vcur = ggml_add(ctx0, Vcur, layer.v_b); + } + + Qcur = ggml_reshape_3d(ctx0, Qcur, d_head, n_head, n_padded); + Kcur = ggml_reshape_3d(ctx0, Kcur, d_head, n_head, n_padded); + Vcur = ggml_reshape_3d(ctx0, Vcur, d_head, n_head, n_padded); + + Qcur = apply_local_rope(Qcur); + Kcur = apply_local_rope(Kcur); + + ggml_tensor * attn_out = build_attn(layer.o_w, nullptr, Qcur, Kcur, Vcur, local_mask, kq_scale, il); + cur_local = ggml_add(ctx0, cur_local, attn_out); + + ggml_tensor * ffn_in = build_norm(cur_local, layer.ln_2_w, nullptr, NORM_TYPE_RMS, eps, il); + ggml_tensor * ffn_out = build_ffn(ffn_in, + layer.ff_up_w, nullptr, + layer.ff_gate_w, nullptr, + layer.ff_down_w, nullptr, + FFN_SILU, il); + cur_local = ggml_add(ctx0, cur_local, ffn_out); + } + + cur_local = build_norm(cur_local, model.mm_a_local_norm_w, nullptr, NORM_TYPE_RMS, eps, -1); + cb(cur_local, "after_local_transformer", -1); + + // flatten each group of `group_size` frames into one (group_size*n_embd)-dim vector + // (matching AudioProjection's flattened input) + ggml_tensor * grouped = ggml_reshape_2d(ctx0, cur_local, n_embd * group_size, n_groups); + + // AudioProjection: Linear (no bias) -> GELU -> Linear (no bias) + projected = build_ffn(grouped, + model.mm_1_w, nullptr, + nullptr, nullptr, + model.mm_2_w, nullptr, + FFN_GELU_ERF, -1); + cb(projected, "after_projection", -1); + } + + ggml_build_forward_expand(gf, projected); + return gf; +} diff --git a/tools/mtmd/models/models.h b/tools/mtmd/models/models.h index 2d7555da41d2..caed438ec513 100644 --- a/tools/mtmd/models/models.h +++ b/tools/mtmd/models/models.h @@ -210,6 +210,11 @@ struct clip_graph_qwen3a : clip_graph { ggml_cgraph * build() override; }; +struct clip_graph_mimo_audio : clip_graph { + clip_graph_mimo_audio(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} + ggml_cgraph * build() override; +}; + struct clip_graph_kimik25 : clip_graph { clip_graph_kimik25(clip_ctx * ctx, const clip_image_f32 & img) : clip_graph(ctx, img) {} ggml_cgraph * build() override; diff --git a/tools/mtmd/mtmd-audio.cpp b/tools/mtmd/mtmd-audio.cpp index b72fd067a508..ed68951c0151 100644 --- a/tools/mtmd/mtmd-audio.cpp +++ b/tools/mtmd/mtmd-audio.cpp @@ -725,6 +725,72 @@ bool mtmd_audio_preprocessor_qwen3a::preprocess(const float * sa return true; } +// +// mtmd_audio_preprocessor_mimo_audio +// +// Matches torchaudio.transforms.MelSpectrogram(power=1.0, center=True) followed by +// log(clip(spec, min=1e-7)): HTK mel scale, no Slaney area norm, magnitude (not power) +// spectrogram, natural log, reflect-padded by n_fft/2 on each side. +// + +void mtmd_audio_preprocessor_mimo_audio::initialize() { + cache.fill_sin_cos_table(hparams.audio_n_fft); + cache.fill_hann_window(hparams.audio_window_len, true); + cache.fill_mel_filterbank_matrix( + hparams.n_mel_bins, hparams.audio_n_fft, hparams.audio_sample_rate, + 0.0f, hparams.audio_sample_rate / 2.0f, + /*slaney_area_norm=*/ false, + /*scale=*/ 1.0f, + /*use_htk=*/ true + ); +} + +bool mtmd_audio_preprocessor_mimo_audio::preprocess(const float * samples, + size_t n_samples, + std::vector & output) { + if (n_samples == 0) { + return false; + } + + GGML_ASSERT(!cache.sin_vals.empty()); + GGML_ASSERT(!cache.cos_vals.empty()); + GGML_ASSERT(!cache.filters.data.empty()); + + const int pad = hparams.audio_n_fft / 2; + + std::vector padded(n_samples + 2 * pad, 0.0f); + for (int i = 0; i < pad; i++) { + int src = pad - i; + padded[i] = (src < (int)n_samples) ? samples[src] : 0.0f; + } + std::copy(samples, samples + n_samples, padded.begin() + pad); + for (int i = 0; i < pad; i++) { + int src = (int)n_samples - 2 - i; + padded[n_samples + pad + i] = (src >= 0) ? samples[src] : 0.0f; + } + + filter_params params; + params.n_mel = hparams.n_mel_bins; + params.n_fft_bins = 1 + (hparams.audio_n_fft / 2); + params.hann_window_size = hparams.audio_window_len; + params.hop_length = hparams.audio_hop_len; + params.sample_rate = hparams.audio_sample_rate; + params.no_padding = true; // reflect padding already applied above + params.use_natural_log = true; + params.use_magnitude = true; + params.mel_floor = 1e-7f; + params.norm_per_feature = false; + + mtmd_audio_mel out; + bool ok = log_mel_spectrogram(padded.data(), (int)padded.size(), 4, params, cache, out); + if (!ok) { + return false; + } + + output.push_back(std::move(out)); + return true; +} + // // mtmd_audio_preprocessor_conformer // diff --git a/tools/mtmd/mtmd-audio.h b/tools/mtmd/mtmd-audio.h index ad96bd847cfc..d8ec72b9d54e 100644 --- a/tools/mtmd/mtmd-audio.h +++ b/tools/mtmd/mtmd-audio.h @@ -111,6 +111,15 @@ struct mtmd_audio_preprocessor_qwen3a : mtmd_audio_preprocessor { mtmd_audio_cache cache; }; +struct mtmd_audio_preprocessor_mimo_audio : mtmd_audio_preprocessor { + mtmd_audio_preprocessor_mimo_audio(const clip_ctx * ctx) : mtmd_audio_preprocessor(ctx) {} + void initialize() override; + bool preprocess(const float * samples, size_t n_samples, std::vector & output) override; + + private: + mtmd_audio_cache cache; +}; + // // streaming ISTFT - converts spectrogram frames back to audio one frame at a time // diff --git a/tools/mtmd/mtmd.cpp b/tools/mtmd/mtmd.cpp index bb49b211efb3..6e61cf3e520b 100644 --- a/tools/mtmd/mtmd.cpp +++ b/tools/mtmd/mtmd.cpp @@ -730,6 +730,12 @@ struct mtmd_context { aud_end = ""; audio_preproc = std::make_unique(ctx_a); } break; + case PROJECTOR_TYPE_MIMO_AUDIO: + { + aud_beg = "<|mimo_audio_start|>"; + aud_end = "<|mimo_audio_end|>"; + audio_preproc = std::make_unique(ctx_a); + } break; default: throw std::runtime_error(string_format("%s: unexpected audio projector type %d\n", __func__, proj)); } diff --git a/tools/server/README-dev.md b/tools/server/README-dev.md index b4ec9f17d3c6..b41d70c63ac8 100644 --- a/tools/server/README-dev.md +++ b/tools/server/README-dev.md @@ -136,13 +136,13 @@ Producer side: `server_res_generator` extends `server_res_spipe`, which keeps al Lifetime safety: the session holds no back reference to the response, so `spipe` is a plain `unique_ptr` touched only by the http worker. `cancel` raises an atomic the producer polls; the producer finalizes the session from its destructor, which also runs `~server_response_reader::stop()` to cancel the generation at the queue level. A `DELETE` stops work by raising the flag and letting the worker unwind. -Consumer side: `GET /v1/stream/?from=N` opens a `text/event-stream` that replays buffered bytes from offset `N` and blocks for live bytes, so the browser reattaches like a fresh EventSource. An offset below the dropped prefix returns 400. +Consumer side: `GET /v1/stream?conv_id=&from=N` opens a `text/event-stream` that replays buffered bytes from offset `N` and blocks for live bytes, so the browser reattaches like a fresh EventSource. An offset below the dropped prefix returns 400. Routes: -- `GET /v1/stream/:conv_id?from=N`: replay or live reattach. +- `GET /v1/stream?conv_id=&from=N`: replay or live reattach. The id travels in the query string because it can embed a model name containing slashes. - `POST /v1/streams/lookup` with `{"conversation_ids": [...]}`: returns session status only for ids the caller already owns. There is no listing route, so live sessions cannot be enumerated (an earlier `GET /v1/streams` was removed for exactly this reason). -- `DELETE /v1/stream/:conv_id`: explicit Stop, idempotent (`evict_and_cancel`). +- `DELETE /v1/stream?conv_id=`: explicit Stop, idempotent (`evict_and_cancel`). Router mode binds the same paths to proxy handlers. A `conv_id -> child` map (`conv_models`), populated when a POST is routed, resolves the owning child in one lookup with no polling. The lookup groups ids per child; GET and DELETE proxy straight to the owner. This loopback REST hop is expected to move to a websocket IPC later, swapping only the transport. @@ -166,8 +166,8 @@ graph TD GC[GC thread] -- drop after TTL --> Sess end Sess -- read_from offset --> Cons[stream_pipe_consumer] - Cons -- "GET /v1/stream/:id?from=N" --> Client - DEL[DELETE /v1/stream/:id] -- evict_and_cancel --> Sess + Cons -- "GET /v1/stream?conv_id=id&from=N" --> Client + DEL[DELETE /v1/stream?conv_id=id] -- evict_and_cancel --> Sess ``` The diagram shows the buffer touch points. The live wire (chunks streamed to the original client during a normal generation) is the producer's default output, described under "Producer side" above. diff --git a/tools/server/README.md b/tools/server/README.md index d34565455455..25aacf9f516f 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -72,10 +72,10 @@ For the full list of features, please refer to [server's changelog](https://gith | `-ctv, --cache-type-v TYPE` | KV cache data type for V
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_CACHE_TYPE_V) | | `-dt, --defrag-thold N` | KV cache defragmentation threshold (DEPRECATED)
(env: LLAMA_ARG_DEFRAG_THOLD) | | `--rpc SERVERS` | comma-separated list of RPC servers (host:port)
(env: LLAMA_ARG_RPC) | -| `--mlock` | DEPRECATED in favor of `--load-mode`: mmap + force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | +| `--mlock` | DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing
(env: LLAMA_ARG_MLOCK) | | `--mmap, --no-mmap` | DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)
(env: LLAMA_ARG_MMAP) | | `-dio, --direct-io, -ndio, --no-direct-io` | DEPRECATED in favor of `--load-mode`: use DirectIO if available
(env: LLAMA_ARG_DIO) | -| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | +| `-lm, --load-mode MODE` | model loading mode (default: mmap)
- none: no special loading mode
- mmap: memory-map model (if mmap disabled, slower load but may reduce pageouts if not using mlock)
- mlock: force system to keep model in RAM rather than swapping or compressing
- mmap+mlock: mmap + force system to keep model in RAM rather than swapping or compressing
- dio: use DirectIO if available

(env: LLAMA_ARG_LOAD_MODE) | | `--numa TYPE` | attempt optimizations that help on some NUMA systems
- distribute: spread execution evenly over all nodes
- isolate: only spawn threads on CPUs on the node that execution started on
- numactl: use the CPU map provided by numactl
if run without this previously, it is recommended to drop the system page cache before using this
see https://github.com/ggml-org/llama.cpp/issues/1437
(env: LLAMA_ARG_NUMA) | | `-dev, --device ` | comma-separated list of devices to use for offloading (none = don't offload)
use --list-devices to see a list of available devices
(env: LLAMA_ARG_DEVICE) | | `--list-devices` | print list of available devices and exit | diff --git a/tools/server/server-models.cpp b/tools/server/server-models.cpp index 923b3533e9cb..188a72a3741c 100644 --- a/tools/server/server-models.cpp +++ b/tools/server/server-models.cpp @@ -1172,7 +1172,7 @@ bool server_models::ensure_model_ready(const std::string & name) { return true; } -server_http_res_ptr server_models::proxy_request(const server_http_req & req, const std::string & method, const std::string & name, bool update_last_used) { +server_http_res_ptr server_models::proxy_request(const server_http_req & req, const std::string & method, const std::string & name, bool update_last_used, bool detached) { auto meta = get_meta(name); if (!meta.has_value()) { throw std::runtime_error("model name=" + name + " is not found"); @@ -1198,7 +1198,10 @@ server_http_res_ptr server_models::proxy_request(const server_http_req & req, co req.headers, req.body, req.files, - req.should_stop, + // a detached request belongs to a replay session that outlives the client socket: + // it reaches the child even when the downstream died during the load wait, the + // session buffer is the recipient and DELETE remains the stop + detached ? std::function([]() { return false; }) : req.should_stop, base_params.timeout_read, base_params.timeout_write ); @@ -1469,13 +1472,9 @@ static bool router_validate_model(std::string & name, server_models & models, bo } // resolve alias to canonical model name name = meta->name; - if (models_autoload) { - models.ensure_model_ready(name); - } else { - if (!meta->is_running()) { - res_err(res, format_error_response("model is not loaded", ERROR_TYPE_INVALID_REQUEST)); - return false; - } + if (!models_autoload && !meta->is_running()) { + res_err(res, format_error_response("model is not loaded", ERROR_TYPE_INVALID_REQUEST)); + return false; } return true; } @@ -1568,6 +1567,9 @@ void server_models_routes::init_routes() { if (!router_validate_model(name, models, autoload, error_res)) { return error_res; } + if (autoload) { + models.ensure_model_ready(name); + } return models.proxy_request(req, method, name, false); }; @@ -1581,12 +1583,23 @@ void server_models_routes::init_routes() { return error_res; } // remember which child serves this conversation so the stream routes can route straight - // to it without polling, keyed on the exact conv id from the header + // to it without polling, keyed on the exact conv id from the header. registered before + // the load wait so a stop issued while the model loads can erase the entry and cancel + // this request instead of leaving an orphan generation std::string conv_id = server_stream_conv_id_from_headers(req.headers); - if (!conv_id.empty()) { - models.conv_models.remember(conv_id, name); + uint64_t ticket = models.conv_models.remember(conv_id, name); + bool waited = autoload && models.ensure_model_ready(name); + if (ticket != 0 && !models.conv_models.alive(conv_id, ticket)) { + SRV_INF("request for conv_id=%s cancelled while model name=%s was loading\n", + conv_id.c_str(), name.c_str()); + res_err(error_res, format_error_response( + "request cancelled by a stop while the model was loading", ERROR_TYPE_INVALID_REQUEST)); + return error_res; } - return models.proxy_request(req, method, name, true); // update last usage for POST request only + // a session request that waited for a load detaches from the client socket: the + // client may have dropped during the wait (page reload) and the session buffer must + // still receive the generation for a later resume + return models.proxy_request(req, method, name, true, waited && ticket != 0); // update last usage for POST request only }; this->post_router_models_load = [this](const server_http_req & req) { @@ -1779,7 +1792,7 @@ void server_models_routes::init_routes() { }; this->router_stream_get = [this](const server_http_req & req) { - // GET /v1/stream/?from=N. resolve the owning child from the conv_id -> model + // GET /v1/stream?conv_id=&from=N. resolve the owning child from the conv_id -> model // map, 404 when nothing maps auto res = std::make_unique(); std::string conv_id = req.get_param("conv_id"); @@ -1789,13 +1802,24 @@ void server_models_routes::init_routes() { } std::optional owner = resolve_child_for_conv(models, conv_id); if (!owner.has_value()) { - res_err(res, format_error_response("Stream not found or expired", ERROR_TYPE_NOT_FOUND)); + // a registered conv whose model is still loading earns a retry: the session appears + // once the load ends and the pending request reaches the child + auto tracked = models.conv_models.lookup(conv_id); + auto meta = tracked.has_value() ? models.get_meta(*tracked) : std::nullopt; + bool transient = meta.has_value() && (meta->status == SERVER_MODEL_STATUS_LOADING || + meta->status == SERVER_MODEL_STATUS_DOWNLOADING || + meta->status == SERVER_MODEL_STATUS_DOWNLOADED); + if (transient) { + res_err(res, format_error_response("Stream owner model is loading, retry later", ERROR_TYPE_UNAVAILABLE)); + } else { + res_err(res, format_error_response("Stream not found or expired", ERROR_TYPE_NOT_FOUND)); + } return res; } std::string from = req.get_param("from"); - std::string child_path = "/v1/stream/" + encode_qs(conv_id); + std::string child_path = "/v1/stream?conv_id=" + encode_qs(conv_id); if (!from.empty()) { - child_path += "?from=" + from; + child_path += "&from=" + from; } SRV_TRC("proxying stream resume to model %s on port %d, path=%s\n", owner->name.c_str(), owner->port, child_path.c_str()); @@ -1875,7 +1899,7 @@ void server_models_routes::init_routes() { }; this->router_stream_delete = [this](const server_http_req & req) { - // DELETE /v1/stream/. resolve the owning child via the map and forward only to + // DELETE /v1/stream?conv_id=. resolve the owning child via the map and forward only to // it, evict_and_cancel is idempotent on the child auto res = std::make_unique(); std::string conv_id = req.get_param("conv_id"); @@ -1883,7 +1907,7 @@ void server_models_routes::init_routes() { res_err(res, format_error_response("Missing conversation id in path", ERROR_TYPE_INVALID_REQUEST)); return res; } - std::string child_path = "/v1/stream/" + encode_qs(conv_id); + std::string child_path = "/v1/stream?conv_id=" + encode_qs(conv_id); auto owner = resolve_child_for_conv(models, conv_id); if (owner.has_value()) { httplib::Client cli(CHILD_ADDR, owner->port); @@ -1892,6 +1916,11 @@ void server_models_routes::init_routes() { cli.set_write_timeout(0, STREAM_LOOKUP_TIMEOUT_MS * 1000); auto resp = cli.Delete(child_path.c_str()); (void) resp; // the child logs its own miss when the session is unknown there + } else if (auto tracked = models.conv_models.lookup(conv_id); tracked.has_value()) { + // the entry exists but its model is still loading: the forget below erases it, + // which cancels the request parked in proxy_post before the generation starts + SRV_INF("router stop for conv_id=%s while model name=%s is loading, cancelling the pending request\n", + conv_id.c_str(), tracked->c_str()); } else { SRV_WRN("router stop for unknown conv_id=%s, no owning child in the conv map\n", conv_id.c_str()); diff --git a/tools/server/server-models.h b/tools/server/server-models.h index 62bed8725b5b..614798186cfc 100644 --- a/tools/server/server-models.h +++ b/tools/server/server-models.h @@ -134,12 +134,24 @@ struct server_models { // proxy_request forwards a POST carrying an X-Conversation-Id. best effort: a stale entry just // makes the child answer not found and the client recovers. owns its lock, one mutex per struct struct conv_model_tracker { - void remember(const std::string & conv_id, const std::string & model) { + // returns the ticket of this registration, 0 when nothing was registered. erasing or + // replacing the entry invalidates the ticket, which is how a stop cancels a request + // parked in the model load wait + uint64_t remember(const std::string & conv_id, const std::string & model) { if (conv_id.empty() || model.empty()) { - return; + return 0; } std::lock_guard lock(mu); - map[conv_id] = model; + uint64_t ticket = next_ticket++; + map[conv_id] = { model, ticket }; + return ticket; + } + + // false means a stop erased the entry or a newer request replaced it + bool alive(const std::string & conv_id, uint64_t ticket) { + std::lock_guard lock(mu); + auto it = map.find(conv_id); + return it != map.end() && it->second.ticket == ticket; } std::optional lookup(const std::string & conv_id) { @@ -151,7 +163,7 @@ struct server_models { if (it == map.end()) { return std::nullopt; } - return it->second; + return it->second.model; } void forget(const std::string & conv_id) { @@ -163,8 +175,13 @@ struct server_models { } private: - std::mutex mu; - std::unordered_map map; + struct entry_t { + std::string model; + uint64_t ticket; + }; + std::mutex mu; + uint64_t next_ticket = 1; + std::unordered_map map; }; common_preset_context ctx_preset; @@ -249,7 +266,7 @@ struct server_models { bool ensure_model_ready(const std::string & name); // proxy an HTTP request to the model instance - server_http_res_ptr proxy_request(const server_http_req & req, const std::string & method, const std::string & name, bool update_last_used); + server_http_res_ptr proxy_request(const server_http_req & req, const std::string & method, const std::string & name, bool update_last_used, bool detached = false); // handle message sent from server_child::notify_to_router() // raw input must starts with CMD_CHILD_TO_ROUTER_STATE, followed by a JSON string diff --git a/tools/server/server-stream.cpp b/tools/server/server-stream.cpp index f0a35b18e525..f6b9b8a9f4cc 100644 --- a/tools/server/server-stream.cpp +++ b/tools/server/server-stream.cpp @@ -453,7 +453,7 @@ static server_http_res_ptr make_error_response(int status, const std::string & m server_http_context::handler_t server_stream_make_get_handler() { return [](const server_http_req & req) -> server_http_res_ptr { - // GET /v1/stream/?from=N replays buffered SSE bytes then blocks for live + // GET /v1/stream?conv_id=&from=N replays buffered SSE bytes then blocks for live // bytes until the session finalizes, streamed as text/event-stream for EventSource std::string conv_id = req.get_param("conv_id"); if (conv_id.empty()) { @@ -560,13 +560,13 @@ server_http_context::handler_t server_stream_make_lookup_handler() { server_http_context::handler_t server_stream_make_delete_handler() { return [](const server_http_req & req) -> server_http_res_ptr { - // DELETE /v1/stream/ is the explicit user Stop, cancels the producer and evicts + // DELETE /v1/stream?conv_id= is the explicit user Stop, cancels the producer and evicts // the buffer. idempotent, returns 204 even if the session was already gone std::string conv_id = req.get_param("conv_id"); if (conv_id.empty()) { return make_error_response(400, "Missing conversation id in path", ERROR_TYPE_INVALID_REQUEST); } - SRV_TRC("DELETE /v1/stream/%s -> evict_and_cancel\n", conv_id.c_str()); + SRV_TRC("DELETE /v1/stream conv_id=%s -> evict_and_cancel\n", conv_id.c_str()); g_stream_sessions.evict_and_cancel(conv_id); auto res = std::make_unique(); res->status = 204; @@ -621,7 +621,7 @@ bool server_res_spipe::conn_alive() { bool server_res_spipe::should_stop() { if (spipe) { - // note: if DELETE /v1/stream/ is called, is_cancelled() will be true + // note: if DELETE /v1/stream is called for this conv, is_cancelled() will be true return spipe->is_cancelled(); } else { return !conn_alive(); diff --git a/tools/server/server-stream.h b/tools/server/server-stream.h index 9753140dd601..1e7461285f43 100644 --- a/tools/server/server-stream.h +++ b/tools/server/server-stream.h @@ -45,7 +45,13 @@ void server_stream_session_manager_start(); void server_stream_session_manager_stop(); // route handler factories wired under /v1/stream/* by server.cpp +// child-side handlers for the resumable stream routes. the conv id travels in the conv_id +// query string because it can embed a model name containing slashes (org/repo), which the +// decoded path would split before the param is captured server_http_context::handler_t server_stream_make_get_handler(); +// POST /v1/streams/lookup with body {"conversation_ids": [...]}: only answers for ids the +// caller already owns (the WebUI passes the convs visible in its sidebar), the server never +// lists ids it has not been asked about, so a random caller cannot enumerate live sessions server_http_context::handler_t server_stream_make_lookup_handler(); server_http_context::handler_t server_stream_make_delete_handler(); diff --git a/tools/server/server.cpp b/tools/server/server.cpp index b6fef99e8747..a3b2a8b0fe1b 100644 --- a/tools/server/server.cpp +++ b/tools/server/server.cpp @@ -272,10 +272,8 @@ int llama_server(common_params & params, int argc, char ** argv) { ctx_http.get ("/slots", ex_wrapper(routes.get_slots)); ctx_http.post("/slots/:id_slot", ex_wrapper(routes.post_slots)); - // resumable streaming, the conversation_id is the session identity end to end. router and - // child wire different handlers under the same paths: a child binds the local session - // factories, the router binds proxies that resolve the owning child through the - // conv_id -> model map + // resumable streaming: a child binds the local session factories, the router binds + // proxies that resolve the owning child, see server-stream.h server_http_context::handler_t stream_get_h; server_http_context::handler_t streams_lookup_h; server_http_context::handler_t stream_delete_h; @@ -288,12 +286,9 @@ int llama_server(common_params & params, int argc, char ** argv) { streams_lookup_h = server_stream_make_lookup_handler(); stream_delete_h = server_stream_make_delete_handler(); } - ctx_http.get ("/v1/stream/:conv_id", ex_wrapper(stream_get_h)); - // POST /v1/streams/lookup with body {"conversation_ids": [...]}. you can only ask for ids - // you already own (the WebUI passes the convs visible in its sidebar). the server never - // lists ids it has not been asked about, so a random caller cannot enumerate live sessions + ctx_http.get ("/v1/stream", ex_wrapper(stream_get_h)); ctx_http.post("/v1/streams/lookup", ex_wrapper(streams_lookup_h)); - ctx_http.del ("/v1/stream/:conv_id", ex_wrapper(stream_delete_h)); + ctx_http.del ("/v1/stream", ex_wrapper(stream_delete_h)); // Google Cloud Platform (Vertex AI) compat ctx_http.register_gcp_compat(); diff --git a/tools/server/tests/unit/test_stream.py b/tools/server/tests/unit/test_stream.py new file mode 100644 index 000000000000..a1ef55567bc7 --- /dev/null +++ b/tools/server/tests/unit/test_stream.py @@ -0,0 +1,153 @@ +import json +import socket +import threading +import time +from urllib.parse import quote +import pytest +from utils import * + +server: ServerProcess + +# a model name with slashes exercises the query string routing of the stream routes: the id +# cannot travel as a path param because the decoded slash would split it before capture +MODEL = "ggml-org/tinygemma3-GGUF:Q8_0" +STREAM_ID = f"conv-stream-test::{MODEL}" +QS = "conv_id=" + quote(STREAM_ID, safe="") + + +@pytest.fixture(autouse=True) +def create_server(): + global server + server = ServerPreset.router() + + +def test_stream_resume_and_stop_with_slashed_model_name(): + global server + server.start() + + content = "" + for data in server.make_stream_request("POST", "/chat/completions", data={ + "model": MODEL, + "stream": True, + "max_tokens": 16, + "messages": [{"role": "user", "content": "hello"}], + }, headers={"X-Conversation-Id": STREAM_ID}): + if data["choices"]: + content += data["choices"][0]["delta"].get("content") or "" + assert len(content) > 0 + + # the finished session replays from the beginning through the router + res = server.make_request("GET", f"/v1/stream?{QS}&from=0") + assert res.status_code == 200 + assert "data: " in str(res.body) + + # the explicit stop reaches the owning child and evicts the session + res = server.make_request("DELETE", f"/v1/stream?{QS}") + assert res.status_code == 204 + res = server.make_request("GET", f"/v1/stream?{QS}&from=0") + assert res.status_code == 404 + + +def test_stream_stop_during_model_load(): + global server + server.start() + + thread_error: list[ServerError] = [] + thread_done = threading.Event() + + def fire_post(): + try: + for _ in server.make_stream_request("POST", "/chat/completions", data={ + "model": MODEL, + "stream": True, + "max_tokens": 512, + "messages": [{"role": "user", "content": "Count from 1 to 1000."}], + }, headers={"X-Conversation-Id": STREAM_ID}): + pass + except ServerError as e: + thread_error.append(e) + finally: + thread_done.set() + + t = threading.Thread(target=fire_post) + t.start() + + # catch the autoload window, tiny models load fast so poll aggressively + saw_loading = False + deadline = time.time() + 5.0 + while time.time() < deadline and not thread_done.is_set(): + res = server.make_request("GET", "/models") + status = next(m["status"]["value"] for m in res.body["data"] if m["id"] == MODEL) + if status == "loading": + saw_loading = True + break + time.sleep(0.002) + if not saw_loading: + t.join() + pytest.skip("load window too short to be observed on this machine") # ty: ignore[too-many-positional-arguments] + + # a stop during the load cancels the parked request instead of leaving an orphan + res = server.make_request("DELETE", f"/v1/stream?{QS}") + assert res.status_code == 204 + assert thread_done.wait(timeout=60) + t.join() + assert len(thread_error) == 1 + assert thread_error[0].code == 400 + assert "cancelled" in json.dumps(thread_error[0].body) + res = server.make_request("GET", f"/v1/stream?{QS}&from=0") + assert res.status_code == 404 + + +def test_stream_resumes_after_reload_during_model_load(): + global server + server.start() + + # raw socket client so the connection can be dropped mid load like a page reload + body = json.dumps({ + "model": MODEL, + "stream": True, + "max_tokens": 16, + "messages": [{"role": "user", "content": "hello"}], + }) + request = ( + f"POST /v1/chat/completions HTTP/1.1\r\n" + f"Host: {server.server_host}:{server.server_port}\r\n" + f"Content-Type: application/json\r\n" + f"X-Conversation-Id: {STREAM_ID}\r\n" + f"Content-Length: {len(body)}\r\n" + f"Connection: close\r\n\r\n{body}" + ) + sock = socket.create_connection((server.server_host, server.server_port)) + sock.sendall(request.encode()) + + # drop the client while the model loads, poll aggressively to catch the window + saw_loading = False + saw_503 = False + deadline = time.time() + 5.0 + while time.time() < deadline: + res = server.make_request("GET", "/models") + status = next(m["status"]["value"] for m in res.body["data"] if m["id"] == MODEL) + if status == "loading": + saw_loading = True + break + if status == "loaded": + break + time.sleep(0.002) + sock.close() + if not saw_loading: + pytest.skip("load window too short to be observed on this machine") # ty: ignore[too-many-positional-arguments] + + # while the model loads the resume route answers retry later, then the session appears, + # receives the whole generation despite the dead client, and replays from the beginning + deadline = time.time() + 60.0 + replay = None + while time.time() < deadline: + res = server.make_request("GET", f"/v1/stream?{QS}&from=0") + if res.status_code == 503: + saw_503 = True + elif res.status_code == 200 and "data: " in str(res.body): + replay = res + break + time.sleep(0.1) + assert saw_503, "resume during the load did not answer 503" + assert replay is not None, "session never became resumable after the client disconnect" diff --git a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte index 00578fcf1a1a..199d75fcec95 100644 --- a/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageAssistant/ChatMessageAssistant.svelte @@ -10,7 +10,7 @@ } from '$lib/components/app'; import { getMessageEditContext } from '$lib/contexts'; import { useProcessingState } from '$lib/hooks/use-processing-state.svelte'; - import { isLoading, isChatStreaming } from '$lib/stores/chat.svelte'; + import { chatStore, isLoading, isChatStreaming } from '$lib/stores/chat.svelte'; import { modelLoadProgressText } from '$lib/utils'; import { MessageRole } from '$lib/enums'; import { config } from '$lib/stores/settings.svelte'; @@ -82,8 +82,11 @@ let hasNoContent = $derived(!message?.content?.trim()); let isActivelyProcessing = $derived(isCurrentlyLoading || isStreaming); - // during a router auto-load the message has no model yet, so target the selected one - let loadTargetModel = $derived(message.model ?? modelsStore.selectedModelName); + // during a router auto-load the message has no model yet: target the model frozen in the + // persisted stream state (survives a reload), then fall back to the dropdown selection + let loadTargetModel = $derived( + message.model ?? chatStore.getResumeModel(message.convId) ?? modelsStore.selectedModelName + ); let modelLoadProgress = $derived( isRouter && loadTargetModel ? modelsStore.getLoadProgress(loadTargetModel) : null ); diff --git a/tools/ui/src/lib/constants/api-endpoints.ts b/tools/ui/src/lib/constants/api-endpoints.ts index 37137c1c7699..ab35708a46da 100644 --- a/tools/ui/src/lib/constants/api-endpoints.ts +++ b/tools/ui/src/lib/constants/api-endpoints.ts @@ -21,7 +21,11 @@ export const API_TOOLS = { EXECUTE: '/tools' }; -// resumable stream routes, the conv::model identity is appended as a path segment +// resumable stream routes, the conv::model identity travels as the conv_id query param +// because model names can contain slashes that a path segment cannot carry +// resume retry cadence while the owning model is still loading (server answers 503) +export const STREAM_RESUME_RETRY_MS = 2000; + export const API_STREAM = { BASE: './v1/stream', LOOKUP: './v1/streams/lookup' diff --git a/tools/ui/src/lib/constants/sandbox.ts b/tools/ui/src/lib/constants/sandbox.ts index 58242678d98b..381621de647e 100644 --- a/tools/ui/src/lib/constants/sandbox.ts +++ b/tools/ui/src/lib/constants/sandbox.ts @@ -14,12 +14,15 @@ export const SANDBOX_EMPTY_OUTPUT = '(no output)'; export const SANDBOX_TRUNCATION_NOTICE = '[output truncated]'; const NERDAMER_DESCRIPTION = ` -Symbolic/numeric math via \`nerdamer\` (pre-loaded, do not require, use it directly). -nerdamer('diff(sin(x)/x,x)') or nerdamer.diff('sin(x)/x','x') → Expression; convert with .toString()/.text()/.toTeX(), or .evaluate() (→ still Expression, then .toString()). -nerdamer(expr,{x:2}) substitutes only; chain .evaluate() or pass 'numer' for numeric result. -solve(expr,var)→Symbol[]; solveEquations([eq1,..])→[[var,val],..] pairs. -Functions: simplify/expand/factor(expr), diff(expr,var[,n]), integrate(expr,var), defint(expr,from,to,var), limit(expr,var,to), laplace(expr,t,s), ilt(expr,s,t), gcd/lcm(a,b), roots/coeffs/partfrac(expr,var), pfactor(n), numer/decimals/erf(expr), product/sum(expr,var,from,to), mean/median/stdev/variance(...vals). -Object.keys(nerdamer).filter(k=>typeof nerdamer[k]==='function') lists all available functions. If you need a function not documented above, list them first — do not guess function names.`; +Symbolic/numeric math via \`nerdamer\` +nerdamer(expr,subs?,opts?)/nerdamer.func(...)→Expression Format via .text(fmt?) (fmt: 'decimals'|'fractions'|'scientific') eval via .evaluate(subs?) +nerdamer(expr,{x:2}) substitutes numeric via opts 'numer' or .evaluate() +simplify/expand/factor(expr) div/gcd/lcm(...) coeffs/partfrac(expr,var) +diff/integrate(expr,var) defint(expr,lo,hi,var?) sum/product(expr,var,lo,hi) limit(expr,var,pt) +solve(expr,var) solveEquations([eq1,eq2],[var1,var2]) +polarform/rectform/arg/realpart/imagpart(z) +set/get Var/Constant(name,val?) setFunction(name,[params],body) +IMPORTANT:Identifier 'nerdamer' has already been declared, use it directly`; /** * Build the sandbox tool definition. When `includeSymbolicMath` is true, diff --git a/tools/ui/src/lib/services/chat.service.ts b/tools/ui/src/lib/services/chat.service.ts index d2455614fbaa..4ce396533de8 100644 --- a/tools/ui/src/lib/services/chat.service.ts +++ b/tools/ui/src/lib/services/chat.service.ts @@ -343,6 +343,9 @@ export class ChatService { // model the ::model suffix keeps the per model session distinct if (stream && conversationId) { headers['X-Conversation-Id'] = streamIdentity(conversationId, options.model); + // persist the pending stream before the fetch: a reload during the model load or + // the prompt processing must still find its way back to the session once it exists + ChatService.saveStreamState(conversationId, 0, options.model ?? null); } const response = await fetch(API_CHAT.COMPLETIONS, { @@ -353,6 +356,11 @@ export class ChatService { }); if (!response.ok) { + // a rejected request (including one cancelled by a stop during the model load) + // leaves nothing to resume + if (conversationId) { + ChatService.clearStreamState(conversationId); + } const error = await ChatService.parseErrorResponse(response); if (onError) { @@ -512,7 +520,7 @@ export class ChatService { if (!conversationId) return; try { const id = streamIdentity(conversationId, model); - await fetch(`${API_STREAM.BASE}/${encodeURIComponent(id)}`, { + await fetch(`${API_STREAM.BASE}?conv_id=${encodeURIComponent(id)}`, { method: 'DELETE', headers: getAuthHeaders() }); @@ -605,6 +613,26 @@ export class ChatService { * existing SSE parser drains it like a fresh stream. The server returns 200 on success, 404 if * no session exists for the conv_id, and 400 if the offset is below the dropped prefix. */ + // probe the resume route status without consuming the stream: the SSE route has no HEAD, + // so issue the GET and abort it right after the status line. 0 on network error + static async probeResumeStatus(streamId: string): Promise { + if (!streamId) return 0; + const ac = new AbortController(); + try { + const resp = await fetch( + `${API_STREAM.BASE}?conv_id=${encodeURIComponent(streamId)}&from=0`, + { + headers: getAuthHeaders(), + signal: ac.signal + } + ); + ac.abort(); + return resp.status; + } catch { + return 0; + } + } + static async resumeStream( conversationId: string, signal?: AbortSignal, @@ -614,7 +642,7 @@ export class ChatService { const state = ChatService.getStreamState(conversationId); const from = state?.bytesReceived ?? 0; const id = streamIdentity(conversationId, model); - const url = `${API_STREAM.BASE}/${encodeURIComponent(id)}?from=${from}`; + const url = `${API_STREAM.BASE}?conv_id=${encodeURIComponent(id)}&from=${from}`; return await fetch(url, { method: 'GET', signal, headers: getAuthHeaders() }); } diff --git a/tools/ui/src/lib/stores/chat.svelte.ts b/tools/ui/src/lib/stores/chat.svelte.ts index 5cbfe213b104..222723ab108d 100644 --- a/tools/ui/src/lib/stores/chat.svelte.ts +++ b/tools/ui/src/lib/stores/chat.svelte.ts @@ -14,6 +14,7 @@ import { SvelteMap, SvelteSet } from 'svelte/reactivity'; import { DatabaseService } from '$lib/services/database.service'; import { ChatService } from '$lib/services/chat.service'; +import { STREAM_RESUME_RETRY_MS } from '$lib/constants/api-endpoints'; import { streamIdentity } from '$lib/utils/stream-identity'; import { getAuthHeaders } from '$lib/utils/api-headers'; import { CONTENT_TYPE_HEADER } from '$lib/constants'; @@ -78,7 +79,7 @@ class ChatStore { // true while the active conversation streams reasoning content but no visible content yet isReasoning = $state(false); // resumable stream connection state for the active conversation - // streaming -> bytes flowing normally, resuming -> waiting on /v1/stream/:id reconnect, lost -> unrecoverable + // streaming -> bytes flowing normally, resuming -> waiting on /v1/stream reconnect, lost -> unrecoverable streamConnectionState = $state(StreamConnectionState.STREAMING); chatLoadingStates = new SvelteMap(); chatReasoningStates = new SvelteMap(); @@ -94,6 +95,11 @@ class ChatStore { // off when one conv finishes while another is still streaming. mirrors chatLoadingStates // in scope but tracks the attach + tee replay path specifically private attachingConvs = new SvelteSet(); + // pending resume retry timers while an owning model loads, one per conv + private resumeRetryTimers = new SvelteMap>(); + // convs whose resume waits on a model load: their loading state belongs to the retry loop, + // so discoverActiveStream must not treat it as a live send and bail + private resumePendingConvs = new SvelteSet(); // in-flight discoverActiveStream guard, keyed by conv id private discoveringConvs = new SvelteSet(); private abortControllers = new SvelteMap(); @@ -263,7 +269,7 @@ class ChatStore { const id = streamId || streamIdentity(convId, selectedModelName()); let response: Response; try { - response = await fetch(`./v1/stream/${encodeURIComponent(id)}?from=0`, { + response = await fetch(`./v1/stream?conv_id=${encodeURIComponent(id)}&from=0`, { headers: getAuthHeaders() }); } catch (e) { @@ -438,13 +444,22 @@ class ChatStore { } } + /** + * Model frozen at send time for a stream awaiting resume, from the persisted stream state. + * The load progress indicator targets it after a reload, when the message row has no model + * yet and the dropdown selection may not be restored. + */ + getResumeModel(convId: string): string | null { + return ChatService.getStreamState(convId)?.model ?? null; + } + async discoverActiveStream(convId: string): Promise { if (!convId) return; if (this.chatStreamingStates.has(convId)) return; - if (this.chatLoadingStates.get(convId)) return; + if (this.chatLoadingStates.get(convId) && !this.resumePendingConvs.has(convId)) return; // concurrency guard: another discover may already be running for this conv (typical race // between mount and visibilitychange on tab switch). a second concurrent fetch on the same - // /v1/stream/ would duplicate every byte into the DB message, this guard bounces it + // /v1/stream would duplicate every byte into the DB message, this guard bounces it if (this.discoveringConvs.has(convId)) return; this.discoveringConvs.add(convId); @@ -470,6 +485,38 @@ class ChatStore { if (!localState) { return; } + // quiet status probe first: a full attach flips the loading UI on every try, probing + // keeps the retry loop invisible while the owning model is still loading (503) + const status = await ChatService.probeResumeStatus(streamId); + if (status === 503) { + // make the wait visible: the empty assistant row persisted at send time renders + // the processing info, whose model load percentage flows from the models feed + this.resumePendingConvs.add(convId); + this.setChatLoading(convId, true); + if (!this.resumeRetryTimers.has(convId)) { + this.resumeRetryTimers.set( + convId, + setTimeout(() => { + this.resumeRetryTimers.delete(convId); + void this.discoverActiveStream(convId); + }, STREAM_RESUME_RETRY_MS) + ); + } + return; + } + if (this.resumePendingConvs.delete(convId) && status !== 200) { + // the wait is over without a session to attach, drop the visible loading state + this.setChatLoading(convId, false); + } + if (status === 0) { + // transient network failure, the next mount or visibility change retries + return; + } + if (status !== 200) { + // the session is gone (stopped, TTL expired), nothing to resume anymore + ChatService.clearStreamState(convId); + return; + } await this.attachServerStream(convId, streamId); // if attachServerStream failed (session gone, TTL expired), clear the local state to avoid retrying forever if (!this.chatStreamingStates.has(convId) && !this.chatLoadingStates.get(convId)) { @@ -1469,8 +1516,16 @@ class ChatStore { // detached drain keeps producing tokens until eos or max_tokens. use the frozen identity // captured when the session started, not the live dropdown const streamStateForStop = this.chatStreamingStates.get(convId); - const modelForStop = streamStateForStop?.model; + const modelForStop = streamStateForStop?.model ?? ChatService.getStreamState(convId)?.model; void ChatService.cancelServerStream(convId, modelForStop); + // an explicit stop leaves nothing to resume and kills a pending resume retry + ChatService.clearStreamState(convId); + const retryTimer = this.resumeRetryTimers.get(convId); + if (retryTimer !== undefined) { + clearTimeout(retryTimer); + this.resumeRetryTimers.delete(convId); + } + this.resumePendingConvs.delete(convId); this.abortRequest(convId); this.setChatLoading(convId, false); this.clearChatStreaming(convId);