diff --git a/.github/workflows/make-release.yml b/.github/workflows/make-release.yml new file mode 100644 index 000000000..fed9c877c --- /dev/null +++ b/.github/workflows/make-release.yml @@ -0,0 +1,46 @@ +name: Make Release + +on: + workflow_dispatch: + inputs: + dry_run: + description: 'Dry run - validate without creating the tag' + required: true + type: boolean + default: true + +env: + GH_TOKEN: ${{ github.token }} + +permissions: + contents: write + +jobs: + make-release: + runs-on: ubuntu-latest + + steps: + - name: Checkout + uses: actions/checkout@v6 + + - name: Run release checks + id: checks + run: bash scripts/make-release-checks.sh ${{ github.event.inputs.dry_run == 'true' && '--dry-run' || '' }} + env: + GITHUB_REPOSITORY: ${{ github.repository }} + + - name: Create release tag + if: ${{ github.event.inputs.dry_run == 'false' }} + run: | + VERSION="${{ steps.checks.outputs.version }}" + git config user.name "github-actions[bot]" + git config user.email "github-actions[bot]@users.noreply.github.com" + git tag -a "${VERSION}" -m "Release ${VERSION}" + git push origin "${VERSION}" + echo "Created and pushed tag ${VERSION}" + + - name: Dry run summary + if: ${{ github.event.inputs.dry_run == 'true' }} + run: | + echo "Dry run complete - all checks passed." + echo "Would have created tag: ${{ steps.checks.outputs.version }}" diff --git a/common/arg.cpp b/common/arg.cpp index 6db1a3559..de60522bb 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -36,6 +36,7 @@ #include #include #include +#include #include // for hardware_concurrency #include @@ -561,6 +562,15 @@ void common_models_handler_apply(common_models_handler & handler, common_params } } + // infer the speculative type from the draft GGUF metadata when none is requested + // note: reads only the first split - sharded drafts need an explicit --spec-type + if (spec_types_is_default(params) && !params.speculative.draft.mparams.path.empty()) { + const auto types_gguf = common_speculative_types_from_gguf(params.speculative.draft.mparams.path); + if (!types_gguf.empty()) { + params.speculative.types = types_gguf; + } + } + // when a sidecar type is requested, the draft repo resolves to its sidecar instead of a full model const bool spec_sidecar_found = !plan_spec.mtp.local_path.empty() || !plan_spec.dflash.local_path.empty() || @@ -705,12 +715,61 @@ void common_models_handler_apply(common_models_handler & handler, common_params // CLI argument parsing functions // +// apply config files (if present), a later file overrides an earlier one: +// 1. system-wide: /etc/llama.cpp/config.ini (%PROGRAMDATA%\llama.cpp\config.ini on windows) +// 2. user-level: ${XDG_CONFIG_HOME:-~/.config}/llama.cpp/config.ini (%APPDATA%\llama.cpp\config.ini on windows) +static void common_params_apply_system_config(common_params & params, llama_example ex) { + std::vector paths; + +#if defined(_WIN32) + const std::string program_data = common_get_env("PROGRAMDATA"); + if (!program_data.empty()) { + paths.push_back(program_data + "\\llama.cpp\\config.ini"); + } +#else + paths.push_back("/etc/llama.cpp/config.ini"); +#endif + + try { + paths.push_back(fs_get_config_directory() + "config.ini"); + } catch (const std::exception & e) { + LOG_DBG("cannot read user-level config file, skipping: %s\n", e.what()); + } + + std::vector found; + for (const auto & path : paths) { + std::error_code ec; + if (std::filesystem::exists(path, ec)) { + found.push_back(path); + } + } + if (found.empty()) { + return; + } + + common_preset_context ctx(ex); + ctx.ignore_unknown_keys = true; // the same config file is shared by all programs + for (const auto & path : found) { + LOG_INF("using config file: %s\n", path.c_str()); + common_preset global; + common_presets presets = ctx.load_from_ini(path, global); + global.apply_to_params(params); + auto it = presets.find(COMMON_PRESET_DEFAULT_NAME); + if (it != presets.end()) { + it->second.apply_to_params(params); + } + } +} + static bool common_params_parse_ex(int argc, char ** argv, common_params_context & ctx_arg) { common_params & params = ctx_arg.params; // setup log directly from params.verbosity: see tools/cli/cli.cpp common_log_set_verbosity_thold(params.verbosity); + // config file applies first, so env variables and CLI arguments override it + common_params_apply_system_config(params, ctx_arg.ex); + std::unordered_map> arg_to_options; for (auto & opt : ctx_arg.options) { for (const auto & arg : opt.args) { @@ -1391,8 +1450,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex {"--version"}, "show version and build info", [](common_params &) { - fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit()); - fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target()); + llama_print_build_info(llama_version()); exit(0); } )); diff --git a/common/build-info.cpp.in b/common/build-info.cpp.in index f888fd079..4ec339708 100644 --- a/common/build-info.cpp.in +++ b/common/build-info.cpp.in @@ -29,7 +29,7 @@ const char * llama_build_info(void) { return s.c_str(); } -void llama_print_build_info(void) { - fprintf(stderr, "%s: build = %d (%s)\n", __func__, llama_build_number(), llama_commit()); - fprintf(stderr, "%s: built with %s for %s\n", __func__, llama_compiler(), llama_build_target()); +void llama_print_build_info(const char * llama_version) { + fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version, llama_build_number(), llama_commit()); + fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target()); } diff --git a/common/build-info.h b/common/build-info.h index afc7b49a1..e6a696acb 100644 --- a/common/build-info.h +++ b/common/build-info.h @@ -2,6 +2,7 @@ #define BUILD_INFO_H #define LLAMA_BUILD_NUMBER 999 +#define LLAMA_VERSION "1.0" #define LLAMA_COMMIT "KOBOLDCPP" #define LLAMA_COMPILER "KCPP" #define LLAMA_TARGET "KCPP" @@ -32,7 +33,7 @@ static inline const char * llama_build_info(void) { return s.c_str(); } -static inline void llama_print_build_info(void) { +static inline void llama_print_build_info(const char *) { fprintf(stderr, "%s: build = %d (%s)\n", __func__, llama_build_number(), llama_commit()); fprintf(stderr, "%s: built with %s for %s\n", __func__, llama_compiler(), llama_build_target()); } diff --git a/common/chat-peg-parser.cpp b/common/chat-peg-parser.cpp index 1910b4f1e..06737b165 100644 --- a/common/chat-peg-parser.cpp +++ b/common/chat-peg-parser.cpp @@ -594,9 +594,7 @@ common_peg_parser common_chat_peg_builder::python_style_tool_calls( // Full argument: name="value" or name=value auto arg_rule = tool_arg( - tool_arg_open(eps()) + - tool_arg_name(arg_name_parser) + - literal("=") + + tool_arg_open(tool_arg_name(arg_name_parser) + literal("=")) + arg_value_parser + tool_arg_close(eps()) ); diff --git a/common/common.cpp b/common/common.cpp index b6cfd16f4..b0d2cc599 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1025,20 +1025,21 @@ std::string fs_get_cache_directory() { std::string cache_directory = ""; auto ensure_trailing_slash = [](std::string p) { // Make sure to add trailing slash - if (p.back() != DIRECTORY_SEPARATOR) { + if (p.empty() || p.back() != DIRECTORY_SEPARATOR) { p += DIRECTORY_SEPARATOR; } return p; }; - if (getenv("LLAMA_CACHE")) { - cache_directory = std::getenv("LLAMA_CACHE"); - } else { + cache_directory = common_get_env("LLAMA_CACHE"); + if (cache_directory.empty()) { #if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \ defined(__OpenBSD__) || defined(__NetBSD__) - if (std::getenv("XDG_CACHE_HOME")) { - cache_directory = std::getenv("XDG_CACHE_HOME"); - } else if (std::getenv("HOME")) { - cache_directory = std::getenv("HOME") + std::string("/.cache/"); + const std::string xdg_cache_home = common_get_env("XDG_CACHE_HOME"); + const std::string home = common_get_env("HOME"); + if (!xdg_cache_home.empty()) { + cache_directory = xdg_cache_home; + } else if (!home.empty()) { + cache_directory = home + "/.cache/"; } else { #if defined(__linux__) /* no $HOME is defined, fallback to getpwuid */ @@ -1053,9 +1054,16 @@ std::string fs_get_cache_directory() { #endif /* defined(__linux__) */ } #elif defined(__APPLE__) - cache_directory = std::getenv("HOME") + std::string("/Library/Caches/"); + cache_directory = common_get_env("HOME"); + if (cache_directory.empty()) { + throw std::runtime_error("Failed to find $HOME directory"); + } + cache_directory += "/Library/Caches/"; #elif defined(_WIN32) - cache_directory = std::getenv("LOCALAPPDATA"); + cache_directory = common_get_env("LOCALAPPDATA"); + if (cache_directory.empty()) { + throw std::runtime_error("Failed to find %LOCALAPPDATA% directory"); + } #elif defined(__EMSCRIPTEN__) GGML_ABORT("not implemented on this platform"); #else @@ -1067,6 +1075,51 @@ std::string fs_get_cache_directory() { return ensure_trailing_slash(cache_directory); } +std::string fs_get_config_directory() { + std::string config_directory = ""; + auto ensure_trailing_slash = [](std::string p) { + if (p.empty() || p.back() != DIRECTORY_SEPARATOR) { + p += DIRECTORY_SEPARATOR; + } + return p; + }; +#if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \ + defined(__OpenBSD__) || defined(__NetBSD__) || defined(__APPLE__) + const std::string xdg_config_home = common_get_env("XDG_CONFIG_HOME"); + const std::string home = common_get_env("HOME"); + if (!xdg_config_home.empty()) { + config_directory = xdg_config_home; + } else if (!home.empty()) { + config_directory = home + "/.config/"; + } else { +#if defined(__linux__) + /* no $HOME is defined, fallback to getpwuid */ + struct passwd *pw = getpwuid(getuid()); + if ((!pw) || (!pw->pw_dir)) { + throw std::runtime_error("Failed to find $HOME directory"); + } + + config_directory = std::string(pw->pw_dir) + std::string("/.config/"); +#else + throw std::runtime_error("Failed to find $HOME directory"); +#endif + } +#elif defined(_WIN32) + config_directory = common_get_env("APPDATA"); + if (config_directory.empty()) { + throw std::runtime_error("Failed to find %APPDATA% directory"); + } +#elif defined(__EMSCRIPTEN__) + // caller decides what to do when there is no config directory + throw std::runtime_error("not implemented on this platform"); +#else +# error Unknown architecture +#endif + config_directory = ensure_trailing_slash(config_directory); + config_directory += "llama.cpp"; + return ensure_trailing_slash(config_directory); +} + std::string fs_get_cache_file(const std::string & filename) { GGML_ASSERT(filename.find(DIRECTORY_SEPARATOR) == std::string::npos); std::string cache_directory = fs_get_cache_directory(); @@ -1228,6 +1281,8 @@ struct common_init_result::impl { // note: the order in which model, context, etc. are declared matters because their destructors will be called bottom-to-top + common_threadpools threadpools; + llama_model_ptr model; llama_context_ptr context; @@ -1329,6 +1384,10 @@ common_init_result::common_init_result(common_params & params, bool model_only) } pimpl->context.reset(lctx); + + set_process_priority(params.cpuparams.priority); + + pimpl->threadpools.init(lctx, params); } llama_model * common_init_result::model() { @@ -1677,6 +1736,10 @@ struct llama_context_params common_context_params_to_llama(const common_params & return cparams; } +// +// Threadpool utils +// + struct ggml_threadpool_params ggml_threadpool_params_from_cpu_params(const common_cpu_params & params) { struct ggml_threadpool_params tpp; @@ -1693,6 +1756,56 @@ struct ggml_threadpool_params ggml_threadpool_params_from_cpu_params(const commo return tpp; } +common_threadpools::~common_threadpools() { + if (!free_fn) { + return; + } + free_fn(threadpool); + free_fn(threadpool_batch); +} + +void common_threadpools::init(llama_context * ctx, const common_params & params) { + GGML_ASSERT(!threadpool); + GGML_ASSERT(!threadpool_batch); + + COM_INF("llama threadpool init, n_threads = %d\n", (int) params.cpuparams.n_threads); + + auto * cpu_dev = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); + if (!cpu_dev) { + COM_WRN("%s", "no CPU backend found\n"); + return; + } + auto * reg = ggml_backend_dev_backend_reg(cpu_dev); + auto * ggml_threadpool_new_fn = (decltype(ggml_threadpool_new) *) ggml_backend_reg_get_proc_address(reg, "ggml_threadpool_new"); + free_fn = (decltype(ggml_threadpool_free) *) ggml_backend_reg_get_proc_address(reg, "ggml_threadpool_free"); + + struct ggml_threadpool_params tpp_batch = + ggml_threadpool_params_from_cpu_params(params.cpuparams_batch); + struct ggml_threadpool_params tpp = + ggml_threadpool_params_from_cpu_params(params.cpuparams); + + if (!ggml_threadpool_params_match(&tpp, &tpp_batch)) { + threadpool_batch = ggml_threadpool_new_fn(&tpp_batch); + if (!threadpool_batch) { + COM_WRN("batch threadpool create failed : n_threads %d\n", tpp_batch.n_threads); + return; + } + + // start the non-batch threadpool in the paused state + tpp.paused = true; + } + + threadpool = ggml_threadpool_new_fn(&tpp); + if (!threadpool) { + COM_WRN("threadpool create failed : n_threads %d\n", tpp.n_threads); + free_fn(threadpool_batch); + threadpool_batch = nullptr; + return; + } + + llama_attach_threadpool(ctx, threadpool, threadpool_batch); +} + // // Batch utils // diff --git a/common/common.h b/common/common.h index 73fffcf1c..ebdc23d18 100644 --- a/common/common.h +++ b/common/common.h @@ -882,6 +882,7 @@ bool fs_is_directory(const std::string & path); std::string fs_get_cache_directory(); std::string fs_get_cache_file(const std::string & filename); +std::string fs_get_config_directory(); struct common_file_info { std::string path; @@ -929,9 +930,8 @@ using common_init_result_ptr = std::unique_ptr; common_init_result_ptr common_init_from_params(common_params & params, bool model_only = false); -struct llama_model_params common_model_params_to_llama ( common_params & params); -struct llama_context_params common_context_params_to_llama(const common_params & params); -struct ggml_threadpool_params ggml_threadpool_params_from_cpu_params(const common_cpu_params & params); +struct llama_model_params common_model_params_to_llama ( common_params & params); +struct llama_context_params common_context_params_to_llama(const common_params & params); // clear LoRA adapters from context, then apply new list of adapters void common_set_adapter_lora(struct llama_context * ctx, std::vector & lora); @@ -942,6 +942,28 @@ std::string common_get_model_endpoint(); // for testing purposes char * common_get_model_or_exit(int, char*[]); +// +// Threadpool utils +// + +struct ggml_threadpool_params ggml_threadpool_params_from_cpu_params(const common_cpu_params & params); + +struct common_threadpools { + common_threadpools() = default; + ~common_threadpools(); + + common_threadpools(const common_threadpools &) = delete; + common_threadpools & operator=(const common_threadpools &) = delete; + + void init(llama_context * ctx, const common_params & params); + +private: + ggml_threadpool * threadpool = nullptr; + ggml_threadpool * threadpool_batch = nullptr; + + decltype(ggml_threadpool_free) * free_fn = nullptr; +}; + // // Context utils // diff --git a/common/preset.cpp b/common/preset.cpp index eb0c60b09..0b29af883 100644 --- a/common/preset.cpp +++ b/common/preset.cpp @@ -322,6 +322,8 @@ common_presets common_preset_context::load_from_ini(const std::string & path, co preset.options[opt] = value; } LOG_DBG("accepted option: %s = %s\n", key.c_str(), preset.options[opt].c_str()); + } else if (ignore_unknown_keys) { + LOG_WRN("ignoring option '%s' from %s: not supported by this program\n", key.c_str(), path.c_str()); } else { throw std::runtime_error(string_format( "option '%s' not recognized in preset '%s'", diff --git a/common/preset.h b/common/preset.h index 52935ebde..d8fc3915b 100644 --- a/common/preset.h +++ b/common/preset.h @@ -59,6 +59,10 @@ struct common_preset_context { bool filter_allowed_keys = false; std::set allowed_keys; + // if true, options unknown to the current example are skipped instead of being an error + // used for config files shared by all binaries, where each binary only knows a subset of options + bool ignore_unknown_keys = false; + // if only_remote_allowed is true, only accept whitelisted keys common_preset_context(llama_example ex); diff --git a/common/speculative.cpp b/common/speculative.cpp index e0342cd43..aec94199b 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -2,6 +2,7 @@ #include "common.h" #include "ggml.h" +#include "ggml-cpp.h" #include "llama.h" #include "log.h" #include "ngram-cache.cpp" @@ -912,6 +913,9 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { std::vector smpls; + // backend sampler chain per seq, attached to ctx_dft + std::vector backend_chains; + int32_t n_embd_dec = 0; // draft hidden size int32_t n_embd_enc = 0; // target_layer_ids_n * target_hidden_size int32_t n_embd_tgt = 0; // target model hidden size @@ -985,6 +989,22 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { s.reset(common_sampler_init(model_dft, sparams)); } + // offload draft sampling to the backend + backend_chains.assign(n_seq, nullptr); + if (this->params.backend_sampling) { + for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { + llama_sampler * chain = llama_sampler_chain_init(llama_sampler_chain_default_params()); + llama_sampler_chain_add(chain, llama_sampler_init_top_k(10)); + + if (!llama_set_sampler(ctx_dft, seq_id, chain)) { + SPC_WRN("backend offload failed for seq_id=%d; using CPU sampler\n", (int) seq_id); + llama_sampler_free(chain); + chain = nullptr; + } + backend_chains[seq_id] = chain; + } + } + // turn on extraction of the target layers' input embeddings for (uint32_t k = 0; k < target_layer_ids_n; ++k) { llama_set_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k], true); @@ -995,6 +1015,18 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { } ~common_speculative_impl_draft_dflash() override { + auto * ctx_dft = this->params.ctx_dft; + for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) backend_chains.size(); ++seq_id) { + if (backend_chains[seq_id] == nullptr) { + continue; + } + if (ctx_dft) { + llama_set_sampler(ctx_dft, seq_id, nullptr); + } + llama_sampler_free(backend_chains[seq_id]); + } + backend_chains.clear(); + llama_batch_free(batch); llama_batch_free(batch_inject); } @@ -2196,6 +2228,43 @@ common_speculative_type common_speculative_type_from_name(const std::string & na return it->second; } +std::vector common_speculative_types_from_gguf(const std::string & path) { + struct gguf_init_params gguf_params = { + /* .no_alloc = */ true, + /* .ctx = */ nullptr, + }; + + gguf_context_ptr gguf_ctx(gguf_init_from_file(path.c_str(), gguf_params)); + if (!gguf_ctx) { + return {}; + } + + const int64_t arch_id = gguf_find_key(gguf_ctx.get(), "general.architecture"); + if (arch_id < 0 || gguf_get_kv_type(gguf_ctx.get(), arch_id) != GGUF_TYPE_STRING) { + return {}; + } + + const std::string arch = gguf_get_val_str(gguf_ctx.get(), arch_id); + if (arch != "dflash") { + const uint32_t block_count = gguf_get_val_u32(gguf_ctx.get(), gguf_find_key(gguf_ctx.get(), (arch + ".block_count").c_str())); + + if (gguf_find_tensor(gguf_ctx.get(), ("blk." + std::to_string(block_count - 1) + ".nextn.eh_proj.weight").c_str()) >= 0) { + return { COMMON_SPECULATIVE_TYPE_DRAFT_MTP }; + } + + return {}; + } + + // the Markov head distinguishes draft-dspark from draft-dflash + const auto type = gguf_find_tensor(gguf_ctx.get(), "markov_w1.weight") >= 0 + ? COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK + : COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH; + + SPC_INF("auto-detected speculative type '%s' from the draft model metadata\n", common_speculative_type_to_str(type).c_str()); + + return { type }; +} + static uint32_t common_get_enabled_speculative_configs(const std::vector & configs) { uint32_t result = 0; for (size_t i = 0; i < configs.size(); i++) { @@ -2263,6 +2332,23 @@ common_params common_base_params_to_speculative(const common_params & params) { result.n_outputs_max = params.n_parallel; result.n_outputs_max_per_seq = 1; + // dflash/dspark decode the whole noise block in a single pass and sample every block position on the backend + // TODO: refactor such properties to be announced by the speculative types + // something like `struct common_speculative_type_props common_speculative_type_get_props(...);` + const bool has_block_draft = std::any_of( + params.speculative.types.begin(), params.speculative.types.end(), + [](common_speculative_type t) { + return t == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH || t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK; + }); + if (has_block_draft) { + // per-seq output positions: DFlash decodes anchor + n_max masks (n_max + 1); DSpark n_max -> +1 covers both + const int32_t per_seq = std::max(1, params_spec.n_max + 1); + result.n_outputs_max = params.n_parallel * per_seq; + if (params_spec.backend_sampling) { + result.n_outputs_max_per_seq = per_seq; + } + } + return result; } diff --git a/common/speculative.h b/common/speculative.h index 06b0992ed..12ae31b7d 100644 --- a/common/speculative.h +++ b/common/speculative.h @@ -14,6 +14,9 @@ const char * common_speculative_all_types_str(); // parse user provided types std::vector common_speculative_types_from_names(const std::vector & names); +// infer the spec types from the GGUF metadata of a draft model; empty if unknown +std::vector common_speculative_types_from_gguf(const std::string & path); + // convert string to type enum common_speculative_type common_speculative_type_from_name(const std::string & name); diff --git a/docs/release.md b/docs/release.md new file mode 100644 index 000000000..4335ef9d4 --- /dev/null +++ b/docs/release.md @@ -0,0 +1,49 @@ +# Release process + +llama.cpp uses [semantic versioning](https://semver.org) (`MAJOR.MINOR.PATCH`). + +## Version bump guidelines + +| Change type | Version component | +|---|---| +| Breaking change to the public C API (`include/llama.h`) | `MAJOR` | +| Backward-compatible features, model support, or API addition | `MINOR` | +| Bug fix with no API change | `PATCH` | + +The version is set in the three variables at the top of the root `CMakeLists.txt`: + +```cmake +set(LLAMA_VERSION_MAJOR 0) +set(LLAMA_VERSION_MINOR 1) +set(LLAMA_VERSION_PATCH 0) +``` + +_A version bump should be included in the PR that introduces the change, or in a +dedicated bump commit merged before the release is cut._ + +_TODO: add PR labels (`semver: patch`, `semver: minor`, `semver: major`) to help +identify which PRs require a version bump before cutting a release._ + +## Making a release + +Releases are created by running the [make-release](.github/workflows/make-release.yml) +which is a manual workflow. + +The workflow creates an annotated git tag (e.g. `v0.1.0`) and pushes it to the +remote. No GitHub Release object is created, the tag is the release artifact. + +## Building a release + +By default, `LLAMA_BUILD_IS_DEV=ON` which appends a `-dev` suffix to `LLAMA_VERSION`, +marking the build as a nightly/development build. Distributors building from a +release tag must pass `-DLLAMA_BUILD_IS_DEV=OFF` to produce a clean version string +(e.g. `0.1.0` instead of `0.1.0-dev`). + +## How releases reach users +Currently releases are not published to github releases, only nightly/development +builds are available there. The way users can access releases are using the following +channels: + +- **llama-install.sh** — downloads pre-built binaries built from the release tag. +- **Package managers** — consume the git tag directly. +- **Build from source** — users clone the repo and check out the tag. diff --git a/examples/test-cmake/.gitignore b/examples/test-cmake/.gitignore new file mode 100644 index 000000000..0ddff317a --- /dev/null +++ b/examples/test-cmake/.gitignore @@ -0,0 +1,3 @@ +llama-build-install +install +build diff --git a/examples/test-cmake/CMakeLists.txt b/examples/test-cmake/CMakeLists.txt new file mode 100644 index 000000000..ed5cb1f3c --- /dev/null +++ b/examples/test-cmake/CMakeLists.txt @@ -0,0 +1,13 @@ +cmake_minimum_required(VERSION 3.14) +project(llama-simple) + +set(CMAKE_CXX_STANDARD 17) + +find_package(llama 0.1.0 REQUIRED) + +add_executable(test-cmake test-cmake.cpp) +target_link_libraries(test-cmake PRIVATE llama) +target_compile_definitions(test-cmake PRIVATE + LLAMA_BUILD_NUMBER=${LLAMA_BUILD_NUMBER} + LLAMA_BUILD_COMMIT="${LLAMA_BUILD_COMMIT}" +) diff --git a/examples/test-cmake/README.md b/examples/test-cmake/README.md new file mode 100644 index 000000000..2f6a2fcfe --- /dev/null +++ b/examples/test-cmake/README.md @@ -0,0 +1,36 @@ +## cmake-test + +This is just for manually testing/developing of a llama.cpp installation to +enable troubleshooting issues and exploration. The idea is that this can be used +after making changes to llama.cpp installation cmake configuration and then +verify it locally. + +### Usage +The following will configure, build, and install llama.cpp + +Configuring/build/install: +```console +./build-install.sh +``` +The above command will create a directory named `install` in the current directory +which will have the follwing files in its lib directory: +```console +(venv) $ ls install/lib/ +cmake libggml.so libllama-common.so.0 libllama.so.0.1.0 llama.cpp +libggml-base.so libggml.so.0 libllama-common.so.0.1.0 libmtmd.so pkgconfig +libggml-base.so.0 libggml.so.0.19.0 libllama.so libmtmd.so.0 +libggml-base.so.0.19.0 libllama-common.so libllama.so.0 libmtmd.so.0.1.0 +``` + +Build/run this project using the installation created above: +```console +(venv) $ ./build.sh +-- Configuring done (0.0s) +-- Generating done (0.0s) +-- Build files have been written to: /path/to/llama.cpp/examples/test-cmake/build +[100%] Built target test-cmake +[test-cmake] Using llama.cpp version 0.1.0-dev-b10335 +[test-cmake] Initializing backend... +load_backend: loaded CPU backend from /path/to/llama.cpp/examples/test-cmake/install/lib/llama.cpp/libggml-cpu-alderlake.so +[test-cmake] Backend initialized. +``` diff --git a/examples/test-cmake/build-install.sh b/examples/test-cmake/build-install.sh new file mode 100755 index 000000000..77a6713d6 --- /dev/null +++ b/examples/test-cmake/build-install.sh @@ -0,0 +1,19 @@ +#!/bin/bash + +set -e + +rm -rf llama-build-install install + +cmake --fresh -S ../../. -B llama-build-install -DCMAKE_BUILD_TYPE=Release \ + -DBUILD_SHARED_LIBS=ON \ + -DGGML_BACKEND_DL=ON \ + -DGGML_CPU_ALL_VARIANTS=ON \ + -DLLAMA_TESTS_INSTALL=OFF \ + -DCMAKE_INSTALL_PREFIX="${PWD}/install" \ + -DGGML_BACKEND_DIR="${PWD}/install/lib/llama.cpp" \ + -DGGML_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \ + -DLLAMA_LIB_INSTALL_DIR="${PWD}/install/lib/llama.cpp" \ + -DLLAMA_TOOLS_INSTALL=OFF + +cmake --build llama-build-install --parallel 12 +cmake --install llama-build-install diff --git a/examples/test-cmake/build.sh b/examples/test-cmake/build.sh new file mode 100755 index 000000000..a212732b8 --- /dev/null +++ b/examples/test-cmake/build.sh @@ -0,0 +1,7 @@ +#!/bin/bash + +set -e + +cmake -S . -B build -DCMAKE_PREFIX_PATH="${PWD}/install" +cmake --build build +LD_LIBRARY_PATH="${PWD}/install/lib/llama.cpp:${PWD}/install/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" ./build/test-cmake diff --git a/examples/test-cmake/test-cmake.cpp b/examples/test-cmake/test-cmake.cpp new file mode 100644 index 000000000..c5c4765b4 --- /dev/null +++ b/examples/test-cmake/test-cmake.cpp @@ -0,0 +1,12 @@ +#include "llama.h" +#include + +int main(void) { + printf("[test-cmake] version: %s, build: %d (%s)\n", + llama_version(), LLAMA_BUILD_NUMBER, LLAMA_BUILD_COMMIT); + printf("[test-cmake] Initializing backend...\n"); + llama_backend_init(); + printf("[test-cmake] Backend initialized.\n"); + llama_backend_free(); + return 0; +} diff --git a/ggml/src/ggml-cpu/arch/arm/cpu-feats.cpp b/ggml/src/ggml-cpu/arch/arm/cpu-feats.cpp index a64ad7a3c..84a11eabd 100644 --- a/ggml/src/ggml-cpu/arch/arm/cpu-feats.cpp +++ b/ggml/src/ggml-cpu/arch/arm/cpu-feats.cpp @@ -1,90 +1,12 @@ #include "ggml-backend-impl.h" +#include "ggml-feats.h" -#if defined(__aarch64__) - -#if defined(__linux__) -#include -#elif defined(__APPLE__) -#include -#endif - -#if !defined(HWCAP_FPHP) -#define HWCAP_FPHP (1 << 9) -#endif - -#if !defined(HWCAP_ASIMDHP) -#define HWCAP_ASIMDHP (1 << 10) -#endif - -#if !defined(HWCAP_ASIMDDP) -#define HWCAP_ASIMDDP (1 << 20) -#endif - -#if !defined(HWCAP_SVE) -#define HWCAP_SVE (1 << 22) -#endif - -#if !defined(HWCAP2_SVE2) -#define HWCAP2_SVE2 (1 << 1) -#endif - -#if !defined(HWCAP2_I8MM) -#define HWCAP2_I8MM (1 << 13) -#endif - -#if !defined(HWCAP2_SME) -#define HWCAP2_SME (1 << 23) -#endif - -struct aarch64_features { - // has_neon not needed, aarch64 has NEON guaranteed - bool has_dotprod = false; - bool has_fp16 = false; - bool has_sve = false; - bool has_sve2 = false; - bool has_i8mm = false; - bool has_sme = false; - bool has_sme2 = false; - - aarch64_features() { -#if defined(__linux__) - uint32_t hwcap = getauxval(AT_HWCAP); - uint32_t hwcap2 = getauxval(AT_HWCAP2); - - has_dotprod = !!(hwcap & HWCAP_ASIMDDP); - has_fp16 = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP); - has_sve = !!(hwcap & HWCAP_SVE); - has_sve2 = !!(hwcap2 & HWCAP2_SVE2); - has_i8mm = !!(hwcap2 & HWCAP2_I8MM); - has_sme = !!(hwcap2 & HWCAP2_SME); -#elif defined(__APPLE__) - int oldp = 0; - size_t size = sizeof(oldp); - - if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, NULL, 0) == 0) { - has_dotprod = static_cast(oldp); - } - - if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, NULL, 0) == 0) { - has_i8mm = static_cast(oldp); - } - - if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, NULL, 0) == 0) { - has_sme = static_cast(oldp); - } - - if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, NULL, 0) == 0) { - has_sme2 = static_cast(oldp); - } - - // Apple apparently does not implement SVE yet -#endif - } -}; +#if defined(__aarch64__) || defined(_M_ARM64) static int ggml_backend_cpu_aarch64_score() { int score = 1; - aarch64_features af; + const ggml_feats_arch64_runtime_t af = ggml_feats_get_arch64_runtime(); + GGML_UNUSED(af); #ifdef GGML_USE_DOTPROD if (!af.has_dotprod) { return 0; } @@ -116,4 +38,4 @@ static int ggml_backend_cpu_aarch64_score() { GGML_BACKEND_DL_SCORE_IMPL(ggml_backend_cpu_aarch64_score) -# endif // defined(__aarch64__) +# endif // defined(__aarch64__) || defined(_M_ARM64) diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 909a8f28e..54dca81e6 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -3630,6 +3630,11 @@ struct ggml_cplan ggml_graph_plan( n_threads = 1; #endif +#if defined(__wasi__) + // WASI doesn't support parallelism yet + n_threads = 1; +#endif + size_t work_size = 0; struct ggml_cplan cplan; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 42ec809ce..25bb74383 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -8941,7 +8941,7 @@ static void ggml_compute_forward_flash_attn_ext_tiled( for (int tk = 0; tk < kv_tile; tk++) { const char * v_data = (const char *)v->data + (ic + tk)*nbv1 + iv2*nbv2 + iv3*nbv3; if (kv_type == GGML_TYPE_F16) { - ggml_fp16_to_fp32_row((const ggml_fp16_t *)v_data, V32 + tk * DV, DV); + ggml_cpu_fp16_to_fp32((const ggml_fp16_t *)v_data, V32 + tk * DV, DV); } else { memcpy(V32 + tk * DV, v_data, DV * sizeof(float)); } diff --git a/ggml/src/ggml-feats.h b/ggml/src/ggml-feats.h new file mode 100644 index 000000000..79a0afd87 --- /dev/null +++ b/ggml/src/ggml-feats.h @@ -0,0 +1,166 @@ +#pragma once + +#if defined(__aarch64__) || defined(_M_ARM64) + +#if defined(__linux__) +#include +#include + +#if !defined(HWCAP2_SVE2) +#define HWCAP2_SVE2 (1ULL << 1) +#endif + +#if !defined(HWCAP_FPHP) +#define HWCAP_FPHP (1 << 9) +#endif + +#if !defined(HWCAP_ASIMDHP) +#define HWCAP_ASIMDHP (1 << 10) +#endif + +#if !defined(HWCAP2_I8MM) +#define HWCAP2_I8MM (1ULL << 13) +#endif + +#if !defined(HWCAP_ASIMDDP) +#define HWCAP_ASIMDDP (1 << 20) +#endif + +#if !defined(HWCAP_SVE) +#define HWCAP_SVE (1 << 22) +#endif + +#if !defined(HWCAP2_SME) +#define HWCAP2_SME (1ULL << 23) +#endif + +#if !defined(HWCAP2_SME2) +#define HWCAP2_SME2 (1ULL << 37) +#endif + +#if !defined(PR_SVE_GET_VL) +#define PR_SVE_GET_VL 51 +#endif + +#if !defined(PR_SVE_VL_LEN_MASK) +#define PR_SVE_VL_LEN_MASK 0xffff +#endif + +#elif defined(__APPLE__) +#include +#elif defined(_WIN32) +#include + +#if !defined(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE 43 +#endif + +#if !defined(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_SVE_INSTRUCTIONS_AVAILABLE 46 +#endif + +#if !defined(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE 47 +#endif + +#if !defined(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE 66 +#endif + +#if !defined(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE 67 +#endif + +#if !defined(PF_ARM_SME_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_SME_INSTRUCTIONS_AVAILABLE 70 +#endif + +#if !defined(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE) +#define PF_ARM_SME2_INSTRUCTIONS_AVAILABLE 71 +#endif + +#endif + +typedef struct ggml_feats_arch64_runtime { + bool has_dotprod; + bool has_fp16; + bool has_sve; + bool has_sve2; + bool has_i8mm; + bool has_sme; + bool has_sme2; + int sve_cnt; +} ggml_feats_arch64_runtime_t; + +static inline ggml_feats_arch64_runtime_t ggml_feats_get_arch64_runtime(void) { + ggml_feats_arch64_runtime_t runtime_feat = {}; + +#if defined(__linux__) + const unsigned long hwcap = getauxval(AT_HWCAP); + const unsigned long hwcap2 = getauxval(AT_HWCAP2); + + runtime_feat.has_dotprod = !!(hwcap & HWCAP_ASIMDDP); + runtime_feat.has_fp16 = !!(hwcap & HWCAP_FPHP) && !!(hwcap & HWCAP_ASIMDHP);; + runtime_feat.has_sve = !!(hwcap & HWCAP_SVE); + runtime_feat.has_sve2 = !!(hwcap2 & HWCAP2_SVE2); + runtime_feat.has_i8mm = !!(hwcap2 & HWCAP2_I8MM); + runtime_feat.has_sme = !!(hwcap2 & HWCAP2_SME); + runtime_feat.has_sme2 = !!(hwcap2 & HWCAP2_SME2); + + if (runtime_feat.has_sve) { + const int vl = prctl(PR_SVE_GET_VL); + if (vl >= 0) { + runtime_feat.sve_cnt = vl & PR_SVE_VL_LEN_MASK; + } + } +#elif defined(__APPLE__) + int oldp = 0; + size_t size = sizeof(oldp); + + if (sysctlbyname("hw.optional.arm.FEAT_DotProd", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_dotprod = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_FP16", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_fp16 = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_SVE", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_sve = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_SVE2", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_sve2 = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_I8MM", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_i8mm = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_SME", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_sme = static_cast(oldp); + } + + if (sysctlbyname("hw.optional.arm.FEAT_SME2", &oldp, &size, nullptr, 0) == 0) { + runtime_feat.has_sme2 = static_cast(oldp); + } + + // Apple does not support userspace non-streaming SVE; keep SVE vector length unknown. + runtime_feat.sve_cnt = 0; +#elif defined (_WIN32) + runtime_feat.has_dotprod = IsProcessorFeaturePresent(PF_ARM_V82_DP_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_fp16 = IsProcessorFeaturePresent(PF_ARM_V82_FP16_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_sve = IsProcessorFeaturePresent(PF_ARM_SVE_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_sve2 = IsProcessorFeaturePresent(PF_ARM_SVE2_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_i8mm = IsProcessorFeaturePresent(PF_ARM_V82_I8MM_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_sme = IsProcessorFeaturePresent(PF_ARM_SME_INSTRUCTIONS_AVAILABLE) != 0; + runtime_feat.has_sme2 = IsProcessorFeaturePresent(PF_ARM_SME2_INSTRUCTIONS_AVAILABLE) != 0; + + // Windows exposes SVE feature presence, but not the runtime SVE vector length here. + runtime_feat.sve_cnt = 0; +#endif + + return runtime_feat; +} + +#endif // defined(__aarch64__) || defined(_M_ARM64) diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index c153bd821..953c75755 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -953,6 +953,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta nr0 = N_R0_IQ4_XS; smem = 32*sizeof(float); } break; + case GGML_TYPE_TQ2_0: + { + nsg = N_SG_TQ2_0; + nr0 = N_R0_TQ2_0; + } break; default: { GGML_LOG_ERROR("Asserting on type %d\n", (int) tsrc0); @@ -1182,6 +1187,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m nr0 = N_R0_IQ4_XS; smem = 32*sizeof(float); } break; + case GGML_TYPE_TQ2_0: + { + nsg = N_SG_TQ2_0; + nr0 = N_R0_TQ2_0; + } break; default: { GGML_LOG_ERROR("Asserting on type %d\n", (int)op->src[2]->type); diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index fc7d2e727..c91aa16de 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1413,6 +1413,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_IQ4_NL: + case GGML_TYPE_TQ2_0: case GGML_TYPE_I32: return true; default: @@ -1441,6 +1442,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_TQ2_0: switch (op->type) { case GGML_TYPE_F32: case GGML_TYPE_F16: @@ -1476,6 +1478,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_IQ4_NL: + case GGML_TYPE_TQ2_0: return true; default: return false; diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index e173b91c0..cf32c5c5b 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -87,6 +87,9 @@ #define N_R0_IQ4_XS 2 #define N_SG_IQ4_XS 2 +#define N_R0_TQ2_0 4 +#define N_SG_TQ2_0 2 + // function constants offsets #define FC_FLASH_ATTN_EXT_PAD 100 #define FC_FLASH_ATTN_EXT_BLK 200 diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 92258b737..b38b23edc 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -468,6 +468,34 @@ void quantize_iq4_nl(device const float * src, device block_iq4_nl & dst) { dst.d = sumq2 > 0 ? sumqx/sumq2 : d; } +void quantize_tq2_0(device const float * src, device block_tq2_0 & dst) { +#pragma METAL fp math_mode(safe) + float amax = 0.0f; // absolute max + + for (int j = 0; j < QK_K; j++) { + const float v = src[j]; + amax = MAX(amax, fabs(v)); + } + + const float d = amax; + const float id = d ? 1.0f/d : 0.0f; + + dst.d = (half) d; + + for (int j = 0; j < QK_K/4; j += 32) { + for (int m = 0; m < 32; ++m) { + uint8_t q = 0; + for (int n = 0; n < 4; ++n) { + // -1, 0, 1 -> 0, 1, 2 + int xi = (int)round(src[m + n*32] * id) + 1; + q += (uint8_t)((xi & 3) << (2*n)); + } + dst.qs[j + m] = q; + } + src += 4*32; + } +} + template void dequantize_q4_1(device const block_q4_1 * xb, short il, thread type4x4 & reg) { device const uint16_t * qs = ((device const uint16_t *)xb + 2); @@ -1021,6 +1049,25 @@ void dequantize_iq4_xs(device const block_iq4_xs * xb, short il, thread type4x4 } } +template +void dequantize_tq2_0(device const block_tq2_0 * xb, short il, thread type4x4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + + float4x4 reg_f; + + // 2 bits per element, 4 elements per byte, 128 elements per 32-byte group + const short base = il * 16; + for (int k = 0; k < 16; k++) { + const int i = base + k; + const int byte = ((i >> 7) & 1) * 32 + (i & 31); + const int l = (i >> 5) & 3; + reg_f[k/4][k%4] = d * (float)(((qs[byte] >> (2*l)) & 3) - 1); + } + + reg = (type4x4) reg_f; +} + enum ggml_sort_order { GGML_SORT_ORDER_ASC, GGML_SORT_ORDER_DESC, @@ -8001,6 +8048,7 @@ template [[host_name("kernel_cpy_f32_q4_1")]] kernel cpy_f_q_t kernel_cpy_f32_ template [[host_name("kernel_cpy_f32_q5_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_q5_1")]] kernel cpy_f_q_t kernel_cpy_f32_q; template [[host_name("kernel_cpy_f32_iq4_nl")]] kernel cpy_f_q_t kernel_cpy_f32_q; +template [[host_name("kernel_cpy_f32_tq2_0")]] kernel cpy_f_q_t kernel_cpy_f32_q; template kernel void kernel_cpy_q_f32( @@ -8048,6 +8096,8 @@ template [[host_name("kernel_cpy_q5_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32< template [[host_name("kernel_cpy_q5_1_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q8_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; +template [[host_name("kernel_cpy_tq2_0_f32")]] kernel cpy_q_f_t kernel_cpy_q_f32; + template [[host_name("kernel_cpy_q1_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q2_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q4_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; @@ -8056,6 +8106,8 @@ template [[host_name("kernel_cpy_q5_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32< template [[host_name("kernel_cpy_q5_1_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; template [[host_name("kernel_cpy_q8_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; +template [[host_name("kernel_cpy_tq2_0_f16")]] kernel cpy_q_f_t kernel_cpy_q_f32; + template kernel void kernel_concat( constant ggml_metal_kargs_concat & args, @@ -9822,6 +9874,121 @@ kernel void kernel_mul_mv_mxfp4_f32( kernel_mul_mv_mxfp4_f32_impl(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg); } +template +void kernel_mul_mv_tq2_0_f32_impl( + args_t args, + device const char * src0, + device const char * src1, + device char * dst, + threadgroup char * shmem, + uint3 tgpig, + ushort tiisg, + ushort sgitg) { + const short NSG = FC_mul_mv_nsg; + + const int nb = args.ne00/QK_K; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * NSG + sgitg) * nr0; + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + device const block_tq2_0 * ax[nr0]; + for (int row = 0; row < nr0; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03; + ax[row] = (device const block_tq2_0 *) ((device char *) src0 + offset0); + } + + float sumf[nr0] = {0.f}; + + // 8 threads per block, NBLOCK blocks per pass, 2 halves per block per pass + constexpr short NBLOCK = 4; + + constexpr short NB = N_SIMDWIDTH/NBLOCK; // threads per block + + const short blk = tiisg / NB; // 0..NBLOCK-1, block handled by this thread + const short htg = tiisg % NB; // 0..NB-1, thread within block (0..7) + + // byte and y base offsets within the block (32 elements per thread, 4 per byte) + device const float4 * yb4 = (device const float4 *)(y + 4*htg + blk*QK_K); + + // hoisted per-byte coefficients (from y) and total y-sum, shared across rows + // ref: https://github.com/ggml-org/llama.cpp/pull/26980 + float4 coef[4]; + + for (int ib = blk; ib < nb; ib += NBLOCK) { + FOR_UNROLL (short h0 = 0; h0 < 2; ++h0) { + const float4 y0 = yb4[ 0 + 32*h0]; + const float4 y1 = yb4[ 8 + 32*h0]; + const float4 y2 = yb4[16 + 32*h0]; + const float4 y3 = yb4[24 + 32*h0]; + + float sumy = 0.f; + FOR_UNROLL (short j = 0; j < 4; ++j) { + coef[j] = float4( + y0[j], + y1[j] - 4.0f*y0[j], + y2[j] - 4.0f*y1[j], + y3[j] - 4.0f*y2[j]); + + sumy += (y0[j] + y1[j]) + (y2[j] + y3[j]); + } + + FOR_UNROLL (short row = 0; row < nr0; ++row) { + device const block_tq2_0 & xb = ax[row][ib]; + device const uchar * qs = xb.qs + 4*htg + 32*h0; + + float sum = -sumy; + FOR_UNROLL (short j = 0; j < 4; ++j) { + // express the 2-bit field shifts (v>>2, v>>4, v>>6) as float floor ops + const float v = (float)qs[j]; + + const float f0 = v; + const float f1 = floor(v*0.25f); // v>>2 + const float f2 = floor(v*0.0625); // v>>4 + const float f3 = floor(v*0.015625); // v>>6 + + sum += coef[j][0]*f0 + coef[j][1]*f1 + coef[j][2]*f2 + coef[j][3]*f3; + } + + sumf[row] += xb.d * sum; + } + } + + yb4 += QK_K * NBLOCK / 4; + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row]); + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + +[[host_name("kernel_mul_mv_tq2_0_f32")]] +kernel void kernel_mul_mv_tq2_0_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + + kernel_mul_mv_tq2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + template kernel void kernel_get_rows_q( constant ggml_metal_kargs_get_rows & args, @@ -9915,6 +10082,38 @@ template [[host_name("kernel_get_rows_iq1_s")]] kernel get_rows_q_t kernel_get template [[host_name("kernel_get_rows_iq1_m")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_iq4_nl")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_iq4_xs")]] kernel get_rows_q_t kernel_get_rows_q; +template [[host_name("kernel_get_rows_tq2_0")]] kernel get_rows_q_t kernel_get_rows_q; + +template +kernel void kernel_set_rows_q( + constant ggml_metal_kargs_set_rows & args, + device const void * src0, + device const void * src1, + device float * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + uint tiitg[[thread_index_in_threadgroup]], + uint3 tptg [[threads_per_threadgroup]]) { + const int32_t i03 = tgpig.z; + const int32_t i02 = tgpig.y; + + const int32_t i12 = i03%args.ne12; + const int32_t i11 = i02%args.ne11; + + const int32_t i01 = tgpig.x*tptg.y + tiitg/tptg.x; + if (i01 >= args.ne01) { + return; + } + + const int32_t i10 = i01; + const TI i1 = ((const device TI *) ((const device char *) src1 + i10*args.nb10 + i11*args.nb11 + i12*args.nb12))[0]; + + device block_q * dst_row = ( device block_q *) (( device char *) dst + i1*args.nb1 + i02*args.nb2 + i03*args.nb3); + const device TS * src_row = (const device TS *) ((const device char *) src0 + i01*args.nb01 + i02*args.nb02 + i03*args.nb03); + + for (int ind = tiitg%tptg.x; ind < args.nk0; ind += tptg.x) { + quantize_func(src_row + QK*ind, dst_row[ind]); + } +} template kernel void kernel_set_rows_q32( @@ -10011,6 +10210,11 @@ template [[host_name("kernel_set_rows_f32_i32_q5_1")]] kernel set_rows_q32_t k template [[host_name("kernel_set_rows_f32_i64_iq4_nl")]] kernel set_rows_q32_t kernel_set_rows_q32; template [[host_name("kernel_set_rows_f32_i32_iq4_nl")]] kernel set_rows_q32_t kernel_set_rows_q32; +typedef decltype(kernel_set_rows_q) set_rows_qK_t; + +template [[host_name("kernel_set_rows_f32_i64_tq2_0")]] kernel set_rows_qK_t kernel_set_rows_q; +template [[host_name("kernel_set_rows_f32_i32_tq2_0")]] kernel set_rows_qK_t kernel_set_rows_q; + kernel void kernel_diag_f32( constant ggml_metal_kargs_diag & args, device const char * src0, @@ -10786,6 +10990,7 @@ template [[host_name("kernel_mul_mm_iq1_s_f32")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_iq1_m_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_iq4_nl_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_iq4_xs_f32")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_tq2_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_f32_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_f16_f16")]] kernel mul_mm_t kernel_mul_mm; @@ -10811,6 +11016,7 @@ template [[host_name("kernel_mul_mm_iq1_s_f16")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_iq1_m_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_iq4_nl_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_iq4_xs_f16")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_tq2_0_f16")]] kernel mul_mm_t kernel_mul_mm; // // indirect matrix-matrix multiplication @@ -10845,6 +11051,7 @@ template [[host_name("kernel_mul_mm_id_iq1_s_f32")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_iq1_m_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_iq4_nl_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_iq4_xs_f32")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_tq2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_f32_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_f16_f16")]] kernel mul_mm_id kernel_mul_mm_id; @@ -10870,6 +11077,7 @@ template [[host_name("kernel_mul_mm_id_iq1_s_f16")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_iq1_m_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_iq4_nl_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_iq4_xs_f16")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_tq2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; // // matrix-vector multiplication @@ -11027,6 +11235,7 @@ template [[host_name("kernel_mul_mv_id_iq3_s_f32")]] kernel kernel_mul_mv_id_t template [[host_name("kernel_mul_mv_id_iq2_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_iq4_nl_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; template [[host_name("kernel_mul_mv_id_iq4_xs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; +template [[host_name("kernel_mul_mv_id_tq2_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id>>; kernel void kernel_pool_2d_max_f32( constant ggml_metal_kargs_pool_2d & args, diff --git a/ggml/src/gguf.cpp b/ggml/src/gguf.cpp index 39e9a2a21..7e212a402 100644 --- a/ggml/src/gguf.cpp +++ b/ggml/src/gguf.cpp @@ -648,6 +648,13 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr GGML_ASSERT(int64_t(ctx->kv.size()) == n_kv); const int alignment_idx = gguf_find_key(ctx, GGUF_KEY_GENERAL_ALIGNMENT); + if (alignment_idx != -1 && gguf_get_kv_type(ctx, alignment_idx) != GGUF_TYPE_UINT32) { + GGML_LOG_ERROR("%s: key '%s' must be of type %s but is %s\n", + __func__, GGUF_KEY_GENERAL_ALIGNMENT, gguf_type_name(GGUF_TYPE_UINT32), + gguf_type_name(gguf_get_kv_type(ctx, alignment_idx))); + gguf_free(ctx); + return nullptr; + } ctx->alignment = alignment_idx == -1 ? GGUF_DEFAULT_ALIGNMENT : gguf_get_val_u32(ctx, alignment_idx); if (ctx->alignment == 0 || (ctx->alignment & (ctx->alignment - 1)) != 0) { @@ -726,9 +733,11 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr } // check that the total number of elements is representable - if (ok && ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) || - (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) || - (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) { + // (a zero-element tensor is trivially representable; the guard also avoids a division by zero below) + if (ok && ggml_nelements(&info.t) > 0 && + ((INT64_MAX/info.t.ne[1] <= info.t.ne[0]) || + (INT64_MAX/info.t.ne[2] <= info.t.ne[0]*info.t.ne[1]) || + (INT64_MAX/info.t.ne[3] <= info.t.ne[0]*info.t.ne[1]*info.t.ne[2]))) { GGML_LOG_ERROR("%s: total number of elements in tensor '%s' with shape " "(%" PRIi64 ", %" PRIi64 ", %" PRIi64 ", %" PRIi64 ") is >= %" PRIi64 "\n", diff --git a/include/llama.h b/include/llama.h index 42f464a9a..e071154bb 100644 --- a/include/llama.h +++ b/include/llama.h @@ -460,6 +460,8 @@ extern "C" { // lora adapter struct llama_adapter_lora; + LLAMA_API const char * llama_version(void); + // Helpers for getting default parameters // TODO: update API to start accepting pointers to params structs (https://github.com/ggml-org/llama.cpp/discussions/9172) LLAMA_API struct llama_model_params llama_model_default_params(void); @@ -886,6 +888,7 @@ extern "C" { const llama_token * tokens, size_t n_token_count); + // If tokens_out is NULL, only the token count is reported through n_token_count_out and no state is loaded LLAMA_API size_t llama_state_seq_load_file( struct llama_context * ctx, const char * filepath, diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 11366fc7f..3f105cfb1 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -3120,6 +3120,17 @@ size_t llama_context::state_seq_load_file(llama_seq_id seq_id, const char * file { const uint32_t n_token_count = file.read_u32(); + if (tokens_out == nullptr) { + const size_t n_token_max = (file.size() - file.tell()) / sizeof(llama_token); + if (n_token_count > n_token_max) { + LLAMA_LOG_ERROR("%s: token count in sequence state file exceeds the file size! %u > %zu\n", __func__, n_token_count, n_token_max); + return 0; + } + + *n_token_count_out = n_token_count; + return file.tell(); + } + if (n_token_count > n_token_capacity) { LLAMA_LOG_ERROR("%s: token count in sequence state file exceeded capacity! %u > %zu\n", __func__, n_token_count, n_token_capacity); return 0; diff --git a/src/llama.cpp b/src/llama.cpp index 2d8b893c0..81eb92e24 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -31,6 +31,7 @@ #include "ggml-cpp.h" #include "ggml-backend.h" #include "gguf.h" +#include "build-info.h" #include #include @@ -135,6 +136,10 @@ bool llama_supports_rpc(void) { return ggml_backend_reg_by_name("RPC") != nullptr; } +const char * llama_version(void) { + return LLAMA_VERSION; +} + void llama_backend_init(void) { ggml_time_init(); diff --git a/src/models/dflash.cpp b/src/models/dflash.cpp index bfbdb28ee..daaa20826 100644 --- a/src/models/dflash.cpp +++ b/src/models/dflash.cpp @@ -14,11 +14,14 @@ void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) { hparams.n_embd_inp_enc_impl = (uint32_t) target_layer_ids.size() * hparams.n_embd; - LLAMA_LOG_INFO("%s: DFlash extract_layers = [", __func__); - for (size_t i = 0; i < target_layer_ids.size(); ++i) { - LLAMA_LOG_INFO("%d%s", target_layer_ids[i], i + 1 < target_layer_ids.size() ? ", " : ""); + std::string layers; + const char * sep = ""; + for (const auto id : target_layer_ids) { + layers += sep; + layers += std::to_string(id); + sep = ", "; } - LLAMA_LOG_INFO("]\n"); + LLAMA_LOG_INFO("%s: DFlash extract_layers = [%s]\n", __func__, layers.c_str()); // DeepSeek-V4 DSpark backbone: stages are full DSV4 blocks, uniform sliding window (the draft KV ring) ml.get_key(LLM_KV_HYPER_CONNECTION_COUNT, hparams.dsv4_hc_mult, false); @@ -66,7 +69,7 @@ void llama_model_dflash::load_arch_hparams(llama_model_loader & ml) { // DFlash has a single rope, so the SWA rope == main rope. if (ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW, hparams.n_swa, false) && hparams.n_swa > 0) { hparams.swa_type = LLAMA_SWA_TYPE_STANDARD; - ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer()); + ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl); hparams.rope_freq_base_train_swa = hparams.rope_freq_base_train; hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train; } diff --git a/tools/completion/completion.cpp b/tools/completion/completion.cpp index ebbd19239..83da454ce 100644 --- a/tools/completion/completion.cpp +++ b/tools/completion/completion.cpp @@ -161,47 +161,6 @@ int llama_completion(int argc, char ** argv) { // start measuring performance timings from here llama_perf_context_reset(ctx); - LOG_INF("%s: llama threadpool init, n_threads = %d\n", __func__, (int) params.cpuparams.n_threads); - - auto * cpu_dev = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU); - if (!cpu_dev) { - LOG_ERR("%s: no CPU backend found\n", __func__); - return 1; - } - auto * reg = ggml_backend_dev_backend_reg(cpu_dev); - auto * ggml_threadpool_new_fn = (decltype(ggml_threadpool_new) *) ggml_backend_reg_get_proc_address(reg, "ggml_threadpool_new"); - auto * ggml_threadpool_free_fn = (decltype(ggml_threadpool_free) *) ggml_backend_reg_get_proc_address(reg, "ggml_threadpool_free"); - - struct ggml_threadpool_params tpp_batch = - ggml_threadpool_params_from_cpu_params(params.cpuparams_batch); - struct ggml_threadpool_params tpp = - ggml_threadpool_params_from_cpu_params(params.cpuparams); - - if (!set_process_priority(params.cpuparams.priority)) { - LOG_ERR("%s: error: failed to set process priority\n", __func__); - return 1; - } - - struct ggml_threadpool * threadpool_batch = NULL; - if (!ggml_threadpool_params_match(&tpp, &tpp_batch)) { - threadpool_batch = ggml_threadpool_new_fn(&tpp_batch); - if (!threadpool_batch) { - LOG_ERR("%s: batch threadpool create failed : n_threads %d\n", __func__, tpp_batch.n_threads); - return 1; - } - - // start the non-batch threadpool in the paused state - tpp.paused = true; - } - - struct ggml_threadpool * threadpool = ggml_threadpool_new_fn(&tpp); - if (!threadpool) { - LOG_ERR("%s: threadpool create failed : n_threads %d\n", __func__, tpp.n_threads); - return 1; - } - - llama_attach_threadpool(ctx, threadpool, threadpool_batch); - const int n_ctx_train = llama_model_n_ctx_train(model); const int n_ctx = llama_n_ctx(ctx); @@ -994,8 +953,5 @@ int llama_completion(int argc, char ** argv) { llama_backend_free(); - ggml_threadpool_free_fn(threadpool); - ggml_threadpool_free_fn(threadpool_batch); - return 0; } diff --git a/tools/gguf-split/gguf-split.cpp b/tools/gguf-split/gguf-split.cpp index 7044e0289..8ce4044b5 100644 --- a/tools/gguf-split/gguf-split.cpp +++ b/tools/gguf-split/gguf-split.cpp @@ -106,7 +106,7 @@ static void split_params_parse_ex(int argc, const char ** argv, split_params & p split_print_usage(argv[0]); exit(0); } else if (arg == "--version") { - fprintf(stderr, "version: %d (%s)\n", llama_build_number(), llama_commit()); + fprintf(stderr, "version: %s (build %d, commit %s)\n", llama_version(), llama_build_number(), llama_commit()); fprintf(stderr, "built with %s for %s\n", llama_compiler(), llama_build_target()); exit(0); } else if (arg == "--dry-run") { diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index 7dce2978e..edca716c5 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -610,7 +610,7 @@ int llama_quantize(int argc, char ** argv) { } } - llama_print_build_info(); + llama_print_build_info(llama_version()); if (params.dry_run) { fprintf(stderr, "%s: calculating quantization size for '%s' as %s", __func__, fname_inp.c_str(), ftype_str.c_str()); diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp index c9109fc96..5ff7685bb 100644 --- a/tools/server/server-common.cpp +++ b/tools/server/server-common.cpp @@ -13,6 +13,8 @@ #include #include #include +#include +#include json format_error_response(const std::string & message, const enum error_type type) { std::string type_str; @@ -58,6 +60,33 @@ json format_error_response(const std::string & message, const enum error_type ty }; } +// +// server_slot_stats +// + +json server_slot_stats::to_json() const { + json base = { + {"cache_n", n_prompt_cached}, + + {"prompt_n", n_prompt_processed}, + {"prompt_ms", t_prompt_ms()}, + {"prompt_per_token_ms", t_prompt_per_token_ms()}, + {"prompt_per_second", n_prompt_tps()}, + + {"predicted_n", n_gen}, + {"predicted_ms", t_gen_ms()}, + {"predicted_per_token_ms", t_gen_per_token_ms()}, + {"predicted_per_second", n_gen_tps()}, + }; + + if (n_draft_tokens > 0) { + base["draft_n"] = n_draft_tokens; + base["draft_n_accepted"] = n_draft_accepted; + } + + return base; +} + // // random string / id // @@ -235,6 +264,102 @@ static inline raw_buffer base64_decode(const std::string & encoded_string) { // server_tokens implementation // +namespace { + +constexpr uint32_t SERVER_TOKENS_STATE_VERSION = 1; + +uint32_t server_tokens_state_u32(size_t value) { + if (value > std::numeric_limits::max()) { + throw std::runtime_error("Server tokens state is too large"); + } + return value; +} + +class server_tokens_state_writer { +public: + template + void write(T value) { + static_assert(std::is_trivially_copyable::value, "T must be trivially copyable"); + const auto * ptr = reinterpret_cast(&value); + data.insert(data.end(), ptr, ptr + sizeof(value)); + } + + template + void write(const std::vector & values) { + static_assert(std::is_trivially_copyable::value, "T must be trivially copyable"); + write(server_tokens_state_u32(values.size())); + if (values.empty()) { + return; + } + const auto * ptr = reinterpret_cast(values.data()); + data.insert(data.end(), ptr, ptr + values.size() * sizeof(T)); + } + + void write_media_chunk(const mtmd_input_chunk * chunk) { + size_t chunk_size = 0; + if (mtmd_input_chunk_save(chunk, nullptr, 0, &chunk_size) != 0 || chunk_size == 0) { + throw std::runtime_error("Cannot serialize media chunk in server tokens"); + } + std::vector chunk_data(server_tokens_state_u32(chunk_size)); + if (mtmd_input_chunk_save(chunk, chunk_data.data(), chunk_data.size(), nullptr) != 0) { + throw std::runtime_error("Cannot serialize media chunk in server tokens"); + } + write(chunk_data); + } + + std::vector take() { + data.resize((data.size() + sizeof(llama_token) - 1) / sizeof(llama_token) * sizeof(llama_token), 0); + return std::move(data); + } + +private: + std::vector data; +}; + +class server_tokens_state_reader { +public: + server_tokens_state_reader(const char * data, size_t size) : data(data), size(size) {} + + template + T read() { + static_assert(std::is_trivially_copyable::value, "T must be trivially copyable"); + if (size - pos < sizeof(T)) { + throw std::runtime_error("Unexpected end of server tokens state"); + } + T value; + std::memcpy(&value, data + pos, sizeof(value)); + pos += sizeof(value); + return value; + } + + template + std::vector read_vector() { + static_assert(std::is_trivially_copyable::value, "T must be trivially copyable"); + const uint32_t n_values = read(); + // reject before resizing, so that a small corrupted payload cannot request a huge allocation + if (n_values > remaining() / sizeof(T)) { + throw std::runtime_error("Unexpected end of server tokens state"); + } + std::vector values(n_values); + if (n_values > 0) { + std::memcpy(values.data(), data + pos, values.size() * sizeof(T)); + pos += values.size() * sizeof(T); + } + return values; + } + + size_t remaining() const { + return size - pos; + } + +private: + const char * data; + size_t size; + size_t pos = 0; +}; + +} // namespace + server_tokens::server_tokens(mtmd::input_chunks & mtmd_chunks, bool has_mtmd) : has_mtmd(has_mtmd) { for (size_t i = 0; i < mtmd_chunks.size(); ++i) { push_back(mtmd_chunks[i]); @@ -408,6 +533,73 @@ const llama_tokens & server_tokens::get_tokens() const { return tokens; } +std::vector server_tokens::serialize() const { + static_assert(sizeof(llama_token) == sizeof(uint32_t), "unexpected llama_token size"); + + server_tokens_state_writer writer; + writer.write((llama_token) LLAMA_TOKEN_NULL); + writer.write(SERVER_TOKENS_STATE_VERSION); + writer.write(tokens); + + std::vector media_keys; + media_keys.reserve(map_idx_to_media.size()); + for (const auto & item : map_idx_to_media) { + media_keys.push_back(server_tokens_state_u32(item.first)); + } + writer.write(media_keys); + + for (const auto & item : map_idx_to_media) { + writer.write_media_chunk(item.second.get()); + } + + return writer.take(); +} + +server_tokens server_tokens::deserialize(const llama_tokens & packed, bool has_mtmd) { + static_assert(sizeof(llama_token) == sizeof(uint32_t), "unexpected llama_token size"); + + if (packed.empty() || packed[0] != LLAMA_TOKEN_NULL) { + // plain token list, as written by older versions + return server_tokens(packed, has_mtmd); + } + + server_tokens_state_reader reader(reinterpret_cast(packed.data()), packed.size() * sizeof(llama_token)); + reader.read(); // format marker + if (reader.read() != SERVER_TOKENS_STATE_VERSION) { + throw std::runtime_error("Unsupported server tokens state version"); + } + + const llama_tokens tokens = reader.read_vector(); + + // the media start indices, followed by the media chunks in the same order + const std::vector media_keys = reader.read_vector(); + if (!media_keys.empty() && !has_mtmd) { + throw std::runtime_error("Cannot restore media tokens without an mmproj"); + } + + server_tokens result(tokens, has_mtmd); + + for (const uint32_t key : media_keys) { + const size_t start_idx = key; + const std::vector chunk_data = reader.read_vector(); + if (chunk_data.empty()) { + throw std::runtime_error("Cannot load media chunk from server tokens state"); + } + + mtmd::input_chunk_ptr chunk(mtmd_input_chunk_load(chunk_data.data(), chunk_data.size())); + if (!chunk) { + throw std::runtime_error("Cannot load media chunk from server tokens state"); + } + result.map_idx_to_media[start_idx] = std::move(chunk); + } + + if (reader.remaining() >= sizeof(llama_token)) { + throw std::runtime_error("Trailing data in server tokens state"); + } + + return result; +} + llama_tokens server_tokens::get_text_tokens() const { llama_tokens res; res.reserve(tokens.size()); @@ -530,14 +722,28 @@ bool server_tokens::validate(const struct llama_context * ctx) const { const llama_model * model = llama_get_model(ctx); const llama_vocab * vocab = llama_model_get_vocab(model); const int32_t n_vocab = llama_vocab_n_tokens(vocab); + size_t n_media = 0; for (size_t i = 0; i < tokens.size(); ++i) { const auto & t = tokens[i]; if (t == LLAMA_TOKEN_NULL) { try { const auto & chunk = find_chunk(i); - size_t n_tokens = mtmd_input_chunk_get_n_tokens(chunk.get()); - i += n_tokens - 1; // will be +1 by the for loop + if (mtmd_input_chunk_get_type(chunk.get()) == MTMD_INPUT_CHUNK_TYPE_TEXT) { + return false; + } + const size_t n_tokens = mtmd_input_chunk_get_n_tokens(chunk.get()); + const llama_pos n_pos = mtmd_input_chunk_get_n_pos(chunk.get()); + if (n_tokens == 0 || n_pos <= 0 || n_tokens > tokens.size() - i) { + return false; + } + for (size_t j = i; j < i + n_tokens; ++j) { + if (tokens[j] != LLAMA_TOKEN_NULL) { + return false; + } + } + ++n_media; + i += n_tokens - 1; } catch (const std::exception & e) { return false; } @@ -545,7 +751,7 @@ bool server_tokens::validate(const struct llama_context * ctx) const { return false; } } - return true; + return n_media == map_idx_to_media.size(); } server_tokens server_tokens::clone() const { diff --git a/tools/server/server-common.h b/tools/server/server-common.h index 6ef797ebb..7082abdd9 100644 --- a/tools/server/server-common.h +++ b/tools/server/server-common.h @@ -201,11 +201,14 @@ public: // for compatibility with context shift and prompt truncation void insert(const llama_tokens & inp_tokens); - // for compatibility with speculative decoding, ctx shift, slot save/load + // for compatibility with speculative decoding, ctx shift const llama_tokens & get_tokens() const; llama_tokens get_text_tokens() const; + std::vector serialize() const; + static server_tokens deserialize(const llama_tokens & packed, bool has_mtmd); + // for compatibility with speculative decoding void set_token(llama_pos pos, llama_token id); @@ -213,9 +216,6 @@ public: bool empty() const { return tokens.empty(); } - // true if the sequence actually contains image/audio chunks. - bool has_media() const { return !map_idx_to_media.empty(); } - void clear() { map_idx_to_media.clear(); tokens.clear(); @@ -230,7 +230,7 @@ public: // split the tokens into message spans, skipping over media chunks common_chat_msg_spans find_message_spans(const common_chat_msg_delimiters & delims) const; - // make sure all text tokens are within the vocab range + // check text token IDs and the mapping between media chunks and token ranges bool validate(const struct llama_context * ctx) const; server_tokens clone() const; @@ -334,6 +334,160 @@ json format_response_rerank( std::vector & texts, int top_n); +// +// stats and metrics +// + +// shared between server_slot and server_task_result_* +struct server_slot_stats { + uint64_t n_prompt_cached = 0; + uint64_t n_prompt_processed = 0; + uint64_t n_gen = 0; + + // speculative decoding stats + // note: the per-position breakdown lives in server_slot, it is not needed in a task result + uint64_t n_draft_tokens = 0; + uint64_t n_draft_accepted = 0; + uint64_t n_draft_verif_steps = 0; + + // these are absolute timestamps (in us) + // note: must be signed - they are subtracted before the later ones are set + int64_t t_start = 0; + int64_t t_prompt_last = 0; + int64_t t_gen_last = 0; + + // can only move one direction: start -> prompt -> gen + void update_prompt_start() { + GGML_ASSERT(t_start == 0); + t_start = ggml_time_us(); + } + void set_prompt_last(int64_t t_us) { + GGML_ASSERT(t_start > 0); + t_prompt_last = t_us; + } + void update_prompt_last() { + set_prompt_last(ggml_time_us()); + } + void update_gen_last() { + GGML_ASSERT(t_prompt_last > 0); + t_gen_last = ggml_time_us(); + } + + // these are time durations + int64_t t_elapsed_us() const { + return ggml_time_us() - t_start; + } + double t_prompt_ms() const { + if (t_prompt_last == 0) { + return 0.0; // the prompt is not processed yet + } + return (t_prompt_last - t_start) / 1000.0; + } + int64_t t_gen_us() const { + if (t_gen_last == 0) { + return 0; // the generation is not started yet + } + // clamp to 1 us, the first token can land in the same us as t_prompt_last + return std::max(1, t_gen_last - t_prompt_last); + } + double t_gen_ms() const { + return t_gen_us() / 1000.0; + } + + // number of decode steps spent on generation + // the first token is free, it comes from the logits of the last prompt batch + uint64_t n_gen_steps() const { + return n_gen > 0 ? n_gen - 1 : 0; + } + + // other derived metrics + // note: all of them return 0.0 if the divisor is not known yet + double t_prompt_per_token_ms() const { + return n_prompt_processed > 0 ? t_prompt_ms() / n_prompt_processed : 0.0; + } + double t_gen_per_token_ms() const { + return n_gen_steps() > 0 ? t_gen_ms() / n_gen_steps() : 0.0; + } + double n_prompt_tps() const { + const double t_ms = t_prompt_ms(); + return t_ms > 0.0 ? 1e3 / t_ms * n_prompt_processed : 0.0; + } + double n_gen_tps() const { + const double t_ms = t_gen_ms(); + return t_ms > 0.0 ? 1e3 / t_ms * n_gen_steps() : 0.0; + } + + // false if the slot never started, i.e. the task result carries no stats + bool is_set() const { + return t_start > 0; + } + + json to_json() const; +}; + +// shared between server_context_impl and server_task_result_* +// unlike server_slot_stats, server_metrics is server-global and cumulative, not tied to a slot +struct server_metrics { + int64_t t_start = 0; + + struct bucket { + uint64_t count = 0; // number of tokens + uint64_t steps = 0; // number of decode steps, + // this excludes first generated token (logits from prompt batch) + uint64_t time = 0; // in microseconds + + // the rate uses the decode steps, so that "free" tokens do not inflate it + double n_per_second() const { + return time > 0 ? (double) steps / (double) time * 1e6 : 0.0; + } + + void add(uint64_t n, uint64_t n_steps, uint64_t t_us) { + count += n; + steps += n_steps; + time += t_us; + } + }; + + // these are reset by reset_bucket(), only the rate is read from them + bucket prompt_bucket; + bucket predict_bucket; + + // metrics below are cumulative since the server started + bucket prompt; // only processed tokens, cached ones are counted separately below + bucket predict; + + // tokens reused from the cache need no decode, so they only have a count + uint64_t n_prompt_cached = 0; + + uint64_t n_tokens_max = 0; + + uint64_t n_decode = 0; + uint64_t n_busy_slots = 0; + + uint64_t n_draft_tokens = 0; // Total draft tokens generated + uint64_t n_draft_accepted = 0; // Draft tokens actually accepted + uint64_t n_draft_verif_steps = 0; // Total draft token verification steps by the target model + std::vector n_accepted_per_pos; // Accepted tokens per draft position + + void init() { + t_start = ggml_time_us(); + } + + void reset_bucket() { + prompt_bucket = {}; + predict_bucket = {}; + } + + void add_prompt(uint64_t n_tokens, uint64_t t_us) { + prompt .add(n_tokens, n_tokens, t_us); + prompt_bucket.add(n_tokens, n_tokens, t_us); + } + + void add_prompt_cached(uint64_t n_tokens) { + n_prompt_cached += n_tokens; + } +}; + // // other utils // diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index d75c6856d..f02a1da68 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -74,6 +74,7 @@ struct server_batch { llama_token token; llama_pos pos; bool output; + bool is_prompt; // for stats tracking }; std::vector tokens; int32_t n_tokens_alloc = 0; @@ -109,22 +110,22 @@ struct server_batch { tokens.reserve(n_tokens_alloc); } - bool add(int32_t id_slot, llama_token token, llama_pos pos, bool output) { + bool add(int32_t id_slot, llama_token token, llama_pos pos, bool output, bool is_prompt) { GGML_ASSERT(!has_embd); // cannot mix tokens + embd in same batch GGML_ASSERT(batch.pos != nullptr); if ((int32_t)tokens.size() >= n_tokens_alloc) { return false; } - tokens.push_back({ id_slot, token, pos, output }); + tokens.push_back({ id_slot, token, pos, output, is_prompt }); return true; } - bool add(int32_t id_slot, const std::vector & embd_in, llama_pos pos, bool output) { + bool add(int32_t id_slot, const std::vector & embd_in, llama_pos pos, bool output, bool is_prompt) { GGML_ASSERT(batch.pos != nullptr); if ((int32_t)tokens.size() >= n_tokens_alloc) { return false; } - tokens.push_back({ id_slot, LLAMA_TOKEN_NULL, pos, output }); + tokens.push_back({ id_slot, LLAMA_TOKEN_NULL, pos, output, is_prompt }); has_embd = true; embd.insert(embd.end(), embd_in.begin(), embd_in.end()); return true; @@ -222,20 +223,19 @@ struct server_slot { int64_t t_last_used = -1; // generation props - int32_t n_ctx = 0; // context size per slot - int32_t n_keep = 0; - int32_t n_decoded = 0; - int32_t n_remaining = -1; - int32_t i_batch = -1; + int32_t n_ctx = 0; // context size per slot + int32_t n_keep = 0; + int32_t i_batch = -1; - int32_t n_prompt_tokens_cache = 0; - int32_t n_prompt_tokens_processed = 0; + // effective generation limit for the current task, -1 means unlimited + int32_t n_predict_max = -1; size_t last_nl_pos = 0; std::string generated_text; std::string debug_generated_text; llama_tokens generated_tokens; + size_t n_sent_text = 0; // number of sent text character (i.e. handle partial UTF-8 on streaming) std::vector generated_token_probs; @@ -309,33 +309,24 @@ struct server_slot { // corresponding to one token position (size = n_embd) std::vector inp_embd; - // stats - size_t n_sent_text = 0; // number of sent text character + server_slot_stats stats; - // TODO @ngxson : move all metrics to a sub-struct for clarity - int64_t t_start_process_prompt; - int64_t t_start_generation; + // accepted tokens per draft position + // not in server_slot_stats to avoid copying to every task result + std::vector n_accepted_per_pos; + + std::function callback_on_release; + std::function callback_on_reset; // called before reset() + + // this is for printing timings with slot progress, not part of metrics int64_t t_print_last = 0; - int32_t n_decoded_last = 0; - - double t_prompt_processing = 0.0; // ms - double t_token_generation = 0.0; // ms - - std::function callback_on_release; - - // Speculative decoding stats - int32_t n_draft_total = 0; // Total draft tokens generated - int32_t n_draft_accepted = 0; // Draft tokens actually accepted - int32_t n_draft_verif_steps = 0; // Total draft token verification steps by the target model - std::vector n_accepted_per_pos; // Accepted tokens per draft position + int32_t n_gen_last = 0; void reset() { SLT_DBG(*this, "%s", "\n"); spec_is_replay = false; - n_prompt_tokens_cache = 0; - last_nl_pos = 0; generated_text = ""; has_new_line = false; @@ -353,15 +344,15 @@ struct server_slot { generated_token_probs.clear(); json_schema = json(); - // clear speculative decoding stats - n_draft_total = 0; - n_draft_accepted = 0; - n_draft_verif_steps = 0; - n_accepted_per_pos.clear(); - task_prev = std::move(task); task.reset(); + // note: callback_on_reset() must have run before this, see release() + stats = {}; + n_accepted_per_pos.clear(); + + n_predict_max = -1; + llama_set_sampler(ctx_tgt, id, nullptr); // clear alora start @@ -419,22 +410,13 @@ struct server_slot { && are_lora_equal(lora, other_slot.lora); } - bool has_budget(const common_params & global_params) { - GGML_ASSERT(task); + // returns -1 if the generation is limitless + int32_t n_remaining() const { + return n_predict_max == -1 ? -1 : n_predict_max - (int32_t) stats.n_gen; + } - if (task->params.n_predict == -1 && global_params.n_predict == -1) { - return true; // limitless - } - - n_remaining = -1; - - if (task->params.n_predict != -1) { - n_remaining = task->params.n_predict - n_decoded; - } else if (global_params.n_predict != -1) { - n_remaining = global_params.n_predict - n_decoded; - } - - return n_remaining > 0; // no budget + bool has_budget() const { + return n_predict_max == -1 || n_remaining() > 0; } bool is_processing() const { @@ -466,8 +448,8 @@ struct server_slot { // also, need to leave space for 1 extra token to allow context shifts int n_draft_max = n_ctx - prompt.n_tokens() - 2; - if (n_remaining > 0) { - n_draft_max = std::min(n_draft_max, n_remaining - 1); + if (n_remaining() > 0) { + n_draft_max = std::min(n_draft_max, n_remaining() - 1); } SLT_DBG(*this, "max possible draft: %d\n", n_draft_max); @@ -483,9 +465,9 @@ struct server_slot { i_batch = batch.size(); if (!inp_embd.empty()) { - add_ok &= batch.add(id, inp_embd, prompt.tokens.pos_next(), true); + add_ok &= batch.add(id, inp_embd, prompt.tokens.pos_next(), true, false); } else { - add_ok &= batch.add(id, sampled, prompt.tokens.pos_next(), true); + add_ok &= batch.add(id, sampled, prompt.tokens.pos_next(), true, false); } SLT_DBG(*this, "slot decode token, id=%d, n_ctx = %d, n_tokens = %d, truncated = %d\n", @@ -503,9 +485,9 @@ struct server_slot { auto pos0 = prompt.tokens.pos_next(); - add_ok &= batch.add(id, sampled, pos0++, true); + add_ok &= batch.add(id, sampled, pos0++, true, false); for (auto token : spec_draft) { - add_ok &= batch.add(this->id, token, pos0++, true); + add_ok &= batch.add(this->id, token, pos0++, true, false); } } @@ -521,8 +503,7 @@ struct server_slot { SLT_INF(*this, "stop processing: n_tokens = %d, truncated = %d\n", prompt.n_tokens(), truncated); - t_last_used = ggml_time_us(); - t_token_generation = (ggml_time_us() - t_start_generation) / 1e3; + t_last_used = ggml_time_us(); state = SLOT_STATE_IDLE; @@ -531,35 +512,14 @@ struct server_slot { prompt_clear(); } + callback_on_reset(*this); + reset(); callback_on_release(id); } } - result_timings get_timings() const { - result_timings timings; - timings.cache_n = n_prompt_tokens_cache; - - timings.prompt_n = n_prompt_tokens_processed; - timings.prompt_ms = t_prompt_processing; - timings.prompt_per_token_ms = t_prompt_processing / n_prompt_tokens_processed; - timings.prompt_per_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed; - - timings.predicted_n = n_decoded; - timings.predicted_ms = t_token_generation; - timings.predicted_per_token_ms = t_token_generation / n_decoded; - timings.predicted_per_second = 1e3 / t_token_generation * n_decoded; - - // Add speculative metrics - if (n_draft_total > 0) { - timings.draft_n = n_draft_total; - timings.draft_n_accepted = n_draft_accepted; - } - - return timings; - } - size_t find_stopping_strings(const std::string & text, const size_t last_token_size, bool is_full_stop) { GGML_ASSERT(task); @@ -592,7 +552,7 @@ struct server_slot { } void print_timings_tg() { - if (n_decoded < 100) { + if (stats.n_gen < 100) { return; } @@ -602,50 +562,59 @@ struct server_slot { return; } - const double n_gen_second = 1e3 / (t_token_generation) * (n_decoded); - const double n_gen_second_win = 1e6 / (t_now - t_print_last) * (n_decoded - n_decoded_last); + const double n_gen_second = stats.n_gen_tps(); + const double n_gen_second_win = 1e6 / (t_now - t_print_last) * (stats.n_gen - n_gen_last); t_print_last = t_now; - n_decoded_last = n_decoded; + n_gen_last = stats.n_gen; - SLT_INF(*this, "n_decoded = %6d, tg = %6.2f t/s, tg_3s = %6.2f t/s\n", n_decoded, n_gen_second, n_gen_second_win); + SLT_INF(*this, "n_gen = %6d, tg = %6.2f t/s, tg_3s = %6.2f t/s\n", (int) stats.n_gen, n_gen_second, n_gen_second_win); } void print_timings_pp() const { - const double n_prompt_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed; - const double f_progress = (float) prompt.n_tokens() / task->n_tokens(); + const double t_prompt_total = stats.t_prompt_ms(); - if (t_prompt_processing < 3000.0) { + if (t_prompt_total < 3000.0) { return; } + const double n_prompt_second = stats.n_prompt_tps(); + const double f_progress = task->n_tokens() > 0 ? (double) prompt.n_tokens() / task->n_tokens() : 0.0; + SLT_INF(*this, "prompt processing, n_tokens = %6d, progress = %.2f, t = %6.2f s / %.2f tokens per second\n", - n_prompt_tokens_processed, f_progress, t_prompt_processing / 1e3, n_prompt_second); + (int) stats.n_prompt_processed, f_progress, t_prompt_total / 1e3, n_prompt_second); } void print_timings() const { - const double t_prompt = t_prompt_processing / n_prompt_tokens_processed; - const double n_prompt_second = 1e3 / t_prompt_processing * n_prompt_tokens_processed; + const double t_prompt_total = stats.t_prompt_ms(); + const double t_gen_total = stats.t_gen_ms(); - const double t_gen = t_token_generation / n_decoded; - const double n_gen_second = 1e3 / t_token_generation * n_decoded; + const double t_prompt = stats.t_prompt_per_token_ms(); + const double n_prompt_second = stats.n_prompt_tps(); + + const double t_gen = stats.t_gen_per_token_ms(); + const double n_gen_second = stats.n_gen_tps(); SLT_INF(*this, "prompt eval time = %10.2f ms / %5d tokens (%8.2f ms per token, %8.2f tokens per second)\n", - t_prompt_processing, n_prompt_tokens_processed, t_prompt, n_prompt_second); + t_prompt_total, (int) stats.n_prompt_processed, t_prompt, n_prompt_second); SLT_INF(*this, " eval time = %10.2f ms / %5d tokens (%8.2f ms per token, %8.2f tokens per second)\n", - t_token_generation, n_decoded, t_gen, n_gen_second); + t_gen_total, (int) stats.n_gen, t_gen, n_gen_second); SLT_INF(*this, " total time = %10.2f ms / %5d tokens\n", - t_prompt_processing + t_token_generation, n_prompt_tokens_processed + n_decoded); + t_prompt_total + t_gen_total, (int) (stats.n_prompt_processed + stats.n_gen)); SLT_INF(*this, " graphs reused = %10d\n", llama_perf_context(ctx_tgt).n_reused); + const int32_t n_draft_total = stats.n_draft_tokens; + const int32_t n_draft_accepted = stats.n_draft_accepted; + const int32_t n_draft_verif_steps = stats.n_draft_verif_steps; + if (n_draft_total > 0) { const float draft_ratio = (float) n_draft_accepted / n_draft_total; const double mean_acc_len = n_draft_verif_steps > 0 ? 1.0 + (double) n_draft_accepted / (double) n_draft_verif_steps : 1.0; @@ -685,15 +654,15 @@ struct server_slot { if (ptask) { res["id_task"] = ptask->id; res["n_prompt_tokens"] = (int32_t) prompt.tokens.size(); - res["n_prompt_tokens_processed"] = n_prompt_tokens_processed; - res["n_prompt_tokens_cache"] = n_prompt_tokens_cache; + res["n_prompt_tokens_processed"] = stats.n_prompt_processed; + res["n_prompt_tokens_cache"] = stats.n_prompt_cached; res["params"] = ptask->params.to_json(only_metrics); res["next_token"] = { { {"has_next_token", has_next_token}, {"has_new_line", has_new_line}, - {"n_remain", n_remaining}, - {"n_decoded", n_decoded}, + {"n_remain", n_remaining()}, + {"n_decoded", stats.n_gen}, } }; @@ -712,14 +681,9 @@ struct server_slot { mem.seq_rm(other.id, -1, -1); mem.seq_cp(id, other.id, -1, -1); - other.n_decoded = n_decoded; - other.n_remaining = n_remaining; - other.i_batch = i_batch; + other.i_batch = i_batch; - other.t_start_process_prompt = t_start_process_prompt; - other.t_prompt_processing = t_prompt_processing; - other.n_prompt_tokens_cache = n_prompt_tokens_cache; - other.n_prompt_tokens_processed = n_prompt_tokens_processed; + other.stats = stats; other.prompt = prompt.clone(); other.init_sampler(); @@ -816,84 +780,6 @@ struct server_slot { -// -// server_metrics -// - -struct server_metrics { - int64_t t_start = 0; - - uint64_t n_prompt_tokens_processed_total = 0; - uint64_t t_prompt_processing_total = 0; - uint64_t n_tokens_predicted_total = 0; - uint64_t t_tokens_generation_total = 0; - - uint64_t n_tokens_max = 0; - - uint64_t n_prompt_tokens_processed = 0; - uint64_t t_prompt_processing = 0; - - uint64_t n_tokens_predicted = 0; - uint64_t t_tokens_generation = 0; - - uint64_t n_decode_total = 0; - uint64_t n_busy_slots_total = 0; - - uint64_t n_draft_tokens_total = 0; - uint64_t n_draft_accepted_total = 0; - uint64_t n_draft_verif_steps_total = 0; - std::vector n_accepted_per_pos_total; - - void init() { - t_start = ggml_time_us(); - } - - void on_prompt_eval(const server_slot & slot) { - n_prompt_tokens_processed_total += slot.n_prompt_tokens_processed; - n_prompt_tokens_processed += slot.n_prompt_tokens_processed; - t_prompt_processing += slot.t_prompt_processing; - t_prompt_processing_total += slot.t_prompt_processing; - - n_tokens_max = std::max(n_tokens_max, (uint64_t) slot.prompt.n_tokens()); - } - - void on_prediction(const server_slot & slot) { - n_tokens_predicted_total += slot.n_decoded; - n_tokens_predicted += slot.n_decoded; - t_tokens_generation += slot.t_token_generation; - t_tokens_generation_total += slot.t_token_generation; - - n_draft_tokens_total += slot.n_draft_total; - n_draft_accepted_total += slot.n_draft_accepted; - n_draft_verif_steps_total += slot.n_draft_verif_steps; - - if (n_accepted_per_pos_total.size() < slot.n_accepted_per_pos.size()) { - n_accepted_per_pos_total.resize(slot.n_accepted_per_pos.size(), 0); - } - for (size_t i = 0; i < slot.n_accepted_per_pos.size(); i++) { - n_accepted_per_pos_total[i] += slot.n_accepted_per_pos[i]; - } - } - - void on_decoded(const std::vector & slots) { - n_decode_total++; - for (const auto & slot : slots) { - if (slot.is_processing()) { - n_busy_slots_total++; - } - n_tokens_max = std::max(n_tokens_max, (uint64_t) slot.prompt.n_tokens()); - } - } - - void reset_bucket() { - n_prompt_tokens_processed = 0; - t_prompt_processing = 0; - n_tokens_predicted = 0; - t_tokens_generation = 0; - } -}; - - // // server_context_impl (private implementation) // @@ -972,6 +858,12 @@ private: server_metrics metrics; + // queued prompt stats - llama_decode() is async, so the timing is only valid after a sync + // note: kept out of server_metrics, which is copied as-is into the task result + int64_t t_decode_start = 0; // start of the last submitted decode + int64_t t_prompt_start = 0; // start of the oldest queued prompt decode + uint64_t n_prompt_queued = 0; + json json_ui_settings = json::object(); // Necessary similarity of prompt for slot selection @@ -1369,6 +1261,13 @@ private: queue_tasks.pop_deferred_task(id_slot); }; + slot.callback_on_reset = [this](const server_slot & slot) { + // flush the generated token stats before reset() + if (slot.stats.n_gen > 0) { + metrics_on_prediction(slot); + } + }; + slot.reset(); } @@ -1846,6 +1745,9 @@ private: slot.smpl.reset(); } + // the per-request limit takes priority over the global one + slot.n_predict_max = task.params.n_predict != -1 ? task.params.n_predict : params_base.n_predict; + slot.task = std::make_unique(std::move(task)); slot.state = slot.task->is_child() @@ -1917,16 +1819,16 @@ private: slot.stop = STOP_TYPE_LIMIT; slot.has_next_token = false; - SLT_DBG(slot, "stopped due to running out of context capacity, prompt.n_tokens() = %d, task.n_tokens = %d, n_decoded = %d, n_ctx = %d\n", - slot.prompt.n_tokens(), slot.task->n_tokens(), slot.n_decoded, slot.n_ctx); + SLT_DBG(slot, "stopped due to running out of context capacity, prompt.n_tokens() = %d, task.n_tokens = %d, n_gen = %d, n_ctx = %d\n", + slot.prompt.n_tokens(), slot.task->n_tokens(), (int) slot.stats.n_gen, slot.n_ctx); } // check the limits - if (slot.n_decoded > 0 && slot.has_next_token && !slot.has_budget(params_base)) { + if (slot.stats.n_gen > 0 && slot.has_next_token && !slot.has_budget()) { slot.stop = STOP_TYPE_LIMIT; slot.has_next_token = false; - SLT_DBG(slot, "stopped by limit, n_decoded = %d, n_predict = %d\n", slot.n_decoded, slot.task->params.n_predict); + SLT_DBG(slot, "stopped by limit, n_gen = %d, n_predict = %d\n", (int) slot.stats.n_gen, slot.task->params.n_predict); } if (slot.has_new_line) { @@ -1950,7 +1852,7 @@ private: // cut the last line slot.generated_text.erase(pos, std::string::npos); - SLT_DBG(slot, "stopped by indentation limit, n_decoded = %d, n_indent = %d\n", slot.n_decoded, n_indent); + SLT_DBG(slot, "stopped by indentation limit, n_gen = %d, n_indent = %d\n", (int) slot.stats.n_gen, n_indent); } } @@ -1970,11 +1872,11 @@ private: slot.has_new_line = true; // if we have seen a new line, we stop after a certain time limit, but only upon another new line - if (slot.task->params.t_max_predict_ms > 0 && (ggml_time_us() - slot.t_start_generation > 1000.0f*slot.task->params.t_max_predict_ms)) { + if (slot.task->params.t_max_predict_ms > 0 && slot.stats.t_gen_ms() > slot.task->params.t_max_predict_ms) { slot.stop = STOP_TYPE_LIMIT; slot.has_next_token = false; - SLT_DBG(slot, "stopped by time limit, n_decoded = %d, t_max_predict_ms = %d ms\n", slot.n_decoded, (int) slot.task->params.t_max_predict_ms); + SLT_DBG(slot, "stopped by time limit, n_gen = %d, t_max_predict_ms = %d ms\n", (int) slot.stats.n_gen, (int) slot.task->params.t_max_predict_ms); } } @@ -1985,7 +1887,7 @@ private: SLT_DBG(slot, "%s", "stopped by EOS\n"); } - SLT_DBG(slot, "n_decoded = %d, n_remaining = %d, next token: %5d '%s'\n", slot.n_decoded, slot.n_remaining, result.tok, token_str.c_str()); + SLT_DBG(slot, "n_gen = %d, n_remaining = %d, next token: %5d '%s'\n", (int) slot.stats.n_gen, slot.n_remaining(), result.tok, token_str.c_str()); return slot.has_next_token; // continue } @@ -2072,18 +1974,6 @@ private: queue_results.send(std::move(res)); } - // Gate slot save/restore/erase on slot content (does it hold media), - // not model capability: a multimodal model may hold a pure-text slot. - bool check_slot_no_media(const server_slot & slot, const int id_task) { - if (slot.prompt.tokens.has_media()) { - send_error(id_task, - "This operation is not supported while the slot holds image/audio tokens (a pure-text prefix is supported)", - ERROR_TYPE_NOT_SUPPORTED); - return false; - } - return true; - } - void send_partial_response(server_slot & slot, const completion_token_output & tkn, bool is_progress, bool is_begin = false) { auto res = std::make_unique(); @@ -2093,9 +1983,9 @@ private: if (is_progress) { res->is_progress = true; res->progress.total = slot.task->n_tokens(); - res->progress.cache = slot.n_prompt_tokens_cache; + res->progress.cache = slot.stats.n_prompt_cached; res->progress.processed = slot.prompt.tokens.size(); - res->progress.time_ms = (ggml_time_us() - slot.t_start_process_prompt) / 1000; + res->progress.time_ms = slot.stats.t_elapsed_us() / 1000; } if (is_begin) { res->is_begin = true; @@ -2104,9 +1994,9 @@ private: res->tokens = { tkn.tok }; } - res->n_decoded = slot.n_decoded; + res->n_decoded = slot.stats.n_gen; res->n_prompt_tokens = slot.task->n_tokens(); - res->n_prompt_tokens_cache = slot.n_prompt_tokens_cache; + res->n_prompt_tokens_cache = slot.stats.n_prompt_cached; res->post_sampling_probs = slot.task->params.post_sampling_probs; res->verbose = slot.task->params.verbose; @@ -2121,7 +2011,7 @@ private: // populate timings if this is final response or timings_per_token is enabled if (slot.stop != STOP_TYPE_NONE || slot.task->params.timings_per_token) { - res->timings = slot.get_timings(); + res->stats = slot.stats; } queue_results.send(std::move(res)); @@ -2148,14 +2038,14 @@ private: res->content = std::move(slot.generated_text); res->tokens = std::move(slot.generated_tokens); } - res->timings = slot.get_timings(); + res->stats = slot.stats; res->prompt = slot.task->tokens.detokenize(ctx_tgt, true); res->response_fields = std::move(slot.task->params.response_fields); res->truncated = slot.truncated; - res->n_decoded = slot.n_decoded; + res->n_decoded = slot.stats.n_gen; res->n_prompt_tokens = slot.task->n_tokens(); - res->n_prompt_tokens_cache = slot.n_prompt_tokens_cache; + res->n_prompt_tokens_cache = slot.stats.n_prompt_cached; res->n_tokens_cached = slot.prompt.n_tokens(); res->has_new_line = slot.has_new_line; res->stopping_word = slot.stopping_word; @@ -2542,27 +2432,7 @@ private: res->n_idle_slots = n_idle_slots; res->n_processing_slots = n_processing_slots; res->n_tasks_deferred = queue_tasks.queue_tasks_deferred_size(); - res->t_start = metrics.t_start; - - res->n_prompt_tokens_processed_total = metrics.n_prompt_tokens_processed_total; - res->t_prompt_processing_total = metrics.t_prompt_processing_total; - res->n_tokens_predicted_total = metrics.n_tokens_predicted_total; - res->t_tokens_generation_total = metrics.t_tokens_generation_total; - - res->n_tokens_max = metrics.n_tokens_max; - - res->n_prompt_tokens_processed = metrics.n_prompt_tokens_processed; - res->t_prompt_processing = metrics.t_prompt_processing; - res->n_tokens_predicted = metrics.n_tokens_predicted; - res->t_tokens_generation = metrics.t_tokens_generation; - - res->n_decode_total = metrics.n_decode_total; - res->n_busy_slots_total = metrics.n_busy_slots_total; - - res->n_draft_tokens_total = metrics.n_draft_tokens_total; - res->n_draft_accepted_total = metrics.n_draft_accepted_total; - res->n_draft_verif_steps_total = metrics.n_draft_verif_steps_total; - res->n_accepted_per_pos_total = metrics.n_accepted_per_pos_total; + res->metrics = metrics; if (task.metrics_reset_bucket) { metrics.reset_bucket(); @@ -2577,9 +2447,6 @@ private: send_error(task, "Invalid slot ID", ERROR_TYPE_INVALID_REQUEST); break; } - if (!check_slot_no_media(*slot, task.id)) { - break; - } if (slot->is_processing()) { // if requested slot is unavailable, we defer this task for processing later SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id); @@ -2592,9 +2459,22 @@ private: std::string filename = task.slot_action.filename; std::string filepath = task.slot_action.filepath; - const llama_tokens tokens = slot->prompt.tokens.get_text_tokens(); - const size_t token_count = tokens.size(); - const size_t nwrite = llama_state_seq_save_file(ctx_tgt, filepath.c_str(), slot->id, tokens.data(), token_count); + std::vector packed; + try { + packed = slot->prompt.tokens.serialize(); + } catch (const std::exception & err) { + send_error(task, err.what(), ERROR_TYPE_NOT_SUPPORTED); + break; + } + + GGML_ASSERT(packed.size() % sizeof(llama_token) == 0); + const size_t nwrite = llama_state_seq_save_file( + ctx_tgt, filepath.c_str(), slot->id, + reinterpret_cast(packed.data()), packed.size() / sizeof(llama_token)); + if (nwrite == 0) { + send_error(task, "Unable to save slot", ERROR_TYPE_SERVER); + break; + } const int64_t t_end = ggml_time_us(); const double t_save_ms = (t_end - t_start) / 1000.0; @@ -2604,7 +2484,7 @@ private: res->id_slot = id_slot; res->filename = filename; res->is_save = true; - res->n_tokens = token_count; + res->n_tokens = slot->prompt.tokens.size(); res->n_bytes = nwrite; res->t_ms = t_save_ms; queue_results.send(std::move(res)); @@ -2629,18 +2509,37 @@ private: std::string filename = task.slot_action.filename; std::string filepath = task.slot_action.filepath; - llama_tokens tokens; - tokens.resize(slot->n_ctx); - size_t token_count = 0; - size_t nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, tokens.data(), tokens.size(), &token_count); - if (nread == 0) { - slot->prompt.clear(); // KV may already been invalidated? - send_error(task, "Unable to restore slot, no available space in KV cache or invalid slot save file", ERROR_TYPE_INVALID_REQUEST); + size_t nread = 0; + try { + size_t n_packed = 0; + llama_tokens packed; + nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, nullptr, 0, &n_packed); + if (nread != 0) { + packed.resize(std::max(1, n_packed)); + nread = llama_state_seq_load_file(ctx_tgt, filepath.c_str(), slot->id, packed.data(), packed.size(), &n_packed); + } + if (nread == 0) { + throw std::runtime_error("No available space in KV cache or invalid slot save file"); + } + packed.resize(n_packed); + + server_tokens restored = server_tokens::deserialize(packed, mctx != nullptr); + + if (restored.size() > (size_t) slot->n_ctx) { + throw std::runtime_error("Restored prompt does not fit in the slot context"); + } + + if (!restored.validate(ctx_tgt)) { + throw std::runtime_error("Invalid tokens in slot save file"); + } + + slot->prompt.clear(); + slot->prompt.tokens = std::move(restored); + } catch (const std::exception & err) { + slot->prompt_clear(); + send_error(task, std::string("Unable to restore slot: ") + err.what(), ERROR_TYPE_INVALID_REQUEST); break; } - tokens.resize(token_count); - slot->prompt.clear(); - slot->prompt.tokens.insert(tokens); const int64_t t_end = ggml_time_us(); const double t_restore_ms = (t_end - t_start) / 1000.0; @@ -2650,7 +2549,7 @@ private: res->id_slot = id_slot; res->filename = filename; res->is_save = false; - res->n_tokens = token_count; + res->n_tokens = slot->prompt.tokens.size(); res->n_bytes = nread; res->t_ms = t_restore_ms; queue_results.send(std::move(res)); @@ -2663,10 +2562,6 @@ private: send_error(task, "Invalid slot ID", ERROR_TYPE_INVALID_REQUEST); break; } - // Gate on slot content, consistent with save/restore. - if (!check_slot_no_media(*slot, task.id)) { - break; - } if (slot->is_processing()) { // if requested slot is unavailable, we defer this task for processing later SRV_DBG("requested slot is unavailable, defer task, id_task = %d\n", task.id); @@ -2817,6 +2712,9 @@ private: if (all_idle) { SRV_TRC("%s", "all slots are idle\n"); + + metrics_flush_idle(); + return; // skip further processing } else { @@ -2835,6 +2733,9 @@ private: } catch (const std::exception & e) { SRV_ERR("pre_decode() failed: %s\n", e.what()); abort_all_slots("pre_decode() failed: " + std::string(e.what())); + + // the batch is half-built and not rendered, skip now to avoid UB + return; } GGML_ASSERT(batch.slot_batched || batch.size() == 0); @@ -3044,7 +2945,7 @@ private: auto & draft = slot.spec_draft; auto & ckpt = slot.spec_ckpt; - slot.n_draft_total += draft.size(); + slot.stats.n_draft_tokens += draft.size(); // TODO: avoid restoring the draft context and re-evaluating the drafted tokens when not needed [TAG_SPEC_AVOID_DRAFT_REEVAL] const bool use_ckpt_dft = ctx_dft_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_FULL; @@ -3132,8 +3033,7 @@ private: // TODO: maybe move branch to outside of this loop in the future if (slot.state == SLOT_STATE_STARTED) { - slot.t_start_process_prompt = ggml_time_us(); - slot.t_start_generation = 0; + slot.stats.update_prompt_start(); slot.state = SLOT_STATE_PROCESSING_PROMPT; @@ -3399,8 +3299,10 @@ private: SLT_WRN(slot, "n_past was set to %d\n", n_past); } - slot.n_prompt_tokens_cache = n_past; - slot.n_prompt_tokens_processed = 0; + slot.stats.n_prompt_cached = n_past; + slot.stats.n_prompt_processed = 0; + + metrics.add_prompt_cached(n_past); slot.prompt.tokens.keep_first(n_past); @@ -3423,8 +3325,8 @@ private: } } - const int64_t t_now = ggml_time_us(); - slot.t_prompt_processing = (t_now - slot.t_start_process_prompt) / 1e3; + // note: the prompt timing is advanced in post_decode(), so it does not cover + // the tokens added to the batch below slot.print_timings_pp(); // truncate any tokens that are beyond n_past for this slot @@ -3465,7 +3367,7 @@ private: bool has_mtmd = false; - // check if we should process the image + // check if we should process the mtmd chunk while (true) { auto cur_token_idx = slot.prompt.n_tokens(); if ( @@ -3475,19 +3377,25 @@ private: break; } - // process the image + // process the mtmd chunk + // note: it submits its own decode, potentially be async + // so the timing is queued and flushed on the next sync + metrics_pre_decode(); + size_t n_tokens_out = 0; int32_t res = slot.process_mtmd_chunk(cur_token_idx, n_tokens_out); if (res != 0) { - SLT_ERR(slot, "failed to process image, res = %d\n", res); - send_error(slot, "failed to process image", ERROR_TYPE_SERVER); + SLT_ERR(slot, "failed to process mtmd chunk, res = %d\n", res); + send_error(slot, "failed to process mtmd chunk", ERROR_TYPE_SERVER); slot.release(); - continue; + return; // the slot is done, skip it entirely } - slot.n_prompt_tokens_processed += n_tokens_out; + metrics_queue_prompt(n_tokens_out); + slot.stats.n_prompt_processed += n_tokens_out; + slot.stats.update_prompt_last(); - // add the image chunk to cache + // add the mtmd chunk to cache { const auto & chunk = input_tokens.find_chunk(cur_token_idx); slot.prompt.tokens.push_back(chunk.get()); // copy @@ -3520,12 +3428,11 @@ private: // streaming hook can mirror t_h_nextn into ctx_dft. add_ok &= batch.add(slot.id, cur_tok, - slot.prompt.tokens.pos_next(), - slot.need_embd()); + /* pos = */ slot.prompt.tokens.pos_next(), + /* output = */ slot.need_embd(), + /* is_prompt = */ true); slot.prompt.tokens.push_back(cur_tok); - slot.n_prompt_tokens_processed++; - // break at the last user message, or at user messages at least min step past the last checkpoint if (do_checkpoint && spans.is_user_start(slot.prompt.n_tokens())) { const auto pos = slot.prompt.n_tokens(); @@ -3577,8 +3484,8 @@ private: // extract the logits only for the last token batch.set_output(batch.size() - 1, true); - slot.n_decoded = 0; - slot.i_batch = batch.size() - 1; + slot.stats.n_gen = 0; + slot.i_batch = batch.size() - 1; slot.init_sampler(); } else { @@ -3627,6 +3534,8 @@ private: bool decode(int32_t & n_batch, int32_t off, llama_batch & batch_view) { SRV_DBG("n_batch (effective) = %d, off = %d\n", n_batch, off); + metrics_pre_decode(); + if (batch.size() == 0) { SRV_WRN("%s", "no tokens to decode\n"); @@ -3650,8 +3559,6 @@ private: const int ret = llama_decode(ctx_tgt, batch_view); - metrics.on_decoded(slots); - if (ret != 0) { { std::string err; @@ -3700,6 +3607,9 @@ private: SRV_WRN("failed to find free space in the KV cache, retrying with smaller batch size, off = %d, n_batch = %d, ret = %d\n", off, n_batch, ret); return false; // retry with the updated n_batch + } else { + // success, apply batch metrics + metrics_post_decode(off, batch_view.n_tokens); } // TODO: avoid restoring the draft context and re-evaluating the drafted tokens when not needed [TAG_SPEC_AVOID_DRAFT_REEVAL] @@ -3820,17 +3730,15 @@ private: // here we have synchronized the llama_context (due to the sampling above), so we can do time measurement const int64_t t_now = ggml_time_us(); - slot.n_decoded += 1; + slot.stats.n_gen += 1; - if (slot.n_decoded == 1) { - slot.t_start_generation = t_now; + if (slot.stats.n_gen == 1) { + slot.stats.update_prompt_last(); slot.t_print_last = t_now; - slot.n_decoded_last = 0; - slot.t_prompt_processing = (slot.t_start_generation - slot.t_start_process_prompt) / 1e3; - metrics.on_prompt_eval(slot); + slot.n_gen_last = 0; } - slot.t_token_generation = std::max(1, t_now - slot.t_start_generation) / 1e3; + slot.stats.update_gen_last(); completion_token_output result; result.tok = id; @@ -3845,7 +3753,6 @@ private: // release slot because of stop condition slot.print_timings(); send_final_response(slot); - metrics.on_prediction(slot); slot.release(); return; @@ -3921,8 +3828,6 @@ private: slot.spec_draft = std::move(accepted); } - const int64_t t_now = ggml_time_us(); - const auto ids = std::move(slot.spec_draft); size_t n_accepted = ids.size() - 1; @@ -3931,17 +3836,18 @@ private: } slot.spec_is_replay = false; - slot.t_token_generation = std::max(1, t_now - slot.t_start_generation) / 1e3; + slot.stats.update_gen_last(); // update how many tokens out of those tested were accepted - slot.n_draft_accepted += n_accepted; - slot.n_draft_verif_steps += 1; + slot.stats.n_draft_accepted += n_accepted; + slot.stats.n_draft_verif_steps += 1; - if (slot.n_accepted_per_pos.empty()) { - slot.n_accepted_per_pos.resize(common_speculative_n_max(¶ms_base.speculative), 0); + auto & n_accepted_per_pos = slot.n_accepted_per_pos; + if (n_accepted_per_pos.empty()) { + n_accepted_per_pos.resize(common_speculative_n_max(¶ms_base.speculative), 0); } - for (size_t i = 0; i < n_accepted && i < slot.n_accepted_per_pos.size(); ++i) { - slot.n_accepted_per_pos[i]++; + for (size_t i = 0; i < n_accepted && i < n_accepted_per_pos.size(); ++i) { + n_accepted_per_pos[i]++; } // add accepted tokens to the prompt @@ -3962,12 +3868,11 @@ private: // TODO: set result.probs - slot.n_decoded += 1; + slot.stats.n_gen += 1; if (!process_token(result, slot)) { slot.print_timings(); send_final_response(slot); - metrics.on_prediction(slot); slot.release(); return; @@ -3987,6 +3892,121 @@ private: server_response_reader get_response_reader() { return server_response_reader(queue_tasks, queue_results, HTTP_POLLING_SECONDS); } + + // + // metrics helpers + // + + // call before submitting a decode, so that the queued prompt stats can be timed + void metrics_pre_decode() { + t_decode_start = ggml_time_us(); + } + + // the batch is submitted, but its compute may not be done yet + void metrics_queue_prompt(uint64_t n_tokens) { + if (n_tokens == 0) { + return; + } + if (n_prompt_queued == 0) { + t_prompt_start = t_decode_start; + } + n_prompt_queued += n_tokens; + } + + // call only after the context is synchronized, otherwise the time is meaningless + void metrics_flush_prompt() { + if (n_prompt_queued == 0) { + return; + } + metrics.add_prompt(n_prompt_queued, ggml_time_us() - t_prompt_start); + n_prompt_queued = 0; + } + + void metrics_post_decode(int32_t off, int32_t n_tokens) { + metrics.n_decode++; + for (const auto & slot : slots) { + if (slot.is_processing()) { + metrics.n_busy_slots++; + } + metrics.n_tokens_max = std::max(metrics.n_tokens_max, (uint64_t) slot.prompt.n_tokens()); + } + + // apply enqueued prompt tokens stats + // note: a slot can be released before we get here, which clears its stats + // the tokens were still computed, counted in the global metrics, not in slot + uint64_t n_prompt_tokens = 0; + bool has_output = false; + + for (int i = off; i < off + n_tokens; ++i) { + const auto & t = batch.tokens[i]; + + has_output |= t.output; + + if (!t.is_prompt) { + continue; // generated tokens are handled after sampling + } + + n_prompt_tokens++; + + auto & slot = slots[t.id_slot]; + if (slot.stats.is_set()) { + slot.stats.n_prompt_processed++; + } + } + + metrics_queue_prompt(n_prompt_tokens); + + if (has_output) { + // sync if we have at least one output in batch + // so that we can calculate the timings correctly + llama_synchronize(ctx_tgt); + metrics_flush_prompt(); + } + + // advance the prompt timing of the slots that had tokens in this batch + // note: a second pass, it must run after the sync above to reflect the compute + const int64_t t_now = ggml_time_us(); + for (int i = off; i < off + n_tokens; ++i) { + const auto & t = batch.tokens[i]; + auto & slot = slots[t.id_slot]; + if (t.is_prompt && slot.stats.is_set()) { + slot.stats.set_prompt_last(t_now); + } + } + } + + // flush any queued prompt metrics if all slots are now idle + void metrics_flush_idle() { + if (n_prompt_queued == 0) { + return; + } + + llama_synchronize(ctx_tgt); + metrics_flush_prompt(); + } + + void metrics_on_prediction(const server_slot & slot) { + const uint64_t t_us = slot.stats.t_gen_us(); + const uint64_t n = slot.stats.n_gen; + const uint64_t n_steps = slot.stats.n_gen_steps(); + + metrics.predict .add(n, n_steps, t_us); + metrics.predict_bucket.add(n, n_steps, t_us); + + metrics.n_draft_tokens += slot.stats.n_draft_tokens; + metrics.n_draft_accepted += slot.stats.n_draft_accepted; + metrics.n_draft_verif_steps += slot.stats.n_draft_verif_steps; + + auto & dst = metrics.n_accepted_per_pos; + const auto & src = slot.n_accepted_per_pos; + + if (dst.size() < src.size()) { + dst.resize(src.size(), 0); + } + for (size_t i = 0; i < src.size(); i++) { + dst[i] += src[i]; + } + } }; // @@ -4406,6 +4426,8 @@ void server_routes::init_routes() { { server_task task(SERVER_TASK_TYPE_METRICS); task.id = res->rd.get_new_id(); + // the gauges are averaged over the window between two scrapes + task.metrics_reset_bucket = true; res->rd.post_task(std::move(task), true); // high-priority task } @@ -4422,104 +4444,13 @@ void server_routes::init_routes() { return res; } - // TODO: get rid of this dynamic_cast auto res_task = dynamic_cast(result.get()); GGML_ASSERT(res_task != nullptr); - // metrics definition: https://prometheus.io/docs/practices/naming/#metric-names - json all_metrics_def = json { - {"counter", {{ - {"name", "prompt_tokens_total"}, - {"help", "Number of prompt tokens processed."}, - {"value", (uint64_t) res_task->n_prompt_tokens_processed_total} - }, { - {"name", "prompt_seconds_total"}, - {"help", "Prompt process time"}, - {"value", (uint64_t) res_task->t_prompt_processing_total / 1.e3} - }, { - {"name", "tokens_predicted_total"}, - {"help", "Number of generation tokens processed."}, - {"value", (uint64_t) res_task->n_tokens_predicted_total} - }, { - {"name", "tokens_predicted_seconds_total"}, - {"help", "Predict process time"}, - {"value", (uint64_t) res_task->t_tokens_generation_total / 1.e3} - }, { - {"name", "n_decode_total"}, - {"help", "Total number of llama_decode() calls"}, - {"value", res_task->n_decode_total} - }, { - {"name", "n_tokens_max"}, - {"help", "Largest observed n_tokens."}, - {"value", res_task->n_tokens_max} - }, { - {"name", "spec_decode_num_draft_tokens_total"}, - {"help", "Total draft tokens generated"}, - {"value", res_task->n_draft_tokens_total} - }, { - {"name", "spec_decode_num_accepted_tokens_total"}, - {"help", "Total draft tokens accepted by the target model"}, - {"value", res_task->n_draft_accepted_total} - }, { - {"name", "spec_decode_num_drafts_total"}, - {"help", "Total speculative decoding verification steps"}, - {"value", res_task->n_draft_verif_steps_total} - }}}, - {"gauge", {{ - {"name", "prompt_tokens_seconds"}, - {"help", "Average prompt throughput in tokens/s."}, - {"value", res_task->n_prompt_tokens_processed ? 1.e3 / res_task->t_prompt_processing * res_task->n_prompt_tokens_processed : 0.} - },{ - {"name", "predicted_tokens_seconds"}, - {"help", "Average generation throughput in tokens/s."}, - {"value", res_task->n_tokens_predicted ? 1.e3 / res_task->t_tokens_generation * res_task->n_tokens_predicted : 0.} - },{ - {"name", "requests_processing"}, - {"help", "Number of requests processing."}, - {"value", (uint64_t) res_task->n_processing_slots} - },{ - {"name", "requests_deferred"}, - {"help", "Number of requests deferred."}, - {"value", (uint64_t) res_task->n_tasks_deferred} - },{ - {"name", "n_busy_slots_per_decode"}, - {"help", "Average number of busy slots per llama_decode() call"}, - {"value", (float) res_task->n_busy_slots_total / std::max((float) res_task->n_decode_total, 1.f)} - }}} - }; - - std::stringstream prometheus; - - for (const auto & el : all_metrics_def.items()) { - const auto & type = el.key(); - const auto & metrics_def = el.value(); - - for (const auto & metric_def : metrics_def) { - const std::string name = metric_def.at("name"); - const std::string help = metric_def.at("help"); - - auto value = json_value(metric_def, "value", 0.); - prometheus << "# HELP llamacpp:" << name << " " << help << "\n" - << "# TYPE llamacpp:" << name << " " << type << "\n" - << "llamacpp:" << name << " " << value << "\n"; - } - } - - // labeled counter: one time series per draft position - if (!res_task->n_accepted_per_pos_total.empty()) { - prometheus << "# HELP llamacpp:spec_decode_num_accepted_tokens_per_pos_total" - " Accepted tokens per draft position\n" - << "# TYPE llamacpp:spec_decode_num_accepted_tokens_per_pos_total counter\n"; - for (size_t i = 0; i < res_task->n_accepted_per_pos_total.size(); i++) { - prometheus << "llamacpp:spec_decode_num_accepted_tokens_per_pos_total{position=\"" - << i << "\"} " << res_task->n_accepted_per_pos_total[i] << "\n"; - } - } - - res->headers["Process-Start-Time-Unix"] = std::to_string(res_task->t_start); + res->headers["Process-Start-Time-Unix"] = std::to_string(res_task->metrics.t_start); res->content_type = "text/plain; version=0.0.4"; res->status = 200; - res->data = prometheus.str(); + res->data = res_task->to_metrics(); return res; }; @@ -4550,7 +4481,6 @@ void server_routes::init_routes() { return res; } - // TODO: get rid of this dynamic_cast auto * res_task = dynamic_cast(result.get()); GGML_ASSERT(res_task != nullptr); @@ -4562,7 +4492,7 @@ void server_routes::init_routes() { } } - res->ok(res_task->slots_data); + res->ok(res_task->to_json()); return res; }; diff --git a/tools/server/server-http.cpp b/tools/server/server-http.cpp index 783b01b82..b11dc09d0 100644 --- a/tools/server/server-http.cpp +++ b/tools/server/server-http.cpp @@ -355,8 +355,15 @@ bool server_http_context::init(const common_params & params) { return true; }; - auto serve_asset_cached = [](const std::string & name, bool isolation) { - return [name, isolation](const httplib::Request & req, httplib::Response & res) { + // Hashed assets never change under a given name, so they can be cached forever. + // `index.html` is the exception: its name is stable while its contents change on + // every build, and it is what names the hashed asset versions the UI loads. + static constexpr auto cache_immutable = "public, max-age=31536000, immutable"; + static constexpr auto cache_revalidate = "no-cache"; + + // Serves an asset with ETag/304 handling, under the given caching policy. + auto serve_asset_cached = [](const std::string & name, bool isolation, const char * cache_control) { + return [name, isolation, cache_control](const httplib::Request & req, httplib::Response & res) { if (!handle_gzip_header(req, res)) { return true; // returns error message } @@ -372,7 +379,7 @@ bool server_http_context::init(const common_params & params) { res.set_header("Cross-Origin-Embedder-Policy", "require-corp"); res.set_header("Cross-Origin-Opener-Policy", "same-origin"); } - res.set_header("Cache-Control", "public, max-age=31536000, immutable"); + res.set_header("Cache-Control", cache_control); res.set_content(reinterpret_cast(a->data), a->size, a->type.c_str()); return false; }; @@ -394,9 +401,9 @@ bool server_http_context::init(const common_params & params) { }; }; - // main index file - srv->Get(params.api_prefix + "/", serve_asset_cached("index.html", true)); - srv->Get(params.api_prefix + "/index.html", serve_asset_cached("index.html", true)); + // main index file -- revalidated, so a new build is picked up on the next load + srv->Get(params.api_prefix + "/", serve_asset_cached("index.html", true, cache_revalidate)); + srv->Get(params.api_prefix + "/index.html", serve_asset_cached("index.html", true, cache_revalidate)); // All remaining assets registered directly from the embedded asset table. // PWA revalidation files (sw.js, manifest, version.json) use no-cache; @@ -414,7 +421,7 @@ bool server_http_context::init(const common_params & params) { SRV_DBG("serve nocache for %s\n", a.name.c_str()); srv->Get(params.api_prefix + "/" + a.name, serve_asset_nocache(a.name)); } else { - srv->Get(params.api_prefix + "/" + a.name, serve_asset_cached(a.name, false)); + srv->Get(params.api_prefix + "/" + a.name, serve_asset_cached(a.name, false, cache_immutable)); } } diff --git a/tools/server/server-task.cpp b/tools/server/server-task.cpp index 1ee677553..64afbc5ed 100644 --- a/tools/server/server-task.cpp +++ b/tools/server/server-task.cpp @@ -10,6 +10,8 @@ #include "speculative.h" #include "server-common.h" +#include + using json = nlohmann::ordered_json; // @@ -236,34 +238,6 @@ common_chat_msg task_result_state::update_chat_msg( return chat_msg; } -// - -// result_timings -// - -json result_timings::to_json() const { - json base = { - {"cache_n", cache_n}, - - {"prompt_n", prompt_n}, - {"prompt_ms", prompt_ms}, - {"prompt_per_token_ms", prompt_per_token_ms}, - {"prompt_per_second", prompt_per_second}, - - {"predicted_n", predicted_n}, - {"predicted_ms", predicted_ms}, - {"predicted_per_token_ms", predicted_per_token_ms}, - {"predicted_per_second", predicted_per_second}, - }; - - if (draft_n > 0) { - base["draft_n"] = draft_n; - base["draft_n_accepted"] = draft_n_accepted; - } - - return base; -} - // // result_prompt_progress // @@ -382,7 +356,7 @@ json server_task_result_cmpl_final::to_json_non_oaicompat() { {"stop_type", stop_type_to_str(stop)}, {"stopping_word", stopping_word}, {"tokens_cached", n_tokens_cached}, - {"timings", timings.to_json()}, + {"timings", stats.to_json()}, }; if (!stream && !probs_output.empty()) { res["completion_probabilities"] = completion_token_output::probs_vector_to_json(probs_output, post_sampling_probs); @@ -432,8 +406,8 @@ json server_task_result_cmpl_final::to_json_oaicompat() { if (verbose) { res["__verbose"] = to_json_non_oaicompat(); } - if (timings.prompt_n >= 0) { - res.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + res.push_back({"timings", stats.to_json()}); } return res; @@ -480,8 +454,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat() { if (verbose) { res["__verbose"] = to_json_non_oaicompat(); } - if (timings.prompt_n >= 0) { - res.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + res.push_back({"timings", stats.to_json()}); } return res; @@ -541,8 +515,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_chat_stream() { }); } - if (timings.prompt_n >= 0) { - deltas.back().push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + deltas.back().push_back({"timings", stats.to_json()}); } // extra fields for debugging purposes @@ -734,8 +708,8 @@ json server_task_result_cmpl_final::to_json_oaicompat_resp_stream() { }} }); - if (timings.prompt_n >= 0) { - server_sent_events.back().at("data").push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + server_sent_events.back().at("data").push_back({"timings", stats.to_json()}); } return server_sent_events; @@ -1086,8 +1060,8 @@ json server_task_result_cmpl_partial::to_json_non_oaicompat() { {"tokens_evaluated", n_prompt_tokens}, }; // populate the timings object when needed (usually for the last response or with timings_per_token enabled) - if (timings.prompt_n > 0) { - res.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + res.push_back({"timings", stats.to_json()}); } if (is_progress) { res.push_back({"prompt_progress", progress.to_json()}); @@ -1126,8 +1100,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat() { if (verbose) { res["__verbose"] = to_json_non_oaicompat(); } - if (timings.prompt_n >= 0) { - res.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + res.push_back({"timings", stats.to_json()}); } if (is_progress) { res.push_back({"prompt_progress", progress.to_json()}); @@ -1180,8 +1154,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat_chat() { }; } - if (timings.prompt_n >= 0) { - last_json.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + last_json.push_back({"timings", stats.to_json()}); } if (is_progress) { last_json.push_back({"prompt_progress", progress.to_json()}); @@ -1330,8 +1304,8 @@ json server_task_result_cmpl_partial::to_json_oaicompat_resp() { if (!events.empty()) { json & data = events.back().at("data"); - if (timings.prompt_n >= 0) { - data.push_back({"timings", timings.to_json()}); + if (stats.is_set()) { + data.push_back({"timings", stats.to_json()}); } if (is_progress) { data.push_back({"prompt_progress", progress.to_json()}); @@ -1539,34 +1513,104 @@ json server_task_result_error::to_json() { // server_task_result_metrics // json server_task_result_metrics::to_json() { - return json { - { "idle", n_idle_slots }, - { "processing", n_processing_slots }, - { "deferred", n_tasks_deferred }, - { "t_start", t_start }, + return slots_data; +} - { "n_prompt_tokens_processed_total", n_prompt_tokens_processed_total }, - { "t_tokens_generation_total", t_tokens_generation_total }, - { "n_tokens_predicted_total", n_tokens_predicted_total }, - { "t_prompt_processing_total", t_prompt_processing_total }, - - { "n_tokens_max", n_tokens_max }, - - { "n_prompt_tokens_processed", n_prompt_tokens_processed }, - { "t_prompt_processing", t_prompt_processing }, - { "n_tokens_predicted", n_tokens_predicted }, - { "t_tokens_generation", t_tokens_generation }, - - { "n_decode_total", n_decode_total }, - { "n_busy_slots_total", n_busy_slots_total }, - - { "n_draft_tokens_total", n_draft_tokens_total }, - { "n_draft_accepted_total", n_draft_accepted_total }, - { "n_draft_verif_steps_total", n_draft_verif_steps_total }, - { "n_accepted_per_pos_total", n_accepted_per_pos_total }, - - { "slots", slots_data }, +// metrics definition: https://prometheus.io/docs/practices/naming/#metric-names +std::string server_task_result_metrics::to_metrics() { + const std::vector counters = { + { + "prompt_tokens_total", + "Number of prompt tokens processed, excluding cached tokens", + (double) metrics.prompt.count + }, { + "prompt_tokens_cached_total", + "Number of prompt tokens reused from the cache", + (double) metrics.n_prompt_cached + }, { + "prompt_seconds_total", + "Total time spent processing prompts", + metrics.prompt.time / 1.e6 + }, { + "tokens_predicted_total", + "Number of generation tokens processed", + (double) metrics.predict.count + }, { + "tokens_predicted_seconds_total", + "Total time spent generating tokens", + metrics.predict.time / 1.e6 + }, { + "n_decode_total", + "Total number of llama_decode() calls, excluding speculative decoding and multimodal decoding", + (double) metrics.n_decode + }, { + "n_tokens_max", + "Largest observed sequence length (prompt + generation)", + (double) metrics.n_tokens_max + }, { + "spec_decode_num_draft_tokens_total", + "Speculative: Total draft tokens generated", + (double) metrics.n_draft_tokens + }, { + "spec_decode_num_accepted_tokens_total", + "Speculative: Total draft tokens accepted by the target model", + (double) metrics.n_draft_accepted + }, { + "spec_decode_num_drafts_total", + "Speculative: Total speculative decoding verification steps", + (double) metrics.n_draft_verif_steps + }, }; + + const std::vector gauges = { + { + "prompt_tokens_seconds", + "Average prompt throughput in tokens/s", + metrics.prompt_bucket.n_per_second() + }, { + "predicted_tokens_seconds", + "Average generation throughput in tokens/s", + metrics.predict_bucket.n_per_second() + }, { + "requests_processing", + "Number of requests processing", + (double) n_processing_slots + }, { + "requests_deferred", + "Number of requests deferred", + (double) n_tasks_deferred + }, { + "n_busy_slots_per_decode", + "Average number of busy slots per llama_decode() call", + (double) metrics.n_busy_slots / std::max((double) metrics.n_decode, 1.0) + }, + }; + + std::stringstream prometheus; + + auto add_items = [&prometheus](const char * type, const std::vector & items) { + for (const auto & item : items) { + prometheus << "# HELP llamacpp:" << item.name << " " << item.description << "\n" + << "# TYPE llamacpp:" << item.name << " " << type << "\n" + << "llamacpp:" << item.name << " " << item.value << "\n"; + } + }; + + add_items("counter", counters); + add_items("gauge", gauges); + + // labeled counter: one time series per draft position + if (!metrics.n_accepted_per_pos.empty()) { + prometheus << "# HELP llamacpp:spec_decode_num_accepted_tokens_per_pos_total" + " Accepted tokens per draft position\n" + << "# TYPE llamacpp:spec_decode_num_accepted_tokens_per_pos_total counter\n"; + for (size_t i = 0; i < metrics.n_accepted_per_pos.size(); i++) { + prometheus << "llamacpp:spec_decode_num_accepted_tokens_per_pos_total{position=\"" + << i << "\"} " << metrics.n_accepted_per_pos[i] << "\n"; + } + } + + return prometheus.str(); } // diff --git a/tools/server/server-task.h b/tools/server/server-task.h index 6275ec760..b6da4d4bd 100644 --- a/tools/server/server-task.h +++ b/tools/server/server-task.h @@ -259,26 +259,6 @@ struct server_task { } }; -struct result_timings { - int32_t cache_n = -1; - - int32_t prompt_n = -1; - double prompt_ms = 0.0; - double prompt_per_token_ms = 0.0; - double prompt_per_second = 0.0; - - int32_t predicted_n = -1; - double predicted_ms = 0.0; - double predicted_per_token_ms = 0.0; - double predicted_per_second = 0.0; - - // Optional speculative metrics - only included when > 0 - int32_t draft_n = 0; - int32_t draft_n_accepted = 0; - - json to_json() const; -}; - struct result_prompt_progress { int32_t total = 0; int32_t cache = 0; @@ -343,7 +323,7 @@ struct server_task_result_cmpl_final : server_task_result { bool stream; bool include_usage; - result_timings timings; + server_slot_stats stats; std::string prompt; bool truncated; @@ -425,7 +405,7 @@ struct server_task_result_cmpl_partial : server_task_result { bool is_begin = false; // whether to send 200 status to HTTP client (begin of SSE stream) // ref: https://github.com/ggml-org/llama.cpp/pull/23884 completion_token_output prob_output; - result_timings timings; + server_slot_stats stats; result_prompt_progress progress; // response formatting @@ -510,38 +490,27 @@ struct server_task_result_error : server_task_result { }; struct server_task_result_metrics : server_task_result { + // these are immediate stats, not accumulated (server_metrics is cumulative) int n_idle_slots; int n_processing_slots; int n_tasks_deferred; - int64_t t_start; - // TODO: somehow reuse server_metrics in the future, instead of duplicating the fields - uint64_t n_prompt_tokens_processed_total = 0; - uint64_t t_prompt_processing_total = 0; - uint64_t n_tokens_predicted_total = 0; - uint64_t t_tokens_generation_total = 0; - - uint64_t n_tokens_max = 0; - - uint64_t n_prompt_tokens_processed = 0; - uint64_t t_prompt_processing = 0; - - uint64_t n_tokens_predicted = 0; - uint64_t t_tokens_generation = 0; - - uint64_t n_decode_total = 0; - uint64_t n_busy_slots_total = 0; - - uint64_t n_draft_tokens_total = 0; - uint64_t n_draft_accepted_total = 0; - uint64_t n_draft_verif_steps_total = 0; - std::vector n_accepted_per_pos_total; + server_metrics metrics; // while we can also use std::vector this requires copying the slot object which can be quite messy // therefore, we use json to temporarily store the slot.to_json() result json slots_data = json::array(); + // used by /slots API virtual json to_json() override; + + // used by /metrics API + struct metric_item { + std::string name; + std::string description; + double value; // prometheus values are always float64 + }; + std::string to_metrics(); }; struct server_task_result_slot_save_load : server_task_result { diff --git a/tools/server/tests/unit/test_metrics.py b/tools/server/tests/unit/test_metrics.py new file mode 100644 index 000000000..10cfc424b --- /dev/null +++ b/tools/server/tests/unit/test_metrics.py @@ -0,0 +1,227 @@ +import pytest +from utils import * + +server = ServerPreset.tinyllama2() + + +@pytest.fixture(autouse=True) +def create_server(): + global server + server = ServerPreset.tinyllama2() + server.server_metrics = True + + +def fetch_metrics(server: ServerProcess) -> str: + """get /metrics as raw prometheus text""" + res = server.make_request("GET", "/metrics") + assert res.status_code == 200 + assert "Process-Start-Time-Unix" in res.headers + assert isinstance(res.body, str) + return res.body + + +def parse_metrics(text: str) -> dict: + """parse the prometheus text format into {name: (type, value)}""" + out = {} + types = {} + for line in text.splitlines(): + if line.startswith("# TYPE "): + _, _, name, kind = line.split(" ", 3) + types[name] = kind + elif line.startswith("llamacpp:") and "{" not in line: + name, value = line.split(" ", 1) + assert name in types, f"{name} has no # TYPE line" + out[name] = (types[name], float(value)) + return out + + +def test_metrics_disabled(): + global server + server.server_metrics = False + server.start() + res = server.make_request("GET", "/metrics") + assert res.status_code == 501 # ERROR_TYPE_NOT_SUPPORTED + + +def test_metrics_prometheus_format(): + global server + server.start() + server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 8}) + + text = fetch_metrics(server) + metrics = parse_metrics(text) + + expected_counters = [ + "llamacpp:prompt_tokens_total", + "llamacpp:prompt_tokens_cached_total", + "llamacpp:prompt_seconds_total", + "llamacpp:tokens_predicted_total", + "llamacpp:tokens_predicted_seconds_total", + "llamacpp:n_decode_total", + "llamacpp:n_tokens_max", + "llamacpp:spec_decode_num_draft_tokens_total", + "llamacpp:spec_decode_num_accepted_tokens_total", + "llamacpp:spec_decode_num_drafts_total", + ] + expected_gauges = [ + "llamacpp:prompt_tokens_seconds", + "llamacpp:predicted_tokens_seconds", + "llamacpp:requests_processing", + "llamacpp:requests_deferred", + "llamacpp:n_busy_slots_per_decode", + ] + + for name in expected_counters: + assert metrics[name][0] == "counter" + for name in expected_gauges: + assert metrics[name][0] == "gauge" + + # every metric must carry a help line + for name in expected_counters + expected_gauges: + assert f"# HELP {name} " in text + + assert metrics["llamacpp:n_decode_total"][1] > 0 + assert metrics["llamacpp:requests_processing"][1] == 0 + + +def test_metrics_prompt_processed_and_cached(): + global server + server.n_slots = 1 # keep the prompt cache on a single slot + server.start() + + prompt = "the quick brown fox jumps over the lazy dog" + + n_processed = 0 + n_cached = 0 + for _ in range(2): + res = server.make_request("POST", "/completion", data={"prompt": prompt, "n_predict": 4}) + assert res.status_code == 200 + n_processed += res.body["timings"]["prompt_n"] + n_cached += res.body["timings"]["cache_n"] + + # the second request must reuse the prompt of the first one + assert n_cached > 0 + + metrics = parse_metrics(fetch_metrics(server)) + + # cached tokens are counted apart, they cost no decode + assert metrics["llamacpp:prompt_tokens_total"][1] == n_processed + assert metrics["llamacpp:prompt_tokens_cached_total"][1] == n_cached + + +def test_metrics_predicted_total_matches_requests(): + global server + server.start() + + n_predicted = 0 + for n_predict in [1, 4, 16]: + res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": n_predict}) + assert res.status_code == 200 + n_predicted += res.body["timings"]["predicted_n"] + + metrics = parse_metrics(fetch_metrics(server)) + assert metrics["llamacpp:tokens_predicted_total"][1] == n_predicted + + +def test_metrics_generation_rate_excludes_first_token(): + global server + server.start() + + # the first token comes from the logits of the last prompt batch, so it costs no decode step + res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 1}) + timings = res.body["timings"] + assert timings["predicted_n"] == 1 + assert timings["predicted_per_second"] == 0.0 + assert timings["predicted_per_token_ms"] == 0.0 + + res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 16}) + timings = res.body["timings"] + assert timings["predicted_n"] == 16 + # the rate is over 15 decode steps, not 16 tokens + expected = 1e3 / timings["predicted_ms"] * 15 + assert abs(timings["predicted_per_second"] - expected) < 1e-6 + + +@pytest.mark.parametrize("n_predict", [1, 8]) +def test_metrics_timings_are_finite(n_predict: int): + global server + server.start() + res = server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": n_predict}) + timings = res.body["timings"] + + # a null here means the server produced inf or nan + for key, value in timings.items(): + assert value is not None, f"{key} is null" + assert value >= 0, f"{key} is negative" + + assert timings["prompt_ms"] > 0 + assert timings["prompt_per_token_ms"] > 0 + + +def test_metrics_timings_on_prompt_progress(): + global server + server.start() + + # a long prompt so that it is split over several batches (n_batch = 32) + prompt = "the quick brown fox jumps over the lazy dog " * 8 + chunks = list(server.make_stream_request("POST", "/completion", data={ + "prompt": prompt, + "n_predict": 4, + "stream": True, + "timings_per_token": True, + "return_progress": True, + })) + + progress = [c for c in chunks if "prompt_progress" in c] + assert len(progress) > 1 # the prompt did not fit in a single batch + + # the very first update is sent before any prompt token is decoded + first = progress[0]["timings"] + assert first["prompt_n"] == 0 + assert first["prompt_ms"] == 0.0 + assert first["predicted_n"] == 0 + assert first["predicted_ms"] == 0.0 + + # timings must never go backwards, nor report bogus values + prompt_ms = 0.0 + for chunk in progress: + timings = chunk["timings"] + for key, value in timings.items(): + assert value is not None, f"{key} is null" + assert value >= 0, f"{key} is negative" + assert timings["prompt_ms"] >= prompt_ms + prompt_ms = timings["prompt_ms"] + + assert prompt_ms > 0 + + +def test_metrics_slots_idle_after_completion(): + global server + server.server_slots = True + server.start() + server.make_request("POST", "/completion", data={"prompt": "I believe", "n_predict": 8}) + + res = server.make_request("GET", "/slots") + assert res.status_code == 200 + for slot in res.body: + assert slot["is_processing"] is False + if "next_token" in slot: + # the budget of the finished task must not leak into the idle slot + assert slot["next_token"][0]["n_remain"] == -1 + assert slot["next_token"][0]["n_decoded"] == 0 + + +def test_metrics_embedding_prompt_is_counted(): + global server + server = ServerPreset.bert_bge_small() + server.server_metrics = True + server.start() + + res = server.make_request("POST", "/v1/embeddings", data={"input": ["hello world", "goodbye world"]}) + assert res.status_code == 200 + + # embedding tasks never sample a token, but their prompt still costs a decode + metrics = parse_metrics(fetch_metrics(server)) + assert metrics["llamacpp:prompt_tokens_total"][1] > 0 + assert metrics["llamacpp:n_decode_total"][1] > 0 + assert metrics["llamacpp:tokens_predicted_total"][1] == 0 diff --git a/tools/server/tests/unit/test_slot_save.py b/tools/server/tests/unit/test_slot_save.py index be22d9859..05acb1be1 100644 --- a/tools/server/tests/unit/test_slot_save.py +++ b/tools/server/tests/unit/test_slot_save.py @@ -2,6 +2,10 @@ import pytest from utils import * import base64 import requests +import struct + +# sequence state file: magic(4) version(4) payload_size(4), then payload_size llama_token words +STATE_FILE_HEADER_SIZE = 12 server = ServerPreset.tinyllama2() @@ -72,6 +76,60 @@ def test_slot_save_restore(): assert res.body["timings"]["prompt_n"] == 1 +def test_slot_restore_legacy_token_list(): + global server + server.start() + + res = server.make_request("POST", "/completion", data={ + "prompt": "What is the capital of France?", + "id_slot": 1, + "cache_prompt": True, + }) + assert res.status_code == 200 + + res = server.make_request("POST", "/slots/1?action=save", data={ + "filename": "slot_legacy.bin", + }) + assert res.status_code == 200 + assert res.body["n_saved"] == 84 + + # rewrite the token payload into a plain token list, as written by servers that predate the packed server_tokens format + path = os.path.join("tmp", "slot_legacy.bin") + with open(path, "rb") as f: + data = bytearray(f.read()) + + # the payload written by this server starts with a packed header: LLAMA_TOKEN_NULL(4) version(4) n_tokens(4) + packed_header_size = 12 + + payload_size = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE - 4)[0] + payload_end = STATE_FILE_HEADER_SIZE + payload_size * 4 + n_tokens = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE + 8)[0] + assert n_tokens == 84 + + tokens_start = STATE_FILE_HEADER_SIZE + packed_header_size + data = data[:STATE_FILE_HEADER_SIZE] + data[tokens_start:tokens_start + n_tokens * 4] + data[payload_end:] + struct.pack_into("=I", data, STATE_FILE_HEADER_SIZE - 4, n_tokens) + + with open(path, "wb") as f: + f.write(data) + + # the plain token list must restore, and the restored KV must be reusable + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "slot_legacy.bin", + }) + assert res.status_code == 200 + assert res.body["n_restored"] == 84 + + res = server.make_request("POST", "/completion", data={ + "prompt": "What is the capital of Germany?", + "id_slot": 0, + "cache_prompt": True, + }) + assert res.status_code == 200 + assert res.body["timings"]["prompt_n"] == 6 # only the different part is processed + + + def test_slot_erase(): global server server.start() @@ -103,14 +161,12 @@ def test_slot_erase(): # # Multimodal server (mmproj loaded) slot save/restore. # -# Regression coverage for issue #21133: slot save/restore/erase must be gated on -# the slot's CONTENT (does it actually hold image/audio tokens) rather than the -# model's CAPABILITY (is an mmproj loaded). A pure-text slot on a multimodal -# server must save/restore/erase normally; a slot that actually holds an image -# must be rejected with ERROR_TYPE_NOT_SUPPORTED (HTTP 501). +# A pure-text slot on a multimodal server and a slot containing images must both support save/restore. +# Erase remains gated on the slot's content. # IMG_URL_CAT = "https://huggingface.co/ggml-org/tinygemma3-GGUF/resolve/main/test/91_cat.png" +IMG_URL_TRUCK = "https://huggingface.co/ggml-org/tinygemma3-GGUF/resolve/main/test/11_truck.png" def _get_img_base64(url: str) -> str: @@ -121,8 +177,7 @@ def _get_img_base64(url: str) -> str: @pytest.fixture def mmproj_server(): - # tinygemma3 is a small multimodal model: the mmproj is provided by the HF - # registry API and auto-downloaded on first run. + # tinygemma3 is a small multimodal model: the mmproj is provided by the HF registry API and auto-downloaded on first run. os.environ['LLAMA_MEDIA_MARKER'] = '<__media__>' mm_server = ServerPreset.tinygemma3() mm_server.slot_save_path = "./tmp" @@ -159,10 +214,7 @@ def test_slot_save_restore_text_only_on_multimodal(mmproj_server): assert res.status_code == 200 assert res.body["n_restored"] == n_saved - # The restored slot is usable for a follow-up completion. We do NOT assert - # prefix reuse here: tinygemma3 is a SWA model, which forces full prompt - # re-processing after a restore (a model property, not the save/restore gate - # under test). + # Prefix reuse is not checked with the default SWA cache. res = server.make_request("POST", "/completion", data={ "prompt": "The quick brown fox jumps over the lazy dog.", "id_slot": 0, @@ -171,54 +223,307 @@ def test_slot_save_restore_text_only_on_multimodal(mmproj_server): assert res.status_code == 200 -def test_slot_save_rejected_when_slot_holds_image(mmproj_server): +def test_slot_save_restore_with_image(mmproj_server): server = mmproj_server + # Use the full SWA cache so the restored image prefix can be reused. + server.swa_full = True server.start() - # Process a prompt that actually contains an image on slot 1. + prompt_cat = { + "prompt_string": "What is this: <__media__>\n", + "multimodal_data": [_get_img_base64(IMG_URL_CAT)], + } res = server.make_request("POST", "/completions", data={ "temperature": 0.0, "top_k": 1, "id_slot": 1, "cache_prompt": True, + "prompt": prompt_cat, + }) + assert res.status_code == 200 + content_cat = res.body["content"] + prompt_n_full = res.body["timings"]["prompt_n"] + assert res.body["timings"]["cache_n"] == 0 + assert prompt_n_full > 32 # text plus image tokens are all processed + + res = server.make_request("POST", "/slots/1?action=save", data={ + "filename": "mm_slot_image.bin", + }) + assert res.status_code == 200 + n_saved = res.body["n_saved"] + n_written = res.body["n_written"] + assert n_saved > 0 + assert n_written > 0 + + res = server.make_request("POST", "/slots/1?action=erase") + assert res.status_code == 200 + + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_image.bin", + }) + assert res.status_code == 200 + assert res.body["n_restored"] == n_saved + assert res.body["n_read"] == n_written + + # a different image must not reuse the restored image tokens; only the text prefix before the image is common + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, "prompt": { "prompt_string": "What is this: <__media__>\n", - "multimodal_data": [ _get_img_base64(IMG_URL_CAT) ], + "multimodal_data": [_get_img_base64(IMG_URL_TRUCK)], + }, + }) + assert res.status_code == 200 + cache_n = res.body["timings"]["cache_n"] + assert cache_n < 16 + assert res.body["timings"]["prompt_n"] == prompt_n_full - cache_n + + # restore again and resend the same image: the image tokens must be reused and greedy sampling must reproduce the original content + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_image.bin", + }) + assert res.status_code == 200 + assert res.body["n_restored"] == n_saved + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt_cat, + }) + assert res.status_code == 200 + assert res.body["timings"]["cache_n"] == prompt_n_full - 1 + assert res.body["timings"]["prompt_n"] == 1 + assert res.body["content"] == content_cat + + +def test_slot_save_restore_with_two_images(mmproj_server): + server = mmproj_server + server.swa_full = True + server.n_ctx = 2048 # two images need more than the default 512 per slot + server.start() + + prompt = { + "prompt_string": "A: <__media__> B: <__media__>\n", + "multimodal_data": [_get_img_base64(IMG_URL_CAT), _get_img_base64(IMG_URL_TRUCK)], + } + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 1, + "cache_prompt": True, + "prompt": prompt, + }) + assert res.status_code == 200 + content = res.body["content"] + prompt_n_full = res.body["timings"]["prompt_n"] + assert prompt_n_full > 64 + + res = server.make_request("POST", "/slots/1?action=save", data={ + "filename": "mm_slot_two_images.bin", + }) + assert res.status_code == 200 + n_saved = res.body["n_saved"] + + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_two_images.bin", + }) + assert res.status_code == 200 + assert res.body["n_restored"] == n_saved + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt, + }) + assert res.status_code == 200 + assert res.body["timings"]["cache_n"] == prompt_n_full - 1 + assert res.body["timings"]["prompt_n"] == 1 + assert res.body["content"] == content + + +def test_slot_save_restore_with_image_across_restart(mmproj_server): + server = mmproj_server + server.swa_full = True + server.start() + + prompt_cat = { + "prompt_string": "What is this: <__media__>\n", + "multimodal_data": [_get_img_base64(IMG_URL_CAT)], + } + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt_cat, + }) + assert res.status_code == 200 + content = res.body["content"] + prompt_n_full = res.body["timings"]["prompt_n"] + + res = server.make_request("POST", "/slots/0?action=save", data={ + "filename": "mm_slot_restart.bin", + }) + assert res.status_code == 200 + n_saved = res.body["n_saved"] + + # restart the server with the same model and mmproj: the saved file must restore in the new process and the image KV must be reused + server.stop() + server.start() + + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_restart.bin", + }) + assert res.status_code == 200 + assert res.body["n_restored"] == n_saved + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt_cat, + }) + assert res.status_code == 200 + assert res.body["timings"]["cache_n"] == prompt_n_full - 1 + assert res.body["timings"]["prompt_n"] == 1 + assert res.body["content"] == content + + +def test_slot_save_restore_image_payload_larger_than_context(mmproj_server): + server = mmproj_server + server.swa_full = True + server.start() + + # the slot context, as the server computed it (n_ctx split across the slots) + res = server.make_request("GET", "/props") + assert res.status_code == 200 + n_ctx_slot = res.body["default_generation_settings"]["n_ctx"] + + # a filler token, used to grow the prompt up to the slot context + res = server.make_request("POST", "/tokenize", data={"content": " hello" * 8}) + assert res.status_code == 200 + assert len(res.body["tokens"]) == 8 + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": { + "prompt_string": "What is this: <__media__>\n", + "multimodal_data": [_get_img_base64(IMG_URL_CAT)], }, }) assert res.status_code == 200 - # Saving a slot that holds image tokens must be rejected (HTTP 501, - # not_supported_error). - res = server.make_request("POST", "/slots/1?action=save", data={ - "filename": "mm_slot_image.bin", + prompt_cat = { + "prompt_string": "What is this: <__media__>\n" + " hello" * (n_ctx_slot - res.body["timings"]["prompt_n"] - 8), + "multimodal_data": [_get_img_base64(IMG_URL_CAT)], + } + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt_cat, }) - assert res.status_code != 200 - assert res.body["error"]["type"] == "not_supported_error" + assert res.status_code == 200 + prompt_n_full = res.body["timings"]["cache_n"] + res.body["timings"]["prompt_n"] + + res = server.make_request("POST", "/slots/0?action=save", data={ + "filename": "mm_slot_large_payload.bin", + }) + assert res.status_code == 200 + + path = os.path.join("tmp", "mm_slot_large_payload.bin") + with open(path, "rb") as f: + data = bytearray(f.read()) + payload_size = struct.unpack_from("=I", data, STATE_FILE_HEADER_SIZE - 4)[0] + assert payload_size > n_ctx_slot # the scenario under test: the payload does not fit in n_ctx + + # drop the image from the slot, then restore it from the file + res = server.make_request("POST", "/completion", data={ + "prompt": "The quick brown fox", + "id_slot": 0, + "cache_prompt": True, + }) + assert res.status_code == 200 + + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_large_payload.bin", + }) + assert res.status_code == 200 + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": prompt_cat, + }) + assert res.status_code == 200 + assert res.body["timings"]["cache_n"] == prompt_n_full - 1 + assert res.body["timings"]["prompt_n"] == 1 -def test_slot_erase_text_only_on_multimodal(mmproj_server): +def test_slot_restore_media_file_without_mmproj(mmproj_server): server = mmproj_server server.start() - res = server.make_request("POST", "/completion", data={ - "prompt": "The quick brown fox jumps over the lazy dog.", - "id_slot": 1, + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, "cache_prompt": True, + "prompt": { + "prompt_string": "What is this: <__media__>\n", + "multimodal_data": [_get_img_base64(IMG_URL_CAT)], + }, }) assert res.status_code == 200 - prompt_n = res.body["timings"]["prompt_n"] - assert prompt_n > 0 # all tokens are processed - # Erasing a pure-text slot must succeed even though an mmproj is loaded. - res = server.make_request("POST", "/slots/1?action=erase") - assert res.status_code == 200 - - # Re-running the same prompt should process all tokens again. - res = server.make_request("POST", "/completion", data={ - "prompt": "The quick brown fox jumps over the lazy dog.", - "id_slot": 1, - "cache_prompt": True, + res = server.make_request("POST", "/slots/0?action=save", data={ + "filename": "mm_slot_no_mmproj.bin", }) assert res.status_code == 200 - assert res.body["timings"]["prompt_n"] == prompt_n # all tokens are processed again + + # restart the same model without the mmproj: restoring the media file must fail gracefully and leave the slot usable + server.stop() + server.no_mmproj = True + server.start() + + res = server.make_request("POST", "/slots/0?action=restore", data={ + "filename": "mm_slot_no_mmproj.bin", + }) + assert res.status_code == 400 + assert "Cannot restore media tokens without an mmproj" in res.body["error"]["message"] + + # A failed restore must leave the slot empty and usable. + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 1, + "cache_prompt": True, + "prompt": "The quick brown fox", + }) + assert res.status_code == 200 + content = res.body["content"] + + res = server.make_request("POST", "/completions", data={ + "temperature": 0.0, + "top_k": 1, + "id_slot": 0, + "cache_prompt": True, + "prompt": "The quick brown fox", + }) + assert res.status_code == 200 + assert res.body["timings"]["cache_n"] == 0 + assert res.body["content"] == content diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py index fffe07a67..9171dbc02 100644 --- a/tools/server/tests/utils.py +++ b/tools/server/tests/utils.py @@ -86,6 +86,7 @@ class ServerProcess: server_reranking: bool | None = False server_metrics: bool | None = False kv_unified: bool | None = False + swa_full: bool | None = False server_slots: bool | None = False pooling: str | None = None api_key: str | None = None @@ -106,6 +107,7 @@ class ServerProcess: chat_template_file: str | None = None server_path: str | None = None mmproj_url: str | None = None + no_mmproj: bool | None = None media_path: str | None = None sleep_idle_seconds: int | None = None cache_ram: int | None = None @@ -198,6 +200,8 @@ class ServerProcess: server_args.append("--metrics") if self.kv_unified: server_args.append("--kv-unified") + if self.swa_full: + server_args.append("--swa-full") if self.server_slots: server_args.append("--slots") else: @@ -259,6 +263,8 @@ class ServerProcess: server_args.extend(["--chat-template-file", self.chat_template_file]) if self.mmproj_url: server_args.extend(["--mmproj-url", self.mmproj_url]) + if self.no_mmproj: + server_args.append("--no-mmproj") if self.media_path: server_args.extend(["--media-path", self.media_path]) if self.sleep_idle_seconds is not None: diff --git a/tools/ui/pwa-assets-dark.config.ts b/tools/ui/pwa-assets-dark.config.ts index 1446b47a8..4d8114ee7 100644 --- a/tools/ui/pwa-assets-dark.config.ts +++ b/tools/ui/pwa-assets-dark.config.ts @@ -1,5 +1,5 @@ import { writeThemeFavicons } from './scripts/favicon-colorize'; -import { FAVICON_COLORS, PWA_ASSET_GENERATOR } from './src/lib/constants/pwa'; +import { FAVICON_COLORS, PWA_ASSET_GENERATOR } from './src/lib/constants/pwa.constants'; import { defineConfig } from '@vite-pwa/assets-generator/config'; writeThemeFavicons(FAVICON_COLORS.LIGHT, FAVICON_COLORS.DARK, { diff --git a/tools/ui/pwa-assets.config.ts b/tools/ui/pwa-assets.config.ts index 5fed0a595..f9f8662a2 100644 --- a/tools/ui/pwa-assets.config.ts +++ b/tools/ui/pwa-assets.config.ts @@ -4,7 +4,7 @@ import { PWA_ASSET_GENERATOR, PWA_GENERATOR_DEVICES, THEME_COLORS -} from './src/lib/constants/pwa'; +} from './src/lib/constants/pwa.constants'; import { SplashOrientation } from './src/lib/enums/splash.enums'; import { combinePresetAndAppleSplashScreens, diff --git a/tools/ui/scripts/vite-plugin-build-info.ts b/tools/ui/scripts/vite-plugin-build-info.ts index 800238630..ec864e8d0 100644 --- a/tools/ui/scripts/vite-plugin-build-info.ts +++ b/tools/ui/scripts/vite-plugin-build-info.ts @@ -1,4 +1,4 @@ -import { BUILD_CONFIG } from '../src/lib/constants/pwa'; +import { BUILD_CONFIG } from '../src/lib/constants/pwa.constants'; import { existsSync, writeFileSync } from 'node:fs'; import { resolve } from 'path'; import type { Plugin } from 'vite'; diff --git a/tools/ui/scripts/vite-plugin-relativize-base.ts b/tools/ui/scripts/vite-plugin-relativize-base.ts index f8eac1d66..0e47741ae 100644 --- a/tools/ui/scripts/vite-plugin-relativize-base.ts +++ b/tools/ui/scripts/vite-plugin-relativize-base.ts @@ -1,4 +1,4 @@ -import { BUILD_CONFIG } from '../src/lib/constants/pwa'; +import { BUILD_CONFIG } from '../src/lib/constants/pwa.constants'; import { existsSync, readFileSync, writeFileSync } from 'node:fs'; import { resolve } from 'path'; import type { Plugin } from 'vite'; diff --git a/tools/ui/scripts/vite-plugin-splash-screen.ts b/tools/ui/scripts/vite-plugin-splash-screen.ts index 45d931bab..62b7a063a 100644 --- a/tools/ui/scripts/vite-plugin-splash-screen.ts +++ b/tools/ui/scripts/vite-plugin-splash-screen.ts @@ -1,5 +1,10 @@ -import { NEWLINE, TAB } from '../src/lib/constants/code'; -import { APPLE_DEVICES, BUILD_CONFIG, REGEX_PATTERNS, SPLASH_LINK } from '../src/lib/constants/pwa'; +import { + APPLE_DEVICES, + BUILD_CONFIG, + REGEX_PATTERNS, + SPLASH_LINK +} from '../src/lib/constants/pwa.constants'; +import { NEWLINE, TAB } from '../src/lib/constants/special-characters.constants'; import { SplashOrientation } from '../src/lib/enums/splash.enums'; import type { SplashDimensions } from '../src/lib/types'; import { existsSync, readdirSync, readFileSync, writeFileSync } from 'node:fs'; diff --git a/tools/ui/src/app.d.ts b/tools/ui/src/app.d.ts index 6f26214e1..5309dce8f 100644 --- a/tools/ui/src/app.d.ts +++ b/tools/ui/src/app.d.ts @@ -32,8 +32,8 @@ import type { ApiRouterModelsStatusResponse, ApiRouterModelsUnloadRequest, ApiRouterModelsUnloadResponse, - // Chat types ChatAttachmentDisplayItem, + // Chat types ChatMessagePromptProgress, ChatMessageSiblingInfo, ChatMessageTimings, diff --git a/tools/ui/src/app.html b/tools/ui/src/app.html index e1de226dc..ef2787ad1 100644 --- a/tools/ui/src/app.html +++ b/tools/ui/src/app.html @@ -2,6 +2,7 @@ + diff --git a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte index 7655f18e8..2d54df89d 100644 --- a/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte +++ b/tools/ui/src/lib/components/app/actions/ActionIconCopyToClipboard.svelte @@ -1,7 +1,7 @@ - +
- {#if useContenteditable} - { - pickers.handleInput(); - onValueChange?.(value); - }} - onPaste={handlePaste} - {disabled} - {placeholder} - /> - {:else} - { - pickers.handleInput(); - onValueChange?.(value); - }} - onPaste={handlePaste} - {disabled} - {placeholder} - /> - {/if} + { + pickers.handleInput(); + onValueChange?.(value); + }} + onPaste={handlePaste} + {disabled} + {placeholder} + {useRichInput} + /> - {#if mcpHasResourceAttachments()} + {#if mcpResourceStore.hasAttachments} { @@ -668,7 +650,7 @@ {#if toolsStore.hasEnabledCwdTools} - void; - onSystemPromptClick?: () => void; - onMcpPromptClick?: () => void; - onMcpSettingsClick?: () => void; - onMcpResourcesClick?: () => void; } - let { - class: className = '', - disabled = false, - hasAudioModality = false, - hasMcpPromptsSupport = false, - hasMcpResourcesSupport = false, - hasVideoModality = false, - hasVisionModality = false, - onFileUpload, - onMcpPromptClick, - onMcpResourcesClick, - onMcpSettingsClick, - onSystemPromptClick - }: Props = $props(); + let { class: className = '' }: Props = $props(); + + const chatFormActions = getChatFormActionsContext(); let dropdownOpen = $state(false); // The system message action moves focus to the message editor, so the menu @@ -54,18 +33,23 @@ function handleMcpSettingsClick() { dropdownOpen = false; - onMcpSettingsClick?.(); + chatFormActions.onMcpSettingsClick?.(); } const attachmentMenu = useAttachmentMenu( () => ({ - hasAudioModality, - hasMcpPromptsSupport, - hasMcpResourcesSupport, - hasVideoModality, - hasVisionModality + hasAudioModality: chatFormActions.hasAudioModality, + hasMcpPromptsSupport: chatFormActions.hasMcpPromptsSupport, + hasMcpResourcesSupport: chatFormActions.hasMcpResourcesSupport, + hasVideoModality: chatFormActions.hasVideoModality, + hasVisionModality: chatFormActions.hasVisionModality + }), + () => ({ + onFileUpload: chatFormActions.onFileUpload, + onMcpPromptClick: chatFormActions.onMcpPromptClick, + onMcpResourcesClick: chatFormActions.onMcpResourcesClick, + onSystemPromptClick: chatFormActions.onSystemPromptClick }), - () => ({ onFileUpload, onMcpPromptClick, onMcpResourcesClick, onSystemPromptClick }), () => { dropdownOpen = false; } @@ -85,7 +69,7 @@ buttonVariants({ variant: 'secondary' }), 'file-upload-button h-8 w-8 cursor-pointer rounded-full p-0' )} - {disabled} + disabled={chatFormActions.disabled} > {ATTACHMENT_TOOLTIP_TEXT} @@ -162,7 +146,7 @@ class="flex cursor-pointer items-center gap-2" onclick={() => { suppressCloseAutoFocus = true; - onSystemPromptClick?.(); + chatFormActions.onSystemPromptClick?.(); }} > @@ -174,12 +158,12 @@ - {#if hasMcpPromptsSupport} + {#if chatFormActions.hasMcpPromptsSupport} @@ -187,10 +171,10 @@ {/if} - {#if hasMcpResourcesSupport} + {#if chatFormActions.hasMcpResourcesSupport} diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte index 852bf3fb6..3d04d14cb 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddMcpServersSubmenu.svelte @@ -4,11 +4,9 @@ import { DropdownMenuSearchable, McpLogo, McpServerIdentity } from '$lib/components/app'; import * as DropdownMenu from '$lib/components/ui/dropdown-menu'; import { Switch } from '$lib/components/ui/switch'; - import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes'; - import { ROUTES } from '$lib/constants/routes'; + import { ICON_CLASS_DEFAULT, ROUTES } from '$lib/constants'; import { HealthCheckStatus } from '$lib/enums'; - import { conversationsStore } from '$lib/stores/conversations.svelte'; - import { mcpStore } from '$lib/stores/mcp.svelte'; + import { conversationsStore, mcpStore } from '$lib/stores'; import type { MCPServerSettingsEntry } from '$lib/types'; interface Props { diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddReasoningSubmenu.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddReasoningSubmenu.svelte index 03b9f9f0a..a3a0b3a20 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddReasoningSubmenu.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddReasoningSubmenu.svelte @@ -2,7 +2,7 @@ import { Check, Info, Lightbulb, LightbulbOff } from '@lucide/svelte'; import * as DropdownMenu from '$lib/components/ui/dropdown-menu'; import * as Tooltip from '$lib/components/ui/tooltip'; - import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes'; + import { ICON_CLASS_DEFAULT } from '$lib/constants'; import { useReasoningMenu } from '$lib/hooks/use-reasoning-menu.svelte'; const reasoning = useReasoningMenu(); diff --git a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte index 73f616f69..63a8c267d 100644 --- a/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddSheet.svelte @@ -14,47 +14,28 @@ import * as Sheet from '$lib/components/ui/sheet'; import { Switch } from '$lib/components/ui/switch'; import * as Tooltip from '$lib/components/ui/tooltip'; - import { TOOLTIP_DELAY_DURATION } from '$lib/constants'; - import { ATTACHMENT_FILE_ITEMS } from '$lib/constants/attachment-menu'; - import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes'; + import { + ATTACHMENT_FILE_ITEMS, + ICON_CLASS_DEFAULT, + TOOLTIP_DELAY_DURATION + } from '$lib/constants'; + import { getChatFormActionsContext } from '$lib/contexts'; import { HealthCheckStatus } from '$lib/enums'; import { AttachmentAction } from '$lib/enums/attachment.enums'; import { useAttachmentMenu } from '$lib/hooks/use-attachment-menu.svelte'; import { useReasoningMenu } from '$lib/hooks/use-reasoning-menu.svelte'; import { useToolsPanel } from '$lib/hooks/use-tools-panel.svelte'; - import { conversationsStore } from '$lib/stores/conversations.svelte'; - import { mcpStore } from '$lib/stores/mcp.svelte'; + import { conversationsStore, mcpStore } from '$lib/stores'; import type { Snippet } from 'svelte'; interface Props { class?: string; - disabled?: boolean; - hasAudioModality?: boolean; - hasVideoModality?: boolean; - hasVisionModality?: boolean; - hasMcpPromptsSupport?: boolean; - hasMcpResourcesSupport?: boolean; - onFileUpload?: () => void; - onSystemPromptClick?: () => void; - onMcpPromptClick?: () => void; - onMcpResourcesClick?: () => void; trigger: Snippet<[{ disabled: boolean; onclick?: () => void }]>; } - let { - class: className = '', - disabled = false, - hasAudioModality = false, - hasMcpPromptsSupport = false, - hasMcpResourcesSupport = false, - hasVideoModality = false, - hasVisionModality = false, - onFileUpload, - onMcpPromptClick, - onMcpResourcesClick, - onSystemPromptClick, - trigger - }: Props = $props(); + let { class: className = '', trigger }: Props = $props(); + + const chatFormActions = getChatFormActionsContext(); let sheetOpen = $state(false); let reasoningExpanded = $state(false); @@ -64,13 +45,18 @@ const attachmentMenu = useAttachmentMenu( () => ({ - hasAudioModality, - hasMcpPromptsSupport, - hasMcpResourcesSupport, - hasVideoModality, - hasVisionModality + hasAudioModality: chatFormActions.hasAudioModality, + hasMcpPromptsSupport: chatFormActions.hasMcpPromptsSupport, + hasMcpResourcesSupport: chatFormActions.hasMcpResourcesSupport, + hasVideoModality: chatFormActions.hasVideoModality, + hasVisionModality: chatFormActions.hasVisionModality + }), + () => ({ + onFileUpload: chatFormActions.onFileUpload, + onMcpPromptClick: chatFormActions.onMcpPromptClick, + onMcpResourcesClick: chatFormActions.onMcpResourcesClick, + onSystemPromptClick: chatFormActions.onSystemPromptClick }), - () => ({ onFileUpload, onMcpPromptClick, onMcpResourcesClick, onSystemPromptClick }), () => { sheetOpen = false; } @@ -90,7 +76,7 @@
- {@render trigger({ disabled, onclick: () => (sheetOpen = true) })} + {@render trigger({ disabled: chatFormActions.disabled, onclick: () => (sheetOpen = true) })} @@ -349,7 +335,7 @@ System Message - {#if hasMcpPromptsSupport} + {#if chatFormActions.hasMcpPromptsSupport} {/if} {#if !isLoadingModel} - {serverError()} + {serverStore.error} {/if}
diff --git a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenStreamResumeStatus.svelte b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenStreamResumeStatus.svelte index 906118610..a8e0dcc19 100644 --- a/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenStreamResumeStatus.svelte +++ b/tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenStreamResumeStatus.svelte @@ -1,7 +1,7 @@ diff --git a/tools/ui/src/lib/components/app/chat/index.ts b/tools/ui/src/lib/components/app/chat/index.ts index 04b1d915f..2de8a6ace 100644 --- a/tools/ui/src/lib/components/app/chat/index.ts +++ b/tools/ui/src/lib/components/app/chat/index.ts @@ -91,7 +91,7 @@ export { default as ChatAttachmentsListItemThumbnailImage } from './ChatAttachme * preview without carousel, or a gallery/carousel view when multiple items exist. * Uses ChatAttachmentPreviewSingle internally for each item's content. */ -export { default as ChatAttachmentsPreview } from './ChatAttachments/ChatAttachmentsPreview.svelte'; +export { default as ChatAttachmentsPreview } from './ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreview.svelte'; export { default as ChatAttachmentsPreviewNavButtons } from './ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreviewNavButtons.svelte'; export { default as ChatAttachmentsPreviewFileInfo } from './ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreviewFileInfo.svelte'; export { default as ChatAttachmentsPreviewThumbnailStrip } from './ChatAttachments/ChatAttachmentsPreview/ChatAttachmentsPreviewThumbnailStrip.svelte'; @@ -120,8 +120,8 @@ export { default as ChatAttachmentsPreviewCurrentItem } from './ChatAttachments/ * Used by ChatScreenForm and ChatMessageEditForm for both new conversations and message editing. * * **Architecture:** - * - Composes ChatFormTextarea (or ChatFormContenteditable for messages with - * file mention links), ChatFormActions, and ChatFormPickerMcpPrompts + * - Composes ChatFormInput (a plain textarea, or a ChatFormInputRich for + * messages with file mention links), ChatFormActions, and ChatFormPickerMcpPrompts * - Manages file upload state via `uploadedFiles` bindable prop * - Integrates with ModelsSelectorDropdown for model selection in router mode * - Communicates with parent via callbacks (onSubmit, onFilesAdd, onStop, etc.) @@ -258,7 +258,7 @@ export { default as ChatFormContextGauge } from './ChatForm/ChatFormContextGauge /** * Hidden file input element for programmatic file selection. */ -export { default as ChatFormFileInputInvisible } from './ChatForm/ChatFormFileInputInvisible.svelte'; +export { default as ChatFormInputFileInputInvisible } from './ChatForm/ChatFormInput/ChatFormInputFileInputInvisible.svelte'; /** * Displays MCP Resource attachments as a horizontal carousel. @@ -267,18 +267,13 @@ export { default as ChatFormFileInputInvisible } from './ChatForm/ChatFormFileIn export { default as ChatFormMcpResourcesList } from './ChatForm/ChatFormMcpResourcesList.svelte'; /** - * Auto-resizing contenteditable input that renders `[name](file://...)` - * mention links as inline chips while keeping the value as the markdown - * source string. ChatForm swaps it in once a mention link lands in the - * buffer. Shares the focus()/resetHeight()/caret handle with the textarea. + * The message editor. Renders a plain auto-resizing textarea by default, + * or a ChatFormInputRich that renders `[name](file://...)` mention links as + * inline chips (keeping the value as the markdown source string) once a + * mention link lands in the buffer. The variant is selected via the + * `useRichInput` prop; both share one imperative handle. */ -export { default as ChatFormContenteditable } from './ChatForm/ChatFormContenteditable.svelte'; - -/** - * Plain auto-resizing textarea with IME composition support. Default input - * renderer inside ChatForm until a file mention lands. - */ -export { default as ChatFormTextarea } from './ChatForm/ChatFormTextarea.svelte'; +export { default as ChatFormInput } from './ChatForm/ChatFormInput/ChatFormInput.svelte'; /** * Working directory selector for agent mode. Renders a chip below the chat @@ -288,7 +283,7 @@ export { default as ChatFormTextarea } from './ChatForm/ChatFormTextarea.svelte' * synthetic "Set working directory to ..." user message into chat history * and is enforced on tool calls via the `x-tool-cwd` request header. */ -export { default as ChatFormWorkingDirectory } from './ChatForm/ChatFormWorkingDirectory.svelte'; +export { default as ChatFormCurrentWorkingDirectory } from './ChatForm/ChatFormCurrentWorkingDirectory/ChatFormCurrentWorkingDirectory.svelte'; /** * **ChatFormPickerMcpPrompts** - MCP prompt selection interface @@ -359,14 +354,14 @@ export { default as ChatFormPickerPopover } from './ChatForm/ChatFormPickers/Cha * Generic scrollable list for picker popovers. Provides search input, * scroll-into-view for keyboard navigation, loading skeletons, empty state, * and optional footer. Uses Svelte 5 snippets for item/skeleton/footer rendering. - * Shared by ChatFormPickerMcpPrompts and ChatFormMentionPicker. + * Shared by ChatFormPickerMcpPrompts and ChatFormPickerMention. */ export { default as ChatFormPickerList } from './ChatForm/ChatFormPickers/ChatFormPicker/ChatFormPickerList.svelte'; /** * Generic button wrapper for picker list items. Provides consistent styling, * hover/selected states, and data-picker-index attribute for scroll-into-view. - * Shared by ChatFormPickerMcpPrompts and ChatFormMentionPicker. + * Shared by ChatFormPickerMcpPrompts and ChatFormPickerMention. */ export { default as ChatFormPickerListItem } from './ChatForm/ChatFormPickers/ChatFormPicker/ChatFormPickerListItem.svelte'; @@ -389,14 +384,14 @@ export { default as ChatFormPickerListItemSkeleton } from './ChatForm/ChatFormPi * tool, scoped to the conversation cwd (or server home when unset). * Selection splices a `[name](file:///)` link into the input. */ -export { default as ChatFormMentionPicker } from './ChatForm/ChatFormPickers/ChatFormMentionPicker.svelte'; +export { default as ChatFormPickerMention } from './ChatForm/ChatFormPickers/ChatFormPickerMention.svelte'; /** * `/`-triggered slash-command picker. Lists the available slash commands * (`/prompt`, `/cwd`, `/model`) filtered by the typed query; selection * hands the command to the parent for dispatch. */ -export { default as ChatFormCommandPicker } from './ChatForm/ChatFormPickers/ChatFormCommandPicker.svelte'; +export { default as ChatFormPickerCommand } from './ChatForm/ChatFormPickers/ChatFormPickerCommand.svelte'; /** * Hosts the chat-form pickers (slash-command, MCP prompt, file mention) diff --git a/tools/ui/src/lib/components/app/content/CollapsibleContentBlock.svelte b/tools/ui/src/lib/components/app/content/CollapsibleContentBlock.svelte index ec57985ac..042a83e0e 100644 --- a/tools/ui/src/lib/components/app/content/CollapsibleContentBlock.svelte +++ b/tools/ui/src/lib/components/app/content/CollapsibleContentBlock.svelte @@ -2,7 +2,7 @@ import ChevronDown from '@lucide/svelte/icons/chevron-down'; import * as Collapsible from '$lib/components/ui/collapsible/index.js'; import { cn } from '$lib/components/ui/utils'; - import { ICON_CLASS_DEFAULT } from '$lib/constants/css-classes'; + import { ICON_CLASS_DEFAULT } from '$lib/constants'; import type { Snippet } from 'svelte'; import type { Component } from 'svelte'; diff --git a/tools/ui/src/lib/components/app/content/MarkdownContent/MarkdownContent.svelte b/tools/ui/src/lib/components/app/content/MarkdownContent/MarkdownContent.svelte index c768a0447..2a1d617ba 100644 --- a/tools/ui/src/lib/components/app/content/MarkdownContent/MarkdownContent.svelte +++ b/tools/ui/src/lib/components/app/content/MarkdownContent/MarkdownContent.svelte @@ -1,5 +1,5 @@ -