mirror of
https://github.com/LostRuins/koboldcpp.git
synced 2026-07-10 01:18:32 +00:00
Merge commit 'ae251b5ff2' into concedo_experimental
# Conflicts: # .github/actions/linux-setup-spacemit/action.yml # .github/actions/unarchive-tar/action.yml # .github/workflows/build-android.yml # .github/workflows/build-cmake-pkg.yml # .github/workflows/build-cross.yml # .github/workflows/build-self-hosted.yml # .github/workflows/build.yml # .github/workflows/check-vendor.yml # .github/workflows/code-style.yml # .github/workflows/editorconfig.yml # .github/workflows/pre-tokenizer-hashes.yml # .github/workflows/python-check-requirements.yml # .github/workflows/python-lint.yml # .github/workflows/python-type-check.yml # .github/workflows/server-self-hosted.yml # .github/workflows/ui-build.yml # .github/workflows/ui.yml # .github/workflows/update-ops-docs.yml # ci/run.sh # docs/build-riscv64-spacemit.md # examples/convert_legacy_llama.py # ggml/cmake/ggml-config.cmake.in # ggml/src/CMakeLists.txt # ggml/src/ggml-cpu/CMakeLists.txt # ggml/src/ggml-opencl/ggml-opencl.cpp # scripts/sync_vendor.py # tests/test-chat-auto-parser.cpp # tests/test-chat.cpp # tests/test-gguf.cpp # tools/cli/README.md # tools/perplexity/perplexity.cpp # tools/server/README.md
This commit is contained in:
Concedo
commit
9204f78926
26 changed files with 1134 additions and 268 deletions
|
|
@ -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)"
|
||||
|
|
|
|||
|
|
@ -310,6 +310,8 @@ std::vector<segment> prune_whitespace_segments(const std::vector<segment> & 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_result> 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;
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
||||
|
|
|
|||
|
|
@ -8,6 +8,9 @@
|
|||
#include "peg-parser.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <cctype>
|
||||
#include <ostream>
|
||||
#include <sstream>
|
||||
|
||||
#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<std::function<void(const common_chat_template & tmpl, autopar
|
|||
analysis.content.end = "<|END_OF_TURN_TOKEN|>";
|
||||
analysis.preserved_tokens.push_back("<|CHATBOT_TOKEN|>");
|
||||
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<std::function<void(const common_chat_template & tmpl, autopar
|
|||
analysis.tools.function.close = "```";
|
||||
LOG_DBG(ANSI_ORANGE "[Patch: DeepSeek-R1-Distill-Qwen]\n" ANSI_RESET);
|
||||
}
|
||||
}
|
||||
},
|
||||
// Nemotron Nano v2
|
||||
[](const common_chat_template & tmpl, autoparser & analysis) -> void {
|
||||
if (tmpl.src.find("<SPECIAL_10>") != std::string::npos && tmpl.src.find("<SPECIAL_11>") != std::string::npos &&
|
||||
tmpl.src.find("<SPECIAL_12>") != std::string::npos && tmpl.src.find("<TOOL_RESPONSE>") != 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 = "<TOOLCALL>";
|
||||
analysis.tools.format.per_call_end = "</TOOLCALL>";
|
||||
analysis.content.mode = content_mode::PLAIN;
|
||||
analysis.content.start = "";
|
||||
analysis.content.end = "";
|
||||
analysis.reasoning.mode = reasoning_mode::TAG_BASED;
|
||||
analysis.reasoning.start = "<think>\n\n";
|
||||
analysis.reasoning.end = "</think>";
|
||||
analysis.assistant_start = "<SPECIAL_11>Assistant";
|
||||
analysis.user_start = "<SPECIAL_11>User";
|
||||
analysis.preserved_tokens.clear();
|
||||
analysis.preserved_tokens.push_back("<SPECIAL_12>");
|
||||
analysis.preserved_tokens.push_back("<SPECIAL_11>");
|
||||
analysis.preserved_tokens.push_back("</think>");
|
||||
analysis.preserved_tokens.push_back("<TOOLCALL>");
|
||||
analysis.preserved_tokens.push_back("</TOOLCALL>");
|
||||
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);
|
||||
|
|
|
|||
|
|
@ -101,6 +101,45 @@ std::string common_chat_msg::render_content(const std::string & delimiter) const
|
|||
return text;
|
||||
}
|
||||
|
||||
std::vector<common_chat_msg_span> common_chat_split_by_role(const std::string & prompt, const std::vector<common_chat_msg_delimiter> & delims) {
|
||||
if (delims.empty() || prompt.empty()) {
|
||||
return {};
|
||||
}
|
||||
|
||||
auto parser = build_peg_parser([&](common_peg_parser_builder & p) {
|
||||
std::vector<std::string> all_delims;
|
||||
std::vector<common_peg_parser> 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<common_chat_msg_span> 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<common_chat_msg_delimiter> 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);
|
||||
|
|
|
|||
|
|
@ -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<std::string> preserved_tokens;
|
||||
std::vector<std::string> additional_stops;
|
||||
std::string parser;
|
||||
std::vector<common_chat_msg_span> message_spans;
|
||||
};
|
||||
|
||||
// per-message parsing syntax
|
||||
|
|
@ -304,6 +316,7 @@ std::optional<common_chat_params> 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_msg_span> common_chat_split_by_role(const std::string & prompt, const std::vector<common_chat_msg_delimiter> & delims);
|
||||
|
||||
|
|
|
|||
|
|
@ -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<std::vector<size_t>> dp(a.size() + 1, std::vector<size_t>(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;
|
||||
|
||||
|
|
|
|||
|
|
@ -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<std::string> & values, const std::string & separator);
|
||||
std::vector<std::string> string_split(const std::string & str, const std::string & delimiter);
|
||||
|
|
|
|||
|
|
@ -26,7 +26,7 @@ class common_params_fit_exception : public std::runtime_error {
|
|||
using std::runtime_error::runtime_error;
|
||||
};
|
||||
|
||||
static std::vector<llama_device_memory_data> common_get_device_memory_data(
|
||||
std::vector<llama_device_memory_data> common_get_device_memory_data(
|
||||
const char * path_model,
|
||||
const llama_model_params * mparams,
|
||||
const llama_context_params * cparams,
|
||||
|
|
|
|||
16
common/fit.h
16
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 <vector>
|
||||
|
||||
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<llama_device_memory_data> common_get_device_memory_data(
|
||||
const char * path_model,
|
||||
const struct llama_model_params * mparams,
|
||||
const struct llama_context_params * cparams,
|
||||
std::vector<ggml_backend_dev_t> & devs,
|
||||
uint32_t & hp_ngl,
|
||||
uint32_t & hp_n_ctx_train,
|
||||
uint32_t & hp_n_expert,
|
||||
enum ggml_log_level log_level);
|
||||
|
|
|
|||
|
|
@ -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);
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -13,6 +13,7 @@
|
|||
#include <cstring>
|
||||
#include <map>
|
||||
#include <memory>
|
||||
#include <set>
|
||||
#include <string>
|
||||
#include <tuple>
|
||||
#include <utility>
|
||||
|
|
@ -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<ggml_context_ptr> ctxs;
|
||||
std::map<const ggml_tensor *, std::vector<ggml_tensor *>> 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<ggml_backend_buffer_ptr> 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<const ggml_tensor *, bool>, std::pair<ggml_backend_meta_split_state, char[nbtc]>> split_state_cache;
|
||||
std::map< const ggml_tensor *, std::vector<ggml_tensor *>> 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<buffer_config> 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<ggml_backend_buffer_t> & 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<ggml_tensor *> 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<ggml_backend_buffer_t> 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<ggml_backend_buffer_t> 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<ggml_backend_buffer_t> 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,
|
||||
|
|
|
|||
|
|
@ -13,6 +13,10 @@
|
|||
#include <stdlib.h> // for qsort
|
||||
#include <stdio.h> // for GGML_ASSERT
|
||||
|
||||
#ifdef GGML_USE_OPENMP
|
||||
#include <omp.h>
|
||||
#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) {
|
||||
|
|
|
|||
|
|
@ -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 <typename T>
|
||||
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 <typename T>
|
||||
|
|
@ -354,24 +361,71 @@ struct gguf_reader {
|
|||
}
|
||||
}
|
||||
dst.resize(static_cast<size_t>(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_t>(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<uint8_t *>(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<void *>(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::vector<struct
|
|||
return true;
|
||||
}
|
||||
|
||||
struct gguf_context * gguf_init_from_file_ptr(FILE * file, struct gguf_init_params params) {
|
||||
if (!file) {
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
const struct gguf_reader gr(file);
|
||||
static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr, struct gguf_init_params params) {
|
||||
struct gguf_context * ctx = new gguf_context;
|
||||
|
||||
bool ok = true;
|
||||
|
|
@ -744,14 +793,14 @@ struct gguf_context * gguf_init_from_file_ptr(FILE * file, struct gguf_init_para
|
|||
GGML_ASSERT(int64_t(ctx->info.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<gguf_file_reader *>(userdata);
|
||||
|
||||
if (reader.offset != offset) {
|
||||
if (offset > INT64_MAX || gguf_fseek(reader.file, static_cast<int64_t>(offset), SEEK_SET) != 0) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
reader.offset = offset;
|
||||
}
|
||||
|
||||
const size_t nread = fread(static_cast<uint8_t *>(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<uint64_t>(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<gguf_buffer_reader *>(userdata);
|
||||
|
||||
if (offset > reader.size || len > reader.size - offset) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
const size_t data_offset = static_cast<size_t>(offset);
|
||||
const size_t nread = std::min(len, reader.size - data_offset);
|
||||
memcpy(static_cast<uint8_t *>(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<const uint8_t *>(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");
|
||||
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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<ggml_backend_dev_t> 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<ggml_backend_dev_t> 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<void(bool)> 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<raw_buffer> & 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<raw_buffer> 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_res_generator> 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<std::string>(),
|
||||
files,
|
||||
ctx_server.vocab,
|
||||
ctx_server.mctx);
|
||||
}
|
||||
|
||||
task.id_slot = json_value(data, "id_slot", -1);
|
||||
|
||||
// OAI-compat
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
||||
|
|
|
|||
|
|
@ -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}
|
||||
)
|
||||
|
|
|
|||
|
|
@ -1 +1 @@
|
|||
export const MEGAPIXELS_TO_PIXELS = 1_000_000;
|
||||
export const MEGAPIXELS_TO_PIXELS = 1_000_000;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -14,9 +14,8 @@ export function capImageDataURLSize(
|
|||
): Promise<string> {
|
||||
return new Promise((resolve, reject) => {
|
||||
try {
|
||||
|
||||
const mimeMatch = base64UrlImage.match(BASE64_IMAGE_URI_REGEX);
|
||||
|
||||
|
||||
if (!mimeMatch) {
|
||||
return reject(new Error('Invalid data URL format.'));
|
||||
}
|
||||
|
|
|
|||
67
vendor/cpp-httplib/httplib.cpp
vendored
67
vendor/cpp-httplib/httplib.cpp
vendored
|
|
@ -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<socket_t> sock(svr_sock_.exchange(INVALID_SOCKET));
|
||||
|
|
@ -12290,9 +12290,18 @@ bool enumerate_windows_system_certs(Callback cb) {
|
|||
template <typename Callback>
|
||||
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<SecCertificateRef>(
|
||||
const_cast<void *>(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<SecCertificateRef>(
|
||||
const_cast<void *>(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;
|
||||
|
|
|
|||
28
vendor/cpp-httplib/httplib.h
vendored
28
vendor/cpp-httplib/httplib.h
vendored
|
|
@ -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 <utility>
|
||||
|
||||
// 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 <Security/Security.h>
|
||||
#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 <Security/Security.h>
|
||||
#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 <Security/Security.h>
|
||||
#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<TaskQueue *(void)> 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<std::string> &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);
|
||||
|
|
|
|||
Loading…
Add table
Add a link
Reference in a new issue