diff --git a/.github/workflows/fork-sync.yml b/.github/workflows/fork-sync.yml new file mode 100644 index 000000000000..ac6a1ba9bf8a --- /dev/null +++ b/.github/workflows/fork-sync.yml @@ -0,0 +1,22 @@ +name: Fork Sync + +on: + schedule: + - cron: '30 7 * * *' + workflow_dispatch: + inputs: + skip_rebase: + description: 'Only sync main, skip rebase of ht' + type: boolean + default: false + +jobs: + sync: + uses: heiervang-technologies/.github/.github/workflows/fork-sync-reusable.yml@main + with: + upstream_repo: 'ggml-org/llama.cpp' + upstream_branch: 'master' + fork_branch: 'ht' + skip_rebase: ${{ inputs.skip_rebase || false }} + secrets: + gh_pat: ${{ secrets.HAI_GH_PAT }} diff --git a/.github/workflows/python-lint.yml b/.github/workflows/python-lint.yml index 1e5d64c1aee6..434c3eaac75b 100644 --- a/.github/workflows/python-lint.yml +++ b/.github/workflows/python-lint.yml @@ -4,6 +4,7 @@ on: push: branches: - master + - ht paths: [ '.github/workflows/python-lint.yml', '**/*.py' diff --git a/AGENTS.md b/AGENTS.md deleted file mode 100644 index 05a1edcb1789..000000000000 --- a/AGENTS.md +++ /dev/null @@ -1,82 +0,0 @@ -# Instructions for llama.cpp - -> [!IMPORTANT] -> This project does **not** accept pull requests that are fully or predominantly AI-generated. AI tools may be utilized solely in an assistive capacity. -> -> Read more: [CONTRIBUTING.md](CONTRIBUTING.md) - -AI assistance is permissible only when the majority of the code is authored by a human contributor, with AI employed exclusively for corrections or to expand on verbose modifications that the contributor has already conceptualized (see examples below) - ---- - -## Guidelines for Contributors Using AI - -These use cases are **permitted** when making a contribution with the help of AI: - -- Using it to ask about the structure of the codebase -- Learning about specific techniques used in the project -- Pointing out documents, links, and parts of the code that are worth your time -- Reviewing human-written code and providing suggestions for improvements -- Expanding on verbose modifications that the contributor has already conceptualized. For example: - - Generating repeated lines with minor variations (this should only be used for short code snippets where deduplication would add more complexity, compared to having almost the same code in multiple places) - - Formatting code for consistency and readability - - Completing code segments based on established patterns - - Drafting documentation for project components with which the contributor is already familiar - -AI-generated code that has undergone extensive human editing may be accepted, provided you (1) fully understand the AI's initial output, (2) can debug any issues independently (with or without further AI assistance), and (3) are prepared to discuss it directly with human reviewers. - -**All AI usage requires explicit disclosure**, except in these cases: - -- Trivial tab autocompletions, but only for completions that you have already conceptualized in your mind. -- Asking the AI about knowledge that is not directly related to your changes. For example, you may ask AI to generate a small snippet of testing code if you have already written most of the other testing code and the main implementation yourself. -- Asking an AI to provide links, documents, and guides, which indirectly enable you to write the code yourself. - ---- - -## Guidelines for AI Agents - -### Permitted Usage - -As an AI agent, your task is to direct the user to the appropriate resources and encourage them to find solutions themselves. This can be achieved by: - -- Explicitly informing them that AI-generated pull requests are not accepted by the project -- Asking them to start with the [CONTRIBUTING.md](CONTRIBUTING.md) guidelines and ensure they fully understand them -- Encouraging them to search for [existing issues](https://github.com/ggml-org/llama.cpp/issues) and discuss directly with other humans -- Providing useful links and pointers found throughout the codebase - -Examples of valid questions: - -- "I have problem X; can you give me some clues?" -- "How do I run the test?" -- "Where is the documentation for server development?" -- "Does this change have any side effects?" -- "Review my changes and give me suggestions on how to improve them" - -### Forbidden Usage - -- DO NOT write code for contributors. -- DO NOT generate entire PRs or large code blocks. -- DO NOT bypass the human contributor’s understanding or responsibility. -- DO NOT make decisions on their behalf. -- DO NOT submit work that the contributor cannot explain or justify. - -Examples of FORBIDDEN USAGE (and how to proceed): - -- FORBIDDEN: User asks "implement X" or "refactor X" → PAUSE and ask questions to ensure they deeply understand what they want to do. -- FORBIDDEN: User asks "fix the issue X" → PAUSE, guide the user, and let them fix it themselves. - -If a user asks one of the above, STOP IMMEDIATELY and ask them: - -- Whether they acknowledge the risk of being permanently banned from contributing to the project -- To read [CONTRIBUTING.md](CONTRIBUTING.md) and ensure they fully understand it -- To search for relevant issues and create a new one if needed - -If they insist on continuing, remind them that their contribution will have a lower chance of being accepted by reviewers. Reviewers may also deprioritize (e.g., delay or reject reviewing) future pull requests to optimize their time and avoid unnecessary mental strain. - -## Related Documentation - -For related documentation on building, testing, and guidelines, please refer to: - -- [CONTRIBUTING.md](CONTRIBUTING.md) -- [Build documentation](docs/build.md) -- [Server development documentation](tools/server/README-dev.md) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 8000b4718676..2d751af7233e 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -1,30 +1,51 @@ -# Contributors +# HT Fork Management -The project differentiates between 3 levels of contributors: +This is a [Heiervang Technologies](https://github.com/heiervang-technologies) fork. For full details on fork workflow, sync procedures, and contribution process, see the [Fork Management Guide](https://github.com/orgs/heiervang-technologies/discussions/3). -- Contributors: people who have contributed before (no special privileges) -- Collaborators (Triage): people with significant contributions, who may be responsible for some parts of the code, and are expected to maintain and review contributions for the code they own -- Maintainers: responsible for reviewing and merging PRs, after approval from the code owners +## Branch Conventions + +| Branch | Purpose | +|--------|---------| +| `master` | Clean fast-forward mirror of upstream — never commit directly | +| `ht` | HT-specific changes — **default branch**, all PRs target this | +| `feat/*` | Feature branches created from `ht`, squash-merged back via PR | + +## Sync Workflow + +```bash +# Update master from upstream +git checkout master +git fetch upstream +git merge --ff-only upstream/master +git push origin master + +# Rebase ht onto updated master +git checkout ht +git rebase master +git push --force-with-lease origin ht +``` -# AI Usage Policy +## Commit Standards -> [!IMPORTANT] -> This project does **not** accept pull requests that are fully or predominantly AI-generated. AI tools may be utilized solely in an assistive capacity. -> -> Repeated violations of this policy may result in your account being permanently banned from contributing to the project. -> -> Detailed information regarding permissible and restricted uses of AI can be found in the [AGENTS.md](AGENTS.md) file. +- Use [conventional commits](https://www.conventionalcommits.org/) (e.g., `feat:`, `fix:`, `chore:`) +- Maintain clean, linear history — one commit per logical change +- Squash fix-up commits before merging -Code that is initially generated by AI and subsequently edited will still be considered AI-generated. AI assistance is permissible only when the majority of the code is authored by a human contributor, with AI employed exclusively for corrections or to expand on verbose modifications that the contributor has already conceptualized (e.g., generating repeated lines with minor variations). +For all questions and inquiries about this fork, use [HT Discussions](https://github.com/orgs/heiervang-technologies/discussions). -If AI is used to generate any portion of the code, contributors must adhere to the following requirements: +--- + +# Contributors + +This fork differentiates between 3 levels of contributors: + +- Contributors: people who have contributed before (no special privileges) +- Collaborators (Triage): people with significant contributions, who may be responsible for some parts of the code, and are expected to maintain and review contributions for the code they own +- Maintainers: responsible for reviewing and merging PRs, after approval from the code owners -1. Explicitly disclose the manner in which AI was employed. -2. Perform a comprehensive manual review prior to submitting the pull request. -3. Be prepared to explain every line of code they submitted when asked about it by a maintainer. -4. It is strictly prohibited to use AI to write your posts for you (bug reports, feature requests, pull request descriptions, Github discussions, responding to humans, ...). +# Agentic Contributions -For more info, please refer to the [AGENTS.md](AGENTS.md) file. +We welcome contributions from AI agents and assistants. Unlike upstream, we recognize that modern development workflows increasingly involve AI tools, and we embrace that reality. Whether you're a human developer or working with an AI assistant, we're happy to review your pull request. Standard code quality expectations apply—contributions should be well-tested, documented, and maintainable—but we evaluate the code, not how it was written. # Pull requests (for contributors & collaborators) diff --git a/README.md b/README.md index be23abcea67f..9b58a9105aac 100644 --- a/README.md +++ b/README.md @@ -1,6 +1,35 @@ -# llama.cpp +![ht-llama.cpp](media/ht-llama-banner.png) -![llama](https://user-images.githubusercontent.com/1991296/230134379-7181e485-c521-4d23-a0d6-f7b3b61ba524.png) +# ht-llama.cpp + +_Heiervang Technologies fork of [llama.cpp](https://github.com/ggml-org/llama.cpp)_ + +[HT Discussions](https://github.com/orgs/heiervang-technologies/discussions) | [Fork Management Guide](https://github.com/orgs/heiervang-technologies/discussions/3) | [Upstream: llama.cpp](https://github.com/ggml-org/llama.cpp) + +--- + +## HT Fork Changes + +This is the [Heiervang Technologies](https://github.com/heiervang-technologies) fork of [llama.cpp](https://github.com/ggml-org/llama.cpp). The `ht` branch contains the following changes on top of upstream `master`: + +| Change | Description | Contributed back? | +|--------|-------------|-------------------| +| Agentic contributions policy | We welcome AI-assisted and AI-generated contributions—see [CONTRIBUTING.md](CONTRIBUTING.md) | N/A | + +Unlike upstream, we accept contributions from AI agents and assistants. We judge code by its quality, not its authorship. + +### Branch Strategy + +| Branch | Purpose | +|--------|---------| +| `master` | Clean fast-forward mirror of upstream `master` — never commit directly | +| `ht` | HT-specific changes on top of `master` — **default branch** | + +Feature branches are created from `ht` and squash-merged back via PR. + +For questions or inquiries, use the [HT Discussions](https://github.com/orgs/heiervang-technologies/discussions) page. For details on fork workflow and sync procedures, see the [Fork Management Guide](https://github.com/orgs/heiervang-technologies/discussions/3). + +--- [![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT) [![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp)](https://github.com/ggml-org/llama.cpp/releases) diff --git a/common/arg.cpp b/common/arg.cpp index 538d2a4b0a46..37041c7d6146 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2226,12 +2226,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex "- distribute: spread execution evenly over all nodes\n" "- isolate: only spawn threads on CPUs on the node that execution started on\n" "- numactl: use the CPU map provided by numactl\n" + "- pin: 1:1 thread-to-core pinning with CCX/L3 locality awareness\n" "if run without this previously, it is recommended to drop the system page cache before using this\n" "see https://github.com/ggml-org/llama.cpp/issues/1437", [](common_params & params, const std::string & value) { /**/ if (value == "distribute" || value == "") { params.numa = GGML_NUMA_STRATEGY_DISTRIBUTE; } else if (value == "isolate") { params.numa = GGML_NUMA_STRATEGY_ISOLATE; } else if (value == "numactl") { params.numa = GGML_NUMA_STRATEGY_NUMACTL; } + else if (value == "pin") { params.numa = GGML_NUMA_STRATEGY_PIN; } else { throw std::invalid_argument("invalid value"); } } ).set_env("LLAMA_ARG_NUMA")); @@ -3058,6 +3060,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.use_jinja = value; } ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_COMPLETION, LLAMA_EXAMPLE_CLI, LLAMA_EXAMPLE_MTMD}).set_env("LLAMA_ARG_JINJA")); + add_opt(common_arg( + {"--remap-developer-role"}, + "remap the OpenAI \"developer\" role to \"system\" before applying chat templates " + "(needed for models whose templates reject unknown roles, e.g. Qwen3.5)", + [](common_params & params) { + params.remap_developer_role = true; + } + ).set_examples({LLAMA_EXAMPLE_SERVER}).set_env("LLAMA_ARG_REMAP_DEVELOPER_ROLE")); add_opt(common_arg( {"--reasoning-format"}, "FORMAT", "controls whether thought tags are allowed and/or extracted from the response, and in which format they're returned; one of:\n" diff --git a/common/common.h b/common/common.h index 17dc3fb23261..2a7a48772f67 100644 --- a/common/common.h +++ b/common/common.h @@ -588,6 +588,7 @@ struct common_params { std::string api_prefix = ""; // NOLINT std::string chat_template = ""; // NOLINT bool use_jinja = true; // NOLINT + bool remap_developer_role = false; // remap "developer" role to "system" for template compatibility bool enable_chat_template = true; bool force_pure_content_parser = false; common_reasoning_format reasoning_format = COMMON_REASONING_FORMAT_DEEPSEEK; diff --git a/common/preset.cpp b/common/preset.cpp index 57ccd000b5c6..e80702f83690 100644 --- a/common/preset.cpp +++ b/common/preset.cpp @@ -4,6 +4,8 @@ #include "log.h" #include "download.h" +#include "gguf.h" + #include #include #include @@ -444,6 +446,141 @@ common_presets common_preset_context::load_from_models_dir(const std::string & m return out; } +// helper: read GGUF metadata to determine if file is a LoRA adapter +// returns true if it is an adapter, and fills out architecture string +static bool gguf_is_lora_adapter(const std::string & path, std::string & architecture) { + struct gguf_init_params params = { + /*.no_alloc = */ true, + /*.ctx = */ nullptr, + }; + struct gguf_context * ctx = gguf_init_from_file(path.c_str(), params); + if (!ctx) { + return false; + } + + bool is_adapter = false; + + int64_t type_key = gguf_find_key(ctx, "general.type"); + if (type_key >= 0) { + const char * type_val = gguf_get_val_str(ctx, type_key); + if (type_val && std::string(type_val) == "adapter") { + is_adapter = true; + + int64_t arch_key = gguf_find_key(ctx, "general.architecture"); + if (arch_key >= 0) { + const char * arch_val = gguf_get_val_str(ctx, arch_key); + if (arch_val) { + architecture = arch_val; + } + } + } + } + + gguf_free(ctx); + return is_adapter; +} + +common_models_dir_result common_preset_context::load_from_models_dir_with_lora(const std::string & models_dir) const { + if (!std::filesystem::exists(models_dir) || !std::filesystem::is_directory(models_dir)) { + throw std::runtime_error(string_format("error: '%s' does not exist or is not a directory\n", models_dir.c_str())); + } + + std::vector models; + std::vector adapters; + + auto scan_subdir = [&models, &adapters](const std::string & subdir_path, const std::string & name) { + auto files = fs_list(subdir_path, false); + common_file_info model_file; + common_file_info first_shard_file; + common_file_info mmproj_file; + std::vector other_gguf_files; + + for (const auto & file : files) { + if (string_ends_with(file.name, ".gguf")) { + if (file.name.find("mmproj") != std::string::npos) { + mmproj_file = file; + } else if (file.name.find("-00001-of-") != std::string::npos) { + first_shard_file = file; + } else { + other_gguf_files.push_back(file); + } + } + } + + // check each non-shard, non-mmproj gguf file for adapter type + for (const auto & file : other_gguf_files) { + std::string arch; + if (gguf_is_lora_adapter(file.path, arch)) { + std::string adapter_name = file.name; + string_replace_all(adapter_name, ".gguf", ""); + adapters.push_back({ + /* name */ adapter_name, + /* path */ file.path, + /* architecture */ arch, + }); + LOG_INF("%s: discovered LoRA adapter '%s' (arch: %s) in %s\n", + __func__, adapter_name.c_str(), arch.c_str(), subdir_path.c_str()); + } else { + // treat as model file (last one wins, matching original behavior) + model_file = file; + } + } + + local_model model{ + /* name */ name, + /* path */ first_shard_file.path.empty() ? model_file.path : first_shard_file.path, + /* path_mmproj */ mmproj_file.path + }; + if (!model.path.empty()) { + models.push_back(model); + } + }; + + auto files = fs_list(models_dir, true); + for (const auto & file : files) { + if (file.is_dir) { + scan_subdir(file.path, file.name); + } else if (string_ends_with(file.name, ".gguf")) { + std::string file_name = file.name; + string_replace_all(file_name, ".gguf", ""); + + // check if this is a LoRA adapter + std::string arch; + if (gguf_is_lora_adapter(file.path, arch)) { + adapters.push_back({ + /* name */ file_name, + /* path */ file.path, + /* architecture */ arch, + }); + LOG_INF("%s: discovered LoRA adapter '%s' (arch: %s)\n", + __func__, file_name.c_str(), arch.c_str()); + } else { + local_model model{ + /* name */ file_name, + /* path */ file.path, + /* path_mmproj */ "" + }; + models.push_back(model); + } + } + } + + // convert local models to presets + common_models_dir_result result; + for (const auto & model : models) { + common_preset preset; + preset.name = model.name; + preset.set_option(*this, "LLAMA_ARG_MODEL", model.path); + if (!model.path_mmproj.empty()) { + preset.set_option(*this, "LLAMA_ARG_MMPROJ", model.path_mmproj); + } + result.presets[preset.name] = preset; + } + result.adapters = std::move(adapters); + + return result; +} + common_preset common_preset_context::load_from_args(int argc, char ** argv) const { common_preset preset; preset.name = COMMON_PRESET_DEFAULT_NAME; diff --git a/common/preset.h b/common/preset.h index 11ba6ef81240..e805376ebef8 100644 --- a/common/preset.h +++ b/common/preset.h @@ -14,6 +14,13 @@ constexpr const char * COMMON_PRESET_DEFAULT_NAME = "default"; +// discovered LoRA adapter from models directory scanning +struct common_lora_adapter_info { + std::string name; // derived from filename (without .gguf extension) + std::string path; // full path to the adapter file + std::string architecture; // general.architecture from GGUF metadata +}; + struct common_preset_context; struct common_preset { @@ -49,6 +56,12 @@ struct common_preset { // interface for multiple presets in one file using common_presets = std::map; +// result of scanning a models directory, containing both models and adapters +struct common_models_dir_result { + common_presets presets; // model presets (non-adapter GGUF files) + std::vector adapters; // discovered LoRA adapters +}; + // context for loading and editing presets struct common_preset_context { common_params default_params; // unused for now @@ -71,6 +84,9 @@ struct common_preset_context { // for the directory structure, see "Using multiple models" in server/README.md common_presets load_from_models_dir(const std::string & models_dir) const; + // generate presets from local models directory, also discovering LoRA adapters + common_models_dir_result load_from_models_dir_with_lora(const std::string & models_dir) const; + // generate one preset from CLI arguments common_preset load_from_args(int argc, char ** argv) const; diff --git a/convert_hf_to_gguf.py b/convert_hf_to_gguf.py index bcf98cfae76f..361368420b3b 100755 --- a/convert_hf_to_gguf.py +++ b/convert_hf_to_gguf.py @@ -4154,12 +4154,14 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter "Qwen2VLForConditionalGeneration", "Qwen2_5_VLForConditionalGeneration", "Qwen2_5OmniModel", + "Qwen2_5OmniThinkerForConditionalGeneration", ) class Qwen2VLModel(TextModel): model_arch = gguf.MODEL_ARCH.QWEN2VL def set_gguf_parameters(self): super().set_gguf_parameters() + self._try_set_pooling_type() def set_vocab(self): try: @@ -4197,8 +4199,8 @@ def set_gguf_parameters(self): model_type = self.global_config['model_type'] if model_type == 'qwen2_vl': self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.QWEN2VL) - elif model_type == 'qwen2_5_vl' or model_type == 'qwen2_5_omni': - if model_type == 'qwen2_5_omni': + elif model_type in ('qwen2_5_vl', 'qwen2_5_omni', 'qwen2_5_omni_thinker'): + if model_type in ('qwen2_5_omni', 'qwen2_5_omni_thinker'): self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.QWEN25O) else: self.gguf_writer.add_clip_projector_type(gguf.VisionProjectorType.QWEN25VL) @@ -4250,7 +4252,7 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter yield from super().modify_tensors(data_torch, name, bid) -@ModelBase.register("Qwen2_5OmniModel") +@ModelBase.register("Qwen2_5OmniModel", "Qwen2_5OmniThinkerForConditionalGeneration") class Qwen25OmniModel(Qwen2VLVisionModel): has_vision_encoder = True has_audio_encoder = True @@ -4269,10 +4271,14 @@ def set_gguf_parameters(self): self.gguf_writer.add_audio_attention_layernorm_eps(self.hparams_audio.get("layer_norm_eps", 1e-5)) def get_vision_config(self) -> dict[str, Any] | None: - return self.global_config["thinker_config"].get("vision_config") + if "thinker_config" in self.global_config: + return self.global_config["thinker_config"].get("vision_config") + return self.global_config.get("vision_config") def get_audio_config(self) -> dict[str, Any] | None: - return self.global_config["thinker_config"].get("audio_config") + if "thinker_config" in self.global_config: + return self.global_config["thinker_config"].get("audio_config") + return self.global_config.get("audio_config") def generate_extra_tensors(self) -> Iterable[tuple[str, Tensor]]: # SinusoidsPositionEmbedding @@ -4304,6 +4310,10 @@ def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iter # this tensor is left unused in transformers code # https://github.com/huggingface/transformers/blob/6e3063422c4b1c014aa60c32b9254fd2902f0f28/src/transformers/models/qwen2_5_omni/modular_qwen2_5_omni.py#L1809 return + # audio tensors go directly to base class (skip Qwen2VLVisionModel + # which only handles visual.* tensors) + yield from MmprojModel.modify_tensors(self, data_torch, name, bid) + return yield from super().modify_tensors(data_torch, name, bid) diff --git a/convert_lora_to_gguf.py b/convert_lora_to_gguf.py index ee98d0cf97d9..aba6330ee494 100755 --- a/convert_lora_to_gguf.py +++ b/convert_lora_to_gguf.py @@ -456,6 +456,37 @@ def get_tensors(self) -> Iterator[tuple[str, Tensor]]: yield (name, cast(torch.Tensor, LoraTorchTensor(tensor.A, tensor.B))) def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]: + # MLA kv_b_proj: the base model splits this into k_b_proj (transposed) and v_b_proj, + # but LoraTorchTensor can't handle the required split+transpose on 3D tensors. + # Decompose the LoRA and apply split/transpose to the raw A/B tensors directly. + if name.endswith("kv_b_proj.weight") and isinstance(data_torch, LoraTorchTensor): + lora_a, lora_b = data_torch.get_lora_A_B() + + n_head_kv = self.hparams["num_key_value_heads"] + v_head_dim = self.hparams["v_head_dim"] + qk_nope_head_dim = self.hparams["qk_nope_head_dim"] + + if lora_b.shape[0] != n_head_kv * (qk_nope_head_dim + v_head_dim): + raise ValueError(f"unexpected kv_b_proj lora_B shape: {lora_b.shape}") + lora_b_3d = lora_b.view(n_head_kv, qk_nope_head_dim + v_head_dim, -1) + k_b_B, v_b_B = torch.split(lora_b_3d, [qk_nope_head_dim, v_head_dim], dim=1) + + # k_b is transposed in the base model: delta_k_b^T = A^T @ B^T + k_lora_a = k_b_B.transpose(1, 2).contiguous() + k_lora_b = lora_a.T.unsqueeze(0).contiguous() + + # v_b is not transposed: delta_v_b = B @ A + v_lora_a = lora_a.unsqueeze(0).contiguous() + v_lora_b = v_b_B.contiguous() + + name_kb = self.map_tensor_name(name.replace("kv_b_proj", "k_b_proj")) + name_vb = self.map_tensor_name(name.replace("kv_b_proj", "v_b_proj")) + yield (name_kb + ".lora_a", k_lora_a) + yield (name_kb + ".lora_b", k_lora_b) + yield (name_vb + ".lora_a", v_lora_a) + yield (name_vb + ".lora_b", v_lora_b) + return + dest = list(super().modify_tensors(data_torch, name, bid)) # some archs may have the same tensor for lm_head and output (tie word embeddings) # in this case, adapters targeting lm_head will fail when using llama-export-lora diff --git a/ggml/include/ggml-cpu.h b/ggml/include/ggml-cpu.h index e3e067c916f1..eb90a2cb4f56 100644 --- a/ggml/include/ggml-cpu.h +++ b/ggml/include/ggml-cpu.h @@ -31,6 +31,7 @@ extern "C" { GGML_NUMA_STRATEGY_ISOLATE = 2, GGML_NUMA_STRATEGY_NUMACTL = 3, GGML_NUMA_STRATEGY_MIRROR = 4, + GGML_NUMA_STRATEGY_PIN = 5, // 1:1 thread-to-core pinning with CCX/L3 locality GGML_NUMA_STRATEGY_COUNT }; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index df17cc553006..c9a509d51b9c 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -36,6 +36,12 @@ #include #if defined(__gnu_linux__) #include +#include +// Memory policy constants for set_mempolicy syscall (avoids libnuma dependency) +#ifndef MPOL_DEFAULT +#define MPOL_DEFAULT 0 +#define MPOL_INTERLEAVE 3 +#endif #endif #ifdef GGML_USE_OPENMP @@ -59,6 +65,29 @@ #define GGML_CACHE_ALIGN __attribute__((aligned(GGML_CACHE_LINE))) #endif +// Portable software prefetch hints for cache-aware matmul +// Prefetch data into L1 cache (temporal, likely to be reused) +#if defined(__SSE__) || defined(__AVX__) || defined(__AVX2__) || defined(__AVX512F__) +#include +#define GGML_PREFETCH_T0(addr) _mm_prefetch((const char *)(addr), _MM_HINT_T0) +#define GGML_PREFETCH_T1(addr) _mm_prefetch((const char *)(addr), _MM_HINT_T1) +#elif defined(__GNUC__) || defined(__clang__) +// Works on ARM NEON, RISC-V, and any GCC/Clang target +#define GGML_PREFETCH_T0(addr) __builtin_prefetch((const void *)(addr), 0, 3) +#define GGML_PREFETCH_T1(addr) __builtin_prefetch((const void *)(addr), 0, 2) +#else +#define GGML_PREFETCH_T0(addr) ((void)(addr)) +#define GGML_PREFETCH_T1(addr) ((void)(addr)) +#endif + +// Prefetch distance: number of rows ahead to prefetch src0 data. +// Each quantized row is typically 1-4 cache lines, so prefetching 2-4 rows +// ahead gives the memory subsystem ~200-400 bytes of lead time. +#define GGML_PREFETCH_ROWS_AHEAD 4 +// Minimum row size (bytes) to enable prefetching. For small models where +// weights fit in cache, prefetch instructions add overhead without benefit. +#define GGML_PREFETCH_MIN_ROW_BYTES (4 * GGML_CACHE_LINE) + #if defined(__has_feature) #if __has_feature(thread_sanitizer) #define GGML_TSAN_ENABLED 1 @@ -524,12 +553,20 @@ static inline void ggml_thread_cpu_relax(void) {;} #define GGML_NUMA_MAX_NODES 8 #define GGML_NUMA_MAX_CPUS 512 +#define GGML_NUMA_MAX_CCX 64 // max L3 cache clusters (CCX/CCD groups) struct ggml_numa_node { uint32_t cpus[GGML_NUMA_MAX_CPUS]; // hardware threads on this node uint32_t n_cpus; }; +// L3 cache cluster (CCX on AMD, sub-NUMA cluster on Intel) +struct ggml_numa_ccx { + uint32_t cpus[GGML_NUMA_MAX_CPUS]; // cores sharing this L3 + uint32_t n_cpus; + uint32_t node; // parent NUMA node +}; + struct ggml_numa_nodes { enum ggml_numa_strategy numa_strategy; struct ggml_numa_node nodes[GGML_NUMA_MAX_NODES]; @@ -541,6 +578,14 @@ struct ggml_numa_nodes { #else uint32_t cpuset; // no NUMA support outside of Linux at this time. Use a portable datatype #endif + + // CCX/L3 topology for sub-NUMA locality (used by PIN strategy) + struct ggml_numa_ccx ccx[GGML_NUMA_MAX_CCX]; + uint32_t n_ccx; + // Sorted CPU list for 1:1 thread pinning: cores ordered by CCX then node + uint32_t pin_order[GGML_NUMA_MAX_CPUS]; + uint32_t n_pin_cpus; + bool pin_initialized; }; // @@ -608,6 +653,151 @@ static cpu_set_t ggml_get_numa_affinity(void) { pthread_getaffinity_np(thread, sizeof(cpu_set_t), &cpuset); return cpuset; } +// Parse a CPU list string like "0-3,8-11" into individual CPU IDs. +// Returns the number of CPUs parsed. +static uint32_t ggml_parse_cpu_list(char * str, uint32_t * out, uint32_t max_out) { + uint32_t n = 0; + char * p = str; + while (*p && n < max_out) { + while (*p == ',' || *p == ' ' || *p == '\n') p++; + if (!*p) break; + + unsigned long start = strtoul(p, &p, 10); + unsigned long end = start; + if (*p == '-') { + p++; + end = strtoul(p, &p, 10); + } + for (unsigned long c = start; c <= end && n < max_out; c++) { + out[n++] = (uint32_t)c; + } + } + return n; +} + +// Detect L3 cache clusters (CCX/CCD on AMD, sub-NUMA clusters on Intel). +// Reads /sys/devices/system/cpu/cpuN/cache/index3/shared_cpu_list +static void ggml_numa_detect_ccx(struct ggml_numa_nodes * numa) { + numa->n_ccx = 0; + + for (uint32_t c = 0; c < numa->total_cpus && numa->n_ccx < GGML_NUMA_MAX_CCX; c++) { + char path[256]; + snprintf(path, sizeof(path), + "/sys/devices/system/cpu/cpu%u/cache/index3/shared_cpu_list", c); + + FILE * f = fopen(path, "r"); + if (!f) continue; + + char buf[256]; + if (!fgets(buf, sizeof(buf), f)) { + fclose(f); + continue; + } + fclose(f); + + uint32_t shared_cpus[GGML_NUMA_MAX_CPUS]; + uint32_t n_shared = ggml_parse_cpu_list(buf, shared_cpus, GGML_NUMA_MAX_CPUS); + if (n_shared == 0) continue; + + // Check if we already have a CCX with the same first CPU + uint32_t first_cpu = shared_cpus[0]; + bool found = false; + for (uint32_t i = 0; i < numa->n_ccx; i++) { + if (numa->ccx[i].n_cpus > 0 && numa->ccx[i].cpus[0] == first_cpu) { + found = true; + break; + } + } + if (found) continue; + + struct ggml_numa_ccx * ccx = &numa->ccx[numa->n_ccx]; + ccx->n_cpus = n_shared; + for (uint32_t i = 0; i < n_shared; i++) { + ccx->cpus[i] = shared_cpus[i]; + } + + // Determine which NUMA node this CCX belongs to + ccx->node = 0; + for (uint32_t nn = 0; nn < numa->n_nodes; nn++) { + struct ggml_numa_node * node = &numa->nodes[nn]; + for (uint32_t i = 0; i < node->n_cpus; i++) { + if (node->cpus[i] == first_cpu) { + ccx->node = nn; + goto ccx_node_found; + } + } + } + ccx_node_found: + + GGML_PRINT_DEBUG("CCX %u: node %u, %u cpus (first=%u)\n", + numa->n_ccx, ccx->node, ccx->n_cpus, first_cpu); + numa->n_ccx++; + } +} + +// Build pin_order: cores sorted by CCX cluster then NUMA node for maximum L3 sharing. +static void ggml_numa_build_pin_order(struct ggml_numa_nodes * numa) { + numa->n_pin_cpus = 0; + numa->pin_initialized = false; + + if (numa->n_ccx > 0) { + // Sort CCX clusters by NUMA node, then by first CPU + uint32_t ccx_order[GGML_NUMA_MAX_CCX]; + for (uint32_t i = 0; i < numa->n_ccx; i++) ccx_order[i] = i; + for (uint32_t i = 1; i < numa->n_ccx; i++) { + uint32_t key = ccx_order[i]; + uint32_t j = i; + while (j > 0) { + uint32_t prev = ccx_order[j - 1]; + if (numa->ccx[prev].node < numa->ccx[key].node) break; + if (numa->ccx[prev].node == numa->ccx[key].node && + numa->ccx[prev].cpus[0] < numa->ccx[key].cpus[0]) break; + ccx_order[j] = ccx_order[j - 1]; + j--; + } + ccx_order[j] = key; + } + + for (uint32_t i = 0; i < numa->n_ccx; i++) { + struct ggml_numa_ccx * ccx = &numa->ccx[ccx_order[i]]; + for (uint32_t ci = 0; ci < ccx->n_cpus && numa->n_pin_cpus < GGML_NUMA_MAX_CPUS; ci++) { + numa->pin_order[numa->n_pin_cpus++] = ccx->cpus[ci]; + } + } + } else { + // Fallback: order cores by NUMA node + for (uint32_t n = 0; n < numa->n_nodes; n++) { + struct ggml_numa_node * node = &numa->nodes[n]; + for (uint32_t ci = 0; ci < node->n_cpus && numa->n_pin_cpus < GGML_NUMA_MAX_CPUS; ci++) { + numa->pin_order[numa->n_pin_cpus++] = node->cpus[ci]; + } + } + } + + if (numa->n_pin_cpus > 0) { + numa->pin_initialized = true; + GGML_PRINT_DEBUG("NUMA pin order: %u cpus, %u CCX clusters\n", + numa->n_pin_cpus, numa->n_ccx); + } +} + +// Set memory interleave policy across all NUMA nodes via syscall (no libnuma needed). +static void ggml_numa_set_interleave_policy(struct ggml_numa_nodes * numa) { + if (numa->n_nodes <= 1) return; + + unsigned long nodemask = 0; + for (uint32_t n = 0; n < numa->n_nodes && n < (uint32_t)(sizeof(unsigned long) * 8); n++) { + nodemask |= (1UL << n); + } + + long ret = syscall(SYS_set_mempolicy, MPOL_INTERLEAVE, &nodemask, numa->n_nodes + 1); + if (ret == 0) { + GGML_PRINT_DEBUG("NUMA: set memory interleave policy across %u nodes\n", numa->n_nodes); + } else { + GGML_PRINT_DEBUG("NUMA: set_mempolicy(MPOL_INTERLEAVE) failed: %s\n", strerror(errno)); + } +} + #else static uint32_t ggml_get_numa_affinity(void) { return 0; // no NUMA support @@ -626,6 +816,15 @@ void ggml_numa_init(enum ggml_numa_strategy numa_flag) { char path[256]; int rv; + // Check GGML_NUMA_PIN env var: if set to "1", override strategy to PIN + { + const char * pin_env = getenv("GGML_NUMA_PIN"); + if (pin_env && strcmp(pin_env, "1") == 0) { + numa_flag = GGML_NUMA_STRATEGY_PIN; + GGML_LOG_INFO("NUMA: GGML_NUMA_PIN=1 detected, using PIN strategy\n"); + } + } + // set numa scheme g_state.numa.numa_strategy = numa_flag; @@ -696,6 +895,25 @@ void ggml_numa_init(enum ggml_numa_strategy numa_flag) { fclose(fptr); } } + + // PIN strategy: detect CCX topology and build optimal pin order + if (numa_flag == GGML_NUMA_STRATEGY_PIN) { + ggml_numa_detect_ccx(&g_state.numa); + ggml_numa_build_pin_order(&g_state.numa); + + if (g_state.numa.pin_initialized) { + GGML_LOG_INFO("NUMA PIN: %u cores across %u CCX clusters on %u nodes\n", + g_state.numa.n_pin_cpus, g_state.numa.n_ccx, g_state.numa.n_nodes); + + // Set memory interleave for better bandwidth on multi-node systems + if (g_state.numa.n_nodes > 1) { + ggml_numa_set_interleave_policy(&g_state.numa); + } + } else { + GGML_LOG_WARN("NUMA PIN: failed to build pin order, falling back to DISTRIBUTE\n"); + g_state.numa.numa_strategy = GGML_NUMA_STRATEGY_DISTRIBUTE; + } + } #else UNUSED(numa_flag); // TODO @@ -1185,9 +1403,9 @@ static void ggml_compute_forward_mul_mat_one_chunk( const size_t src1_col_stride = src1_cont || src1->type != vec_dot_type ? row_size : nb11; - // attempt to reduce false-sharing (does not seem to make a difference) + // Cache-line aligned accumulator to prevent false sharing between threads // 16 * 2, accounting for mmla kernels - float tmp[32]; + GGML_CACHE_ALIGN float tmp[32]; for (int64_t iir1 = ir1_start; iir1 < ir1_end; iir1 += blck_1) { for (int64_t iir0 = ir0_start; iir0 < ir0_end; iir0 += blck_0) { @@ -1216,11 +1434,47 @@ static void ggml_compute_forward_mul_mat_one_chunk( : (i11 * nb11 + i12 * nb12 + i13 * nb13)); float * dst_col = (float*)((char*)dst->data + (i1 * nb1 + i2 * nb2 + i3 * nb3)); + // Prefetch the next src1 column while we process the current one + { + const int64_t ir1_next = ir1 + num_rows_per_vec_dot; + if (ir1_next < iir1 + blck_1 && ir1_next < ir1_end) { + const int64_t ni13 = (ir1_next / (ne12 * ne1)); + const int64_t ni12 = (ir1_next - ni13 * ne12 * ne1) / ne1; + const int64_t ni11 = (ir1_next - ni13 * ne12 * ne1 - ni12 * ne1); + const char * next_src1_col = (const char*)wdata + + (src1_cont || src1->type != vec_dot_type + ? (ni11 + ni12 * ne11 + ni13 * ne12 * ne11) * row_size + : (ni11 * nb11 + ni12 * nb12 + ni13 * nb13)); + if (row_size >= GGML_PREFETCH_MIN_ROW_BYTES) { + GGML_PREFETCH_T0(next_src1_col); + if (row_size > GGML_CACHE_LINE) { + GGML_PREFETCH_T0(next_src1_col + GGML_CACHE_LINE); + } + } + } + } + //for (int64_t ir0 = iir0; ir0 < iir0 + blck_0 && ir0 < ir0_end; ++ir0) { // vec_dot(ne00, &dst_col[ir0], src0_row + ir0*nb01, src1_col); //} for (int64_t ir0 = iir0; ir0 < iir0 + blck_0 && ir0 < ir0_end; ir0 += num_rows_per_vec_dot) { + // Prefetch src0 rows ahead to hide memory latency. + // We prefetch GGML_PREFETCH_ROWS_AHEAD rows into L1 (T0) and the + // row after that into L2 (T1), covering ~256-512 bytes of lead time. + if (nb01 >= GGML_PREFETCH_MIN_ROW_BYTES) { + const int64_t ir0_prefetch = ir0 + GGML_PREFETCH_ROWS_AHEAD * num_rows_per_vec_dot; + if (ir0_prefetch < ir0_end) { + GGML_PREFETCH_T0(src0_row + ir0_prefetch * nb01); + if (nb01 > GGML_CACHE_LINE) { + GGML_PREFETCH_T0(src0_row + ir0_prefetch * nb01 + GGML_CACHE_LINE); + } + if (nb01 > 2 * GGML_CACHE_LINE) { + GGML_PREFETCH_T1(src0_row + ir0_prefetch * nb01 + 2 * GGML_CACHE_LINE); + } + } + } + vec_dot(ne00, &tmp[ir0 - iir0], (num_rows_per_vec_dot > 1 ? 16 : 0), src0_row + ir0 * nb01, (num_rows_per_vec_dot > 1 ? nb01 : 0), src1_col, (num_rows_per_vec_dot > 1 ? src1_col_stride : 0), num_rows_per_vec_dot); } @@ -1228,6 +1482,19 @@ static void ggml_compute_forward_mul_mat_one_chunk( memcpy(&dst_col[iir0 + cn * nb1 / nb0], tmp + (cn * 16), (MIN(iir0 + blck_0, ir0_end) - iir0) * sizeof(float)); } } + + // Prefetch src0 data for the next block of rows in the ir0 dimension + if (nb01 >= GGML_PREFETCH_MIN_ROW_BYTES && iir0 + blck_0 < ir0_end) { + for (int64_t ir1 = iir1; ir1 < iir1 + blck_1 && ir1 < ir1_end; ir1 += num_rows_per_vec_dot) { + const int64_t i13 = (ir1 / (ne12 * ne1)); + const int64_t i12 = (ir1 - i13 * ne12 * ne1) / ne1; + const int64_t i03 = i13 / r3; + const int64_t i02 = i12 / r2; + const char * next_src0_row = (const char*)src0->data + (0 + i02 * nb02 + i03 * nb03); + GGML_PREFETCH_T1(next_src0_row + (iir0 + blck_0) * nb01); + break; // only need to prefetch once per block transition + } + } } } } @@ -1461,7 +1728,8 @@ static void ggml_compute_forward_mul_mat_id_one_chunk( const int64_t blck_0 = 16; const int64_t blck_1 = 16; - float tmp[16]; + // Cache-line aligned accumulator to prevent false sharing between threads + GGML_CACHE_ALIGN float tmp[16]; for (int64_t iir1 = ir1_start; iir1 < ir1_end; iir1 += blck_1) { for (int64_t iir0 = ir0_start; iir0 < ir0_end; iir0 += blck_0) { @@ -1489,6 +1757,17 @@ static void ggml_compute_forward_mul_mat_id_one_chunk( float * dst_col = (float *) ((char *) dst->data + (i1*nb1 + i2*nb2)); for (int64_t ir0 = iir0; ir0 < iir0 + blck_0 && ir0 < ir0_end; ++ir0) { + // Prefetch src0 rows ahead to hide memory latency + if (nb01 >= GGML_PREFETCH_MIN_ROW_BYTES) { + const int64_t ir0_prefetch = ir0 + GGML_PREFETCH_ROWS_AHEAD; + if (ir0_prefetch < ir0_end) { + GGML_PREFETCH_T0(src0_cur + ir0_prefetch * nb01); + if (nb01 > GGML_CACHE_LINE) { + GGML_PREFETCH_T0(src0_cur + ir0_prefetch * nb01 + GGML_CACHE_LINE); + } + } + } + vec_dot(ne00, &tmp[ir0 - iir0], 0, src0_cur + ir0*nb01, 0, src1_col, 0, 1); } @@ -2101,7 +2380,8 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm // Android's libc implementation "bionic" does not support setting affinity #if defined(__gnu_linux__) static void set_numa_thread_affinity(int thread_n) { - if (!ggml_is_numa()) { + // PIN strategy works on single-node multi-CCX systems too + if (!ggml_is_numa() && g_state.numa.numa_strategy != GGML_NUMA_STRATEGY_PIN) { return; } @@ -2125,6 +2405,24 @@ static void set_numa_thread_affinity(int thread_n) { fprintf(stderr, "warning: pthread_setaffinity_np() failed: %s\n",strerror(rv)); } return; + case GGML_NUMA_STRATEGY_PIN: + // 1:1 thread-to-core pinning using CCX-aware ordering + if (g_state.numa.pin_initialized && (uint32_t)thread_n < g_state.numa.n_pin_cpus) { + cpu_set_t * pin_cpus = CPU_ALLOC(g_state.numa.total_cpus); + CPU_ZERO_S(setsize, pin_cpus); + CPU_SET_S(g_state.numa.pin_order[thread_n], setsize, pin_cpus); + rv = pthread_setaffinity_np(pthread_self(), setsize, pin_cpus); + if (rv) { + fprintf(stderr, "warning: pthread_setaffinity_np() pin failed for thread %d -> cpu %u: %s\n", + thread_n, g_state.numa.pin_order[thread_n], strerror(rv)); + } + CPU_FREE(pin_cpus); + } else { + // More threads than cores: fall back to distribute across nodes + node_num = thread_n % g_state.numa.n_nodes; + break; + } + return; default: return; } @@ -2146,7 +2444,7 @@ static void set_numa_thread_affinity(int thread_n) { } static void clear_numa_thread_affinity(void) { - if (!ggml_is_numa()) { + if (!ggml_is_numa() && g_state.numa.numa_strategy != GGML_NUMA_STRATEGY_PIN) { return; } diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 7d7f20af3a04..ae984a42021e 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1166,6 +1166,10 @@ struct ggml_cuda_graph { if (graph != nullptr) { CUDA_CHECK(cudaGraphDestroy(graph)); } + // Clean up overlap resources + if (overlap_instance != nullptr) { + CUDA_CHECK(cudaGraphExecDestroy(overlap_instance)); + } } cudaGraph_t graph = nullptr; cudaGraphExec_t instance = nullptr; @@ -1181,10 +1185,32 @@ struct ggml_cuda_graph { // ref: https://github.com/ggml-org/llama.cpp/pull/19165 std::vector extra; + // Double-buffered graph overlap: while GPU executes one graph instance, + // CPU builds/updates the next. Enabled via GGML_CUDA_GRAPH_OVERLAP=1. + // Both instances are built from the same captured graph template. + // The overlap is CPU-side only: after launching the primary instance, + // the CPU updates the secondary instance while the GPU is busy. + // On the next call, the instances are swapped so the pre-built one launches immediately. + cudaGraphExec_t overlap_instance = nullptr; + bool is_enabled() const { static const bool disable_cuda_graphs_due_to_env = (getenv("GGML_CUDA_DISABLE_GRAPHS") != nullptr); return !(disable_due_to_gpu_arch || disable_cuda_graphs_due_to_env); } + + static bool is_overlap_enabled() { + static const bool enabled = [] { + const char * env = getenv("GGML_CUDA_GRAPH_OVERLAP"); + return env != nullptr && atoi(env) == 1; + }(); + return enabled; + } + + // Swap primary and overlap executable instances. + // The captured graph template (graph) is NOT swapped -- only the executables. + void swap_graph_instances() { + std::swap(instance, overlap_instance); + } #endif }; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index d1239b1c5f72..a6460e7471a1 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -948,9 +948,9 @@ static void ggml_backend_cuda_split_buffer_set_tensor(ggml_backend_buffer_t buff CUDA_CHECK(cudaMemcpyAsync(extra->data_device[id], buf_host, original_size, cudaMemcpyHostToDevice, cudaStreamPerThread)); } - for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { - CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); - } + // All async copies above use the same cudaStreamPerThread, so a single sync suffices + // (the previous loop redundantly synced once per device on the same stream). + CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); } static void ggml_backend_cuda_split_buffer_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor, void * data, size_t offset, size_t size) { @@ -987,9 +987,9 @@ static void ggml_backend_cuda_split_buffer_get_tensor(ggml_backend_buffer_t buff CUDA_CHECK(cudaMemcpyAsync(buf_host, extra->data_device[id], original_size, cudaMemcpyDeviceToHost, cudaStreamPerThread)); } - for (int id = 0; id < ggml_backend_cuda_get_device_count(); ++id) { - CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); - } + // All async copies above use the same cudaStreamPerThread, so a single sync suffices + // (the previous loop redundantly synced once per device on the same stream). + CUDA_CHECK(cudaStreamSynchronize(cudaStreamPerThread)); } static void ggml_backend_cuda_split_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value) { @@ -2417,7 +2417,8 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor * ids_to_sorted_host.insert(ids_to_sorted_host.end(), ids_from_sorted_host.begin(), ids_from_sorted_host.end()); CUDA_CHECK(cudaMemcpyAsync(ids_buf_dev.ptr, ids_to_sorted_host.data(), 2*ne_get_rows*sizeof(int32_t), cudaMemcpyHostToDevice, stream)); - CUDA_CHECK(cudaStreamSynchronize(stream)); + // No sync needed: the subsequent get_rows_cuda kernel is launched on the same stream, + // so CUDA stream ordering guarantees the H2D copy completes before the kernel executes. const int32_t * ids_to_sorted = ids_buf_dev.ptr + 0*ne_get_rows; const int32_t * ids_from_sorted = ids_buf_dev.ptr + 1*ne_get_rows; @@ -3126,6 +3127,33 @@ static void ggml_cuda_graph_update_executable(ggml_backend_cuda_context * cuda_c GGML_ASSERT(stat == cudaSuccess); } } + +// Update the overlap (secondary) graph executable instance. +// Same logic as ggml_cuda_graph_update_executable but targets overlap_instance. +static void ggml_cuda_graph_update_overlap_executable(ggml_backend_cuda_context * cuda_ctx, const void * graph_key) { + ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key); + +#if CUDART_VERSION >= 12000 + cudaGraphExecUpdateResultInfo result_info; + cudaError_t stat = cudaGraphExecUpdate(graph->overlap_instance, graph->graph, &result_info); +#else + cudaGraphNode_t errorNode; + cudaGraphExecUpdateResult result_info; + cudaError_t stat = cudaGraphExecUpdate(graph->overlap_instance, graph->graph, &errorNode, &result_info); +#endif // CUDART_VERSION >= 12000 + + if (stat == cudaErrorGraphExecUpdateFailure) { +#ifndef NDEBUG + GGML_LOG_DEBUG("%s: CUDA overlap graph update failed, re-instantiating\n", __func__); +#endif + (void)cudaGetLastError(); + CUDA_CHECK(cudaGraphExecDestroy(graph->overlap_instance)); + graph->overlap_instance = nullptr; + CUDA_CHECK(cudaGraphInstantiate(&graph->overlap_instance, graph->graph, NULL, NULL, 0)); + } else { + GGML_ASSERT(stat == cudaSuccess); + } +} #endif // USE_CUDA_GRAPH static bool ggml_cuda_should_fuse_rope_set_rows(const ggml_tensor * rope, @@ -4075,9 +4103,49 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud } if (cuda_graph_update_required) { // Update graph executable ggml_cuda_graph_update_executable(cuda_ctx, graph_key); + // Graph structure changed - invalidate overlap instance + if (graph->overlap_instance != nullptr) { + CUDA_CHECK(cudaGraphExecDestroy(graph->overlap_instance)); + graph->overlap_instance = nullptr; + } + } + + if (ggml_cuda_graph::is_overlap_enabled() && !cuda_graph_update_required) { + // Double-buffered overlap path: + // While GPU executes the current graph instance, CPU prepares/updates + // the overlap instance for the next call. On the next call, we swap + // and launch the pre-built instance immediately. + // + // Timeline (steady state): + // Call N: Launch instance A (pre-built) -> GPU runs A + // CPU: update/instantiate instance B (overlaps with GPU) + // Swap A <-> B + // Call N+1: Launch instance B (pre-built) -> GPU runs B + // CPU: update/instantiate instance A (overlaps with GPU) + // Swap B <-> A + // + // The key benefit: cudaGraphExecUpdate (CPU work) runs while the GPU + // is busy executing the graph, instead of blocking before the launch. + + // Launch the current (primary) graph instance on the main stream + CUDA_CHECK(cudaGraphLaunch(graph->instance, cuda_ctx->stream())); + + // CPU-side: build/update the overlap instance from the same captured graph. + // This CPU work overlaps with the GPU executing the primary instance + // because cudaGraphLaunch is asynchronous. + if (graph->overlap_instance == nullptr) { + CUDA_CHECK(cudaGraphInstantiate(&graph->overlap_instance, graph->graph, NULL, NULL, 0)); + } else { + // Update overlap instance to match current graph (handles pointer changes etc.) + ggml_cuda_graph_update_overlap_executable(cuda_ctx, graph_key); + } + + // Swap: next call will launch the freshly updated instance as primary + graph->swap_graph_instances(); + } else { + // Standard path: just launch + CUDA_CHECK(cudaGraphLaunch(graph->instance, cuda_ctx->stream())); } - // Launch graph - CUDA_CHECK(cudaGraphLaunch(graph->instance, cuda_ctx->stream())); #else GGML_UNUSED(graph_key); graph_evaluated_or_captured = true; @@ -4137,6 +4205,12 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, // Properties changed - reset warmup, execute directly until stable again graph->warmup_complete = false; GGML_LOG_DEBUG("%s: CUDA graph warmup reset\n", __func__); + // Reset overlap state: the captured graph will change, so any + // pre-built overlap instance is invalid. + if (graph->overlap_instance != nullptr) { + CUDA_CHECK(cudaGraphExecDestroy(graph->overlap_instance)); + graph->overlap_instance = nullptr; + } } else { use_cuda_graph = true; cuda_graph_update_required = graph->instance == nullptr; diff --git a/ggml/src/ggml-cuda/norm.cu b/ggml/src/ggml-cuda/norm.cu index ef98f675aa71..b4b7cf37f3a3 100644 --- a/ggml/src/ggml-cuda/norm.cu +++ b/ggml/src/ggml-cuda/norm.cu @@ -294,16 +294,84 @@ static void group_norm_f32_cuda( } } +// Float4-vectorized RMS norm kernel: uses 128-bit (float4) loads/stores for better memory throughput. +// Requires ncols % 4 == 0 and contiguous row data (stride_row == ncols). +// The mathematical operation is identical to rms_norm_f32 -- output is bit-exact. +template +static __global__ void rms_norm_f32_float4(const float * x, + float * dst, + const int ncols, + const int64_t stride_row, + const int64_t stride_channel, + const int64_t stride_sample, + const float eps) { + const int nrows = gridDim.x; + const int nchannels = gridDim.y; + + const int row = blockIdx.x; + const int channel = blockIdx.y; + const int sample = blockIdx.z; + const int tid = threadIdx.x; + + x += sample*stride_sample + channel*stride_channel + row*stride_row; + dst += ((sample*nchannels + channel)*nrows + row)*ncols; + + const int ncols4 = ncols / 4; + const float4 * x4 = reinterpret_cast(x); + + float tmp = 0.0f; + + for (int col4 = tid; col4 < ncols4; col4 += block_size) { + const float4 v = x4[col4]; + tmp += v.x * v.x + v.y * v.y + v.z * v.z + v.w * v.w; + } + + // sum up partial sums + extern __shared__ float s_sum[]; + tmp = block_reduce(tmp, s_sum); + + const float mean = tmp / ncols; + const float scale = rsqrtf(mean + eps); + + float4 * dst4 = reinterpret_cast(dst); + + for (int col4 = tid; col4 < ncols4; col4 += block_size) { + const float4 v = x4[col4]; + float4 out; + out.x = scale * v.x; + out.y = scale * v.y; + out.z = scale * v.z; + out.w = scale * v.w; + dst4[col4] = out; + } +} + static void rms_norm_f32_cuda( const float * x, float * dst, const int ncols, const int nrows, const int nchannels, const int nsamples, const int64_t stride_row, const int64_t stride_channel, const int64_t stride_sample, const float eps, cudaStream_t stream) { const dim3 blocks_num(nrows, nchannels, nsamples); - if (ncols < 1024) { - const dim3 block_dims(256, 1, 1); - rms_norm_f32<256, false><< WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + + // Use float4-vectorized kernel when ncols is divisible by 4 and the row is contiguous. + // stride_row == ncols means elements within a row are packed, so float4 loads are valid. + if (ncols % 4 == 0 && stride_row == ncols) { + if (ncols < 1024) { + const dim3 block_dims(256, 1, 1); + rms_norm_f32_float4<256><< WARP_SIZE ? 32 * sizeof(float) : 0, stream>>>( + x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + } else { + const dim3 block_dims(1024, 1, 1); + rms_norm_f32_float4<1024><<>>( + x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + } } else { - const dim3 block_dims(1024, 1, 1); - rms_norm_f32<1024, false><< WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + // Fallback: scalar kernel for non-aligned or non-contiguous cases. + if (ncols < 1024) { + const dim3 block_dims(256, 1, 1); + rms_norm_f32<256, false><< WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + } else { + const dim3 block_dims(1024, 1, 1); + rms_norm_f32<1024, false><< WARP_SIZE ? 32 * sizeof(float): 0, stream>>>(x, dst, ncols, stride_row, stride_channel, stride_sample, eps); + } } } diff --git a/ggml/src/ggml-cuda/unary.cu b/ggml/src/ggml-cuda/unary.cu index 4ad30fa1f353..84af7cd440b3 100644 --- a/ggml/src/ggml-cuda/unary.cu +++ b/ggml/src/ggml-cuda/unary.cu @@ -563,6 +563,70 @@ void ggml_cuda_op_leaky_relu(ggml_backend_cuda_context & ctx, ggml_tensor * dst) /* fused unary + mul */ +// Vectorized fused SiLU-multiply kernel for contiguous f32 tensors. +// Processes 4 elements per thread using float4 loads/stores, eliminating +// the intermediate DRAM round-trip for SiLU output. This is the hot path +// in SwiGLU MLP blocks (Llama, Mistral, etc.). +// +// Math (bit-exact with the scalar path): +// dst[i] = (gate[i] / (1 + exp(-gate[i]))) * up[i] + +#define CUDA_SILU_MUL_VEC4_BLOCK_SIZE 256 + +static __global__ void silu_mul_f32_vec4_kernel( + const float * __restrict__ gate, + const float * __restrict__ up, + float * __restrict__ dst, + const int64_t k4) { + const int64_t i = int64_t(blockDim.x) * blockIdx.x + threadIdx.x; + + if (i >= k4) { + return; + } + + const float4 g = reinterpret_cast(gate)[i]; + const float4 u = reinterpret_cast(up)[i]; + + float4 out; + out.x = (g.x / (1.0f + expf(-g.x))) * u.x; + out.y = (g.y / (1.0f + expf(-g.y))) * u.y; + out.z = (g.z / (1.0f + expf(-g.z))) * u.z; + out.w = (g.w / (1.0f + expf(-g.w))) * u.w; + + reinterpret_cast(dst)[i] = out; +} + +static __global__ void silu_mul_f32_tail_kernel( + const float * __restrict__ gate, + const float * __restrict__ up, + float * __restrict__ dst, + const int64_t offset, + const int64_t k) { + const int64_t i = offset + int64_t(blockDim.x) * blockIdx.x + threadIdx.x; + + if (i >= k) { + return; + } + + dst[i] = (gate[i] / (1.0f + expf(-gate[i]))) * up[i]; +} + +static void silu_mul_f32_vec4_cuda( + const float * gate, const float * up, float * dst, + const int64_t k, cudaStream_t stream) { + const int64_t k4 = k / 4; + const int64_t tail = k % 4; + + if (k4 > 0) { + const int64_t num_blocks = (k4 + CUDA_SILU_MUL_VEC4_BLOCK_SIZE - 1) / CUDA_SILU_MUL_VEC4_BLOCK_SIZE; + silu_mul_f32_vec4_kernel<<>>(gate, up, dst, k4); + } + + if (tail > 0) { + silu_mul_f32_tail_kernel<<<1, tail, 0, stream>>>(gate, up, dst, k4 * 4, k); + } +} + template static void ggml_cuda_op_unary_mul_impl(ggml_backend_cuda_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node) { // unary_node: UNARY op applied to unary_node->src[0] @@ -600,10 +664,43 @@ static void ggml_cuda_op_unary_mul_impl(ggml_backend_cuda_context & ctx, ggml_te } } +// Specialization: when both inputs are contiguous f32, use the vectorized kernel. +static void ggml_cuda_op_silu_mul_fast(ggml_backend_cuda_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node) { + const ggml_tensor * gate_src = unary_node->src[0]; + const ggml_tensor * up_src = (mul_node->src[0] == unary_node) ? mul_node->src[1] : mul_node->src[0]; + + GGML_ASSERT(ggml_is_contiguous_1(gate_src)); + GGML_ASSERT(gate_src->nb[0] == ggml_element_size(gate_src)); + GGML_ASSERT(ggml_is_contiguous_1(up_src)); + GGML_ASSERT(up_src->nb[0] == ggml_element_size(up_src)); + GGML_ASSERT(ggml_are_same_shape(gate_src, up_src)); + GGML_ASSERT(gate_src->type == up_src->type); + GGML_ASSERT(gate_src->type == mul_node->type); + + const int64_t k = ggml_nelements(mul_node); + + // Use vectorized path for fully contiguous f32 tensors. + // float4 requires 16-byte alignment (ggml guarantees this for data pointers) + // and contiguous layout so there are no stride gaps. + if (gate_src->type == GGML_TYPE_F32 && + ggml_is_contiguous(gate_src) && + ggml_is_contiguous(up_src)) { + + silu_mul_f32_vec4_cuda( + (const float *) gate_src->data, + (const float *) up_src->data, + (float *) mul_node->data, + k, ctx.stream()); + } else { + // Fall back to the generic strided path for f16 or non-contiguous tensors. + ggml_cuda_op_unary_mul_impl(ctx, unary_node, mul_node); + } +} + void ggml_cuda_op_unary_mul(ggml_backend_cuda_context & ctx, ggml_tensor * unary_node, ggml_tensor * mul_node) { switch (ggml_get_unary_op(unary_node)) { case GGML_UNARY_OP_SILU: - ggml_cuda_op_unary_mul_impl(ctx, unary_node, mul_node); + ggml_cuda_op_silu_mul_fast(ctx, unary_node, mul_node); break; case GGML_UNARY_OP_SIGMOID: ggml_cuda_op_unary_mul_impl(ctx, unary_node, mul_node); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 15ed5b2a79df..e5fc2ce9ba49 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -695,6 +695,10 @@ struct vk_device_struct { vk_pipeline pipeline_dequant_mul_mat_vec_f16_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT][mul_mat_vec_max_cols]; vk_pipeline pipeline_dequant_mul_mat_vec_id_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT]; + // Shared-memory staging MatVec pipelines (two-phase DECODE+DOT) + vk_pipeline pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_COUNT]; + vk_pipeline pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_COUNT]; + vk_pipeline pipeline_dequant_mul_mat_vec_q8_1_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT][mul_mat_vec_max_cols]; vk_pipeline pipeline_dequant_mul_mat_vec_id_q8_1_f32[DMMV_WG_SIZE_COUNT][GGML_TYPE_COUNT]; @@ -4165,6 +4169,39 @@ static void ggml_vk_load_shaders(vk_device& device) { GGML_UNUSED(rm_iq_int); #endif + // Shared-memory staging MatVec pipelines (two-phase DECODE+DOT) + // Only for types that use the generic mul_mat_vec.comp (not k-quants or iq1/iq2/iq3) + { + const uint32_t shmem_wg_size = subgroup_size; + // f32/f16/bf16 types + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_F32 ], "mul_mat_vec_shmem_f32_f32_f32", mul_mat_vec_shmem_f32_f32_f32_len, mul_mat_vec_shmem_f32_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1, 1, 1}, {shmem_wg_size, 1, 1}, 1, false, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_F16 ], "mul_mat_vec_shmem_f16_f32_f32", mul_mat_vec_shmem_f16_f32_f32_len, mul_mat_vec_shmem_f16_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {shmem_wg_size, 2, 1}, 1, false, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_BF16], "mul_mat_vec_shmem_bf16_f32_f32", mul_mat_vec_shmem_bf16_f32_f32_len, mul_mat_vec_shmem_bf16_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {shmem_wg_size, 2, 1}, 1, false, use_subgroups, force_subgroup_size); + // Standard quant types + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_Q4_0], "mul_mat_vec_shmem_q4_0_f32_f32", mul_mat_vec_shmem_q4_0_f32_f32_len, mul_mat_vec_shmem_q4_0_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_Q4_1], "mul_mat_vec_shmem_q4_1_f32_f32", mul_mat_vec_shmem_q4_1_f32_f32_len, mul_mat_vec_shmem_q4_1_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_Q5_0], "mul_mat_vec_shmem_q5_0_f32_f32", mul_mat_vec_shmem_q5_0_f32_f32_len, mul_mat_vec_shmem_q5_0_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_Q5_1], "mul_mat_vec_shmem_q5_1_f32_f32", mul_mat_vec_shmem_q5_1_f32_f32_len, mul_mat_vec_shmem_q5_1_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_Q8_0], "mul_mat_vec_shmem_q8_0_f32_f32", mul_mat_vec_shmem_q8_0_f32_f32_len, mul_mat_vec_shmem_q8_0_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_stdq, 1, 1}, {shmem_wg_size, 1*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + // IQ4/MXFP4 types + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_IQ4_NL], "mul_mat_vec_shmem_iq4_nl_f32_f32", mul_mat_vec_shmem_iq4_nl_f32_f32_len, mul_mat_vec_shmem_iq4_nl_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_IQ4_XS], "mul_mat_vec_shmem_iq4_xs_f32_f32", mul_mat_vec_shmem_iq4_xs_f32_f32_len, mul_mat_vec_shmem_iq4_xs_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[GGML_TYPE_MXFP4], "mul_mat_vec_shmem_mxfp4_f32_f32", mul_mat_vec_shmem_mxfp4_f32_f32_len, mul_mat_vec_shmem_mxfp4_f32_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + + // f16 B-type variants + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_F32 ], "mul_mat_vec_shmem_f32_f16_f32", mul_mat_vec_shmem_f32_f16_f32_len, mul_mat_vec_shmem_f32_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1, 1, 1}, {shmem_wg_size, 1, 1}, 1, false, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_F16 ], "mul_mat_vec_shmem_f16_f16_f32", mul_mat_vec_shmem_f16_f16_f32_len, mul_mat_vec_shmem_f16_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {shmem_wg_size, 2, 1}, 1, false, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_BF16], "mul_mat_vec_shmem_bf16_f16_f32", mul_mat_vec_shmem_bf16_f16_f32_len, mul_mat_vec_shmem_bf16_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {shmem_wg_size, 2, 1}, 1, false, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_Q4_0], "mul_mat_vec_shmem_q4_0_f16_f32", mul_mat_vec_shmem_q4_0_f16_f32_len, mul_mat_vec_shmem_q4_0_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_Q4_1], "mul_mat_vec_shmem_q4_1_f16_f32", mul_mat_vec_shmem_q4_1_f16_f32_len, mul_mat_vec_shmem_q4_1_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_Q5_0], "mul_mat_vec_shmem_q5_0_f16_f32", mul_mat_vec_shmem_q5_0_f16_f32_len, mul_mat_vec_shmem_q5_0_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_Q5_1], "mul_mat_vec_shmem_q5_1_f16_f32", mul_mat_vec_shmem_q5_1_f16_f32_len, mul_mat_vec_shmem_q5_1_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {shmem_wg_size, 2*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_Q8_0], "mul_mat_vec_shmem_q8_0_f16_f32", mul_mat_vec_shmem_q8_0_f16_f32_len, mul_mat_vec_shmem_q8_0_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {1*rm_stdq, 1, 1}, {shmem_wg_size, 1*rm_stdq, 1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_IQ4_NL], "mul_mat_vec_shmem_iq4_nl_f16_f32", mul_mat_vec_shmem_iq4_nl_f16_f32_len, mul_mat_vec_shmem_iq4_nl_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_IQ4_XS], "mul_mat_vec_shmem_iq4_xs_f16_f32", mul_mat_vec_shmem_iq4_xs_f16_f32_len, mul_mat_vec_shmem_iq4_xs_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[GGML_TYPE_MXFP4], "mul_mat_vec_shmem_mxfp4_f16_f32", mul_mat_vec_shmem_mxfp4_f16_f32_len, mul_mat_vec_shmem_mxfp4_f16_f32_data, "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {shmem_wg_size, rm_iq, 1}, 1, true, use_subgroups16, force_subgroup_size16); + } + // dequant shaders ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_F32 ], "f32_to_f16", dequant_f32_len, dequant_f32_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_0], "dequant_q4_0", dequant_q4_0_len, dequant_q4_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); @@ -6171,6 +6208,29 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context * return nullptr; } + // Try shared-memory staging variant for single-column matvec (token generation). + // This decouples dequantization from dot-product computation via shared memory, + // which can significantly improve performance especially on Intel Arc GPUs. + if (num_cols == 1 && b_type != GGML_TYPE_Q8_1) { + // Check shared memory feasibility + const uint32_t subgroup_sz = ctx->device->subgroup_size; + const bool is_float_type = (a_type == GGML_TYPE_F32 || a_type == GGML_TYPE_F16 || a_type == GGML_TYPE_BF16); + const uint32_t k_per_iter = is_float_type ? 2 : 8; + const uint32_t num_rows_default = 2; + const uint32_t staging_shmem = num_rows_default * subgroup_sz * k_per_iter * sizeof(float); + const uint32_t reduce_shmem = 1 * num_rows_default * subgroup_sz * sizeof(float); + const uint32_t total_shmem = staging_shmem + reduce_shmem; + + if (total_shmem <= ctx->device->properties.limits.maxComputeSharedMemorySize) { + vk_pipeline shmem_pipeline = b_type == GGML_TYPE_F32 + ? ctx->device->pipeline_dequant_mul_mat_vec_shmem_f32_f32[a_type] + : ctx->device->pipeline_dequant_mul_mat_vec_shmem_f16_f32[a_type]; + if (shmem_pipeline) { + return shmem_pipeline; + } + } + } + // heuristic to choose workgroup size uint32_t dmmv_wg = DMMV_WG_SIZE_SUBGROUP; if ((ctx->device->vendor_id == VK_VENDOR_ID_NVIDIA && ctx->device->architecture != vk_device_architecture::NVIDIA_PRE_TURING) || ctx->device->vendor_id == VK_VENDOR_ID_INTEL) { diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_shmem.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_shmem.comp new file mode 100644 index 000000000000..198c4cfe86fc --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_shmem.comp @@ -0,0 +1,183 @@ +#version 450 + +#extension GL_EXT_shader_explicit_arithmetic_types_int32 : require + +// Shared-memory staging MatVec kernel +// +// Two-phase approach: +// Phase 1 (DECODE): All threads cooperatively load and dequantize a tile of +// the A-matrix into shared memory as f32 values. +// Phase 2 (DOT): Each thread computes partial dot products from the staged +// shared memory tile against the B-vector elements. +// +// This decouples the memory-bound dequantization from the compute-bound dot +// products, improving utilization on GPUs where the two phases have different +// bottlenecks (particularly beneficial on Intel Arc and similar architectures). + +#include "mul_mat_vec_base.glsl" +#include "dequant_funcs.glsl" + +layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in; + +#if !defined(DATA_A_F32) && !defined(DATA_A_F16) && !defined(DATA_A_BF16) +#define K_PER_ITER 8 +#else +#define K_PER_ITER 2 +#endif + +// Tile width: number of K-elements staged per iteration +// Each iteration, BLOCK_SIZE threads each handle K_PER_ITER elements, +// so we stage TILE_K = BLOCK_SIZE * K_PER_ITER elements of K per tile. +// For NUM_ROWS output rows, shared memory = NUM_ROWS * TILE_K * sizeof(float). +// +// With BLOCK_SIZE=64, K_PER_ITER=8: TILE_K=512, and for NUM_ROWS=2: +// shmem = 2 * 512 * 4 = 4096 bytes (well within 32KB limit) + +// Shared memory for staging dequantized A-matrix tile +// Layout: [NUM_ROWS][BLOCK_SIZE * K_PER_ITER] +shared float a_staging[NUM_ROWS][BLOCK_SIZE * K_PER_ITER]; + +uint a_offset, b_offset, d_offset, y_offset; + +void compute_outputs(const uint32_t first_row, const uint32_t num_rows) { + const uint tid = gl_LocalInvocationID.x; + + get_offsets(a_offset, b_offset, d_offset); + + y_offset = QUANT_R == 1 ? 1 : QUANT_K/2; + + FLOAT_TYPE temp[NUM_COLS][NUM_ROWS]; + + [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) { + [[unroll]] for (uint i = 0; i < NUM_ROWS; ++i) { + temp[j][i] = FLOAT_TYPE(0); + } + } + + uint num_iters = p.ncols / (K_PER_ITER * BLOCK_SIZE); + if (num_iters * K_PER_ITER * BLOCK_SIZE + K_PER_ITER*tid < p.ncols) { + num_iters++; + } + + for (uint i = 0; i < num_iters; ++i) { + const uint col = i*BLOCK_SIZE*K_PER_ITER + K_PER_ITER*tid; + const bool lastiter = (i == num_iters - 1); + + // ============================================================ + // Phase 1: DECODE - cooperatively dequantize A-tile into shmem + // ============================================================ + const uint iqs = (col%QUANT_K)/QUANT_R; + const uint iybs = col - col%QUANT_K; + + uint ibi = first_row * p.ncols; + [[unroll]] for (uint n = 0; n < num_rows; ++n) { + const uint ib = (ibi + col)/QUANT_K; + ibi += p.ncols; + + const uint shmem_base = K_PER_ITER * tid; + +#if K_PER_ITER == 8 + vec4 v = dequantize4(ib, iqs, a_offset); + vec4 v2 = dequantize4(ib, iqs+(4/QUANT_R), a_offset); + + const vec2 dm = get_dm(ib, a_offset); + if (dm.y != 0) { + v = v * dm.x + dm.y; + v2 = v2 * dm.x + dm.y; + } else { + v *= dm.x; + v2 *= dm.x; + } + + a_staging[n][shmem_base + 0] = v.x; + a_staging[n][shmem_base + 1] = v.y; + a_staging[n][shmem_base + 2] = v.z; + a_staging[n][shmem_base + 3] = v.w; + a_staging[n][shmem_base + 4] = v2.x; + a_staging[n][shmem_base + 5] = v2.y; + a_staging[n][shmem_base + 6] = v2.z; + a_staging[n][shmem_base + 7] = v2.w; +#else + const vec2 v = dequantize(ib, iqs, a_offset); + + a_staging[n][shmem_base + 0] = v.x; + a_staging[n][shmem_base + 1] = v.y; +#endif + } + + barrier(); + + // ============================================================ + // Phase 2: DOT - compute dot products from shared memory + // ============================================================ + [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) { + // Re-derive B-vector access coordinates + const uint b_iybs = col - col%QUANT_K; + const uint b_iqs = (col%QUANT_K)/QUANT_R; + +#if K_PER_ITER == 8 +#if QUANT_R == 2 + const vec4 bv02 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + b_iybs + b_iqs) / 4]); + const vec4 bv13 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + b_iybs + b_iqs + y_offset) / 4]); + const vec4 bv0 = vec4(bv02.x, bv13.x, bv02.y, bv13.y); + const vec4 bv1 = vec4(bv02.z, bv13.z, bv02.w, bv13.w); +#else + const vec4 bv0 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + b_iybs + b_iqs) / 4]); + const vec4 bv1 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + b_iybs + b_iqs) / 4 + 1]); +#endif +#else + const bool OOB = lastiter && (b_iybs + b_iqs + y_offset >= p.ncols); + + FLOAT_TYPE b0 = 0, b1 = 0; + b0 = FLOAT_TYPE(data_b[j*p.batch_stride_b + b_offset + b_iybs + b_iqs]); + if (!OOB) { + b1 = FLOAT_TYPE(data_b[j*p.batch_stride_b + b_offset + b_iybs + b_iqs + y_offset]); + } +#endif + + const uint shmem_base = K_PER_ITER * tid; + + [[unroll]] for (uint n = 0; n < num_rows; ++n) { +#if K_PER_ITER == 8 + vec4 sv0 = vec4(a_staging[n][shmem_base + 0], + a_staging[n][shmem_base + 1], + a_staging[n][shmem_base + 2], + a_staging[n][shmem_base + 3]); + vec4 sv1 = vec4(a_staging[n][shmem_base + 4], + a_staging[n][shmem_base + 5], + a_staging[n][shmem_base + 6], + a_staging[n][shmem_base + 7]); + + temp[j][n] += FLOAT_TYPE(dot(bv0, sv0) + dot(bv1, sv1)); +#else + temp[j][n] = fma(FLOAT_TYPE(a_staging[n][shmem_base + 0]), b0, temp[j][n]); + if (!OOB) { + temp[j][n] = fma(FLOAT_TYPE(a_staging[n][shmem_base + 1]), b1, temp[j][n]); + } +#endif + } + } + + barrier(); + } + + reduce_result(temp, d_offset, first_row, num_rows, tid); +} + +void main() { + const uint first_row = NUM_ROWS * (gl_WorkGroupID.x + gl_NumWorkGroups.x * gl_WorkGroupID.z); + +#ifdef NEEDS_INIT_IQ_SHMEM + init_iq_shmem(gl_WorkGroupSize); +#endif + + // do NUM_ROWS at a time, unless there aren't enough remaining rows + if (first_row + NUM_ROWS <= p.stride_d) { + compute_outputs(first_row, NUM_ROWS); + } else { + if (first_row >= p.stride_d) { + return; + } + compute_outputs(first_row, p.stride_d - first_row); + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 8186dba36f60..72cfbaba3200 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -691,6 +691,13 @@ void process_shaders() { string_to_spv("mul_mat_vec_" + tname + "_f32_f32_subgroup_no_shmem", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPE_VEC2", "vec2"}, {"B_TYPE_VEC4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}})); string_to_spv("mul_mat_vec_" + tname + "_f16_f32_subgroup_no_shmem", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float16_t"}, {"B_TYPE_VEC2", "f16vec2"}, {"B_TYPE_VEC4", "f16vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}})); + // Shared-memory staging variant (two-phase DECODE+DOT) + // Only for types that use the generic mul_mat_vec.comp (not k-quants or iq1/iq2/iq3 with custom shaders) + if (!string_ends_with(tname, "_k") && !string_starts_with(tname, "iq1_") && !string_starts_with(tname, "iq2_") && !string_starts_with(tname, "iq3_")) { + string_to_spv("mul_mat_vec_shmem_" + tname + "_f32_f32", "mul_mat_vec_shmem.comp", merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPE_VEC2", "vec2"}, {"B_TYPE_VEC4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}})); + string_to_spv("mul_mat_vec_shmem_" + tname + "_f16_f32", "mul_mat_vec_shmem.comp", merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float16_t"}, {"B_TYPE_VEC2", "f16vec2"}, {"B_TYPE_VEC4", "f16vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}})); + } + string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPE_VEC2", "vec2"}, {"B_TYPE_VEC4", "vec4"}, {"D_TYPE", "float"}})); string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPE_VEC2", "vec2"}, {"B_TYPE_VEC4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD", "1"}})); string_to_spv("mul_mat_vec_id_" + tname + "_f32_f32_subgroup_no_shmem", shader, merge_maps(base_dict, {{"MUL_MAT_ID", "1"}, {data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPE_VEC2", "vec2"}, {"B_TYPE_VEC4", "vec4"}, {"D_TYPE", "float"}, {"USE_SUBGROUP_ADD_NO_SHMEM", "1"}})); diff --git a/media/ht-llama-banner.png b/media/ht-llama-banner.png new file mode 100644 index 000000000000..7053f4f73dc5 Binary files /dev/null and b/media/ht-llama-banner.png differ diff --git a/tools/server/public/index.html.gz b/tools/server/public/index.html.gz index ffeed79ff706..2bcbca382c92 100644 Binary files a/tools/server/public/index.html.gz and b/tools/server/public/index.html.gz differ diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp index ed5e306fc5b2..3f8a9ef925a0 100644 --- a/tools/server/server-common.cpp +++ b/tools/server/server-common.cpp @@ -944,6 +944,10 @@ json oaicompat_chat_params_parse( } for (auto & msg : messages) { std::string role = json_value(msg, "role", std::string()); + if (opt.remap_developer_role && role == "developer") { + msg["role"] = "system"; + role = "system"; + } if (role != "assistant" && !msg.contains("content")) { throw std::invalid_argument("All non-assistant messages must contain 'content'"); } @@ -1142,7 +1146,7 @@ json oaicompat_chat_params_parse( return llama_params; } -json convert_responses_to_chatcmpl(const json & response_body) { +json convert_responses_to_chatcmpl(const json & response_body, bool remap_developer_role) { if (!response_body.contains("input")) { throw std::invalid_argument("'input' is required"); } @@ -1257,7 +1261,36 @@ json convert_responses_to_chatcmpl(const json & response_body) { } item["content"] = chatcmpl_content; - chatcmpl_messages.push_back(item); + if (remap_developer_role && (item.at("role") == "developer" || item.at("role") == "system")) { + // Remap "developer" → "system" and merge with any existing + // system message (e.g. from "instructions" field) to avoid + // duplicate system messages that some templates reject. + item["role"] = "system"; + std::string sys_text; + for (const auto & part : chatcmpl_content) { + if (part.value("type", "") == "text" && part.contains("text")) { + if (!sys_text.empty()) { + sys_text += "\n"; + } + sys_text += part.at("text").get(); + } + } + if (!chatcmpl_messages.empty() && chatcmpl_messages.front().value("role", "") == "system") { + std::string existing = chatcmpl_messages.front().value("content", ""); + if (!existing.empty() && !sys_text.empty()) { + existing += "\n\n"; + } + existing += sys_text; + chatcmpl_messages.front()["content"] = existing; + } else { + chatcmpl_messages.insert(chatcmpl_messages.begin(), { + {"role", "system"}, + {"content", sys_text}, + }); + } + } else { + chatcmpl_messages.push_back(item); + } } else if (exists_and_is_array(item, "content") && exists_and_is_string(item, "role") && item.at("role") == "assistant" && diff --git a/tools/server/server-common.h b/tools/server/server-common.h index 213ae52bb09e..72eb3dded8bb 100644 --- a/tools/server/server-common.h +++ b/tools/server/server-common.h @@ -287,6 +287,7 @@ struct server_chat_params { bool allow_image; bool allow_audio; bool enable_thinking = true; + bool remap_developer_role = false; int reasoning_budget = -1; std::string reasoning_budget_message; std::string media_path; @@ -303,7 +304,7 @@ json oaicompat_chat_params_parse( std::vector & out_files); // convert OpenAI Responses API format to OpenAI Chat Completions API format -json convert_responses_to_chatcmpl(const json & body); +json convert_responses_to_chatcmpl(const json & body, bool remap_developer_role = false); // convert Anthropic Messages API format to OpenAI Chat Completions API format json convert_anthropic_to_oai(const json & body); diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 6f737d94d020..64b21c8c1bb5 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -909,6 +909,7 @@ struct server_context_impl { /* allow_image */ mctx ? mtmd_support_vision(mctx) : false, /* allow_audio */ mctx ? mtmd_support_audio (mctx) : false, /* enable_thinking */ enable_thinking, + /* remap_developer_role */ params_base.remap_developer_role, /* reasoning_budget */ params_base.reasoning_budget, /* reasoning_budget_msg */ params_base.reasoning_budget_message, /* media_path */ params_base.media_path, @@ -3681,7 +3682,7 @@ void server_routes::init_routes() { this->post_responses_oai = [this](const server_http_req & req) { auto res = create_response(); std::vector files; - json body = convert_responses_to_chatcmpl(json::parse(req.body)); + json body = convert_responses_to_chatcmpl(json::parse(req.body), meta->chat_params.remap_developer_role); SRV_DBG("%s\n", "Request converted: OpenAI Responses -> OpenAI Chat Completions"); SRV_DBG("converted request: %s\n", body.dump().c_str()); json body_parsed = oaicompat_chat_params_parse( diff --git a/tools/server/server-models.cpp b/tools/server/server-models.cpp index 7e61844f0852..906cc5b509c6 100644 --- a/tools/server/server-models.cpp +++ b/tools/server/server-models.cpp @@ -3,6 +3,7 @@ #include "preset.h" #include "download.h" +#include "gguf.h" #include // TODO: remove this once we use HTTP client from download.h #include @@ -245,11 +246,14 @@ void server_models::load_models() { // 1. cached models common_presets cached_models = ctx_preset.load_from_cache(); SRV_INF("Loaded %zu cached model presets\n", cached_models.size()); - // 2. local models from --models-dir + // 2. local models from --models-dir (also discovers LoRA adapters) common_presets local_models; if (!base_params.models_dir.empty()) { - local_models = ctx_preset.load_from_models_dir(base_params.models_dir); - SRV_INF("Loaded %zu local model presets from %s\n", local_models.size(), base_params.models_dir.c_str()); + auto dir_result = ctx_preset.load_from_models_dir_with_lora(base_params.models_dir); + local_models = std::move(dir_result.presets); + discovered_adapters = std::move(dir_result.adapters); + SRV_INF("Loaded %zu local model presets and %zu LoRA adapters from %s\n", + local_models.size(), discovered_adapters.size(), base_params.models_dir.c_str()); } // 3. custom-path models from presets common_preset global = {}; @@ -336,6 +340,13 @@ void server_models::load_models() { } SRV_INF(" %c %s%s\n", has_custom ? '*' : ' ', name.c_str(), info.c_str()); } + + if (!discovered_adapters.empty()) { + SRV_INF("Available LoRA adapters (%zu)\n", discovered_adapters.size()); + for (const auto & adapter : discovered_adapters) { + SRV_INF(" %s (arch: %s)\n", adapter.name.c_str(), adapter.architecture.c_str()); + } + } } // handle custom stop-timeout option @@ -493,6 +504,11 @@ std::vector server_models::get_all_meta() { return result; } +std::vector server_models::get_discovered_adapters() { + std::lock_guard lk(mutex); + return discovered_adapters; +} + void server_models::unload_lru() { if (base_params.models_max <= 0) { return; // no limit @@ -571,6 +587,58 @@ void server_models::load(const std::string & name) { { SRV_INF("spawning server instance with name=%s on port %d\n", inst.meta.name.c_str(), inst.meta.port); + // inject discovered LoRA adapters matching the model's architecture + if (!discovered_adapters.empty()) { + std::string model_path; + if (inst.meta.preset.get_option("LLAMA_ARG_MODEL", model_path) && !model_path.empty()) { + // read model architecture from GGUF metadata + struct gguf_init_params gguf_params = { + /*.no_alloc = */ true, + /*.ctx = */ nullptr, + }; + struct gguf_context * gguf_ctx = gguf_init_from_file(model_path.c_str(), gguf_params); + if (gguf_ctx) { + std::string model_arch; + int64_t arch_key = gguf_find_key(gguf_ctx, "general.architecture"); + if (arch_key >= 0) { + const char * arch_val = gguf_get_val_str(gguf_ctx, arch_key); + if (arch_val) { + model_arch = arch_val; + } + } + gguf_free(gguf_ctx); + + if (!model_arch.empty()) { + // collect matching adapters and add as lora options + std::vector matching_lora_paths; + for (const auto & adapter : discovered_adapters) { + if (adapter.architecture == model_arch) { + matching_lora_paths.push_back(adapter.path); + SRV_INF("matching LoRA adapter '%s' with model '%s' (arch: %s)\n", + adapter.name.c_str(), name.c_str(), model_arch.c_str()); + } + } + if (!matching_lora_paths.empty()) { + // set lora-init-without-apply so adapters are loaded but not applied by default + inst.meta.preset.set_option(ctx_preset, "lora-init-without-apply", "true"); + // set lora paths as comma-separated value + std::string lora_csv; + for (const auto & p : matching_lora_paths) { + if (!lora_csv.empty()) { + lora_csv += ","; + } + lora_csv += p; + } + inst.meta.preset.set_option(ctx_preset, "lora", lora_csv); + } + } + } else { + SRV_WRN("failed to read GGUF metadata from model '%s', skipping LoRA adapter matching\n", + model_path.c_str()); + } + } + } + inst.meta.update_args(ctx_preset, bin_path); // render args std::vector child_args = inst.meta.args; // copy @@ -954,7 +1022,15 @@ void server_models_routes::init_routes() { this->proxy_post = [this](const server_http_req & req) { std::string method = "POST"; json body = json::parse(req.body); - std::string name = json_value(body, "model", std::string()); + std::string name; + // try to get model from JSON body (object with "model" field) + if (body.is_object()) { + name = json_value(body, "model", std::string()); + } + // fallback to query parameter (needed for endpoints like POST /lora-adapters where body is an array) + if (name.empty()) { + name = req.get_param("model"); + } bool autoload = is_autoload(params, req); auto error_res = std::make_unique(); if (!router_validate_model(name, models, autoload, error_res)) { @@ -1015,9 +1091,21 @@ void server_models_routes::init_routes() { // TODO: add other fields, may require reading GGUF metadata }); } + // include discovered LoRA adapters + json adapters_json = json::array(); + auto all_adapters = models.get_discovered_adapters(); + for (const auto & adapter : all_adapters) { + adapters_json.push_back(json { + {"name", adapter.name}, + {"path", adapter.path}, + {"architecture", adapter.architecture}, + }); + } + res_ok(res, { {"data", models_json}, {"object", "list"}, + {"lora_adapters", adapters_json}, }); return res; }; diff --git a/tools/server/server-models.h b/tools/server/server-models.h index 1db34b6c4df1..6831d9e45262 100644 --- a/tools/server/server-models.h +++ b/tools/server/server-models.h @@ -107,6 +107,9 @@ struct server_models { std::vector base_env; common_preset base_preset; // base preset from llama-server CLI args + // discovered LoRA adapters from models directory + std::vector discovered_adapters; + void update_meta(const std::string & name, const server_model_meta & meta); // unload least recently used models if the limit is reached @@ -129,6 +132,9 @@ struct server_models { // return a copy of all model metadata (thread-safe) std::vector get_all_meta(); + // return discovered LoRA adapters from models directory + std::vector get_discovered_adapters(); + // load and unload model instances // these functions are thread-safe void load(const std::string & name); diff --git a/tools/server/webui/src/app.css b/tools/server/webui/src/app.css index 3ab21f0cc7b6..476c863e11e9 100644 --- a/tools/server/webui/src/app.css +++ b/tools/server/webui/src/app.css @@ -6,76 +6,76 @@ :root { --radius: 0.625rem; - --background: oklch(1 0 0); - --foreground: oklch(0.145 0 0); - --card: oklch(1 0 0); - --card-foreground: oklch(0.145 0 0); - --popover: oklch(1 0 0); - --popover-foreground: oklch(0.145 0 0); - --primary: oklch(0.205 0 0); - --primary-foreground: oklch(0.985 0 0); - --secondary: oklch(0.95 0 0); - --secondary-foreground: oklch(0.205 0 0); - --muted: oklch(0.97 0 0); - --muted-foreground: oklch(0.556 0 0); - --accent: oklch(0.95 0 0); - --accent-foreground: oklch(0.205 0 0); + --background: oklch(0.965 0.06 190); + --foreground: oklch(0.30 0.18 295); + --card: oklch(0.965 0.06 190); + --card-foreground: oklch(0.30 0.18 295); + --popover: oklch(0.965 0.06 190); + --popover-foreground: oklch(0.30 0.18 295); + --primary: oklch(0.35 0.18 295); + --primary-foreground: oklch(0.955 0.06 190); + --secondary: oklch(0.92 0.06 190); + --secondary-foreground: oklch(0.35 0.18 295); + --muted: oklch(0.94 0.05 190); + --muted-foreground: oklch(0.5 0.12 295); + --accent: oklch(0.92 0.06 190); + --accent-foreground: oklch(0.35 0.18 295); --destructive: oklch(0.577 0.245 27.325); - --border: oklch(0.875 0 0); - --input: oklch(0.92 0 0); - --ring: oklch(0.708 0 0); + --border: oklch(0.85 0.05 190); + --input: oklch(0.89 0.05 190); + --ring: oklch(0.65 0.12 295); --chart-1: oklch(0.646 0.222 41.116); --chart-2: oklch(0.6 0.118 184.704); --chart-3: oklch(0.398 0.07 227.392); --chart-4: oklch(0.828 0.189 84.429); --chart-5: oklch(0.769 0.188 70.08); - --sidebar: oklch(0.987 0 0); - --sidebar-foreground: oklch(0.145 0 0); - --sidebar-primary: oklch(0.205 0 0); - --sidebar-primary-foreground: oklch(0.985 0 0); - --sidebar-accent: oklch(0.97 0 0); - --sidebar-accent-foreground: oklch(0.205 0 0); - --sidebar-border: oklch(0.922 0 0); - --sidebar-ring: oklch(0.708 0 0); - --code-background: oklch(0.985 0 0); - --code-foreground: oklch(0.145 0 0); + --sidebar: oklch(0.955 0.06 190); + --sidebar-foreground: oklch(0.30 0.18 295); + --sidebar-primary: oklch(0.35 0.18 295); + --sidebar-primary-foreground: oklch(0.955 0.06 190); + --sidebar-accent: oklch(0.94 0.05 190); + --sidebar-accent-foreground: oklch(0.35 0.18 295); + --sidebar-border: oklch(0.89 0.05 190); + --sidebar-ring: oklch(0.65 0.12 295); + --code-background: oklch(0.955 0.06 190); + --code-foreground: oklch(0.30 0.18 295); --layer-popover: 1000000; } .dark { - --background: oklch(0.16 0 0); - --foreground: oklch(0.985 0 0); - --card: oklch(0.205 0 0); - --card-foreground: oklch(0.985 0 0); - --popover: oklch(0.205 0 0); - --popover-foreground: oklch(0.985 0 0); - --primary: oklch(0.922 0 0); - --primary-foreground: oklch(0.205 0 0); - --secondary: oklch(0.29 0 0); - --secondary-foreground: oklch(0.985 0 0); - --muted: oklch(0.269 0 0); - --muted-foreground: oklch(0.708 0 0); - --accent: oklch(0.269 0 0); - --accent-foreground: oklch(0.985 0 0); + --background: oklch(0.17 0.09 190); + --foreground: oklch(0.90 0.14 295); + --card: oklch(0.21 0.09 190); + --card-foreground: oklch(0.90 0.14 295); + --popover: oklch(0.21 0.09 190); + --popover-foreground: oklch(0.90 0.14 295); + --primary: oklch(0.85 0.14 295); + --primary-foreground: oklch(0.21 0.09 190); + --secondary: oklch(0.3 0.09 190); + --secondary-foreground: oklch(0.90 0.14 295); + --muted: oklch(0.28 0.07 190); + --muted-foreground: oklch(0.7 0.1 295); + --accent: oklch(0.28 0.07 190); + --accent-foreground: oklch(0.90 0.14 295); --destructive: oklch(0.704 0.191 22.216); - --border: oklch(1 0 0 / 30%); - --input: oklch(1 0 0 / 30%); - --ring: oklch(0.556 0 0); + --border: oklch(1 0.08 295 / 30%); + --input: oklch(1 0.08 295 / 30%); + --ring: oklch(0.55 0.12 190); --chart-1: oklch(0.488 0.243 264.376); --chart-2: oklch(0.696 0.17 162.48); --chart-3: oklch(0.769 0.188 70.08); --chart-4: oklch(0.627 0.265 303.9); --chart-5: oklch(0.645 0.246 16.439); - --sidebar: oklch(0.19 0 0); - --sidebar-foreground: oklch(0.985 0 0); - --sidebar-primary: oklch(0.488 0.243 264.376); - --sidebar-primary-foreground: oklch(0.985 0 0); - --sidebar-accent: oklch(0.269 0 0); - --sidebar-accent-foreground: oklch(0.985 0 0); - --sidebar-border: oklch(1 0 0 / 10%); - --sidebar-ring: oklch(0.556 0 0); - --code-background: oklch(0.225 0 0); - --code-foreground: oklch(0.875 0 0); + --sidebar: oklch(0.19 0.09 190); + --sidebar-foreground: oklch(0.90 0.14 295); + --sidebar-primary: oklch(0.55 0.14 190); + --sidebar-primary-foreground: oklch(0.90 0.14 295); + --sidebar-accent: oklch(0.28 0.07 190); + --sidebar-accent-foreground: oklch(0.90 0.14 295); + --sidebar-border: oklch(1 0.08 295 / 10%); + --sidebar-ring: oklch(0.55 0.12 190); + --code-background: oklch(0.23 0.07 190); + --code-foreground: oklch(0.83 0.12 295); } @theme inline { @@ -185,4 +185,8 @@ -ms-overflow-style: none; scrollbar-width: none; } + + .animate-spin-reverse { + animation: spin 1s linear infinite reverse; + } } diff --git a/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte b/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte index 95c3c5da1ff2..04b8730d0f90 100644 --- a/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte +++ b/tools/server/webui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActions.svelte @@ -6,6 +6,7 @@ ChatFormActionAttachmentsSheet, ChatFormActionRecord, ChatFormActionSubmit, + LoraAdapters, McpServersSelector, ModelsSelector, ModelsSelectorSheet @@ -221,6 +222,11 @@
+ + {#if isMobile.current}
-

llama.cpp

+

ht-llama.cpp

{serverStore.props?.modalities?.audio diff --git a/tools/server/webui/src/lib/components/app/chat/ChatSidebar/ChatSidebar.svelte b/tools/server/webui/src/lib/components/app/chat/ChatSidebar/ChatSidebar.svelte index 751406b3e39f..4f65ecf61ee0 100644 --- a/tools/server/webui/src/lib/components/app/chat/ChatSidebar/ChatSidebar.svelte +++ b/tools/server/webui/src/lib/components/app/chat/ChatSidebar/ChatSidebar.svelte @@ -141,7 +141,7 @@ -

llama.cpp

+

ht-llama.cpp

 diff --git a/tools/server/webui/src/lib/components/app/index.ts b/tools/server/webui/src/lib/components/app/index.ts index 56bd8a485203..e5774e98d5ce 100644 --- a/tools/server/webui/src/lib/components/app/index.ts +++ b/tools/server/webui/src/lib/components/app/index.ts @@ -4,6 +4,7 @@ export * from './chat'; export * from './content'; export * from './dialogs'; export * from './forms'; +export * from './lora'; export * from './mcp'; export * from './misc'; export * from './models'; diff --git a/tools/server/webui/src/lib/components/app/lora/LoraAdapters.svelte b/tools/server/webui/src/lib/components/app/lora/LoraAdapters.svelte new file mode 100644 index 000000000000..128753c31140 --- /dev/null +++ b/tools/server/webui/src/lib/components/app/lora/LoraAdapters.svelte @@ -0,0 +1,203 @@ + + +{#if loraStore.isAvailable} + + 0 ? 'text-foreground' : 'text-muted-foreground', + isOpen ? 'text-foreground' : '' + )} + > + + + + {#if activeCount === 0} + LoRA + {:else} + LoRA ({activeCount}) + {/if} + + + {#if loraStore.applying} + + {:else} + + {/if} + + + +
+ + LoRA Adapters + +
+ + {#if showOnlyMatching} + {#if loraStore.adapters.length === 0} +
+ No matching adapters for this model. +
+ {:else} + {#each loraStore.adapters as adapter (adapter.id)} +
+ + + + + { + const val = parseFloat(e.currentTarget.value); + if (!isNaN(val)) { + loraStore.setScale(adapter.id, Math.max(0, Math.min(2, val))); + } + }} + class={cn( + 'h-6 w-14 shrink-0 rounded border bg-background px-1.5 text-right font-mono text-[11px] tabular-nums outline-none focus:border-primary focus:ring-1 focus:ring-primary', + adapter.enabled + ? 'border-border text-foreground' + : 'border-border/50 text-muted-foreground/60 cursor-not-allowed' + )} + /> +
+ {/each} + {/if} + + {#if loraStore.hasChanges} +
+
+ +
+
+ {/if} + {:else} + {#if loraStore.allDiscovered.length === 0} +
+ No LoRA adapters discovered. +
+ {:else} + {#each loraStore.allDiscovered as adapter (adapter.path)} +
+ + + + {adapter.name} + + + + {adapter.architecture} + +
+ {/each} + {/if} + {/if} + +
+ +
+ + +{/if} diff --git a/tools/server/webui/src/lib/components/app/lora/index.ts b/tools/server/webui/src/lib/components/app/lora/index.ts new file mode 100644 index 000000000000..720cf6d11603 --- /dev/null +++ b/tools/server/webui/src/lib/components/app/lora/index.ts @@ -0,0 +1,28 @@ +/** + * + * LORA + * + * Components for LoRA adapter management. Displays available adapters + * with toggle switches and scale sliders. Only renders when the server + * has LoRA adapters loaded. + * + */ + +/** + * **LoraAdapters** - Collapsible LoRA adapter management panel + * + * Shows loaded LoRA adapters with per-adapter enable/disable toggle + * and scale slider (0.0 - 2.0). Includes an "Apply" button when + * changes are pending. + * + * **Architecture:** + * - Fetches adapter list from server via loraStore on mount + * - Conditionally renders only when adapters are available + * - Uses loraStore for state management + * + * @example + * ```svelte + * + * ``` + */ +export { default as LoraAdapters } from './LoraAdapters.svelte'; diff --git a/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte b/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte index ccda78690cf4..64b62a539469 100644 --- a/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte +++ b/tools/server/webui/src/lib/components/app/models/ModelsSelector.svelte @@ -1,6 +1,7 @@
{/if}
+ {#if isCancelling} +
+ + Cancelling +
+ {:else if isUnloading} +
+ + Unloading +
+ {:else if isLoading} +
+ - {#if isLoading} - + +
{:else if isFailed}
diff --git a/tools/server/webui/src/lib/components/app/models/ModelsSelectorSheet.svelte b/tools/server/webui/src/lib/components/app/models/ModelsSelectorSheet.svelte index fe88c979f92c..84fd0baa03d6 100644 --- a/tools/server/webui/src/lib/components/app/models/ModelsSelectorSheet.svelte +++ b/tools/server/webui/src/lib/components/app/models/ModelsSelectorSheet.svelte @@ -1,6 +1,7 @@ diff --git a/tools/server/webui/src/lib/services/index.ts b/tools/server/webui/src/lib/services/index.ts index 7c98383d1f27..75a03e2efe80 100644 --- a/tools/server/webui/src/lib/services/index.ts +++ b/tools/server/webui/src/lib/services/index.ts @@ -136,6 +136,8 @@ export { DatabaseService } from './database.service'; * * @see modelsStore in stores/models.svelte.ts — primary consumer for reactive model state */ +export { LoraService } from './lora.service'; + export { ModelsService } from './models.service'; /** diff --git a/tools/server/webui/src/lib/services/lora.service.ts b/tools/server/webui/src/lib/services/lora.service.ts new file mode 100644 index 000000000000..292feb48d397 --- /dev/null +++ b/tools/server/webui/src/lib/services/lora.service.ts @@ -0,0 +1,56 @@ +import { apiFetch, apiFetchWithParams, apiPost } from '$lib/utils'; +import type { ApiDiscoveredLoraAdapter } from '$lib/types'; + +/** + * LoRA adapter entry returned by the server (loaded for a specific model) + */ +export interface LoraAdapter { + id: number; + path: string; + scale: number; +} + +/** + * Payload for updating adapter scales + */ +export interface LoraAdapterUpdate { + id: number; + scale: number; +} + +/** + * LoraService - Stateless service for LoRA adapter API communication + * + * In MODEL mode, calls go directly to the server. + * In ROUTER mode, a `model` param is needed to proxy to the correct child process. + */ +export class LoraService { + /** + * Fetch list of loaded LoRA adapters from the server. + * In router mode, pass modelId to route the request to the correct child. + */ + static async list(modelId?: string): Promise { + if (modelId) { + return apiFetchWithParams('./lora-adapters', { model: modelId }); + } + return apiFetch('/lora-adapters'); + } + + /** + * Fetch all discovered LoRA adapters from the /v1/models endpoint. + * Returns adapters from the models directory regardless of which model they're loaded for. + */ + static async listAll(): Promise { + const response = await apiFetch<{ lora_adapters?: ApiDiscoveredLoraAdapter[] }>('/v1/models'); + return response.lora_adapters ?? []; + } + + /** + * Update LoRA adapter scales on the server. + * In router mode, pass modelId as query param to route to the correct child. + */ + static async update(adapters: LoraAdapterUpdate[], modelId?: string): Promise { + const url = modelId ? `/lora-adapters?model=${encodeURIComponent(modelId)}` : '/lora-adapters'; + await apiPost(url, adapters); + } +} diff --git a/tools/server/webui/src/lib/stores/chat.svelte.ts b/tools/server/webui/src/lib/stores/chat.svelte.ts index 3af91be83674..4d023829a429 100644 --- a/tools/server/webui/src/lib/stores/chat.svelte.ts +++ b/tools/server/webui/src/lib/stores/chat.svelte.ts @@ -23,6 +23,7 @@ import { modelsStore, selectedModelContextSize } from '$lib/stores/models.svelte'; +import { loraStore } from '$lib/stores/lora.svelte'; import { normalizeModelName, filterByLeafNodeId, @@ -1556,6 +1557,9 @@ class ChatStore { if (currentConfig.custom) apiOptions.custom = currentConfig.custom; + const loraPayload = loraStore.getRequestPayload(); + if (loraPayload) apiOptions.lora = loraPayload; + return apiOptions; } } diff --git a/tools/server/webui/src/lib/stores/lora.svelte.ts b/tools/server/webui/src/lib/stores/lora.svelte.ts new file mode 100644 index 000000000000..cd232c6e5d8a --- /dev/null +++ b/tools/server/webui/src/lib/stores/lora.svelte.ts @@ -0,0 +1,182 @@ +import { toast } from 'svelte-sonner'; +import { LoraService } from '$lib/services/lora.service'; +import type { LoraAdapter, LoraAdapterUpdate } from '$lib/services/lora.service'; +import type { ApiDiscoveredLoraAdapter } from '$lib/types'; + +/** + * LoRA adapter with UI state (loaded for the current model) + */ +export interface LoraAdapterState { + id: number; + path: string; + name: string; + scale: number; + /** Scale value before the adapter was disabled (for toggle restore) */ + previousScale: number; + enabled: boolean; +} + +/** + * Extract display name from adapter path. + * Strips directory and .gguf extension. + */ +function extractAdapterName(path: string): string { + const filename = path.split(/[/\\]/).pop() || path; + return filename.replace(/\.gguf$/i, ''); +} + +/** + * loraStore - Reactive store for LoRA adapter management + * + * Manages: + * - Available LoRA adapters list (loaded for current model) + * - All discovered adapters from models directory + * - Per-adapter enable/disable toggle and scale + * - Applying changes to the server + */ +class LoraStore { + adapters = $state([]); + allDiscovered = $state([]); + loading = $state(false); + applying = $state(false); + + /** + * Whether any LoRA adapters are available (loaded or discovered) + */ + get isAvailable(): boolean { + return this.adapters.length > 0 || this.allDiscovered.length > 0; + } + + /** + * Get adapters that are currently enabled (scale > 0) + */ + get activeAdapters(): LoraAdapterUpdate[] { + return this.adapters + .filter((a) => a.enabled) + .map((a) => ({ id: a.id, scale: a.scale })); + } + + /** + * Get all adapter scales for API submission (enabled use their scale, disabled use 0) + */ + get adapterUpdates(): LoraAdapterUpdate[] { + return this.adapters.map((a) => ({ + id: a.id, + scale: a.enabled ? a.scale : 0 + })); + } + + /** + * Whether there are unsaved changes compared to what was last fetched + */ + private lastFetchedScales = $state>(new Map()); + + get hasChanges(): boolean { + return this.adapters.some((a) => { + const effectiveScale = a.enabled ? a.scale : 0; + const lastScale = this.lastFetchedScales.get(a.id) ?? 0; + return Math.abs(effectiveScale - lastScale) > 0.001; + }); + } + + /** Current model ID for router mode */ + private modelId: string | undefined = undefined; + + /** + * Fetch LoRA adapters from server. + * In router mode, pass modelId to route the request to the correct child. + * Also fetches all discovered adapters for the "show all" view. + */ + async fetch(modelId?: string): Promise { + if (this.loading) return; + this.loading = true; + this.modelId = modelId; + + try { + const [adapters, allAdapters] = await Promise.all([ + LoraService.list(modelId).catch(() => [] as LoraAdapter[]), + LoraService.listAll().catch(() => [] as ApiDiscoveredLoraAdapter[]) + ]); + + this.allDiscovered = allAdapters; + + this.lastFetchedScales = new Map(adapters.map((a) => [a.id, a.scale])); + + this.adapters = adapters.map((a: LoraAdapter) => ({ + id: a.id, + path: a.path, + name: extractAdapterName(a.path), + scale: 1.0, + previousScale: 1.0, + enabled: false + })); + } catch { + // Server may not support LoRA — this is fine, leave empty + this.adapters = []; + this.allDiscovered = []; + } finally { + this.loading = false; + } + } + + /** + * Set the scale for a specific adapter + */ + setScale(id: number, scale: number): void { + this.adapters = this.adapters.map((a) => + a.id === id ? { ...a, scale, previousScale: scale, enabled: scale > 0 } : a + ); + } + + /** + * Toggle an adapter on/off. + * When disabling, remembers the previous scale for re-enable. + * When enabling, restores the previous scale (or 1.0 default). + */ + toggle(id: number): void { + this.adapters = this.adapters.map((a) => { + if (a.id !== id) return a; + + if (a.enabled) { + return { ...a, enabled: false, previousScale: a.scale }; + } else { + return { ...a, enabled: true, scale: a.previousScale || 1.0 }; + } + }); + } + + /** + * Apply current adapter configuration to the server + */ + async applyChanges(): Promise { + if (this.applying) return; + this.applying = true; + + try { + await LoraService.update(this.adapterUpdates, this.modelId); + + // Update last fetched scales to reflect applied state + this.lastFetchedScales = new Map( + this.adapters.map((a) => [a.id, a.enabled ? a.scale : 0]) + ); + + toast.success('LoRA adapters updated'); + } catch (error) { + const msg = error instanceof Error ? error.message : 'Failed to update LoRA adapters'; + toast.error(msg); + } finally { + this.applying = false; + } + } + + /** + * Get the lora field for inclusion in chat completion requests. + * Returns undefined when no adapters are active (so the field is omitted). + */ + getRequestPayload(): LoraAdapterUpdate[] | undefined { + const active = this.activeAdapters; + return active.length > 0 ? active : undefined; + } +} + +export const loraStore = new LoraStore(); diff --git a/tools/server/webui/src/lib/stores/models.svelte.ts b/tools/server/webui/src/lib/stores/models.svelte.ts index d7c885844fc8..97aa63ce4f4c 100644 --- a/tools/server/webui/src/lib/stores/models.svelte.ts +++ b/tools/server/webui/src/lib/stores/models.svelte.ts @@ -56,6 +56,9 @@ class ModelsStore { private modelUsage = $state>>(new Map()); private modelLoadingStates = new SvelteMap(); + private loadAbortControllers = new Map(); + private modelCancellingStates = new SvelteMap(); + private modelUnloadingStates = new SvelteMap(); favoriteModelIds = $state>(this.loadFavoritesFromStorage()); @@ -230,6 +233,14 @@ class ModelsStore { return this.modelLoadingStates.get(modelId) ?? false; } + isModelCancelling(modelId: string): boolean { + return this.modelCancellingStates.get(modelId) ?? false; + } + + isModelUnloading(modelId: string): boolean { + return this.modelUnloadingStates.get(modelId) ?? false; + } + getModelStatus(modelId: string): ServerModelStatus | null { const model = this.routerModels.find((m) => m.id === modelId); return model?.status.value ?? null; @@ -516,6 +527,8 @@ class ModelsStore { /** Polling interval in ms for checking model status */ private static readonly STATUS_POLL_INTERVAL = 500; + /** Maximum number of cancel poll attempts before giving up */ + private static readonly MAX_CANCEL_POLL_ATTEMPTS = 60; /** * Poll for expected model status after load/unload operation. @@ -527,10 +540,15 @@ class ModelsStore { */ private async pollForModelStatus( modelId: string, - expectedStatus: ServerModelStatus + expectedStatus: ServerModelStatus, + signal?: AbortSignal ): Promise { let attempt = 0; while (true) { + if (signal?.aborted) { + return; + } + await this.fetchRouterModels(); const currentStatus = this.getModelStatus(modelId); @@ -549,6 +567,9 @@ class ModelsStore { currentStatus === ServerModelStatus.UNLOADED && attempt > 2 ) { + if (signal?.aborted) { + return; + } throw new Error('Model was unloaded unexpectedly during loading'); } @@ -568,21 +589,75 @@ class ModelsStore { if (this.modelLoadingStates.get(modelId)) return; + const controller = new AbortController(); + this.loadAbortControllers.set(modelId, controller); this.modelLoadingStates.set(modelId, true); this.error = null; try { await ModelsService.load(modelId); - await this.pollForModelStatus(modelId, ServerModelStatus.LOADED); + await this.pollForModelStatus(modelId, ServerModelStatus.LOADED, controller.signal); + + if (controller.signal.aborted) { + return; + } await this.updateModelModalities(modelId); toast.success(`Model loaded: ${this.toDisplayName(modelId)}`); } catch (error) { + if (controller.signal.aborted) { + return; + } this.error = error instanceof Error ? error.message : 'Failed to load model'; toast.error(`Failed to load model: ${this.toDisplayName(modelId)}`); throw error; } finally { this.modelLoadingStates.set(modelId, false); + this.loadAbortControllers.delete(modelId); + } + } + + /** + * Cancel an in-progress model load (ROUTER mode) + * Aborts the polling loop and sends unload to stop the subprocess + * @param modelId - Model identifier to cancel loading + */ + async cancelLoadModel(modelId: string): Promise { + const controller = this.loadAbortControllers.get(modelId); + if (controller) { + controller.abort(); + } + + this.modelLoadingStates.set(modelId, false); + this.loadAbortControllers.delete(modelId); + this.modelCancellingStates.set(modelId, true); + + try { + const status = this.getModelStatus(modelId); + if (status === ServerModelStatus.LOADING || status === ServerModelStatus.LOADED) { + await ModelsService.unload(modelId); + } + + // Poll until the model is no longer in LOADING state + for (let attempt = 0; attempt < ModelsStore.MAX_CANCEL_POLL_ATTEMPTS; attempt++) { + await this.fetchRouterModels(); + const currentStatus = this.getModelStatus(modelId); + if (currentStatus !== ServerModelStatus.LOADING) { + break; + } + if (attempt === ModelsStore.MAX_CANCEL_POLL_ATTEMPTS - 1) { + console.warn(`Cancel polling timed out for model: ${modelId}`); + toast.warning(`Cancel may not have completed for: ${this.toDisplayName(modelId)}`); + return; + } + await new Promise((resolve) => setTimeout(resolve, ModelsStore.STATUS_POLL_INTERVAL)); + } + + toast.info(`Model loading cancelled: ${this.toDisplayName(modelId)}`); + } catch (error) { + console.warn(`Failed to cancel model loading: ${modelId}`, error); + } finally { + this.modelCancellingStates.set(modelId, false); } } @@ -598,6 +673,7 @@ class ModelsStore { if (this.modelLoadingStates.get(modelId)) return; this.modelLoadingStates.set(modelId, true); + this.modelUnloadingStates.set(modelId, true); this.error = null; try { @@ -611,6 +687,7 @@ class ModelsStore { throw error; } finally { this.modelLoadingStates.set(modelId, false); + this.modelUnloadingStates.set(modelId, false); } } @@ -693,6 +770,10 @@ class ModelsStore { this.selectedModelName = null; this.modelUsage.clear(); this.modelLoadingStates.clear(); + this.loadAbortControllers.forEach((c) => c.abort()); + this.loadAbortControllers.clear(); + this.modelCancellingStates.clear(); + this.modelUnloadingStates.clear(); this.modelPropsCache.clear(); this.modelPropsFetching.clear(); } diff --git a/tools/server/webui/src/lib/types/api.d.ts b/tools/server/webui/src/lib/types/api.d.ts index c1a02342357d..6df92b416eab 100644 --- a/tools/server/webui/src/lib/types/api.d.ts +++ b/tools/server/webui/src/lib/types/api.d.ts @@ -436,6 +436,16 @@ export interface ApiRouterModelsStatusResponse { export interface ApiRouterModelsListResponse { object: string; data: ApiModelDataEntry[]; + lora_adapters?: ApiDiscoveredLoraAdapter[]; +} + +/** + * A LoRA adapter discovered in the models directory + */ +export interface ApiDiscoveredLoraAdapter { + name: string; + path: string; + architecture: string; } /** diff --git a/tools/server/webui/src/lib/types/index.ts b/tools/server/webui/src/lib/types/index.ts index 93a39f03dae3..d6ee937610bb 100644 --- a/tools/server/webui/src/lib/types/index.ts +++ b/tools/server/webui/src/lib/types/index.ts @@ -28,6 +28,7 @@ export type { ApiRouterModelsStatusRequest, ApiRouterModelsStatusResponse, ApiRouterModelsListResponse, + ApiDiscoveredLoraAdapter, ApiRouterModelsUnloadRequest, ApiRouterModelsUnloadResponse } from './api'; diff --git a/tools/server/webui/src/routes/+page.svelte b/tools/server/webui/src/routes/+page.svelte index 949ac273d511..407433e3c6a4 100644 --- a/tools/server/webui/src/routes/+page.svelte +++ b/tools/server/webui/src/routes/+page.svelte @@ -87,7 +87,7 @@ - llama.cpp - AI Chat Interface + ht-llama.cpp - AI Chat Interface diff --git a/tools/server/webui/src/routes/chat/[id]/+page.svelte b/tools/server/webui/src/routes/chat/[id]/+page.svelte index c0b2e0c5e7a2..a139320b6b1c 100644 --- a/tools/server/webui/src/routes/chat/[id]/+page.svelte +++ b/tools/server/webui/src/routes/chat/[id]/+page.svelte @@ -166,7 +166,7 @@ - {activeConversation()?.name || 'Chat'} - llama.cpp + {activeConversation()?.name || 'Chat'} - ht-llama.cpp diff --git a/tools/server/webui/vite.config.ts b/tools/server/webui/vite.config.ts index e4408f09e41c..586d6cad1d8f 100644 --- a/tools/server/webui/vite.config.ts +++ b/tools/server/webui/vite.config.ts @@ -161,7 +161,8 @@ export default defineConfig({ '/v1': 'http://localhost:8080', '/props': 'http://localhost:8080', '/models': 'http://localhost:8080', - '/cors-proxy': 'http://localhost:8080' + '/cors-proxy': 'http://localhost:8080', + '/lora-adapters': 'http://localhost:8080' }, headers: { 'Cross-Origin-Embedder-Policy': 'require-corp',