diff --git a/common/arg.cpp b/common/arg.cpp index be561114c..7e21ed4bc 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -1336,12 +1336,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } ).set_env("LLAMA_ARG_CTX_CHECKPOINTS").set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI})); add_opt(common_arg( - {"-cpent", "--checkpoint-every-n-tokens"}, "N", - string_format("create a checkpoint every n tokens during prefill (processing), -1 to disable (default: %d)", params.checkpoint_every_nt), + {"-cms", "--checkpoint-min-step"}, "N", + string_format("minimum spacing between context checkpoints in tokens (default: %d, 0 = no minimum)", params.checkpoint_min_step), [](common_params & params, int value) { - params.checkpoint_every_nt = value; + if (value < 0) { + throw std::invalid_argument("checkpoint-min-step must be non-negative"); + } + params.checkpoint_min_step = value; } - ).set_env("LLAMA_ARG_CHECKPOINT_EVERY_NT").set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI})); + ).set_env("LLAMA_ARG_CHECKPOINT_MIN_SPACING_NT").set_examples({LLAMA_EXAMPLE_SERVER})); add_opt(common_arg( {"-cram", "--cache-ram"}, "N", string_format("set the maximum cache size in MiB (default: %d, -1 - no limit, 0 - disable)" diff --git a/common/chat-auto-parser-helpers.cpp b/common/chat-auto-parser-helpers.cpp index 2499464cd..81b17e5e1 100644 --- a/common/chat-auto-parser-helpers.cpp +++ b/common/chat-auto-parser-helpers.cpp @@ -310,6 +310,8 @@ std::vector prune_whitespace_segments(const std::vector & segm namespace autoparser { +static const std::string ERR_TMPL = "#**ERROR**#"; + std::string apply_template(const common_chat_template & tmpl, const template_params & params) { generation_params tmpl_params; tmpl_params.messages = params.messages; @@ -326,7 +328,7 @@ std::string apply_template(const common_chat_template & tmpl, const template_par return common_chat_template_direct_apply(tmpl, tmpl_params); } catch (const std::exception & e) { LOG_DBG("Template application failed: %s\n", e.what()); - return ""; + return ERR_TMPL; } } @@ -347,7 +349,7 @@ std::optional compare_variants( std::string output_B = apply_template(tmpl, params_B); // Check for template application failures - if (output_A.empty() || output_B.empty()) { + if (output_A == ERR_TMPL || output_B == ERR_TMPL) { return std::nullopt; } diff --git a/common/chat-auto-parser.h b/common/chat-auto-parser.h index c680e6868..7858f6572 100644 --- a/common/chat-auto-parser.h +++ b/common/chat-auto-parser.h @@ -377,6 +377,8 @@ struct analyze_tools : analyze_base { struct autoparser { jinja::caps jinja_caps; + std::string user_start; + std::string assistant_start; analyze_reasoning reasoning; analyze_content content; analyze_tools tools; @@ -387,6 +389,10 @@ struct autoparser { autoparser() = default; + // Find the starting marker for the user message and assistant message + std::string detect_user_start_marker(const common_chat_template & tmpl); + std::string detect_assistant_start_marker(const common_chat_template & tmpl); + // Run full differential analysis on a template void analyze_template(const common_chat_template & tmpl); diff --git a/common/chat-diff-analyzer.cpp b/common/chat-diff-analyzer.cpp index 9c7c9678a..0875c5347 100644 --- a/common/chat-diff-analyzer.cpp +++ b/common/chat-diff-analyzer.cpp @@ -8,6 +8,9 @@ #include "peg-parser.h" #include +#include +#include +#include #define ANSI_RESET "\033[0m" #define ANSI_PURPLE "\033[1m\x1b[38;5;126m" @@ -23,6 +26,7 @@ static const std::string FUN_SECOND = "SSS_SECOND_FUN_S"; static const std::string ARG_FIRST = "AA_ARG_FST_AA"; static const std::string ARG_SECOND = "BB_ARG_SND_BB"; static const std::string USER_MSG = "U_USER_MSG Hello END_U"; +static const std::string USER_MSG_TWO = "V_USER_MSG Hello END_V"; static const std::string ASSISTANT_MSG = "A_ASST_MSG I can help END_A"; static const std::string THINKING_CONTENT = "REASON_PART I am thinking END_R"; static const std::string CALL_ID_001 = "call00001"; @@ -71,6 +75,7 @@ static std::vector"); analysis.preserved_tokens.push_back("<|END_OF_TURN_TOKEN|>"); + analysis.user_start = "<|START_OF_TURN_TOKEN|><|USER_TOKEN|>"; LOG_DBG(ANSI_ORANGE "[Patch: Cohere Command R+]\n" ANSI_RESET); } }, @@ -108,7 +113,59 @@ static std::vector void { + if (tmpl.src.find("") != std::string::npos && tmpl.src.find("") != std::string::npos && + tmpl.src.find("") != std::string::npos && tmpl.src.find("") != std::string::npos) { + + analysis.tools.format.mode = tool_format::JSON_NATIVE; + analysis.tools.format.section_start = ""; + analysis.tools.format.section_end = ""; + analysis.tools.format.per_call_start = ""; + analysis.tools.format.per_call_end = ""; + analysis.content.mode = content_mode::PLAIN; + analysis.content.start = ""; + analysis.content.end = ""; + analysis.reasoning.mode = reasoning_mode::TAG_BASED; + analysis.reasoning.start = "\n\n"; + analysis.reasoning.end = ""; + analysis.assistant_start = "Assistant"; + analysis.user_start = "User"; + analysis.preserved_tokens.clear(); + analysis.preserved_tokens.push_back(""); + analysis.preserved_tokens.push_back(""); + analysis.preserved_tokens.push_back(""); + analysis.preserved_tokens.push_back(""); + analysis.preserved_tokens.push_back(""); + LOG_DBG(ANSI_ORANGE "[Patch: Nemotron Nano v2]\n" ANSI_RESET); + } + }, + // Fireworks + [](const common_chat_template & tmpl, autoparser & analysis) -> void { + if (tmpl.src.find("{%- set system_prompt = '<|start_header_id|>' + 'system' + '<|end_header_id|>\\n\\n'" + " + message['content'] | trim + '\\n' + system_prompt_suffix + '<|eot_id|>' -%}") != std::string::npos) { + analysis.assistant_start = "<|start_header_id|>assistant<|end_header_id|>"; + analysis.user_start = "<|start_header_id|>user<|end_header_id|>"; + LOG_DBG(ANSI_ORANGE "[Patch: Fireworks v2]\n" ANSI_RESET); + } + }, + // Solar Open + [](const common_chat_template & tmpl, autoparser & analysis) -> void { + if (tmpl.src.find("<|begin|>assistant<|think|><|end|>") != std::string::npos) { + analysis.assistant_start = "<|begin|>assistant"; + LOG_DBG(ANSI_ORANGE "[Patch: Solar Open]\n" ANSI_RESET); + } + }, + // Apriel 1.6 + [](const common_chat_template & tmpl, autoparser & analysis) -> void { + if (tmpl.src.find("if not loop.last and '[BEGIN FINAL RESPONSE]' in asst_text") != std::string::npos) { + analysis.user_start = "<|begin_user|>"; + analysis.assistant_start = "<|begin_assistant|>"; + LOG_DBG(ANSI_ORANGE "[Patch: Apriel 1.6]\n" ANSI_RESET); + } + }, + }); // Common JSON structures @@ -166,6 +223,8 @@ void autoparser::analyze_template(const common_chat_template & tmpl) { reasoning = analyze_reasoning(tmpl, jinja_caps.supports_tool_calls); content = analyze_content(tmpl, reasoning); tools = analyze_tools(jinja_caps.supports_tool_calls ? analyze_tools(tmpl, jinja_caps, reasoning) : analyze_tools()); + assistant_start = detect_assistant_start_marker(tmpl); + user_start = detect_user_start_marker(tmpl); collect_preserved_tokens(); for (auto & workaround : workarounds) { @@ -173,6 +232,8 @@ void autoparser::analyze_template(const common_chat_template & tmpl) { } LOG_DBG("\n--- Reasoning & Content Structure ---\n"); + LOG_DBG("user_msg_start: %s\n", user_start.c_str()); + LOG_DBG("assistant_msg_start: %s\n", assistant_start.c_str()); LOG_DBG("reasoning_mode: %s\n", mode_to_str(reasoning.mode).c_str()); LOG_DBG("reasoning_start: '%s'\n", reasoning.start.c_str()); LOG_DBG("reasoning_end: '%s'\n", reasoning.end.c_str()); @@ -245,6 +306,120 @@ void autoparser::collect_preserved_tokens() { add_token(tools.call_id.suffix); } +std::string autoparser::detect_assistant_start_marker(const common_chat_template & tmpl) { + json user_msg = json{ + { "role", "user" }, + { "content", USER_MSG } + }; + + json assistant_no_reasoning = json{ + { "role", "assistant" }, + { "content", ASSISTANT_MSG } + }; + + template_params params; + params.messages = json::array({ user_msg }); + params.add_generation_prompt = false; + params.enable_thinking = true; + + auto comparison = compare_variants( + tmpl, params, [&](template_params & p) { + p.messages = json::array({ user_msg, assistant_no_reasoning }); + } + ); + + if (!comparison) { + LOG_DBG(ANSI_ORANGE "%s: Template application failed, skipping assistant start detection\n" ANSI_RESET, __func__); + return ""; + } + + auto usermsg = comparison->diff.right; + if (usermsg.find(ASSISTANT_MSG) == std::string::npos) { + LOG_DBG(ANSI_ORANGE "%s: Did not find assistant message in assistant message block, skipping detection\n" ANSI_RESET, __func__); + } + + auto ast_prefix = usermsg.substr(0, usermsg.find(ASSISTANT_MSG)); + if (!reasoning.start.empty() && ast_prefix.find(trim_whitespace(reasoning.start)) != std::string::npos) { + ast_prefix = ast_prefix.substr(0, ast_prefix.find(trim_whitespace(reasoning.start))); + } + if (!reasoning.end.empty() && ast_prefix.find(trim_whitespace(reasoning.end)) != std::string::npos) { + ast_prefix = ast_prefix.substr(0, ast_prefix.find(trim_whitespace(reasoning.end))); + } + return trim_whitespace(ast_prefix); +} + +std::string autoparser::detect_user_start_marker(const common_chat_template & tmpl) { + json user_msg = json{ + { "role", "user" }, + { "content", USER_MSG } + }; + + json assistant = json{ + { "role", "assistant" }, + { "content", ASSISTANT_MSG } + }; + + json user_msg_two = json{ + { "role", "user" }, + { "content", USER_MSG_TWO } + }; + + template_params params; + params.messages = json::array({}); + params.add_generation_prompt = false; + params.enable_thinking = true; + + auto comparison = compare_variants( + tmpl, params, [&](template_params & p) { + p.messages = json::array({ user_msg }); + } + ); + + if (!comparison) { + LOG_DBG(ANSI_ORANGE "%s: Template application failed, unsupported empty messages? trying complex variant\n" ANSI_RESET, __func__); + params.messages = json::array({ user_msg_two, assistant }); + comparison = compare_variants( + tmpl, params, [&](template_params & p) { + p.messages = json::array({ user_msg_two, assistant, user_msg }); + } + ); + if (!comparison) { + LOG_DBG(ANSI_ORANGE "%s: Template application failed for reserve variant, aborting\n" ANSI_RESET, __func__); + return ""; + } + } + + auto usermsg = comparison->diff.right; + if (usermsg.find(USER_MSG) == std::string::npos) { + LOG_DBG(ANSI_ORANGE "%s: Did not find user message in user message block, aborting detection\n" ANSI_RESET, __func__); + } + + if (usermsg.find(ASSISTANT_MSG) != std::string::npos) { + usermsg = usermsg.substr(usermsg.find(ASSISTANT_MSG) + ASSISTANT_MSG.size()); + } + + auto candidate = usermsg.substr(0, usermsg.find(USER_MSG)); + auto candidate_split = segmentize_markers(candidate); + std::stringstream result; + bool encountered_marker = false; + for (const auto & mrk : candidate_split) { + std::string lower_mrk = std::string(mrk.value); + std::transform(lower_mrk.begin(), lower_mrk.end(), lower_mrk.begin(), + [](unsigned char c) { return std::tolower(c); }); + // heuristic to weed out potential end markers, but only at the start + if (mrk.type == segment_type::MARKER && !encountered_marker && + (lower_mrk.find("end") != std::string::npos || lower_mrk.find("close") != std::string::npos)) { + continue; + } + if (mrk.type == segment_type::TEXT && !encountered_marker && trim_whitespace(mrk.value).empty()) { + continue; + } + encountered_marker |= mrk.type == segment_type::MARKER; + result << mrk.value; + } + return trim_whitespace(result.str()); +} + analyze_reasoning::analyze_reasoning(const common_chat_template & tmpl, bool supports_tools) : analyze_base(tmpl) { LOG_DBG(ANSI_PURPLE "=== Starting differential analysis ===\n" ANSI_RESET); diff --git a/common/chat.cpp b/common/chat.cpp index 1d4ded6d8..aeefa83e3 100644 --- a/common/chat.cpp +++ b/common/chat.cpp @@ -101,6 +101,45 @@ std::string common_chat_msg::render_content(const std::string & delimiter) const return text; } +std::vector common_chat_split_by_role(const std::string & prompt, const std::vector & delims) { + if (delims.empty() || prompt.empty()) { + return {}; + } + + auto parser = build_peg_parser([&](common_peg_parser_builder & p) { + std::vector all_delims; + std::vector tagged_messages; + + all_delims.reserve(delims.size()); + tagged_messages.reserve(delims.size()); + for (const auto & d : delims) { + all_delims.push_back(d.delimiter); + } + + auto any_delim = p.until_one_of(all_delims); + for (const auto & d : delims) { + tagged_messages.push_back(p.tag(d.role, p.literal(d.delimiter) + any_delim)); + } + + return any_delim + p.zero_or_more(p.choice(tagged_messages)) + p.end(); + }); + + common_peg_parse_context ctx(prompt); + const auto result = parser.parse(ctx); + if (!result.success()) { + return {}; + } + + std::vector spans; + ctx.ast.visit(result, [&](const common_peg_ast_node & node) { + if (!node.tag.empty()) { + spans.push_back({ node.tag, node.start, node.end - node.start }); + } + }); + + return spans; +} + json common_chat_msg::to_json_oaicompat(bool concat_typed_text) const { if (!content.empty() && !content_parts.empty()) { throw std::runtime_error("Cannot specify both content and content_parts"); @@ -1057,6 +1096,14 @@ static common_chat_params common_chat_params_init_gpt_oss(const common_chat_temp data.prompt = prompt; data.generation_prompt = common_chat_template_generation_prompt_impl(tmpl, inputs, /* messages_override= */ adjusted_messages); + data.message_spans = common_chat_split_by_role(prompt, { + { "assistant", "<|start|>assistant" }, + { "user", "<|start|>user" }, + { "system", "<|start|>developer" }, + { "system", "<|start|>system" }, + { "tool", "<|start|>functions" }, + }); + data.format = COMMON_CHAT_FORMAT_PEG_NATIVE; data.supports_thinking = true; @@ -1196,6 +1243,11 @@ static common_chat_params common_chat_params_init_gemma4(const common_chat_templ data.prompt += data.generation_prompt; } + data.message_spans = common_chat_split_by_role(data.prompt, { + { "user", "<|turn>user\n" }, + { "assistant", "<|turn>model\n" }, + }); + data.format = COMMON_CHAT_FORMAT_PEG_GEMMA4; data.supports_thinking = true; data.thinking_start_tag = "<|channel>thought"; @@ -2408,6 +2460,19 @@ static common_chat_params common_chat_templates_apply_jinja(const struct common_ struct autoparser::autoparser autoparser; autoparser.analyze_template(tmpl); auto auto_params = autoparser::peg_generator::generate_parser(tmpl, params, autoparser); + + std::vector delimiters; + if (!autoparser.assistant_start.empty()) { + delimiters.push_back({ "assistant", autoparser.assistant_start }); + } + if (!autoparser.user_start.empty()) { + delimiters.push_back({ "user", autoparser.user_start }); + } + + if (!delimiters.empty()) { + auto_params.message_spans = common_chat_split_by_role(auto_params.prompt, delimiters); + } + auto_params.supports_thinking = autoparser.reasoning.mode != autoparser::reasoning_mode::NONE; if (auto_params.supports_thinking) { auto_params.thinking_start_tag = trim_whitespace(autoparser.reasoning.start); diff --git a/common/chat.h b/common/chat.h index b29c627e6..5659cd42a 100644 --- a/common/chat.h +++ b/common/chat.h @@ -143,6 +143,17 @@ struct common_chat_msg_diff { } }; +struct common_chat_msg_span { + std::string role; + std::size_t pos = 0; + std::size_t len = 0; +}; + +struct common_chat_msg_delimiter { + std::string role; + std::string delimiter; +}; + struct common_chat_tool { std::string name; std::string description; @@ -208,6 +219,7 @@ struct common_chat_params { std::vector preserved_tokens; std::vector additional_stops; std::string parser; + std::vector message_spans; }; // per-message parsing syntax @@ -304,6 +316,7 @@ std::optional common_chat_try_specialized_template( const std::string & src, autoparser::generation_params & params); + // specialized per-task preset struct common_chat_prompt_preset { std::string system; @@ -311,3 +324,6 @@ struct common_chat_prompt_preset { }; common_chat_prompt_preset common_chat_get_asr_prompt(const common_chat_templates * chat_templates); + +std::vector common_chat_split_by_role(const std::string & prompt, const std::vector & delims); + diff --git a/common/common.cpp b/common/common.cpp index 452460609..b3c10e8bb 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -451,6 +451,27 @@ std::string string_strip(const std::string & str) { return str.substr(start, end - start); } +std::string string_lcs(std::string_view a, std::string_view b) { + if (a.empty() || b.empty()) return {}; + + std::vector> dp(a.size() + 1, std::vector(b.size() + 1, 0)); + size_t best_len = 0; + size_t best_end_a = 0; + + for (size_t i = 1; i <= a.size(); ++i) { + for (size_t j = 1; j <= b.size(); ++j) { + if (a[i - 1] == b[j - 1]) { + dp[i][j] = dp[i - 1][j - 1] + 1; + if (dp[i][j] > best_len) { + best_len = dp[i][j]; + best_end_a = i; + } + } + } + } + return std::string(a.substr(best_end_a - best_len, best_len)); +} + std::string string_get_sortable_timestamp() { using clock = std::chrono::system_clock; diff --git a/common/common.h b/common/common.h index 2fb986f15..cf97a7484 100644 --- a/common/common.h +++ b/common/common.h @@ -595,7 +595,7 @@ struct common_params { bool cache_prompt = true; // whether to enable prompt caching bool cache_idle_slots = true; // save and clear idle slots upon starting a new task int32_t n_ctx_checkpoints = 32; // max number of context checkpoints per slot - int32_t checkpoint_every_nt = 8192; // make a checkpoint every n tokens during prefill + int32_t checkpoint_min_step = 256; // minimum spacing between context checkpoints int32_t cache_ram_mib = 8192; // -1 = no limit, 0 - disable, 1 = 1 MiB, etc. std::string hostname = "127.0.0.1"; @@ -732,6 +732,7 @@ std::string string_format(const char * fmt, ...); std::string string_strip(const std::string & str); std::string string_get_sortable_timestamp(); +std::string string_lcs(std::string_view a, std::string_view b); std::string string_join(const std::vector & values, const std::string & separator); std::vector string_split(const std::string & str, const std::string & delimiter); diff --git a/common/fit.cpp b/common/fit.cpp index c10cb7f08..668d892e9 100644 --- a/common/fit.cpp +++ b/common/fit.cpp @@ -26,7 +26,7 @@ class common_params_fit_exception : public std::runtime_error { using std::runtime_error::runtime_error; }; -static std::vector common_get_device_memory_data( +std::vector common_get_device_memory_data( const char * path_model, const llama_model_params * mparams, const llama_context_params * cparams, diff --git a/common/fit.h b/common/fit.h index e066092ec..643d34200 100644 --- a/common/fit.h +++ b/common/fit.h @@ -1,6 +1,11 @@ #pragma once #include "ggml.h" +#include "ggml-backend.h" +#include "llama.h" +#include "../src/llama-ext.h" + +#include enum common_params_fit_status { COMMON_PARAMS_FIT_STATUS_SUCCESS = 0, // found allocations that are projected to fit @@ -30,3 +35,14 @@ void common_fit_print( struct llama_context_params * cparams); void common_memory_breakdown_print(const struct llama_context * ctx); + +// Load a model + context with no_alloc and return the per-device memory breakdown. +std::vector common_get_device_memory_data( + const char * path_model, + const struct llama_model_params * mparams, + const struct llama_context_params * cparams, + std::vector & devs, + uint32_t & hp_ngl, + uint32_t & hp_n_ctx_train, + uint32_t & hp_n_expert, + enum ggml_log_level log_level); diff --git a/ggml/include/ggml-alloc.h b/ggml/include/ggml-alloc.h index 78aa059dd..a7926a21a 100644 --- a/ggml/include/ggml-alloc.h +++ b/ggml/include/ggml-alloc.h @@ -76,6 +76,7 @@ GGML_API size_t ggml_gallocr_get_buffer_size(ggml_gallocr_t galloc, int buffer_i // Utils // Create a buffer and allocate all the tensors in a ggml_context // ggml_backend_alloc_ctx_tensors_from_buft_size returns the size of the buffer that would be allocated by ggml_backend_alloc_ctx_tensors_from_buft +// ggml_backend_alloc_ctx_tensors_from_buft returns NULL on failure or if all tensors in ctx are already allocated or zero-sized GGML_API size_t ggml_backend_alloc_ctx_tensors_from_buft_size(struct ggml_context * ctx, ggml_backend_buffer_type_t buft); GGML_API struct ggml_backend_buffer * ggml_backend_alloc_ctx_tensors_from_buft(struct ggml_context * ctx, ggml_backend_buffer_type_t buft); GGML_API struct ggml_backend_buffer * ggml_backend_alloc_ctx_tensors(struct ggml_context * ctx, ggml_backend_t backend); diff --git a/ggml/include/gguf.h b/ggml/include/gguf.h index 02d5f221c..67851ba6f 100644 --- a/ggml/include/gguf.h +++ b/ggml/include/gguf.h @@ -76,10 +76,16 @@ extern "C" { struct ggml_context ** ctx; }; + // callback to simulate or wrap a FILE pointer - read up to `len` bytes at `offset` into `output` and return the number of bytes read + typedef size_t (*gguf_reader_callback_t)(void * userdata, void * output, uint64_t offset, size_t len); + GGML_API struct gguf_context * gguf_init_empty(void); GGML_API struct gguf_context * gguf_init_from_file_ptr(FILE * file, struct gguf_init_params params); GGML_API struct gguf_context * gguf_init_from_file(const char * fname, struct gguf_init_params params); - //GGML_API struct gguf_context * gguf_init_from_buffer(..); + GGML_API struct gguf_context * gguf_init_from_buffer(const void * data, size_t size, struct gguf_init_params params); + + // max_chunk_read is the maximum number of bytes that the GGUF code will read at once from the callback, a value of 0 means no limit + GGML_API struct gguf_context * gguf_init_from_callback(gguf_reader_callback_t callback, void * userdata, size_t max_chunk_read, uint64_t max_expected_size, struct gguf_init_params params); GGML_API void gguf_free(struct gguf_context * ctx); @@ -87,7 +93,7 @@ extern "C" { GGML_API uint32_t gguf_get_version (const struct gguf_context * ctx); GGML_API size_t gguf_get_alignment (const struct gguf_context * ctx); - GGML_API size_t gguf_get_data_offset(const struct gguf_context * ctx); + GGML_API size_t gguf_get_data_offset(const struct gguf_context * ctx); // padded to gguf_get_alignment if and only if the gguf_context contains at least one tensor GGML_API int64_t gguf_get_n_kv(const struct gguf_context * ctx); GGML_API int64_t gguf_find_key(const struct gguf_context * ctx, const char * key); // returns -1 if key is not found diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index df0f405ed..d0d64523b 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include #include #include @@ -392,64 +393,100 @@ static ggml_backend_buffer_type_t ggml_backend_meta_device_get_host_buffer_type( // meta backend buffer // +// Container to hold the tensor slices per simple ggml backend buffer. +struct ggml_backend_meta_simple_tensor_container { + std::vector ctxs; + std::map> simple_tensors; + + ggml_backend_meta_simple_tensor_container(const ggml_init_params & params, const int n_simple) { + ctxs.reserve(n_simple); + for (int i = 0; i < n_simple; i++) { + ctxs.emplace_back(ggml_init(params)); + } + } + ggml_backend_meta_simple_tensor_container() {} +}; + struct ggml_backend_meta_buffer_context { + // FIXME + // Most tensors can simply be stored statically in their own buffer. + // Externally created views however also need a mapping to simple tensors but they use the buffer of the view source. + // If external views are simply using that buffer they will slowly deplete its memory. + // Current solution: rotating set of 2 "compute" containers to hold external views, works correctly for llama.cpp. + // Long-term: tie the lifetime of external views to the meta backend executing the graph instead, + // currently not possible due to graph-external operations in the backend scheduler. + ggml_backend_meta_simple_tensor_container stc_static; + ggml_backend_meta_simple_tensor_container stc_compute[2]; + int stc_compute_index = 0; + int stc_compute_index_next = 0; + std::vector bufs; + + // FIXME + // The size of the split state cache is unbounded and can theoretically grow infinitely large. + // However, it is also expensive to build and clearing it on every rebuild in ggml_backend_meta_graph_compute is too expensive. static constexpr size_t nbtc = GGML_TENSOR_SIZE - sizeof(ggml_tensor::padding); - std::map, std::pair> split_state_cache; - std::map< const ggml_tensor *, std::vector> simple_tensors; - - struct buffer_config { - ggml_context * ctx; - ggml_backend_buffer_t buf; - - buffer_config(ggml_context * ctx, ggml_backend_buffer_t buf) : ctx(ctx), buf(buf) {} - }; - std::vector buf_configs; int debug; - ggml_backend_meta_buffer_context() { + ggml_backend_meta_buffer_context( + ggml_backend_meta_simple_tensor_container & stc_static, + ggml_backend_meta_simple_tensor_container & stc_compute_0, + ggml_backend_meta_simple_tensor_container & stc_compute_1, + const std::vector & bufs) + : stc_static(std::move(stc_static)), stc_compute{std::move(stc_compute_0), std::move(stc_compute_1)} { + this->bufs.reserve(bufs.size()); + for (ggml_backend_buffer_t buf : bufs) { + this->bufs.emplace_back(buf); + } const char * GGML_META_DEBUG = getenv("GGML_META_DEBUG"); debug = GGML_META_DEBUG ? atoi(GGML_META_DEBUG) : 0; } + + ggml_backend_meta_simple_tensor_container & get_simple_tensor_container(const ggml_tensor * tensor) { + if (stc_static.simple_tensors.find(tensor) != stc_static.simple_tensors.end()) { + return stc_static; + } + return stc_compute[stc_compute_index]; + } }; static void ggml_backend_meta_buffer_free_buffer(ggml_backend_buffer_t buffer) { GGML_ASSERT(ggml_backend_buffer_is_meta(buffer)); ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buffer->context; - for (auto & [ctx, buf] : buf_ctx->buf_configs) { - ggml_backend_buffer_free(buf); - ggml_free(ctx); - } delete buf_ctx; } static size_t ggml_backend_meta_buffer_n_bufs(ggml_backend_buffer_t meta_buf) { GGML_ASSERT(ggml_backend_buffer_is_meta(meta_buf)); ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) meta_buf->context; - return buf_ctx->buf_configs.size(); + return buf_ctx->bufs.size(); } static ggml_backend_buffer_t ggml_backend_meta_buffer_simple_buffer(ggml_backend_buffer_t meta_buf, size_t index) { GGML_ASSERT(ggml_backend_buffer_is_meta(meta_buf)); ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) meta_buf->context; - GGML_ASSERT(index < buf_ctx->buf_configs.size()); - return buf_ctx->buf_configs[index].buf; + GGML_ASSERT(index < buf_ctx->bufs.size()); + return buf_ctx->bufs[index].get(); } static struct ggml_tensor * ggml_backend_meta_buffer_simple_tensor(const struct ggml_tensor * tensor, size_t index) { GGML_ASSERT(ggml_backend_buffer_is_meta(tensor->buffer)); ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) tensor->buffer->context; - GGML_ASSERT(index < buf_ctx->buf_configs.size()); + GGML_ASSERT(index < buf_ctx->bufs.size()); - auto it = buf_ctx->simple_tensors.find(tensor); - if (it == buf_ctx->simple_tensors.end()) { + ggml_backend_meta_simple_tensor_container & stc = buf_ctx->get_simple_tensor_container(tensor); + auto it = stc.simple_tensors.find(tensor); + if (it == stc.simple_tensors.end()) { return nullptr; } return it->second[index]; } -static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(const struct ggml_tensor * tensor, bool assume_sync) { +static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(const struct ggml_tensor * tensor, bool assume_sync); + +static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state( + ggml_backend_meta_simple_tensor_container & stc, const struct ggml_tensor * tensor, bool assume_sync) { const size_t n_bufs = ggml_backend_meta_buffer_n_bufs(tensor->buffer); ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) tensor->buffer->context; @@ -785,7 +822,7 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(co src_ss[i] = {GGML_BACKEND_SPLIT_AXIS_UNKNOWN, {0}, 1}; continue; } - src_ss[i] = ggml_backend_meta_get_split_state(tensor->src[i], /*assume_sync =*/ true); + src_ss[i] = ggml_backend_meta_get_split_state(stc, tensor->src[i], /*assume_sync =*/ true); GGML_ASSERT(src_ss[i].axis != GGML_BACKEND_SPLIT_AXIS_UNKNOWN); } @@ -1079,17 +1116,23 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(co return ret; } +static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(const struct ggml_tensor * tensor, bool assume_sync) { + GGML_ASSERT(ggml_backend_buffer_is_meta(tensor->buffer)); + ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) tensor->buffer->context; + return ggml_backend_meta_get_split_state(buf_ctx->get_simple_tensor_container(tensor), tensor, assume_sync); +} + static void * ggml_backend_meta_buffer_get_base(ggml_backend_buffer_t buffer) { GGML_UNUSED(buffer); return (void *) 0x1000000000000000; // FIXME } -static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor) { - GGML_ASSERT(ggml_backend_buffer_is_meta(buffer)); - ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buffer->context; - const size_t n_simple_bufs = ggml_backend_meta_buffer_n_bufs(buffer); +static enum ggml_status ggml_backend_meta_buffer_init_tensor_impl(ggml_backend_meta_simple_tensor_container & stc, ggml_tensor * tensor) { + GGML_ASSERT(ggml_backend_buffer_is_meta(tensor->buffer)); + ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) tensor->buffer->context; + const size_t n_simple_bufs = ggml_backend_meta_buffer_n_bufs(tensor->buffer); - const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(tensor, /*assume_sync =*/ true); + const ggml_backend_meta_split_state split_state = ggml_backend_meta_get_split_state(stc, tensor, /*assume_sync =*/ true); GGML_ASSERT(ggml_nelements(tensor) == 0 || split_state.axis != GGML_BACKEND_SPLIT_AXIS_UNKNOWN); GGML_ASSERT(split_state.n_segments <= 16); @@ -1104,8 +1147,8 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer std::vector simple_tensors; simple_tensors.reserve(n_simple_bufs); for (size_t j = 0; j < n_simple_bufs; j++) { - ggml_context * simple_ctx = buf_ctx->buf_configs[j].ctx; - ggml_backend_buffer_t simple_buf = buf_ctx->buf_configs[j].buf; + ggml_context * simple_ctx = stc.ctxs[j].get(); + ggml_backend_buffer_t simple_buf = buf_ctx->bufs[j].get(); if (split_dim >= 0 && split_dim < GGML_MAX_DIMS) { // TODO: the following assert fails for llama-parallel even though the results are correct: @@ -1158,7 +1201,7 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer t_ij->data = (char *) t_ij->view_src->data + t_ij->view_offs; } else if (simple_buf != nullptr) { t_ij->data = (char *) ggml_backend_buffer_get_base(simple_buf) - + size_t(tensor->data) - size_t(ggml_backend_buffer_get_base(buffer)); + + size_t(tensor->data) - size_t(ggml_backend_buffer_get_base(tensor->buffer)); } t_ij->extra = tensor->extra; for (int i = 0; i < GGML_MAX_SRC; i++) { @@ -1194,11 +1237,18 @@ static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer } } - buf_ctx->simple_tensors[tensor] = simple_tensors; + stc.simple_tensors[tensor] = simple_tensors; return GGML_STATUS_SUCCESS; } +static enum ggml_status ggml_backend_meta_buffer_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor) { + GGML_ASSERT(ggml_backend_buffer_is_meta(buffer)); + ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buffer->context; + buf_ctx->stc_compute_index = buf_ctx->stc_compute_index_next; + return ggml_backend_meta_buffer_init_tensor_impl(buf_ctx->get_simple_tensor_container(tensor), tensor); +} + static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor, const void * data, size_t offset, size_t size) { const size_t n_bufs = ggml_backend_meta_buffer_n_bufs(buffer); GGML_ASSERT(ggml_is_contiguous(tensor)); @@ -1275,6 +1325,9 @@ static void ggml_backend_meta_buffer_set_tensor(ggml_backend_buffer_t buffer, gg for (size_t j = 0; j < n_bufs; j++) { ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j); const size_t chunk_size_j = simple_tensor->nb[split_state.axis + 1]; + if (chunk_size_j == 0) { + continue; + } const size_t simple_offset = i_start * chunk_size_j; ggml_backend_tensor_set_2d(simple_tensor, (const char *) data + offset_j, simple_offset, chunk_size_j, i_stop - i_start, chunk_size_j, chunk_size_full); offset_j += chunk_size_j; @@ -1382,6 +1435,9 @@ static void ggml_backend_meta_buffer_get_tensor(ggml_backend_buffer_t buffer, co for (size_t j = 0; j < n_bufs; j++){ const ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j); const size_t chunk_size_j = simple_tensor->nb[split_state.axis + 1]; + if (chunk_size_j == 0) { + continue; + } const size_t simple_offset = i_start * chunk_size_j; ggml_backend_tensor_get_2d(simple_tensor, (char *) data + offset_j, simple_offset, chunk_size_j, i_stop - i_start, chunk_size_j, chunk_size_full); offset_j += chunk_size_j; @@ -1407,8 +1463,9 @@ static void ggml_backend_meta_buffer_clear(ggml_backend_buffer_t buffer, uint8_t } static void ggml_backend_meta_buffer_reset(ggml_backend_buffer_t buffer) { - const size_t n_buffers = ggml_backend_meta_buffer_n_bufs(buffer); - for (size_t i = 0; i < n_buffers; i++) { + GGML_ASSERT(ggml_backend_buffer_is_meta(buffer)); + ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buffer->context; + for (size_t i = 0; i < buf_ctx->bufs.size(); i++) { ggml_backend_buffer_reset(ggml_backend_meta_buffer_simple_buffer(buffer, i)); } } @@ -1434,20 +1491,24 @@ bool ggml_backend_buffer_is_meta(ggml_backend_buffer_t buf) { static ggml_backend_buffer_t ggml_backend_meta_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) { const size_t n_simple_bufts = ggml_backend_meta_buft_n_bufts(buft); - ggml_init_params params = { - /*.mem_size =*/ 1024*1024*1024, // FIXME + const ggml_init_params params = { + /*.mem_size =*/ 1024*1024*ggml_tensor_overhead(), // FIXME /*.mem_buffer =*/ nullptr, /*.no_alloc =*/ true, }; + ggml_backend_meta_simple_tensor_container stc_static; + ggml_backend_meta_simple_tensor_container stc_compute_0(params, n_simple_bufts); + ggml_backend_meta_simple_tensor_container stc_compute_1(params, n_simple_bufts); - ggml_backend_meta_buffer_context * buf_ctx = new ggml_backend_meta_buffer_context(); size_t max_size = 0; - buf_ctx->buf_configs.reserve(n_simple_bufts); + std::vector bufs; + bufs.reserve(n_simple_bufts); for (size_t i = 0; i < n_simple_bufts; i++) { - ggml_backend_buffer_t simple_buf = ggml_backend_buft_alloc_buffer(ggml_backend_meta_buft_simple_buft(buft, i), size); - max_size = std::max(max_size, ggml_backend_buffer_get_size(simple_buf)); - buf_ctx->buf_configs.emplace_back(ggml_init(params), simple_buf); + bufs.push_back(ggml_backend_buft_alloc_buffer(ggml_backend_meta_buft_simple_buft(buft, i), size)); + GGML_ASSERT(bufs.back() != nullptr); + max_size = std::max(max_size, ggml_backend_buffer_get_size(bufs.back())); } + ggml_backend_meta_buffer_context * buf_ctx = new ggml_backend_meta_buffer_context(stc_static, stc_compute_0, stc_compute_1, bufs); return ggml_backend_buffer_init(buft, ggml_backend_meta_buffer_iface, buf_ctx, max_size); } @@ -1455,28 +1516,53 @@ static ggml_backend_buffer_t ggml_backend_meta_buffer_type_alloc_buffer(ggml_bac struct ggml_backend_buffer * ggml_backend_meta_alloc_ctx_tensors_from_buft(struct ggml_context * ctx, ggml_backend_buffer_type_t buft) { const size_t n_simple_bufts = ggml_backend_meta_buft_n_bufts(buft); - ggml_init_params params = { - /*.mem_size =*/ 1024*1024*1024, // FIXME + constexpr size_t compute_headroom = 16; // Maximum number of views per statically allocated tensor that can be created between evals. + const ggml_init_params params_static = { + /*.mem_size =*/ ggml_get_mem_size(ctx), /*.mem_buffer =*/ nullptr, /*.no_alloc =*/ true, }; + const ggml_init_params params_compute = { + /*.mem_size =*/ compute_headroom*ggml_get_mem_size(ctx), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ggml_backend_meta_simple_tensor_container stc_static (params_static, n_simple_bufts); + ggml_backend_meta_simple_tensor_container stc_compute_0(params_compute, n_simple_bufts); + ggml_backend_meta_simple_tensor_container stc_compute_1(params_compute, n_simple_bufts); - ggml_backend_meta_buffer_context * meta_buf_ctx = new ggml_backend_meta_buffer_context(); - meta_buf_ctx->buf_configs.reserve(n_simple_bufts); - for (size_t i = 0; i < n_simple_bufts; i++) { - meta_buf_ctx->buf_configs.emplace_back(ggml_init(params), nullptr); - } + std::vector bufs(n_simple_bufts, nullptr); + ggml_backend_meta_buffer_context * meta_buf_ctx = new ggml_backend_meta_buffer_context(stc_static, stc_compute_0, stc_compute_1, bufs); ggml_backend_buffer_t meta_buf = ggml_backend_buffer_init(buft, ggml_backend_meta_buffer_iface, meta_buf_ctx, 0); for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) { t->buffer = meta_buf; - ggml_backend_meta_buffer_init_tensor(meta_buf, t); + ggml_backend_meta_buffer_init_tensor_impl(meta_buf_ctx->stc_static, t); t->data = (void *) 0x2000000000000000; // FIXME } for (size_t i = 0; i < n_simple_bufts; i++) { - meta_buf_ctx->buf_configs[i].buf = ggml_backend_alloc_ctx_tensors_from_buft( - meta_buf_ctx->buf_configs[i].ctx, ggml_backend_meta_buft_simple_buft(buft, i)); - meta_buf->size = std::max(meta_buf->size, ggml_backend_buffer_get_size(meta_buf_ctx->buf_configs[i].buf)); + ggml_context * ctx = meta_buf_ctx->stc_static.ctxs[i].get(); + ggml_backend_buffer_type_t simple_buft = ggml_backend_meta_buft_simple_buft(buft, i); + + // If a ggml_context only has zero-sized tensors, ggml_backend_alloc_ctx_tensors_from_buft returns NULL. + // For those edge cases, allocate a dummy buffer instead. + bool any_nonzero_slice = false; + for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) { + if (ggml_nelements(t) != 0) { + any_nonzero_slice = true; + break; + } + } + if (any_nonzero_slice) { + meta_buf_ctx->bufs[i].reset(ggml_backend_alloc_ctx_tensors_from_buft(ctx, simple_buft)); + } else { + meta_buf_ctx->bufs[i].reset(ggml_backend_buft_alloc_buffer(simple_buft, 0)); + for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) { + t->buffer = meta_buf_ctx->bufs[i].get(); + } + } + GGML_ASSERT(meta_buf_ctx->bufs[i]); + meta_buf->size = std::max(meta_buf->size, ggml_backend_buffer_get_size(meta_buf_ctx->bufs[i].get())); } return meta_buf; } @@ -1605,6 +1691,9 @@ static void ggml_backend_meta_set_tensor_async(ggml_backend_t backend, ggml_tens ggml_backend_t simple_backend = ggml_backend_meta_simple_backend(backend, j); ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j); const size_t chunk_size_j = simple_tensor->nb[split_state.axis + 1]; + if (chunk_size_j == 0) { + continue; + } ggml_backend_tensor_set_2d_async(simple_backend, simple_tensor, (const char *) data + offset_j, offset, chunk_size_j, i_stop - i_start, chunk_size_j, chunk_size_full); offset_j += chunk_size_j; @@ -1646,6 +1735,9 @@ static void ggml_backend_meta_get_tensor_async(ggml_backend_t backend, const ggm ggml_backend_t simple_backend = ggml_backend_meta_simple_backend(backend, j); const ggml_tensor * simple_tensor = ggml_backend_meta_buffer_simple_tensor(tensor, j); const size_t chunk_size_j = simple_tensor->nb[split_state.axis + 1]; + if (chunk_size_j == 0) { + continue; + } ggml_backend_tensor_get_2d_async(simple_backend, simple_tensor, (char *) data + offset_j, offset, chunk_size_j, i_stop - i_start, chunk_size_j, chunk_size_full); offset_j += chunk_size_j; @@ -1692,6 +1784,26 @@ static enum ggml_status ggml_backend_meta_graph_compute(ggml_backend_t backend, } if (needs_rebuild) { + std::set used_buffers; + for (int i = 0; i < cgraph->n_leafs; i++) { + if (ggml_backend_buffer_is_meta(cgraph->leafs[i]->buffer)) { + used_buffers.emplace(cgraph->leafs[i]->buffer); + } + } + for (int i = 0; i < cgraph->n_nodes; i++) { + if (ggml_backend_buffer_is_meta(cgraph->nodes[i]->buffer)) { + used_buffers.emplace(cgraph->nodes[i]->buffer); + } + } + for (ggml_backend_buffer_t buf : used_buffers) { + ggml_backend_meta_buffer_context * buf_ctx = (ggml_backend_meta_buffer_context *) buf->context; + buf_ctx->stc_compute_index_next = buf_ctx->stc_compute_index ^ 1; + ggml_backend_meta_simple_tensor_container & stc = buf_ctx->stc_compute[buf_ctx->stc_compute_index_next]; + for (ggml_context_ptr & ctx : stc.ctxs) { + ggml_reset(ctx.get()); + } + stc.simple_tensors.clear(); + } size_t n_subgraphs = 0; size_t max_tmp_size = 0; @@ -1877,7 +1989,7 @@ static enum ggml_status ggml_backend_meta_graph_compute(ggml_backend_t backend, const size_t mem_per_device_graphs_main = backend_ctx->max_subgraphs*ggml_graph_overhead_custom(backend_ctx->max_nnodes, cgraph->grads); const size_t mem_per_device_graphs_aux = n_cgraphs_per_device*backend_ctx->max_subgraphs*ggml_graph_overhead_custom(1, cgraph->grads); const size_t mem_per_device_nodes_aux = n_nodes_per_device*backend_ctx->max_subgraphs*ggml_tensor_overhead(); - ggml_init_params params = { + const ggml_init_params params = { /*.mem_size =*/ n_backends * (mem_per_device_graphs_main + mem_per_device_graphs_aux + mem_per_device_nodes_aux), /*.mem_buffer =*/ nullptr, /*.no_alloc =*/ true, diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 15443aa55..15d231f70 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -13,6 +13,10 @@ #include // for qsort #include // for GGML_ASSERT +#ifdef GGML_USE_OPENMP +#include +#endif + #define GROUP_MAX_EPS 1e-15f #define GROUP_MAX_EPS_IQ3_XXS 1e-8f #define GROUP_MAX_EPS_IQ2_S 1e-8f @@ -3064,70 +3068,121 @@ void iq2xs_init_impl(enum ggml_type type) { } kmap_q2xs[index] = i; } - int8_t pos[8]; - int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + // The neighbour search runs in three passes: + // 1. Parallel: for each i, qsort and count its neighbours into n_per_i, + // and reduce the totals (num_neighbors, num_not_in_map). + // 2. Serial: prefix-sum n_per_i into offsets[], so each i has a + // pre-assigned slice of kneighbors_q2xs to write into. + // 3. Parallel: redo the qsort and write each i's neighbour list at + // offsets[i]. + int * n_per_i = (int *)malloc(kmap_size*sizeof(int)); + GGML_ASSERT(n_per_i); int num_neighbors = 0, num_not_in_map = 0; - for (int i = 0; i < kmap_size; ++i) { - if (kmap_q2xs[i] >= 0) continue; - ++num_not_in_map; - for (int k = 0; k < 8; ++k) { - int l = (i >> 2*k) & 0x3; - pos[k] = 2*l + 1; - } - for (int j = 0; j < grid_size; ++j) { - const int8_t * pg = (const int8_t *)(kgrid_q2xs + j); - int d2 = 0; - for (int k = 0; k < 8; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); - dist2[2*j+0] = d2; - dist2[2*j+1] = j; - } - qsort(dist2, grid_size, 2*sizeof(int), iq2_compare_func); - int n = 0; int d2 = dist2[0]; - int nhave = 1; - for (int j = 0; j < grid_size; ++j) { - if (dist2[2*j] > d2) { - if (nhave == nwant) break; - d2 = dist2[2*j]; - ++nhave; +#ifdef GGML_USE_OPENMP + #pragma omp parallel reduction(+:num_neighbors,num_not_in_map) +#endif + { + int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + GGML_ASSERT(dist2); + int8_t pos[8]; + int i; +#ifdef GGML_USE_OPENMP + #pragma omp for schedule(dynamic, 64) +#endif + for (i = 0; i < kmap_size; ++i) { + if (kmap_q2xs[i] >= 0) { + n_per_i[i] = 0; + continue; } - ++n; + ++num_not_in_map; + for (int k = 0; k < 8; ++k) { + int l = (i >> 2*k) & 0x3; + pos[k] = 2*l + 1; + } + for (int j = 0; j < grid_size; ++j) { + const int8_t * pg = (const int8_t *)(kgrid_q2xs + j); + int d2 = 0; + for (int k = 0; k < 8; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); + dist2[2*j+0] = d2; + dist2[2*j+1] = j; + } + qsort(dist2, grid_size, 2*sizeof(int), iq2_compare_func); + int n = 0; int d2 = dist2[0]; + int nhave = 1; + for (int j = 0; j < grid_size; ++j) { + if (dist2[2*j] > d2) { + if (nhave == nwant) break; + d2 = dist2[2*j]; + ++nhave; + } + ++n; + } + n_per_i[i] = n; + num_neighbors += n; } - num_neighbors += n; + free(dist2); } //printf("%s: %d neighbours in total\n", __func__, num_neighbors); kneighbors_q2xs = (uint16_t *)malloc((num_neighbors + num_not_in_map)*sizeof(uint16_t)); iq2_data[gindex].neighbours = kneighbors_q2xs; + + int * offsets = (int *)malloc(kmap_size*sizeof(int)); + GGML_ASSERT(offsets); int counter = 0; for (int i = 0; i < kmap_size; ++i) { - if (kmap_q2xs[i] >= 0) continue; - for (int k = 0; k < 8; ++k) { - int l = (i >> 2*k) & 0x3; - pos[k] = 2*l + 1; + if (kmap_q2xs[i] >= 0) { + offsets[i] = -1; + continue; } - for (int j = 0; j < grid_size; ++j) { - const int8_t * pg = (const int8_t *)(kgrid_q2xs + j); - int d2 = 0; - for (int k = 0; k < 8; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); - dist2[2*j+0] = d2; - dist2[2*j+1] = j; - } - qsort(dist2, grid_size, 2*sizeof(int), iq2_compare_func); - kmap_q2xs[i] = -(counter + 1); - int d2 = dist2[0]; - uint16_t * start = &kneighbors_q2xs[counter++]; - int n = 0, nhave = 1; - for (int j = 0; j < grid_size; ++j) { - if (dist2[2*j] > d2) { - if (nhave == nwant) break; - d2 = dist2[2*j]; - ++nhave; - } - kneighbors_q2xs[counter++] = dist2[2*j+1]; - ++n; - } - *start = n; + offsets[i] = counter; + counter += 1 + n_per_i[i]; } - free(dist2); + +#ifdef GGML_USE_OPENMP + #pragma omp parallel +#endif + { + int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + GGML_ASSERT(dist2); + int8_t pos[8]; + int i; +#ifdef GGML_USE_OPENMP + #pragma omp for schedule(dynamic, 64) +#endif + for (i = 0; i < kmap_size; ++i) { + if (kmap_q2xs[i] >= 0) continue; + for (int k = 0; k < 8; ++k) { + int l = (i >> 2*k) & 0x3; + pos[k] = 2*l + 1; + } + for (int j = 0; j < grid_size; ++j) { + const int8_t * pg = (const int8_t *)(kgrid_q2xs + j); + int d2 = 0; + for (int k = 0; k < 8; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); + dist2[2*j+0] = d2; + dist2[2*j+1] = j; + } + qsort(dist2, grid_size, 2*sizeof(int), iq2_compare_func); + int local_counter = offsets[i]; + kmap_q2xs[i] = -(local_counter + 1); + int d2 = dist2[0]; + uint16_t * start = &kneighbors_q2xs[local_counter++]; + int n = 0, nhave = 1; + for (int j = 0; j < grid_size; ++j) { + if (dist2[2*j] > d2) { + if (nhave == nwant) break; + d2 = dist2[2*j]; + ++nhave; + } + kneighbors_q2xs[local_counter++] = dist2[2*j+1]; + ++n; + } + *start = n; + } + free(dist2); + } + free(offsets); + free(n_per_i); } void iq2xs_free_impl(enum ggml_type type) { @@ -3663,70 +3718,115 @@ void iq3xs_init_impl(int grid_size) { } kmap_q3xs[index] = i; } - int8_t pos[4]; - int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + // See explanation of parallelism in iq2xs_init_impl + int * n_per_i = (int *)malloc(kmap_size*sizeof(int)); + GGML_ASSERT(n_per_i); int num_neighbors = 0, num_not_in_map = 0; - for (int i = 0; i < kmap_size; ++i) { - if (kmap_q3xs[i] >= 0) continue; - ++num_not_in_map; - for (int k = 0; k < 4; ++k) { - int l = (i >> 3*k) & 0x7; - pos[k] = 2*l + 1; - } - for (int j = 0; j < grid_size; ++j) { - const int8_t * pg = (const int8_t *)(kgrid_q3xs + j); - int d2 = 0; - for (int k = 0; k < 4; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); - dist2[2*j+0] = d2; - dist2[2*j+1] = j; - } - qsort(dist2, grid_size, 2*sizeof(int), iq3_compare_func); - int n = 0; int d2 = dist2[0]; - int nhave = 1; - for (int j = 0; j < grid_size; ++j) { - if (dist2[2*j] > d2) { - if (nhave == nwant) break; - d2 = dist2[2*j]; - ++nhave; +#ifdef GGML_USE_OPENMP + #pragma omp parallel reduction(+:num_neighbors,num_not_in_map) +#endif + { + int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + GGML_ASSERT(dist2); + int8_t pos[4]; + int i; +#ifdef GGML_USE_OPENMP + #pragma omp for schedule(dynamic, 64) +#endif + for (i = 0; i < kmap_size; ++i) { + if (kmap_q3xs[i] >= 0) { + n_per_i[i] = 0; + continue; } - ++n; + ++num_not_in_map; + for (int k = 0; k < 4; ++k) { + int l = (i >> 3*k) & 0x7; + pos[k] = 2*l + 1; + } + for (int j = 0; j < grid_size; ++j) { + const int8_t * pg = (const int8_t *)(kgrid_q3xs + j); + int d2 = 0; + for (int k = 0; k < 4; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); + dist2[2*j+0] = d2; + dist2[2*j+1] = j; + } + qsort(dist2, grid_size, 2*sizeof(int), iq3_compare_func); + int n = 0; int d2 = dist2[0]; + int nhave = 1; + for (int j = 0; j < grid_size; ++j) { + if (dist2[2*j] > d2) { + if (nhave == nwant) break; + d2 = dist2[2*j]; + ++nhave; + } + ++n; + } + n_per_i[i] = n; + num_neighbors += n; } - num_neighbors += n; + free(dist2); } //printf("%s: %d neighbours in total\n", __func__, num_neighbors); kneighbors_q3xs = (uint16_t *)malloc((num_neighbors + num_not_in_map)*sizeof(uint16_t)); iq3_data[gindex].neighbours = kneighbors_q3xs; + + int * offsets = (int *)malloc(kmap_size*sizeof(int)); + GGML_ASSERT(offsets); int counter = 0; for (int i = 0; i < kmap_size; ++i) { - if (kmap_q3xs[i] >= 0) continue; - for (int k = 0; k < 4; ++k) { - int l = (i >> 3*k) & 0x7; - pos[k] = 2*l + 1; + if (kmap_q3xs[i] >= 0) { + offsets[i] = -1; + continue; } - for (int j = 0; j < grid_size; ++j) { - const int8_t * pg = (const int8_t *)(kgrid_q3xs + j); - int d2 = 0; - for (int k = 0; k < 4; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); - dist2[2*j+0] = d2; - dist2[2*j+1] = j; - } - qsort(dist2, grid_size, 2*sizeof(int), iq3_compare_func); - kmap_q3xs[i] = -(counter + 1); - int d2 = dist2[0]; - uint16_t * start = &kneighbors_q3xs[counter++]; - int n = 0, nhave = 1; - for (int j = 0; j < grid_size; ++j) { - if (dist2[2*j] > d2) { - if (nhave == nwant) break; - d2 = dist2[2*j]; - ++nhave; - } - kneighbors_q3xs[counter++] = dist2[2*j+1]; - ++n; - } - *start = n; + offsets[i] = counter; + counter += 1 + n_per_i[i]; } - free(dist2); + +#ifdef GGML_USE_OPENMP + #pragma omp parallel +#endif + { + int * dist2 = (int *)malloc(2*grid_size*sizeof(int)); + GGML_ASSERT(dist2); + int8_t pos[4]; + int i; +#ifdef GGML_USE_OPENMP + #pragma omp for schedule(dynamic, 64) +#endif + for (i = 0; i < kmap_size; ++i) { + if (kmap_q3xs[i] >= 0) continue; + for (int k = 0; k < 4; ++k) { + int l = (i >> 3*k) & 0x7; + pos[k] = 2*l + 1; + } + for (int j = 0; j < grid_size; ++j) { + const int8_t * pg = (const int8_t *)(kgrid_q3xs + j); + int d2 = 0; + for (int k = 0; k < 4; ++k) d2 += (pg[k] - pos[k])*(pg[k] - pos[k]); + dist2[2*j+0] = d2; + dist2[2*j+1] = j; + } + qsort(dist2, grid_size, 2*sizeof(int), iq3_compare_func); + int local_counter = offsets[i]; + kmap_q3xs[i] = -(local_counter + 1); + int d2 = dist2[0]; + uint16_t * start = &kneighbors_q3xs[local_counter++]; + int n = 0, nhave = 1; + for (int j = 0; j < grid_size; ++j) { + if (dist2[2*j] > d2) { + if (nhave == nwant) break; + d2 = dist2[2*j]; + ++nhave; + } + kneighbors_q3xs[local_counter++] = dist2[2*j+1]; + ++n; + } + *start = n; + } + free(dist2); + } + free(offsets); + free(n_per_i); } void iq3xs_free_impl(int grid_size) { diff --git a/ggml/src/gguf.cpp b/ggml/src/gguf.cpp index daf45a478..04878e9e8 100644 --- a/ggml/src/gguf.cpp +++ b/ggml/src/gguf.cpp @@ -229,9 +229,18 @@ struct gguf_context { static bool isggufv1 = false; struct gguf_reader { - gguf_reader(FILE * file) : file(file) { - // read the remaining bytes once and update on each read - nbytes_remain = file_remain(file); + gguf_reader( + gguf_reader_callback_t callback, + void * userdata, + size_t max_chunk_read, + uint64_t data_offset = 0, + uint64_t nbytes_remain = 0) + : callback(callback), + userdata(userdata), + max_chunk_read(max_chunk_read), + data_offset(data_offset), + nbytes_remain(nbytes_remain) { + GGML_ASSERT(max_chunk_read > 0); } // helper for remaining bytes in a file @@ -258,12 +267,10 @@ struct gguf_reader { template bool read(T & dst) const { const size_t size = sizeof(dst); - if (nbytes_remain < size) { + if (size > nbytes_remain) { return false; } - const size_t nread = fread(&dst, 1, size, file); - nbytes_remain -= nread; - return nread == size; + return read_raw(&dst, size) == size; } template @@ -354,24 +361,71 @@ struct gguf_reader { } } dst.resize(static_cast(size)); - const size_t nread = fread(dst.data(), 1, size, file); - nbytes_remain -= nread; - return nread == size; + return read_raw(dst.data(), static_cast(size)) == size; } bool read(void * dst, const size_t size) const { if (size > nbytes_remain) { return false; } - const size_t nread = fread(dst, 1, size, file); - nbytes_remain -= nread; - return nread == size; + return read_raw(dst, size) == size; + } + + uint64_t tell() const { + return data_offset; + } + + bool seek(uint64_t absolute_offset) const { + const uint64_t end_offset = uint64_t(data_offset) + nbytes_remain; + if (absolute_offset > end_offset) { + return false; + } + + data_offset = absolute_offset; + nbytes_remain = end_offset - absolute_offset; + + return true; } private: - FILE * file; + size_t read_raw(void * dst, size_t size) const { + if (callback == nullptr || size == 0) { + return 0; + } - mutable uint64_t nbytes_remain; + uint8_t * data = static_cast(dst); + size_t total_nread = 0; + bool reached_eof = false; + + while (total_nread < size) { + const size_t chunk_size = std::min(max_chunk_read, size - total_nread); + if (data_offset + total_nread < data_offset) { + break; + } + const size_t nread = callback(userdata, static_cast(data + total_nread), data_offset + total_nread, chunk_size); + total_nread += nread; + if (nread != chunk_size) { + reached_eof = true; + break; + } + } + + data_offset += total_nread; + GGML_ASSERT(total_nread <= nbytes_remain); + nbytes_remain -= total_nread; + + if (reached_eof) { + nbytes_remain = 0; + } + + return total_nread; + } + + gguf_reader_callback_t callback = nullptr; + void * userdata = nullptr; + size_t max_chunk_read = 0; + mutable uint64_t data_offset = 0; + mutable uint64_t nbytes_remain = 0; }; struct gguf_context * gguf_init_empty(void) { @@ -404,12 +458,7 @@ bool gguf_read_emplace_helper(const struct gguf_reader & gr, std::vectorinfo.size()) == n_tensors); // we require the data section to be aligned, so take into account any padding - if (gguf_fseek(file, GGML_PAD(gguf_ftell(file), ctx->alignment), SEEK_SET) != 0) { + if (n_tensors > 0 && !gr.seek(GGML_PAD(gr.tell(), ctx->alignment))) { GGML_LOG_ERROR("%s: failed to seek to beginning of data section\n", __func__); gguf_free(ctx); return nullptr; } // store the current file offset - this is where the data section starts - ctx->offset = gguf_ftell(file); + ctx->offset = gr.tell(); // compute the total size of the data section, taking into account the alignment { @@ -888,6 +937,89 @@ struct gguf_context * gguf_init_from_file_ptr(FILE * file, struct gguf_init_para return ctx; } +struct gguf_context * gguf_init_from_callback(gguf_reader_callback_t callback, void * userdata, size_t max_chunk_read, uint64_t max_expected_size, struct gguf_init_params params) { + if (callback == nullptr) { + return nullptr; + } + + const struct gguf_reader gr(callback, userdata, max_chunk_read == 0 ? SIZE_MAX : max_chunk_read, 0, max_expected_size); + return gguf_init_from_reader(gr, params); +} + +struct gguf_file_reader { + FILE * file; + uint64_t offset; +}; + +static size_t gguf_file_reader_callback(void * userdata, void * output, uint64_t offset, size_t len) { + GGML_ASSERT(len > 0); + + gguf_file_reader & reader = *static_cast(userdata); + + if (reader.offset != offset) { + if (offset > INT64_MAX || gguf_fseek(reader.file, static_cast(offset), SEEK_SET) != 0) { + return 0; + } + + reader.offset = offset; + } + + const size_t nread = fread(static_cast(output), 1, len, reader.file); + reader.offset += nread; + return nread; +} + +struct gguf_context * gguf_init_from_file_ptr(FILE * file, struct gguf_init_params params) { + if (!file) { + return nullptr; + } + + const int64_t cur = gguf_ftell(file); + if (cur < 0) { + return nullptr; + } + + gguf_file_reader reader = { + /*.file = */ file, + /*.offset = */ static_cast(cur), + }; + const struct gguf_reader gr(gguf_file_reader_callback, &reader, SIZE_MAX, reader.offset, gguf_reader::file_remain(file)); + return gguf_init_from_reader(gr, params); +} + +struct gguf_buffer_reader { + const uint8_t * data; + size_t size; +}; + +static size_t gguf_buffer_reader_callback(void * userdata, void * output, uint64_t offset, size_t len) { + GGML_ASSERT(len > 0); + + const gguf_buffer_reader & reader = *static_cast(userdata); + + if (offset > reader.size || len > reader.size - offset) { + return 0; + } + + const size_t data_offset = static_cast(offset); + const size_t nread = std::min(len, reader.size - data_offset); + memcpy(static_cast(output), reader.data + data_offset, nread); + return nread; +} + +struct gguf_context * gguf_init_from_buffer(const void * data, size_t size, struct gguf_init_params params) { + if (data == nullptr || size == 0) { + return nullptr; + } + + gguf_buffer_reader reader = { + /*.data = */ static_cast(data), + /*.size = */ size, + }; + const struct gguf_reader gr(gguf_buffer_reader_callback, &reader, SIZE_MAX, 0, size); + return gguf_init_from_reader(gr, params); +} + struct gguf_context * gguf_init_from_file(const char * fname, struct gguf_init_params params) { FILE * file = ggml_fopen(fname, "rb"); diff --git a/gguf-py/gguf/quants.py b/gguf-py/gguf/quants.py index 1d9d9ab7d..80966b6ef 100644 --- a/gguf-py/gguf/quants.py +++ b/gguf-py/gguf/quants.py @@ -28,6 +28,7 @@ def quant_shape_from_byte_shape(shape: Sequence[int], quant_type: GGMLQuantizati # This is faster than np.vectorize and np.apply_along_axis because it works on more than one row at a time def _apply_over_grouped_rows(func: Callable[[np.ndarray], np.ndarray], arr: np.ndarray, otype: DTypeLike, oshape: tuple[int, ...]) -> np.ndarray: rows = arr.reshape((-1, arr.shape[-1])) + assert len(rows.shape) osize = 1 for dim in oshape: osize *= dim diff --git a/include/llama.h b/include/llama.h index c2f466a0d..e2b6a9bad 100644 --- a/include/llama.h +++ b/include/llama.h @@ -877,7 +877,8 @@ extern "C" { // work only with partial states, such as SWA KV cache or recurrent cache (e.g. Mamba) #define LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY 1 -// keeps the tensor data on device buffers (i.e. not accessible in host memory, but faster save/load) +// Keeps the tensor data on device buffers (i.e. not accessible in host memory, but faster save/load). +// Getting the state for a seq_id with this flag invalidates all prior states gotten for that seq_id with this flag. #define LLAMA_STATE_SEQ_FLAGS_ON_DEVICE 2 typedef uint32_t llama_state_seq_flags; diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp index dc00edfa8..fb71792fe 100644 --- a/tools/server/server-common.cpp +++ b/tools/server/server-common.cpp @@ -1110,6 +1110,16 @@ json oaicompat_chat_params_parse( llama_params["chat_parser"] = chat_params.parser; } + llama_params["message_spans"] = json::array(); + + for (const auto & span : chat_params.message_spans) { + llama_params["message_spans"].push_back({ + { "role", span.role }, + { "pos", span.pos }, + { "len", span.len }, + }); + } + // Reasoning budget: pass parameters through to sampling layer { int reasoning_budget = opt.reasoning_budget; diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index b939e3b75..ae9e0bf60 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -8,6 +8,7 @@ #include "build-info.h" #include "common.h" +#include "fit.h" #include "llama.h" #include "log.h" #include "sampling.h" @@ -775,7 +776,7 @@ private: for (auto & [dev, size] : mmproj_mem) { total += size; } - SRV_INF("[mtmd] estimated memory usage of mmproj is %.2f MiB\n", total / (1024.0 * 1024.0)); + SRV_INF("[mtmd] estimated worst-case memory usage of mmproj is %.2f MiB\n", total / (1024.0 * 1024.0)); GGML_ASSERT(!params_base.fit_params_target.empty()); for (auto & [dev, size] : mmproj_mem) { for (size_t i = 0; i < ggml_backend_dev_count(); i++) { @@ -793,6 +794,84 @@ private: } } + // optionally reserve VRAM for the draft / MTP context before fitting the target model + if (params_base.fit_params) { + const bool spec_mtp = std::find(params_base.speculative.types.begin(), + params_base.speculative.types.end(), + COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != params_base.speculative.types.end(); + const bool has_draft = params_base.speculative.has_dft(); + + if (has_draft || spec_mtp) { + common_params params_dft = params_base; + bool measure_model_bytes = true; + + if (has_draft) { + const auto & params_spec = params_base.speculative.draft; + params_dft.devices = params_spec.devices; + params_dft.model = params_spec.mparams; + params_dft.n_gpu_layers = params_spec.n_gpu_layers; + params_dft.cache_type_k = params_spec.cache_type_k; + params_dft.cache_type_v = params_spec.cache_type_v; + params_dft.tensor_buft_overrides = params_spec.tensor_buft_overrides; + } else { + // MTP draft context lives on the target model, only context+compute are new + measure_model_bytes = false; + } + + auto mparams_dft = common_model_params_to_llama(params_dft); + auto cparams_dft = common_context_params_to_llama(params_dft); + if (spec_mtp) { + cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP; + cparams_dft.type_k = params_base.speculative.draft.cache_type_k; + cparams_dft.type_v = params_base.speculative.draft.cache_type_v; + } + cparams_dft.n_rs_seq = 0; + + std::vector devs; + uint32_t hp_ngl = 0; + uint32_t hp_nct = 0; + uint32_t hp_nex = 0; + try { + auto dmd = common_get_device_memory_data( + params_dft.model.path.c_str(), &mparams_dft, &cparams_dft, + devs, hp_ngl, hp_nct, hp_nex, GGML_LOG_LEVEL_ERROR); + + GGML_ASSERT(!params_base.fit_params_target.empty()); + size_t total = 0; + + std::vector tgt_devices = params.devices; + + if (tgt_devices.empty()) { + for(size_t i = 0; i < ggml_backend_dev_count(); ++i) { + tgt_devices.push_back(ggml_backend_dev_get(i)); + } + } + + for (size_t j = 0; j < devs.size(); ++j) { + const size_t bytes = + (measure_model_bytes ? dmd[j].mb.model : 0) + + dmd[j].mb.context + + dmd[j].mb.compute; + total += bytes; + for (size_t i = 0; i < tgt_devices.size(); i++) { + if (tgt_devices[i] == devs[j]) { + SRV_DBG("[spec] adding %.2f MiB to fit_params_target for device %s\n", + bytes / (1024.0 * 1024.0), ggml_backend_dev_name(devs[j])); + params_base.fit_params_target[i] += bytes; + break; + } + } + } + SRV_INF("[spec] estimated memory usage of %s is %.2f MiB\n", + has_draft ? "draft model" : "MTP context", + total / (1024.0 * 1024.0)); + } catch (const std::exception & e) { + SRV_ERR("[spec] failed to measure %s memory: %s\n", + has_draft ? "draft model" : "MTP context", e.what()); + } + } + } + llama_init = common_init_from_params(params_base); model_tgt = llama_init->model(); @@ -863,6 +942,8 @@ private: auto cparams_mtp = common_context_params_to_llama(params_base); cparams_mtp.ctx_type = LLAMA_CONTEXT_TYPE_MTP; + cparams_mtp.type_k = params_base.speculative.draft.cache_type_k; + cparams_mtp.type_v = params_base.speculative.draft.cache_type_v; cparams_mtp.n_rs_seq = 0; ctx_dft.reset(llama_init_from_model(model_tgt, cparams_mtp)); @@ -1026,6 +1107,13 @@ private: } SRV_INF("%s", "for more info see https://github.com/ggml-org/llama.cpp/pull/16391\n"); + if (params_base.n_ctx_checkpoints > 0) { + SRV_INF("context checkpoints enabled, max = %d, min spacing = %d\n", + params_base.n_ctx_checkpoints, params_base.checkpoint_min_step); + } else { + SRV_INF("%s", "context checkpoints disabled\n"); + } + if (!params_base.model_alias.empty()) { // backward compat: use first alias as model name model_name = *params_base.model_alias.begin(); @@ -2681,8 +2769,6 @@ private: } if (pos_min >= pos_min_thold) { - SLT_WRN(slot, "n_past = %d, slot.prompt.tokens.size() = %d, seq_id = %d, pos_min = %d, n_swa = %d\n", n_past, (int) slot.prompt.tokens.size(), slot.id, pos_min, n_swa); - // search for a context checkpoint const auto it = std::find_if( slot.prompt.checkpoints.rbegin(), @@ -2699,7 +2785,6 @@ private: if (!do_reset) { // restore the context checkpoint - it->load_tgt(ctx_tgt, slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY); it->load_dft(ctx_dft.get(), slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY); @@ -2835,6 +2920,9 @@ private: has_mtmd = true; } + const int32_t n_before_user = slot.task->params.n_before_user; + const bool n_before_user_known = n_before_user > 0; + // add prompt tokens for processing in the current batch while (slot.prompt.n_tokens() < slot.task->n_tokens() && batch.n_tokens < n_batch) { // get next token to process @@ -2863,6 +2951,13 @@ private: slot.n_prompt_tokens_processed++; + // stop the prompt batch exactly before the latest user input, so a checkpoint + // can be created after the previous messages + if (n_before_user_known && + slot.prompt.n_tokens() == n_before_user) { + break; + } + // process the last few tokens of the prompt separately in order to allow for a checkpoint to be created. // create checkpoints that many tokens before the end of the prompt: // - 4 + n_ubatch @@ -2888,6 +2983,8 @@ private: // the number of tokens added to the batch for the current slot const auto n_tokens_cur = batch.n_tokens - n_tokens_prev; + const bool near_prompt_end = slot.task->n_tokens() < slot.prompt.n_tokens() + n_ubatch; + // entire prompt has been processed if (slot.prompt.n_tokens() == slot.task->n_tokens()) { slot.state = SLOT_STATE_DONE_PROMPT; @@ -2902,39 +2999,49 @@ private: slot.init_sampler(); } else { - if (slot.task->n_tokens() < slot.prompt.n_tokens() + n_ubatch) { - // near the end of the prompt - do_checkpoint = do_checkpoint && true; - } else { - // only do non-end checkpoints if the "checkpoint every n tokens" option is set - do_checkpoint = do_checkpoint && params_base.checkpoint_every_nt > 0; - - if (do_checkpoint) { - llama_pos last_checkpoint = 0; - if (!slot.prompt.checkpoints.empty()) { - last_checkpoint = slot.prompt.checkpoints.back().n_tokens; - } - - do_checkpoint = do_checkpoint && slot.prompt.n_tokens() - batch.n_tokens - last_checkpoint >= params_base.checkpoint_every_nt; - - if (do_checkpoint) { - SLT_INF(slot, "%d tokens since last checkpoint at %d, creating new checkpoint during processing at position %d\n", params_base.checkpoint_every_nt, last_checkpoint, slot.prompt.n_tokens()); - } - } + // skip ordinary mid-prompt checkpoints + if (!n_before_user_known && !near_prompt_end) { + do_checkpoint = false; } } const auto pos_min = llama_memory_seq_pos_min(llama_get_memory(ctx_tgt), slot.id); const auto pos_max = llama_memory_seq_pos_max(llama_get_memory(ctx_tgt), slot.id); - // no need for empty or small checkpoints - do_checkpoint = do_checkpoint && (pos_min >= 0 && slot.prompt.n_tokens() >= 64); + // checkpoints are created before the current batch is decoded, so + // their token position is the batch start rather than the prompt end + const int32_t n_tokens_start = slot.prompt.n_tokens() - n_tokens_cur; + + { + const bool is_on_user = + n_before_user_known && + n_tokens_start == n_before_user; + + const bool is_after_user = + n_before_user_known && + n_tokens_start > n_before_user; + + const bool is_allowed = + !n_before_user_known || + is_on_user || + (is_after_user && near_prompt_end); + + if (do_checkpoint && !is_allowed) { + do_checkpoint = false; + } + } + + // nothing to checkpoint yet + // TODO: is this check needed? + if (do_checkpoint && pos_min < 0) { + do_checkpoint = false; + } // do not checkpoint after mtmd chunks do_checkpoint = do_checkpoint && !has_mtmd; // no need to create checkpoints that are too close together - do_checkpoint = do_checkpoint && (slot.prompt.checkpoints.empty() || slot.prompt.n_tokens() - n_tokens_cur > slot.prompt.checkpoints.back().n_tokens + 64); + do_checkpoint = do_checkpoint && (slot.prompt.checkpoints.empty() || n_tokens_start > slot.prompt.checkpoints.back().n_tokens + params_base.checkpoint_min_step); SLT_DBG(slot, "main/do_checkpoint = %s, pos_min = %d, pos_max = %d\n", do_checkpoint ? "yes" : "no", pos_min, pos_max); // note: we create the checkpoint before calling llama_decode(), so the current batch is not @@ -3451,6 +3558,53 @@ void server_context::on_sleeping_changed(std::function callback) { impl->queue_tasks.on_sleeping_state(std::move(callback)); } +// compute the number of tokens before the last user message in the prompt +static int32_t prompt_get_n_before_user( + const json & message_spans, + const std::string & prompt, + const std::vector & files, + const llama_vocab * vocab, + mtmd_context * mctx) { + int32_t result = -1; + int32_t byte_pos = -1; + + for (const auto & span : message_spans) { + const std::string role = json_value(span, "role", std::string()); + + if (role == "user") { + byte_pos = json_value(span, "pos", -1); + } + } + + if (byte_pos >= 0) { + GGML_ASSERT((size_t) byte_pos <= prompt.size()); + + const std::string prefix = prompt.substr(0, (size_t) byte_pos); + + const std::string marker = get_media_marker(); + size_t n_prefix_media = 0; + for (size_t pos = 0; (pos = prefix.find(marker, pos)) != std::string::npos; pos += marker.size()) { + n_prefix_media++; + } + + GGML_ASSERT(n_prefix_media <= files.size()); + + if (mctx != nullptr && n_prefix_media > 0) { + // TODO: this makes a copy - avoid it + std::vector prefix_files(files.begin(), files.begin() + n_prefix_media); + + result = (int32_t) process_mtmd_prompt(mctx, prefix, prefix_files).size(); + } else { + result = (int32_t) tokenize_input_prompts(vocab, nullptr, prefix, true, true)[0].size(); + } + + SRV_TRC("message_spans: last user message: byte_pos=%d, media=%zu, n_before_user=%d\n", + byte_pos, n_prefix_media, result); + } + + return result; +} + // // server_routes @@ -3500,6 +3654,18 @@ std::unique_ptr server_routes::handle_completions_impl( meta->slot_n_ctx, meta->logit_bias_eog, data); + + const auto message_spans = json_value(data, "message_spans", json::array()); + if (prompt.is_string() && message_spans.is_array()) { + task.params.n_before_user = + prompt_get_n_before_user( + message_spans, + prompt.get(), + files, + ctx_server.vocab, + ctx_server.mctx); + } + task.id_slot = json_value(data, "id_slot", -1); // OAI-compat diff --git a/tools/server/server-task.h b/tools/server/server-task.h index 0978bb6ff..60e216e79 100644 --- a/tools/server/server-task.h +++ b/tools/server/server-task.h @@ -61,6 +61,9 @@ struct task_params { int32_t n_cache_reuse = 0; // min chunk size to attempt reusing from the cache via KV shifting (0 = disabled) + // number of prompt tokens before the latest user message + int32_t n_before_user = -1; + int64_t t_max_prompt_ms = -1; // TODO: implement int64_t t_max_predict_ms = -1; // if positive, limit the generation phase to this time limit diff --git a/tools/ui/CMakeLists.txt b/tools/ui/CMakeLists.txt index d4cf35802..60d9020da 100644 --- a/tools/ui/CMakeLists.txt +++ b/tools/ui/CMakeLists.txt @@ -43,9 +43,9 @@ if(CMAKE_CROSSCOMPILING) message(STATUS "UI: building llama-ui-embed with host compiler ${HOST_CXX_COMPILER}") if(CMAKE_HOST_WIN32) - set(LLAMA_UI_EMBED_EXE "${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed.exe") + set(LLAMA_UI_EMBED_EXE "${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed-host.exe") else() - set(LLAMA_UI_EMBED_EXE "${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed") + set(LLAMA_UI_EMBED_EXE "${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed-host") endif() add_custom_command( @@ -56,6 +56,8 @@ if(CMAKE_CROSSCOMPILING) COMMENT "Building llama-ui-embed (host)" VERBATIM ) + + # phony target to tie it into the dependency graph add_custom_target(llama-ui-embed DEPENDS "${LLAMA_UI_EMBED_EXE}") else() add_executable(llama-ui-embed embed.cpp) @@ -93,6 +95,10 @@ add_library(${TARGET} STATIC ${UI_CPP} ${UI_H}) target_compile_features(${TARGET} PRIVATE cxx_std_17) add_dependencies(${TARGET} llama-ui-assets) +if (BUILD_SHARED_LIBS) + set_target_properties(${TARGET} PROPERTIES POSITION_INDEPENDENT_CODE ON) +endif() + target_include_directories(${TARGET} PUBLIC ${CMAKE_CURRENT_BINARY_DIR} ) diff --git a/tools/ui/src/lib/constants/image-size.ts b/tools/ui/src/lib/constants/image-size.ts index 0d0c4e484..bcc1da08b 100644 --- a/tools/ui/src/lib/constants/image-size.ts +++ b/tools/ui/src/lib/constants/image-size.ts @@ -1 +1 @@ -export const MEGAPIXELS_TO_PIXELS = 1_000_000; \ No newline at end of file +export const MEGAPIXELS_TO_PIXELS = 1_000_000; diff --git a/tools/ui/src/lib/services/chat.service.ts b/tools/ui/src/lib/services/chat.service.ts index ae29206e2..b7c4a6307 100644 --- a/tools/ui/src/lib/services/chat.service.ts +++ b/tools/ui/src/lib/services/chat.service.ts @@ -883,14 +883,6 @@ export class ChatService { }); } - if (message.content) { - contentParts.push({ - type: ContentPartType.TEXT, - text: message.content - }); - } - - // Include images from all messages const imageFiles = message.extra.filter( (extra: DatabaseMessageExtra): extra is DatabaseMessageExtraImageFile => extra.type === AttachmentType.IMAGE @@ -923,6 +915,13 @@ export class ChatService { }); } + if (message.content) { + contentParts.push({ + type: ContentPartType.TEXT, + text: message.content + }); + } + const videoFiles = message.extra.filter( (extra: DatabaseMessageExtra): extra is DatabaseMessageExtraVideoFile => extra.type === AttachmentType.VIDEO diff --git a/tools/ui/src/lib/utils/cap-img-size.ts b/tools/ui/src/lib/utils/cap-img-size.ts index fa2af53f6..1ba218316 100644 --- a/tools/ui/src/lib/utils/cap-img-size.ts +++ b/tools/ui/src/lib/utils/cap-img-size.ts @@ -14,9 +14,8 @@ export function capImageDataURLSize( ): Promise { return new Promise((resolve, reject) => { try { - const mimeMatch = base64UrlImage.match(BASE64_IMAGE_URI_REGEX); - + if (!mimeMatch) { return reject(new Error('Invalid data URL format.')); } diff --git a/vendor/cpp-httplib/httplib.cpp b/vendor/cpp-httplib/httplib.cpp index b28549607..4ac497d03 100644 --- a/vendor/cpp-httplib/httplib.cpp +++ b/vendor/cpp-httplib/httplib.cpp @@ -1567,7 +1567,7 @@ void mmap::close() { #endif size_ = 0; } -int close_socket(socket_t sock) { +int close_socket(socket_t sock) noexcept { #ifdef _WIN32 return closesocket(sock); #else @@ -1794,7 +1794,7 @@ bool process_client_socket( return callback(strm); } -int shutdown_socket(socket_t sock) { +int shutdown_socket(socket_t sock) noexcept { #ifdef _WIN32 return shutdown(sock, SD_BOTH); #else @@ -7149,7 +7149,7 @@ void Server::wait_until_ready() const { } } -void Server::stop() { +void Server::stop() noexcept { if (is_running_) { assert(svr_sock_ != INVALID_SOCKET); std::atomic sock(svr_sock_.exchange(INVALID_SOCKET)); @@ -12290,9 +12290,18 @@ bool enumerate_windows_system_certs(Callback cb) { template bool enumerate_macos_keychain_certs(Callback cb) { bool loaded = false; - CFArrayRef certs = nullptr; - OSStatus status = SecTrustCopyAnchorCertificates(&certs); - if (status == errSecSuccess && certs) { + const SecTrustSettingsDomain domains[] = { + kSecTrustSettingsDomainSystem, + kSecTrustSettingsDomainAdmin, + kSecTrustSettingsDomainUser, + }; + for (auto domain : domains) { + CFArrayRef certs = nullptr; + OSStatus status = SecTrustSettingsCopyCertificates(domain, &certs); + if (status != errSecSuccess || !certs) { + if (certs) CFRelease(certs); + continue; + } CFIndex count = CFArrayGetCount(certs); for (CFIndex i = 0; i < count; i++) { SecCertificateRef cert = @@ -12655,28 +12664,36 @@ bool load_system_certs(ctx_t ctx) { auto store = SSL_CTX_get_cert_store(ssl_ctx); if (!store) return false; - CFArrayRef certs = nullptr; - if (SecTrustCopyAnchorCertificates(&certs) != errSecSuccess || !certs) { - return SSL_CTX_set_default_verify_paths(ssl_ctx) == 1; - } - bool loaded_any = false; - auto count = CFArrayGetCount(certs); - for (CFIndex i = 0; i < count; i++) { - auto cert = reinterpret_cast( - const_cast(CFArrayGetValueAtIndex(certs, i))); - CFDataRef der = SecCertificateCopyData(cert); - if (der) { - const unsigned char *data = CFDataGetBytePtr(der); - auto x509 = d2i_X509(nullptr, &data, CFDataGetLength(der)); - if (x509) { - if (X509_STORE_add_cert(store, x509) == 1) { loaded_any = true; } - X509_free(x509); - } - CFRelease(der); + const SecTrustSettingsDomain domains[] = { + kSecTrustSettingsDomainSystem, + kSecTrustSettingsDomainAdmin, + kSecTrustSettingsDomainUser, + }; + for (auto domain : domains) { + CFArrayRef certs = nullptr; + if (SecTrustSettingsCopyCertificates(domain, &certs) != errSecSuccess || + !certs) { + if (certs) CFRelease(certs); + continue; } + auto count = CFArrayGetCount(certs); + for (CFIndex i = 0; i < count; i++) { + auto cert = reinterpret_cast( + const_cast(CFArrayGetValueAtIndex(certs, i))); + CFDataRef der = SecCertificateCopyData(cert); + if (der) { + const unsigned char *data = CFDataGetBytePtr(der); + auto x509 = d2i_X509(nullptr, &data, CFDataGetLength(der)); + if (x509) { + if (X509_STORE_add_cert(store, x509) == 1) { loaded_any = true; } + X509_free(x509); + } + CFRelease(der); + } + } + CFRelease(certs); } - CFRelease(certs); return loaded_any || SSL_CTX_set_default_verify_paths(ssl_ctx) == 1; #else return SSL_CTX_set_default_verify_paths(ssl_ctx) == 1; diff --git a/vendor/cpp-httplib/httplib.h b/vendor/cpp-httplib/httplib.h index af750cdd9..536f0cb4d 100644 --- a/vendor/cpp-httplib/httplib.h +++ b/vendor/cpp-httplib/httplib.h @@ -8,8 +8,8 @@ #ifndef CPPHTTPLIB_HTTPLIB_H #define CPPHTTPLIB_HTTPLIB_H -#define CPPHTTPLIB_VERSION "0.45.0" -#define CPPHTTPLIB_VERSION_NUM "0x002d00" +#define CPPHTTPLIB_VERSION "0.45.1" +#define CPPHTTPLIB_VERSION_NUM "0x002d01" #ifdef _WIN32 #if defined(_WIN32_WINNT) && _WIN32_WINNT < 0x0A00 @@ -339,16 +339,26 @@ using socket_t = int; #include // On macOS with a TLS backend, enable Keychain root certificates by default -// unless the user explicitly opts out. +// unless the user explicitly opts out. Not enabled on iOS/tvOS/watchOS since +// the SecTrustSettings APIs used to enumerate anchor certificates are macOS +// only; on those platforms the user must provide a CA bundle explicitly. #if defined(__APPLE__) && defined(__clang__) && \ !defined(CPPHTTPLIB_DISABLE_MACOSX_AUTOMATIC_ROOT_CERTIFICATES) && \ (defined(CPPHTTPLIB_OPENSSL_SUPPORT) || \ defined(CPPHTTPLIB_MBEDTLS_SUPPORT) || \ defined(CPPHTTPLIB_WOLFSSL_SUPPORT)) +#if TARGET_OS_OSX #ifndef CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN #define CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN #endif #endif +#endif + +#if defined(CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN) && \ + defined(__APPLE__) && !TARGET_OS_OSX +#error \ + "CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN is only supported on macOS. On iOS/tvOS/watchOS, supply a CA bundle via set_ca_cert_path()." +#endif // On Windows, enable Schannel certificate verification by default // unless the user explicitly opts out. @@ -382,7 +392,7 @@ using socket_t = int; #endif // _WIN32 #ifdef CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN -#if TARGET_OS_MAC +#if TARGET_OS_OSX #include #endif #endif @@ -430,7 +440,7 @@ using socket_t = int; #endif #endif // _WIN32 #ifdef CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN -#if TARGET_OS_MAC +#if TARGET_OS_OSX #include #endif #endif @@ -473,7 +483,7 @@ using socket_t = int; #endif #endif // _WIN32 #ifdef CPPHTTPLIB_USE_CERTS_FROM_MACOSX_KEYCHAIN -#if TARGET_OS_MAC +#if TARGET_OS_OSX #include #endif #endif @@ -1597,7 +1607,7 @@ private: std::regex regex_; }; -int close_socket(socket_t sock); +int close_socket(socket_t sock) noexcept; ssize_t write_headers(Stream &strm, const Headers &headers); @@ -1734,7 +1744,7 @@ public: bool is_running() const; void wait_until_ready() const; - void stop(); + void stop() noexcept; void decommission(); std::function new_task_queue; @@ -3028,8 +3038,6 @@ bool parse_range_header(const std::string &s, Ranges &ranges); bool parse_accept_header(const std::string &s, std::vector &content_types); -int close_socket(socket_t sock); - ssize_t send_socket(socket_t sock, const void *ptr, size_t size, int flags); ssize_t read_socket(socket_t sock, void *ptr, size_t size, int flags);